diff --git a/src/util/cuda/merge_ofp.cuh b/src/util/cuda/merge_ofp.cuh index 2c9250053ffe7d0b05b93cf788248703baebd36d..367da42881f5bfc7d5ffaa708b18d7c8f43d88de 100644 --- a/src/util/cuda/merge_ofp.cuh +++ b/src/util/cuda/merge_ofp.cuh @@ -24,11 +24,14 @@ #define __CUDACC__ #define __CUDA__ #else - #include "util/cuda/moderngpu/kernel_merge.hxx" + #include <thrust/merge.h> + #include <thrust/execution_policy.h> #endif #endif #else - #include "util/cuda/moderngpu/kernel_merge.hxx" + #include <thrust/merge.h> + #include <thrust/execution_policy.h> +// #include "util/cuda/moderngpu/kernel_merge.hxx" #endif #include "util/cuda/ofp_context.hxx" @@ -98,7 +101,14 @@ #else - mgpu::merge(a_keys,a_vals,a_count,b_keys,b_vals,b_count,c_keys,c_vals,comp,context); +// It seems broken on some CUDA on some hardware. Anyway is not anymore supported +// on some hardware ... we move to thrust +// mgpu::merge(a_keys,a_vals,a_count,b_keys,b_vals,b_count,c_keys,c_vals,comp,context); + + thrust::merge_by_key(thrust::device, a_keys,a_keys + a_count, + b_keys,b_keys + b_count, + a_vals,b_vals, + c_keys,c_vals,comp); #endif