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 "<