From caa5f9772339ea5124759e4dc4c47af4ef0b01cd Mon Sep 17 00:00:00 2001 From: dbollweg Date: Wed, 31 Jan 2024 16:50:06 -0500 Subject: [PATCH 01/14] Add sliceSum gpu using cub/hipcub --- Grid/lattice/Lattice_reduction.h | 1 + Grid/lattice/Lattice_slicesum_gpu.h | 144 ++++++++++++++++++++++++++++ tests/core/Test_sliceSum.cc | 56 +++++++++++ 3 files changed, 201 insertions(+) create mode 100644 Grid/lattice/Lattice_slicesum_gpu.h create mode 100644 tests/core/Test_sliceSum.cc diff --git a/Grid/lattice/Lattice_reduction.h b/Grid/lattice/Lattice_reduction.h index a6bbcf15..f85ed7e3 100644 --- a/Grid/lattice/Lattice_reduction.h +++ b/Grid/lattice/Lattice_reduction.h @@ -27,6 +27,7 @@ Author: Christoph Lehner #if defined(GRID_CUDA)||defined(GRID_HIP) #include +#include #endif #if defined(GRID_SYCL) #include diff --git a/Grid/lattice/Lattice_slicesum_gpu.h b/Grid/lattice/Lattice_slicesum_gpu.h new file mode 100644 index 00000000..53d0afae --- /dev/null +++ b/Grid/lattice/Lattice_slicesum_gpu.h @@ -0,0 +1,144 @@ +#pragma once +#if defined(GRID_CUDA) + +#include +#define gpucub cub +#define gpuMalloc cudaMalloc +#define gpuMemcpy cudaMemcpy +#define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost +#define gpuError_t cudaError_t +#define gpuSuccess cudaSuccess + +#elif defined(GRID_HIP) + +#include +#define gpucub hipcub +#define gpuMalloc hipMalloc +#define gpuMemcpy hipMemcpy +#define gpuMemcpyDeviceToHost hipMemcpyDeviceToHost +#define gpuError_t hipError_t +#define gpuSuccess hipSuccess + +// extern hipStream_t computeStream; +#endif + + +NAMESPACE_BEGIN(Grid); + +template inline void sliceSumGpu(const Lattice &Data,std::vector &result,int orthogdim) +{ + + typedef typename vobj::scalar_object sobj; + typedef typename vobj::scalar_object::scalar_type scalar_type; + GridBase *grid = Data.Grid(); + assert(grid!=NULL); + + const int Nd = grid->_ndimension; + const int Nsimd = grid->Nsimd(); + + assert(orthogdim >= 0); + assert(orthogdim < Nd); + + int fd=grid->_fdimensions[orthogdim]; + int ld=grid->_ldimensions[orthogdim]; + int rd=grid->_rdimensions[orthogdim]; + + int e1= grid->_slice_nblock[orthogdim]; + int e2= grid->_slice_block [orthogdim]; + int stride=grid->_slice_stride[orthogdim]; + int ostride=grid->_ostride[orthogdim]; + Vector lvSum(rd); + Vector lsSum(ld,Zero()); + commVector reduction_buffer(e1*e2); + ExtractBuffer extracted(Nsimd); + + result.resize(fd); + for(int r=0;rNsimd(),{ //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream) + + int n = s / e2; + int b = s % e2; + int so=r*ostride; // base offset for start of plane + int ss= so+n*stride+b; + + coalescedWrite(rb_p[s], coalescedRead(Data_v[ss])); + + }); + + //issue reductions in computeStream + gpuErr =gpucub::DeviceReduce::Sum(helperArray, temp_storage_bytes, rb_p, &d_out[r], size, computeStream); + if (gpuErr!=gpuSuccess) { + std::cout << "Encountered error during cub::DeviceReduce::Sum(2)! Error: " << gpuErr <iCoorFromIindex(icoor,idx); + + int ldx =rt+icoor[orthogdim]*rd; + + lsSum[ldx]=lsSum[ldx]+extracted[idx]; + + } + } + + // sum over nodes. + for(int t=0;t_processor_coor[orthogdim] ) { + result[t]=lsSum[lt]; + } else { + result[t]=Zero(); + } + + } + scalar_type * ptr = (scalar_type *) &result[0]; + int words = fd*sizeof(sobj)/sizeof(scalar_type); + grid->GlobalSumVector(ptr, words); +} + +template inline +std::vector +sliceSumGpu(const Lattice &Data,int orthogdim) +{ + std::vector result; + sliceSumGpu(Data,result,orthogdim); + return result; +} + +NAMESPACE_END(Grid); \ No newline at end of file diff --git a/tests/core/Test_sliceSum.cc b/tests/core/Test_sliceSum.cc new file mode 100644 index 00000000..4a04b41c --- /dev/null +++ b/tests/core/Test_sliceSum.cc @@ -0,0 +1,56 @@ +#include + + +int main (int argc, char ** argv) { + + using namespace Grid; + + Grid_init(&argc,&argv); + + + Coordinate latt_size({64,64,64,16}); + auto simd_layout = GridDefaultSimd(Nd, vComplexD::Nsimd()); + auto mpi_layout = GridDefaultMpi(); + GridCartesian Grid(latt_size, simd_layout, mpi_layout); + + std::vector seeds({1, 2, 3, 4}); + + GridParallelRNG pRNG(&Grid); + pRNG.SeedFixedIntegers(seeds); + + LatticeComplexD test_data(&Grid); + gaussian(pRNG,test_data); + + std::vector reduction_reference; + std::vector reduction_result; + + //warmup + for (int sweeps = 0; sweeps < 5; sweeps++) { + sliceSumGpu(test_data,reduction_result,0); + } + + + for (int i = 0; i < Nd; i++) { + RealD t=-usecond(); + sliceSum(test_data,reduction_reference,i); + t+=usecond(); + std::cout << " sliceSum took "< Date: Thu, 1 Feb 2024 16:41:03 -0500 Subject: [PATCH 02/14] Use accelerator_for2d and DeviceSegmentedRecude to avoid kernel launch latencies --- Grid/lattice/Lattice_slicesum_gpu.h | 71 +++++++++++++++++++---------- Grid/perfmon/Tracing.h | 2 +- tests/core/Test_sliceSum.cc | 3 +- 3 files changed, 50 insertions(+), 26 deletions(-) diff --git a/Grid/lattice/Lattice_slicesum_gpu.h b/Grid/lattice/Lattice_slicesum_gpu.h index 53d0afae..d8927708 100644 --- a/Grid/lattice/Lattice_slicesum_gpu.h +++ b/Grid/lattice/Lattice_slicesum_gpu.h @@ -6,6 +6,7 @@ #define gpuMalloc cudaMalloc #define gpuMemcpy cudaMemcpy #define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost +#define gpuMemcpyHostToDevice cudaMemcpyHostToDevice #define gpuError_t cudaError_t #define gpuSuccess cudaSuccess @@ -16,6 +17,7 @@ #define gpuMalloc hipMalloc #define gpuMemcpy hipMemcpy #define gpuMemcpyDeviceToHost hipMemcpyDeviceToHost +#define gpuMemcpyHostToDevice hipMemcpyHostToDevice #define gpuError_t hipError_t #define gpuSuccess hipSuccess @@ -49,14 +51,15 @@ template inline void sliceSumGpu(const Lattice &Data,std::vect int ostride=grid->_ostride[orthogdim]; Vector lvSum(rd); Vector lsSum(ld,Zero()); - commVector reduction_buffer(e1*e2); + commVector reduction_buffer(rd*e1*e2); ExtractBuffer extracted(Nsimd); result.resize(fd); for(int r=0;r inline void sliceSumGpu(const Lattice &Data,std::vect vobj *d_out; size_t temp_storage_bytes = 0; size_t size = e1*e2; - gpuMalloc(&d_out,rd*sizeof(vobj)); - gpuError_t gpuErr =gpucub::DeviceReduce::Sum(helperArray, temp_storage_bytes, rb_p,d_out, size, computeStream); + std::vector offsets(rd+1,0); + for (int i = 0; i < offsets.size(); i++) { + offsets[i] = i*size; + } + int* d_offsets; + + gpuError_t gpuErr = gpuMalloc(&d_out,rd*sizeof(vobj)); + if (gpuErr != gpuSuccess) { + std::cout << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc(1) Error: " << gpuErr <Nsimd(),{ //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream) - - int n = s / e2; - int b = s % e2; - int so=r*ostride; // base offset for start of plane - int ss= so+n*stride+b; - coalescedWrite(rb_p[s], coalescedRead(Data_v[ss])); + //prepare buffer for reduction + accelerator_for2dNB( s,e1*e2, r,rd, grid->Nsimd(),{ //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream) + //use 2d accelerator_for to avoid launch latencies found when looping over rd + int n = s / e2; + int b = s % e2; + int so=r*ostride; // base offset for start of plane + int ss= so+n*stride+b; - }); - - //issue reductions in computeStream - gpuErr =gpucub::DeviceReduce::Sum(helperArray, temp_storage_bytes, rb_p, &d_out[r], size, computeStream); - if (gpuErr!=gpuSuccess) { - std::cout << "Encountered error during cub::DeviceReduce::Sum(2)! Error: " << gpuErr < offsets(rd+1,0); + + for (int i = 0; i < offsets.size(); i++) { + offsets[i] = i*subvol_size; } + //Allocate memory for output and offset arrays on device + gpuError_t gpuErr = gpuMalloc(&d_out,rd*sizeof(vobj)); + if (gpuErr != gpuSuccess) { + std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc (d_out)! Error: " << gpuErr <Nsimd(),{ //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream) - //use 2d accelerator_for to avoid launch latencies found when looping over rd + //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream) + //use 2d accelerator_for to avoid launch latencies found when serially looping over rd + + accelerator_for2dNB( s,subvol_size, r,rd, grid->Nsimd(),{ + int n = s / e2; int b = s % e2; int so=r*ostride; // base offset for start of plane int ss= so+n*stride+b; - coalescedWrite(rb_p[r*e1*e2+s], coalescedRead(Data_v[ss])); + coalescedWrite(rb_p[r*subvol_size+s], coalescedRead(Data_v[ss])); }); - //issue reductions in computeStream - gpuErr =gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p, d_out, rd, d_offsets, d_offsets+1,::gpucub::Sum(), identity, computeStream); + //issue segmented reductions in computeStream + gpuErr = gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p, d_out, rd, d_offsets, d_offsets+1,::gpucub::Sum(), vobj_zero, computeStream); if (gpuErr!=gpuSuccess) { - std::cout << "Lattice_slicesum_gpu.h: Encountered error during cub::DeviceReduce::Sum(2)! Error: " << gpuErr < Date: Thu, 1 Feb 2024 18:02:30 -0500 Subject: [PATCH 04/14] Fix cuda compilation of Lattice_slicesum_gpu.h --- Grid/lattice/Lattice_slicesum_gpu.h | 16 +++++++++------- 1 file changed, 9 insertions(+), 7 deletions(-) diff --git a/Grid/lattice/Lattice_slicesum_gpu.h b/Grid/lattice/Lattice_slicesum_gpu.h index 196956d1..8e13808f 100644 --- a/Grid/lattice/Lattice_slicesum_gpu.h +++ b/Grid/lattice/Lattice_slicesum_gpu.h @@ -4,8 +4,9 @@ #include #define gpucub cub #define gpuMalloc cudaMalloc -#define gpuMemcpyDtoHAsync cudaMemcpyDtoHAsync -#define gpuMemcpyHtoDAsync cudaMemcpyHtoDAsync +#define gpuMemcpyAsync cudaMemcpyAsync +#define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost +#define gpuMemcpyHostToDevice cudaMemcpyHostToDevice #define gpuError_t cudaError_t #define gpuSuccess cudaSuccess @@ -14,8 +15,9 @@ #include #define gpucub hipcub #define gpuMalloc hipMalloc -#define gpuMemcpyDtoHAsync hipMemcpyDtoHAsync -#define gpuMemcpyHtoDAsync hipMemcpyHtoDAsync +#define gpuMemcpyAsync hipMemcpyAsync +#define gpuMemcpyDeviceToHost hipMemcpyDeviceToHost +#define gpuMemcpyHostToDevice hipMemcpyHostToDevice #define gpuError_t hipError_t #define gpuSuccess hipSuccess @@ -71,7 +73,7 @@ template inline void sliceSumGpu(const Lattice &Data,std::vect int* d_offsets; std::vector offsets(rd+1,0); - + for (int i = 0; i < offsets.size(); i++) { offsets[i] = i*subvol_size; } @@ -90,7 +92,7 @@ template inline void sliceSumGpu(const Lattice &Data,std::vect } //copy offsets to device - gpuErr = gpuMemcpyHtoDAsync(d_offsets,&offsets[0],sizeof(int)*(rd+1),computeStream); + gpuErr = gpuMemcpyAsync(d_offsets,&offsets[0],sizeof(int)*(rd+1),gpuMemcpyHostToDevice,computeStream); if (gpuErr != gpuSuccess) { std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy (d_offsets)! Error: " << gpuErr < inline void sliceSumGpu(const Lattice &Data,std::vect exit(EXIT_FAILURE); } - gpuErr = gpuMemcpyDtoHAsync(&lvSum[0],d_out,rd*sizeof(vobj),computeStream); + gpuErr = gpuMemcpyAsync(&lvSum[0],d_out,rd*sizeof(vobj),gpuMemcpyDeviceToHost,computeStream); if (gpuErr!=gpuSuccess) { std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy (d_out)! Error: " << gpuErr < Date: Tue, 6 Feb 2024 13:24:45 -0500 Subject: [PATCH 05/14] work towards sliceSum for sycl backend --- Grid/lattice/Lattice_reduction.h | 15 ++++ Grid/lattice/Lattice_slicesum_gpu.h | 9 --- Grid/lattice/Lattice_slicesum_sycl.h | 115 +++++++++++++++++++++++++++ Grid/threads/Accelerator.h | 4 +- tests/core/Test_sliceSum.cc | 4 +- 5 files changed, 134 insertions(+), 13 deletions(-) create mode 100644 Grid/lattice/Lattice_slicesum_sycl.h diff --git a/Grid/lattice/Lattice_reduction.h b/Grid/lattice/Lattice_reduction.h index f85ed7e3..bfd41b6c 100644 --- a/Grid/lattice/Lattice_reduction.h +++ b/Grid/lattice/Lattice_reduction.h @@ -31,6 +31,7 @@ Author: Christoph Lehner #endif #if defined(GRID_SYCL) #include +#include #endif NAMESPACE_BEGIN(Grid); @@ -505,6 +506,20 @@ sliceSum(const Lattice &Data,int orthogdim) return result; } +template inline +std::vector +sliceSumGpu(const Lattice &Data,int orthogdim) +{ + std::vector result; + #if defined(GRID_CUDA) || defined(GRID_HIP) + sliceSumGpu(Data,result,orthogdim); + #elif defined(GRID_SYCL) + sliceSum_sycl(Data,result,orthogdim); + #endif + return result; +} + + template static void sliceInnerProductVector( std::vector & result, const Lattice &lhs,const Lattice &rhs,int orthogdim) { diff --git a/Grid/lattice/Lattice_slicesum_gpu.h b/Grid/lattice/Lattice_slicesum_gpu.h index 8e13808f..5d2ad049 100644 --- a/Grid/lattice/Lattice_slicesum_gpu.h +++ b/Grid/lattice/Lattice_slicesum_gpu.h @@ -177,13 +177,4 @@ template inline void sliceSumGpu(const Lattice &Data,std::vect grid->GlobalSumVector(ptr, words); } -template inline -std::vector -sliceSumGpu(const Lattice &Data,int orthogdim) -{ - std::vector result; - sliceSumGpu(Data,result,orthogdim); - return result; -} - NAMESPACE_END(Grid); \ No newline at end of file diff --git a/Grid/lattice/Lattice_slicesum_sycl.h b/Grid/lattice/Lattice_slicesum_sycl.h new file mode 100644 index 00000000..034e9dd3 --- /dev/null +++ b/Grid/lattice/Lattice_slicesum_sycl.h @@ -0,0 +1,115 @@ +#pragma once + +NAMESPACE_BEGIN(Grid); + +template +inline void sliceSum_sycl(const Lattice &Data, std::vector &result, int orthogdim) +{ + typedef typename vobj::scalar_object sobj; + typedef typename vobj::scalar_object::scalar_type scalar_type; + + GridBase *grid = Data.Grid(); + assert(grid!=NULL); + + const int Nd = grid->_ndimension; + const size_t Nsimd = grid->Nsimd(); + + + assert(orthogdim >= 0); + assert(orthogdim < Nd); + + int fd=grid->_fdimensions[orthogdim]; + int ld=grid->_ldimensions[orthogdim]; + int rd=grid->_rdimensions[orthogdim]; + + int e1= grid->_slice_nblock[orthogdim]; + int e2= grid->_slice_block [orthogdim]; + int stride=grid->_slice_stride[orthogdim]; + int ostride=grid->_ostride[orthogdim]; + size_t subvol_size = e1*e2; + + vobj *mysum = (vobj *) malloc_shared(sizeof(vobj),*theGridAccelerator); + vobj vobj_zero; + zeroit(vobj_zero); + + + result.resize(fd); + + Vector lvSum(rd); + Vector lsSum(ld,Zero()); + commVector reduction_buffer(rd*subvol_size); + ExtractBuffer extracted(Nsimd); + + for(int r=0;rsubmit([&](cl::sycl::handler &cgh) { + auto Reduction = cl::sycl::reduction(mysum,vobj_zero,std::plus<>()); + cgh.parallel_for(cl::sycl::range<1>{subvol_size}, + Reduction, + [=](cl::sycl::id<1> item, auto &sum) { + auto s = item[0]; + sum += rb_p[r*subvol_size+s]; + }); + }); + theGridAccelerator->wait(); + lvSum[r] = mysum[0]; + } + + Coordinate icoor(Nd); + + for(int rt=0;rtiCoorFromIindex(icoor,idx); + + int ldx =rt+icoor[orthogdim]*rd; + + lsSum[ldx]=lsSum[ldx]+extracted[idx]; + + } + } + + // sum over nodes. + for(int t=0;t_processor_coor[orthogdim] ) { + result[t]=lsSum[lt]; + } else { + result[t]=Zero(); + } + + } + scalar_type * ptr = (scalar_type *) &result[0]; + int words = fd*sizeof(sobj)/sizeof(scalar_type); + grid->GlobalSumVector(ptr, words); + +} + +NAMESPACE_END(Grid); \ No newline at end of file diff --git a/Grid/threads/Accelerator.h b/Grid/threads/Accelerator.h index ff5ccd7a..eaafea5d 100644 --- a/Grid/threads/Accelerator.h +++ b/Grid/threads/Accelerator.h @@ -256,12 +256,12 @@ NAMESPACE_END(Grid); #if 0 #include #include -#include +#include #include #else #include #include -#include +#include #include #endif diff --git a/tests/core/Test_sliceSum.cc b/tests/core/Test_sliceSum.cc index 7dd69cdc..399ab899 100644 --- a/tests/core/Test_sliceSum.cc +++ b/tests/core/Test_sliceSum.cc @@ -26,7 +26,7 @@ int main (int argc, char ** argv) { //warmup for (int sweeps = 0; sweeps < 5; sweeps++) { - sliceSumGpu(test_data,reduction_result,0); + reduction_result = sliceSumGpu(test_data,0); } int trace_id = traceStart("sliceSum benchmark"); @@ -46,7 +46,7 @@ int main (int argc, char ** argv) { RealD tgpu=-usecond(); tracePush("sliceSumGpu"); - sliceSumGpu(test_data,reduction_result,i); + reduction_result = sliceSumGpu(test_data,i); tracePop("sliceSumGpu"); tgpu+=usecond(); From 1514b4f13733f8952b232d1b95c00d7abb01d74a Mon Sep 17 00:00:00 2001 From: dbollweg Date: Tue, 6 Feb 2024 19:08:44 -0500 Subject: [PATCH 06/14] slicesum_sycl passes test --- Grid/lattice/Lattice_slicesum_sycl.h | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/Grid/lattice/Lattice_slicesum_sycl.h b/Grid/lattice/Lattice_slicesum_sycl.h index 034e9dd3..04ec8a6a 100644 --- a/Grid/lattice/Lattice_slicesum_sycl.h +++ b/Grid/lattice/Lattice_slicesum_sycl.h @@ -14,7 +14,6 @@ inline void sliceSum_sycl(const Lattice &Data, std::vector_ndimension; const size_t Nsimd = grid->Nsimd(); - assert(orthogdim >= 0); assert(orthogdim < Nd); @@ -29,9 +28,6 @@ inline void sliceSum_sycl(const Lattice &Data, std::vector &Data, std::vector lsSum(ld,Zero()); commVector reduction_buffer(rd*subvol_size); ExtractBuffer extracted(Nsimd); + vobj vobj_zero; + zeroit(vobj_zero); for(int r=0;r &Data, std::vector &Data, std::vectorsubmit([&](cl::sycl::handler &cgh) { - auto Reduction = cl::sycl::reduction(mysum,vobj_zero,std::plus<>()); + auto Reduction = cl::sycl::reduction(mysum,std::plus<>()); cgh.parallel_for(cl::sycl::range<1>{subvol_size}, Reduction, [=](cl::sycl::id<1> item, auto &sum) { From 9514035b87b09e5f6d4a52a8724a144621456a20 Mon Sep 17 00:00:00 2001 From: dbollweg Date: Fri, 9 Feb 2024 13:02:28 -0500 Subject: [PATCH 07/14] refactor slicesum: slicesum uses GPU version by default now --- Grid/lattice/Lattice_reduction.h | 33 +---- Grid/lattice/Lattice_slicesum_core.h | 204 +++++++++++++++++++++++++++ Grid/lattice/Lattice_slicesum_gpu.h | 180 ----------------------- Grid/lattice/Lattice_slicesum_sycl.h | 110 --------------- tests/core/Test_sliceSum.cc | 86 ++++++++++- 5 files changed, 289 insertions(+), 324 deletions(-) create mode 100644 Grid/lattice/Lattice_slicesum_core.h delete mode 100644 Grid/lattice/Lattice_slicesum_gpu.h delete mode 100644 Grid/lattice/Lattice_slicesum_sycl.h diff --git a/Grid/lattice/Lattice_reduction.h b/Grid/lattice/Lattice_reduction.h index bfd41b6c..66788a4c 100644 --- a/Grid/lattice/Lattice_reduction.h +++ b/Grid/lattice/Lattice_reduction.h @@ -27,12 +27,11 @@ Author: Christoph Lehner #if defined(GRID_CUDA)||defined(GRID_HIP) #include -#include #endif #if defined(GRID_SYCL) #include -#include #endif +#include NAMESPACE_BEGIN(Grid); @@ -450,19 +449,10 @@ template inline void sliceSum(const Lattice &Data,std::vector< int e1= grid->_slice_nblock[orthogdim]; int e2= grid->_slice_block [orthogdim]; int stride=grid->_slice_stride[orthogdim]; - - // sum over reduced dimension planes, breaking out orthog dir - // Parallel over orthog direction - autoView( Data_v, Data, CpuRead); - thread_for( r,rd, { - int so=r*grid->_ostride[orthogdim]; // base offset for start of plane - for(int n=0;n_ostride[orthogdim]; + + //Reduce Data down to lvSum + sliceSumReduction(Data,lvSum,rd, e1,e2,stride,ostride,Nsimd); // Sum across simd lanes in the plane, breaking out orthog dir. Coordinate icoor(Nd); @@ -506,19 +496,6 @@ sliceSum(const Lattice &Data,int orthogdim) return result; } -template inline -std::vector -sliceSumGpu(const Lattice &Data,int orthogdim) -{ - std::vector result; - #if defined(GRID_CUDA) || defined(GRID_HIP) - sliceSumGpu(Data,result,orthogdim); - #elif defined(GRID_SYCL) - sliceSum_sycl(Data,result,orthogdim); - #endif - return result; -} - template static void sliceInnerProductVector( std::vector & result, const Lattice &lhs,const Lattice &rhs,int orthogdim) diff --git a/Grid/lattice/Lattice_slicesum_core.h b/Grid/lattice/Lattice_slicesum_core.h new file mode 100644 index 00000000..2548884a --- /dev/null +++ b/Grid/lattice/Lattice_slicesum_core.h @@ -0,0 +1,204 @@ +#pragma once +#if defined(GRID_CUDA) + +#include +#define gpucub cub +#define gpuMalloc cudaMalloc +#define gpuMemcpyAsync cudaMemcpyAsync +#define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost +#define gpuMemcpyHostToDevice cudaMemcpyHostToDevice +#define gpuError_t cudaError_t +#define gpuSuccess cudaSuccess + +#elif defined(GRID_HIP) + +#include +#define gpucub hipcub +#define gpuMalloc hipMalloc +#define gpuMemcpyAsync hipMemcpyAsync +#define gpuMemcpyDeviceToHost hipMemcpyDeviceToHost +#define gpuMemcpyHostToDevice hipMemcpyHostToDevice +#define gpuError_t hipError_t +#define gpuSuccess hipSuccess + +#endif + + +NAMESPACE_BEGIN(Grid); + +#if defined(GRID_CUDA) || defined(GRID_HIP) +template inline void sliceSumReduction_cub(const Lattice &Data, Vector &lvSum, const int rd, const int e1, const int e2, const int stride, const int ostride, const int Nsimd) +{ + typedef typename vobj::scalar_object sobj; + + size_t subvol_size = e1*e2; + + commVector reduction_buffer(rd*subvol_size); + auto rb_p = &reduction_buffer[0]; + + vobj vobj_zero; //Need to provide initial value for reduction operation + zeroit(vobj_zero); + + + void *temp_storage_array = NULL; + size_t temp_storage_bytes = 0; + vobj *d_out; + int* d_offsets; + + std::vector offsets(rd+1,0); + + for (int i = 0; i < offsets.size(); i++) { + offsets[i] = i*subvol_size; + } + + //Allocate memory for output and offset arrays on device + gpuError_t gpuErr = gpuMalloc(&d_out,rd*sizeof(vobj)); + if (gpuErr != gpuSuccess) { + std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc (d_out)! Error: " << gpuErr < inline void sliceSumReduction_sycl(const Lattice &Data, Vector &lvSum, const int &rd, const int &e1, const int &e2, const int &stride, const int &ostride, const int &Nsimd) +{ + typedef typename vobj::scalar_object sobj; + size_t subvol_size = e1*e2; + + vobj *mysum = (vobj *) malloc_shared(sizeof(vobj),*theGridAccelerator); + vobj vobj_zero; + zeroit(vobj_zero); + + commVector reduction_buffer(rd*subvol_size); + + auto rb_p = &reduction_buffer[0]; + + autoView(Data_v, Data, AcceleratorRead); + + //prepare reduction buffer + accelerator_for2d( s,subvol_size, r,rd, (size_t)Nsimd,{ + + int n = s / e2; + int b = s % e2; + int so=r*ostride; // base offset for start of plane + int ss= so+n*stride+b; + + coalescedWrite(rb_p[r*subvol_size+s], coalescedRead(Data_v[ss])); + + }); + + for (int r = 0; r < rd; r++) { + mysum[0] = vobj_zero; //dirty hack: cannot pass vobj_zero as identity to sycl::reduction as its not device_copyable + theGridAccelerator->submit([&](cl::sycl::handler &cgh) { + auto Reduction = cl::sycl::reduction(mysum,std::plus<>()); + cgh.parallel_for(cl::sycl::range<1>{subvol_size}, + Reduction, + [=](cl::sycl::id<1> item, auto &sum) { + auto s = item[0]; + sum += rb_p[r*subvol_size+s]; + }); + }); + theGridAccelerator->wait(); + lvSum[r] = mysum[0]; + } + +} +#endif + +template inline void sliceSumReduction_cpu(const Lattice &Data, Vector &lvSum, const int &rd, const int &e1, const int &e2, const int &stride, const int &ostride, const int &Nsimd) +{ + // sum over reduced dimension planes, breaking out orthog dir + // Parallel over orthog direction + autoView( Data_v, Data, CpuRead); + thread_for( r,rd, { + int so=r*ostride; // base offset for start of plane + for(int n=0;n inline void sliceSumReduction(const Lattice &Data, Vector &lvSum, const int &rd, const int &e1, const int &e2, const int &stride, const int &ostride, const int &Nsimd) +{ + #if defined(GRID_CUDA) || defined(GRID_HIP) + + sliceSumReduction_cub(Data, lvSum, rd, e1, e2, stride, ostride, Nsimd); + + #elif defined(GRID_SYCL) + + sliceSumReduction_sycl(Data, lvSum, rd, e1, e2, stride, ostride, Nsimd); + + #else + sliceSumReduction_cpu(Data, lvSum, rd, e1, e2, stride, ostride, Nsimd); + + #endif +} + + +NAMESPACE_END(Grid); \ No newline at end of file diff --git a/Grid/lattice/Lattice_slicesum_gpu.h b/Grid/lattice/Lattice_slicesum_gpu.h deleted file mode 100644 index 5d2ad049..00000000 --- a/Grid/lattice/Lattice_slicesum_gpu.h +++ /dev/null @@ -1,180 +0,0 @@ -#pragma once -#if defined(GRID_CUDA) - -#include -#define gpucub cub -#define gpuMalloc cudaMalloc -#define gpuMemcpyAsync cudaMemcpyAsync -#define gpuMemcpyDeviceToHost cudaMemcpyDeviceToHost -#define gpuMemcpyHostToDevice cudaMemcpyHostToDevice -#define gpuError_t cudaError_t -#define gpuSuccess cudaSuccess - -#elif defined(GRID_HIP) - -#include -#define gpucub hipcub -#define gpuMalloc hipMalloc -#define gpuMemcpyAsync hipMemcpyAsync -#define gpuMemcpyDeviceToHost hipMemcpyDeviceToHost -#define gpuMemcpyHostToDevice hipMemcpyHostToDevice -#define gpuError_t hipError_t -#define gpuSuccess hipSuccess - -#endif - - -NAMESPACE_BEGIN(Grid); - -template inline void sliceSumGpu(const Lattice &Data,std::vector &result,int orthogdim) -{ - - typedef typename vobj::scalar_object sobj; - typedef typename vobj::scalar_object::scalar_type scalar_type; - GridBase *grid = Data.Grid(); - assert(grid!=NULL); - - const int Nd = grid->_ndimension; - const int Nsimd = grid->Nsimd(); - - assert(orthogdim >= 0); - assert(orthogdim < Nd); - - int fd=grid->_fdimensions[orthogdim]; - int ld=grid->_ldimensions[orthogdim]; - int rd=grid->_rdimensions[orthogdim]; - - int e1= grid->_slice_nblock[orthogdim]; - int e2= grid->_slice_block [orthogdim]; - int stride=grid->_slice_stride[orthogdim]; - int ostride=grid->_ostride[orthogdim]; - size_t subvol_size = e1*e2; - - Vector lvSum(rd); - Vector lsSum(ld,Zero()); - commVector reduction_buffer(rd*e1*e2); - ExtractBuffer extracted(Nsimd); - - result.resize(fd); - - for(int r=0;r offsets(rd+1,0); - - for (int i = 0; i < offsets.size(); i++) { - offsets[i] = i*subvol_size; - } - - //Allocate memory for output and offset arrays on device - gpuError_t gpuErr = gpuMalloc(&d_out,rd*sizeof(vobj)); - if (gpuErr != gpuSuccess) { - std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMalloc (d_out)! Error: " << gpuErr <Nsimd(),{ - - int n = s / e2; - int b = s % e2; - int so=r*ostride; // base offset for start of plane - int ss= so+n*stride+b; - - coalescedWrite(rb_p[r*subvol_size+s], coalescedRead(Data_v[ss])); - - }); - - //issue segmented reductions in computeStream - gpuErr = gpucub::DeviceSegmentedReduce::Reduce(helperArray, temp_storage_bytes, rb_p, d_out, rd, d_offsets, d_offsets+1,::gpucub::Sum(), vobj_zero, computeStream); - if (gpuErr!=gpuSuccess) { - std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce! Error: " << gpuErr <iCoorFromIindex(icoor,idx); - - int ldx =rt+icoor[orthogdim]*rd; - - lsSum[ldx]=lsSum[ldx]+extracted[idx]; - - } - } - - // sum over nodes. - for(int t=0;t_processor_coor[orthogdim] ) { - result[t]=lsSum[lt]; - } else { - result[t]=Zero(); - } - - } - scalar_type * ptr = (scalar_type *) &result[0]; - int words = fd*sizeof(sobj)/sizeof(scalar_type); - grid->GlobalSumVector(ptr, words); -} - -NAMESPACE_END(Grid); \ No newline at end of file diff --git a/Grid/lattice/Lattice_slicesum_sycl.h b/Grid/lattice/Lattice_slicesum_sycl.h deleted file mode 100644 index 04ec8a6a..00000000 --- a/Grid/lattice/Lattice_slicesum_sycl.h +++ /dev/null @@ -1,110 +0,0 @@ -#pragma once - -NAMESPACE_BEGIN(Grid); - -template -inline void sliceSum_sycl(const Lattice &Data, std::vector &result, int orthogdim) -{ - typedef typename vobj::scalar_object sobj; - typedef typename vobj::scalar_object::scalar_type scalar_type; - - GridBase *grid = Data.Grid(); - assert(grid!=NULL); - - const int Nd = grid->_ndimension; - const size_t Nsimd = grid->Nsimd(); - - assert(orthogdim >= 0); - assert(orthogdim < Nd); - - int fd=grid->_fdimensions[orthogdim]; - int ld=grid->_ldimensions[orthogdim]; - int rd=grid->_rdimensions[orthogdim]; - - int e1= grid->_slice_nblock[orthogdim]; - int e2= grid->_slice_block [orthogdim]; - int stride=grid->_slice_stride[orthogdim]; - int ostride=grid->_ostride[orthogdim]; - size_t subvol_size = e1*e2; - - vobj *mysum = (vobj *) malloc_shared(sizeof(vobj),*theGridAccelerator); - - result.resize(fd); - - Vector lvSum(rd); - Vector lsSum(ld,Zero()); - commVector reduction_buffer(rd*subvol_size); - ExtractBuffer extracted(Nsimd); - vobj vobj_zero; - zeroit(vobj_zero); - - for(int r=0;rsubmit([&](cl::sycl::handler &cgh) { - auto Reduction = cl::sycl::reduction(mysum,std::plus<>()); - cgh.parallel_for(cl::sycl::range<1>{subvol_size}, - Reduction, - [=](cl::sycl::id<1> item, auto &sum) { - auto s = item[0]; - sum += rb_p[r*subvol_size+s]; - }); - }); - theGridAccelerator->wait(); - lvSum[r] = mysum[0]; - } - - Coordinate icoor(Nd); - - for(int rt=0;rtiCoorFromIindex(icoor,idx); - - int ldx =rt+icoor[orthogdim]*rd; - - lsSum[ldx]=lsSum[ldx]+extracted[idx]; - - } - } - - // sum over nodes. - for(int t=0;t_processor_coor[orthogdim] ) { - result[t]=lsSum[lt]; - } else { - result[t]=Zero(); - } - - } - scalar_type * ptr = (scalar_type *) &result[0]; - int words = fd*sizeof(sobj)/sizeof(scalar_type); - grid->GlobalSumVector(ptr, words); - -} - -NAMESPACE_END(Grid); \ No newline at end of file diff --git a/tests/core/Test_sliceSum.cc b/tests/core/Test_sliceSum.cc index 399ab899..e0e0c1ae 100644 --- a/tests/core/Test_sliceSum.cc +++ b/tests/core/Test_sliceSum.cc @@ -1,5 +1,79 @@ #include +template inline void sliceSumCPU(const Grid::Lattice &Data,std::vector &result,int orthogdim) +{ + using namespace Grid; + /////////////////////////////////////////////////////// + // FIXME precision promoted summation + // may be important for correlation functions + // But easily avoided by using double precision fields + /////////////////////////////////////////////////////// + typedef typename vobj::scalar_object sobj; + typedef typename vobj::scalar_object::scalar_type scalar_type; + GridBase *grid = Data.Grid(); + assert(grid!=NULL); + + const int Nd = grid->_ndimension; + const int Nsimd = grid->Nsimd(); + + assert(orthogdim >= 0); + assert(orthogdim < Nd); + + int fd=grid->_fdimensions[orthogdim]; + int ld=grid->_ldimensions[orthogdim]; + int rd=grid->_rdimensions[orthogdim]; + + Vector lvSum(rd); // will locally sum vectors first + Vector lsSum(ld,Zero()); // sum across these down to scalars + ExtractBuffer extracted(Nsimd); // splitting the SIMD + + result.resize(fd); // And then global sum to return the same vector to every node + for(int r=0;r_slice_nblock[orthogdim]; + int e2= grid->_slice_block [orthogdim]; + int stride=grid->_slice_stride[orthogdim]; + int ostride=grid->_ostride[orthogdim]; + + //Reduce Data down to lvSum + sliceSumReduction_cpu(Data,lvSum,rd, e1,e2,stride,ostride,Nsimd); + + // Sum across simd lanes in the plane, breaking out orthog dir. + Coordinate icoor(Nd); + + for(int rt=0;rtiCoorFromIindex(icoor,idx); + + int ldx =rt+icoor[orthogdim]*rd; + + lsSum[ldx]=lsSum[ldx]+extracted[idx]; + + } + } + + // sum over nodes. + for(int t=0;t_processor_coor[orthogdim] ) { + result[t]=lsSum[lt]; + } else { + result[t]=Zero(); + } + + } + scalar_type * ptr = (scalar_type *) &result[0]; + int words = fd*sizeof(sobj)/sizeof(scalar_type); + grid->GlobalSumVector(ptr, words); +} + int main (int argc, char ** argv) { @@ -26,7 +100,7 @@ int main (int argc, char ** argv) { //warmup for (int sweeps = 0; sweeps < 5; sweeps++) { - reduction_result = sliceSumGpu(test_data,0); + reduction_result = sliceSum(test_data,0); } int trace_id = traceStart("sliceSum benchmark"); @@ -35,23 +109,23 @@ int main (int argc, char ** argv) { RealD t=-usecond(); tracePush("sliceSum"); - sliceSum(test_data,reduction_reference,i); + sliceSumCPU(test_data,reduction_reference,i); tracePop("sliceSum"); t+=usecond(); - - std::cout << GridLogMessage << " sliceSum took "< Date: Fri, 9 Feb 2024 13:07:56 -0500 Subject: [PATCH 08/14] Undo include path changes for level zero api header --- Grid/threads/Accelerator.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Grid/threads/Accelerator.h b/Grid/threads/Accelerator.h index eaafea5d..ff5ccd7a 100644 --- a/Grid/threads/Accelerator.h +++ b/Grid/threads/Accelerator.h @@ -256,12 +256,12 @@ NAMESPACE_END(Grid); #if 0 #include #include -#include +#include #include #else #include #include -#include +#include #include #endif From b5659d106ec5997631caef6ca6d80e39cc69cfc6 Mon Sep 17 00:00:00 2001 From: dbollweg Date: Fri, 9 Feb 2024 13:37:14 -0500 Subject: [PATCH 09/14] more test cases --- tests/core/Test_sliceSum.cc | 113 +++++++++++++++++++++++++++++++++++- 1 file changed, 112 insertions(+), 1 deletion(-) diff --git a/tests/core/Test_sliceSum.cc b/tests/core/Test_sliceSum.cc index e0e0c1ae..0ca2881b 100644 --- a/tests/core/Test_sliceSum.cc +++ b/tests/core/Test_sliceSum.cc @@ -103,7 +103,8 @@ int main (int argc, char ** argv) { reduction_result = sliceSum(test_data,0); } - int trace_id = traceStart("sliceSum benchmark"); + int trace_id = traceStart("sliceSum benchmark - ComplexD"); + std::cout << GridLogMessage << "Testing ComplexD" << std::endl; for (int i = 0; i < Nd; i++) { RealD t=-usecond(); @@ -138,6 +139,116 @@ int main (int argc, char ** argv) { } traceStop(trace_id); + + LatticeColourVectorD test_data_cv(&Grid); + gaussian(pRNG,test_data_cv); + + std::vector reduction_reference_cv; + std::vector reduction_result_cv; + + //warmup + for (int sweeps = 0; sweeps < 5; sweeps++) { + reduction_result_cv = sliceSum(test_data_cv,0); + } + trace_id = traceStart("sliceSum benchmark - ColourVectorD"); + + std::cout << GridLogMessage << "Testing ColourVectorD" << std::endl; + for (int i = 0; i < Nd; i++) { + + RealD t=-usecond(); + + tracePush("sliceSum"); + sliceSumCPU(test_data_cv,reduction_reference_cv,i); + tracePop("sliceSum"); + + t+=usecond(); + std::cout << GridLogMessage << "Orthog. dir. = " << i << std::endl; + std::cout << GridLogMessage << "CPU sliceSum took "< reduction_reference_scv; + std::vector reduction_result_scv; + + //warmup + for (int sweeps = 0; sweeps < 5; sweeps++) { + reduction_result_scv = sliceSum(test_data_scv,0); + } + trace_id = traceStart("sliceSum benchmark - SpinColourVectorD"); + + std::cout << GridLogMessage << "Testing SpinColourVectorD" << std::endl; + for (int i = 0; i < Nd; i++) { + + RealD t=-usecond(); + + tracePush("sliceSum"); + sliceSumCPU(test_data_scv,reduction_reference_scv,i); + tracePop("sliceSum"); + + t+=usecond(); + std::cout << GridLogMessage << "Orthog. dir. = " << i << std::endl; + std::cout << GridLogMessage << "CPU sliceSum took "< Date: Fri, 16 Feb 2024 13:15:02 -0500 Subject: [PATCH 10/14] Adding sliceSumReduction_cub_small/large since hipcub cannot deal with arb. large vobjs --- Grid/lattice/Lattice_slicesum_core.h | 65 ++++++++++++++++++++++------ tests/core/Test_sliceSum.cc | 22 ++++++---- 2 files changed, 65 insertions(+), 22 deletions(-) diff --git a/Grid/lattice/Lattice_slicesum_core.h b/Grid/lattice/Lattice_slicesum_core.h index 2548884a..f4301c33 100644 --- a/Grid/lattice/Lattice_slicesum_core.h +++ b/Grid/lattice/Lattice_slicesum_core.h @@ -1,4 +1,5 @@ #pragma once +#include #if defined(GRID_CUDA) #include @@ -26,20 +27,16 @@ NAMESPACE_BEGIN(Grid); + #if defined(GRID_CUDA) || defined(GRID_HIP) -template inline void sliceSumReduction_cub(const Lattice &Data, Vector &lvSum, const int rd, const int e1, const int e2, const int stride, const int ostride, const int Nsimd) -{ - typedef typename vobj::scalar_object sobj; - +template inline void sliceSumReduction_cub_small(const vobj *Data, Vector &lvSum, const int rd, const int e1, const int e2, const int stride, const int ostride, const int Nsimd) { size_t subvol_size = e1*e2; - commVector reduction_buffer(rd*subvol_size); auto rb_p = &reduction_buffer[0]; + vobj zero_init; + zeroit(zero_init); - vobj vobj_zero; //Need to provide initial value for reduction operation - zeroit(vobj_zero); - void *temp_storage_array = NULL; size_t temp_storage_bytes = 0; vobj *d_out; @@ -71,8 +68,8 @@ template inline void sliceSumReduction_cub(const Lattice &Data exit(EXIT_FAILURE); } - //determine temp_storage_array size - gpuErr = gpucub::DeviceSegmentedReduce::Reduce(temp_storage_array, temp_storage_bytes, rb_p,d_out, rd, d_offsets, d_offsets+1, ::gpucub::Sum(), vobj_zero, computeStream); + + 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); if (gpuErr!=gpuSuccess) { std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce (setup)! Error: " << gpuErr < inline void sliceSumReduction_cub(const Lattice &Data exit(EXIT_FAILURE); } - autoView( Data_v, Data, AcceleratorRead); //prepare buffer for reduction //use non-blocking accelerator_for to avoid syncs (ok because we submit to same computeStream) //use 2d accelerator_for to avoid launch latencies found when serially looping over rd - accelerator_for2dNB( s,subvol_size, r,rd, Nsimd,{ int n = s / e2; @@ -97,12 +92,12 @@ template inline void sliceSumReduction_cub(const Lattice &Data int so=r*ostride; // base offset for start of plane int ss= so+n*stride+b; - coalescedWrite(rb_p[r*subvol_size+s], coalescedRead(Data_v[ss])); + coalescedWrite(rb_p[r*subvol_size+s], coalescedRead(Data[ss])); }); //issue segmented reductions in computeStream - gpuErr = gpucub::DeviceSegmentedReduce::Reduce(temp_storage_array, temp_storage_bytes, rb_p, d_out, rd, d_offsets, d_offsets+1,::gpucub::Sum(), vobj_zero, computeStream); + 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); if (gpuErr!=gpuSuccess) { std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce! Error: " << gpuErr < inline void sliceSumReduction_cub(const Lattice &Data } + +template inline void sliceSumReduction_cub_large(const vobj *Data, Vector &lvSum, const int rd, const int e1, const int e2, const int stride, const int ostride, const int Nsimd) { + typedef typename vobj::vector_type vector; + const int words = sizeof(vobj)/sizeof(vector); + const int osites = rd*e1*e2; + Vectorbuffer(osites); + vector *dat = (vector *)Data; + vector *buf = &buffer[0]; + Vector lvSum_small(rd); + vector *lvSum_ptr = (vector *)&lvSum[0]; + + for (int w = 0; w < words; w++) { + accelerator_for(ss,osites,1,{ + buf[ss] = dat[ss*words+w]; + }); + + sliceSumReduction_cub_small(buf,lvSum_small,rd,e1,e2,stride, ostride,Nsimd); + + for (int r = 0; r < rd; r++) { + lvSum_ptr[w+words*r]=lvSum_small[r]; + } + + } + + +} + +template inline void sliceSumReduction_cub(const Lattice &Data, Vector &lvSum, const int rd, const int e1, const int e2, const int stride, const int ostride, const int Nsimd) +{ + autoView(Data_v, Data, AcceleratorRead); + #if defined(GRID_CUDA) + sliceSumReduction_cub_small(&Data_v[0], lvSum, rd, e1, e2, stride, ostride, Nsimd); + + #elif defined (GRID_HIP) //hipcub cannot deal with large vobjs that don't fit in shared memory, therefore separate into _small/_large. + if constexpr (sizeof(vobj) <= 256) { + sliceSumReduction_cub_small(&Data_v[0], lvSum, rd, e1, e2, stride, ostride, Nsimd); + } + else { + sliceSumReduction_cub_large(&Data_v[0], lvSum, rd, e1, e2, stride, ostride, Nsimd); + } + #endif +} #endif diff --git a/tests/core/Test_sliceSum.cc b/tests/core/Test_sliceSum.cc index 0ca2881b..f86f96f8 100644 --- a/tests/core/Test_sliceSum.cc +++ b/tests/core/Test_sliceSum.cc @@ -140,19 +140,21 @@ int main (int argc, char ** argv) { } traceStop(trace_id); - LatticeColourVectorD test_data_cv(&Grid); + LatticeSpinVectorD test_data_cv(&Grid); gaussian(pRNG,test_data_cv); - std::vector reduction_reference_cv; - std::vector reduction_result_cv; + std::vector reduction_reference_cv; + std::vector reduction_result_cv; //warmup for (int sweeps = 0; sweeps < 5; sweeps++) { reduction_result_cv = sliceSum(test_data_cv,0); } - trace_id = traceStart("sliceSum benchmark - ColourVectorD"); + trace_id = traceStart("sliceSum benchmark - SpinVectorD"); - std::cout << GridLogMessage << "Testing ColourVectorD" << std::endl; + std::cout << GridLogMessage << "Testing SpinVectorD" << std::endl; + std::cout << GridLogMessage << "sizeof(SpinVectorD) = " << sizeof(SpinVectorD) << std::endl; + std::cout << GridLogMessage << "sizeof(vSpinVectorD) = " << sizeof(vSpinVectorD) << std::endl; for (int i = 0; i < Nd; i++) { RealD t=-usecond(); @@ -180,9 +182,10 @@ int main (int argc, char ** argv) { for(int t=0;t(acceleratorAllocDevice(rd*sizeof(vobj))); - 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); + d_offsets = static_cast(acceleratorAllocDevice((rd+1)*sizeof(int))); + + //copy offsets to device + acceleratorCopyToDeviceAsync(&offsets[0],d_offsets,sizeof(int)*(rd+1),computeStream); + + + 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); if (gpuErr!=gpuSuccess) { std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpucub::DeviceSegmentedReduce::Reduce (setup)! Error: " << gpuErr < inline void sliceSumReduction_cub_small(const vobj *Data, V exit(EXIT_FAILURE); } - gpuErr = gpuMemcpyAsync(&lvSum[0],d_out,rd*sizeof(vobj),gpuMemcpyDeviceToHost,computeStream); - if (gpuErr!=gpuSuccess) { - std::cout << GridLogError << "Lattice_slicesum_gpu.h: Encountered error during gpuMemcpy (d_out)! Error: " << gpuErr < Date: Tue, 27 Feb 2024 11:28:32 -0500 Subject: [PATCH 14/14] Added SpinColourMatrix case to sliceSum Test --- tests/core/Test_sliceSum.cc | 58 +++++++++++++++++++++++++++++++++++++ 1 file changed, 58 insertions(+) diff --git a/tests/core/Test_sliceSum.cc b/tests/core/Test_sliceSum.cc index 5275ddc2..e366f1f3 100644 --- a/tests/core/Test_sliceSum.cc +++ b/tests/core/Test_sliceSum.cc @@ -257,7 +257,65 @@ int main (int argc, char ** argv) { } traceStop(trace_id); + LatticeSpinColourMatrixD test_data_scm(&Grid); + gaussian(pRNG,test_data_scm); + + std::vector reduction_reference_scm; + std::vector reduction_result_scm; + + //warmup + for (int sweeps = 0; sweeps < 5; sweeps++) { + reduction_result_scm = sliceSum(test_data_scm,0); + } + trace_id = traceStart("sliceSum benchmark - SpinColourMatrixD"); + + std::cout << GridLogMessage << "Testing SpinColourMatrixD" << std::endl; + std::cout << GridLogMessage << "sizeof(SpinColourMatrixD) = " << sizeof(SpinColourMatrixD) << std::endl; + std::cout << GridLogMessage << "sizeof(vSpinColourMatrixD) = " << sizeof(vSpinColourMatrixD) << std::endl; + for (int i = 0; i < Nd; i++) { + + RealD t=-usecond(); + + tracePush("sliceSum"); + sliceSumCPU(test_data_scm,reduction_reference_scm,i); + tracePop("sliceSum"); + + t+=usecond(); + std::cout << GridLogMessage << "Orthog. dir. = " << i << std::endl; + std::cout << GridLogMessage << "CPU sliceSum took "<