diff --git a/Grid/lattice/Lattice_slicesum_core.h b/Grid/lattice/Lattice_slicesum_core.h index 5d29e64f..5d96c4fb 100644 --- a/Grid/lattice/Lattice_slicesum_core.h +++ b/Grid/lattice/Lattice_slicesum_core.h @@ -1,7 +1,6 @@ #pragma once #if defined(GRID_CUDA) - #include #define gpucub cub #define gpuError_t cudaError_t @@ -57,8 +56,13 @@ inline void sliceSumReduction_cub_small(const vobj *Data, //copy offsets to device acceleratorCopyToDeviceAsynch(&offsets[0],d_offsets,sizeof(int)*(rd+1),computeStream); +#if defined(__CUDACC__) && (__CUDACC_VER_MAJOR__ >= 13) + #define GRID_CUB_SUM_OP ::cuda::std::plus<>{} +#else + #define GRID_CUB_SUM_OP ::cub::Sum() +#endif - gpuError_t gpuErr = gpucub::DeviceSegmentedReduce::Reduce(temp_storage_array, temp_storage_bytes, rb_p,d_out, rd, d_offsets, d_offsets+1, ::gpucub::Sum(), zero_init, computeStream); + gpuError_t gpuErr = gpucub::DeviceSegmentedReduce::Reduce(temp_storage_array, temp_storage_bytes, rb_p,d_out, rd, d_offsets, d_offsets+1, GRID_CUB_SUM_OP, zero_init, computeStream); if (gpuErr!=gpuSuccess) { std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce (setup)! Error: " << gpuErr < +NAMESPACE_BEGIN(Grid); #ifdef __CUDA_ARCH__ #define GRID_SIMT