portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2025-04-11 22:50:45 +01:00
Code Releases Activity

slicesum_sycl passes test

Browse Source
This commit is contained in:
dbollweg 2024-02-06 19:08:44 -05:00
parent ab2de131bd
commit 1514b4f137
1 changed files with 5 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
Grid/lattice/Lattice_slicesum_sycl.h
View File

@ -14,7 +14,6 @@ inline void sliceSum_sycl(const Lattice<vobj> &Data, std::vector<typename vobj::
const int Nd = grid->_ndimension; const int Nd = grid->_ndimension;
const size_t Nsimd = grid->Nsimd(); const size_t Nsimd = grid->Nsimd();
assert(orthogdim >= 0); assert(orthogdim >= 0);
assert(orthogdim < Nd); assert(orthogdim < Nd);
@ -29,9 +28,6 @@ inline void sliceSum_sycl(const Lattice<vobj> &Data, std::vector<typename vobj::
size_t subvol_size = e1*e2; size_t subvol_size = e1*e2;
vobj *mysum = (vobj *) malloc_shared(sizeof(vobj),*theGridAccelerator); vobj *mysum = (vobj *) malloc_shared(sizeof(vobj),*theGridAccelerator);
vobj vobj_zero;
zeroit(vobj_zero);
result.resize(fd); result.resize(fd);
@ -39,6 +35,8 @@ inline void sliceSum_sycl(const Lattice<vobj> &Data, std::vector<typename vobj::
Vector<sobj> lsSum(ld,Zero()); Vector<sobj> lsSum(ld,Zero());
commVector<vobj> reduction_buffer(rd*subvol_size); commVector<vobj> reduction_buffer(rd*subvol_size);
ExtractBuffer<sobj> extracted(Nsimd); ExtractBuffer<sobj> extracted(Nsimd);
vobj vobj_zero;
zeroit(vobj_zero);
for(int r=0;r<rd;r++){ for(int r=0;r<rd;r++){
lvSum[r]=Zero(); lvSum[r]=Zero();
@ -46,12 +44,9 @@ inline void sliceSum_sycl(const Lattice<vobj> &Data, std::vector<typename vobj::
auto rb_p = &reduction_buffer[0]; auto rb_p = &reduction_buffer[0];
autoView(Data_v, Data, AcceleratorRead); autoView(Data_v, Data, AcceleratorRead);
//prepare reduction buffer (can i use this with sycl backend?) //prepare reduction buffer
accelerator_for2d( s,subvol_size, r,rd, Nsimd,{ accelerator_for2d( s,subvol_size, r,rd, Nsimd,{
int n = s / e2; int n = s / e2;
@ -64,9 +59,9 @@ inline void sliceSum_sycl(const Lattice<vobj> &Data, std::vector<typename vobj::
}); });
for (int r = 0; r < rd; r++) { 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) { theGridAccelerator->submit([&](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}, cgh.parallel_for(cl::sycl::range<1>{subvol_size},
Reduction, Reduction,
[=](cl::sycl::id<1> item, auto &sum) { [=](cl::sycl::id<1> item, auto &sum) {