portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2024-11-09 23:45:36 +00:00
Code Releases Activity

Xor csum for repro testing

Browse Source
This commit is contained in:
Peter Boyle 2024-03-22 15:42:57 +00:00
parent d01e5fa838
commit d1e9fe50d2
1 changed files with 11 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
Grid/lattice/Lattice_reduction_sycl.h
View File

@ -69,29 +69,30 @@ inline typename vobj::scalar_object sum_gpu_large(const vobj *lat, Integer osite
return result; return result;
} }
NAMESPACE_END(Grid);
/* template<class Word> Word svm_xor(Word *vec,uint64_t L)
template<class Double> Double svm_reduce(Double *vec,uint64_t L)
{ {
Double sumResult; zeroit(sumResult); Word xorResult; xorResult = 0;
Double *d_sum =(Double *)cl::sycl::malloc_shared(sizeof(Double),*theGridAccelerator); Word *d_sum =(Word *)cl::sycl::malloc_shared(sizeof(Word),*theGridAccelerator);
Double identity; zeroit(identity); Word identity; identity=0;
theGridAccelerator->submit([&](cl::sycl::handler &cgh) { theGridAccelerator->submit([&](cl::sycl::handler &cgh) {
auto Reduction = cl::sycl::reduction(d_sum,identity,std::plus<>()); auto Reduction = cl::sycl::reduction(d_sum,identity,std::bit_xor<>());
cgh.parallel_for(cl::sycl::range<1>{L}, cgh.parallel_for(cl::sycl::range<1>{L},
Reduction, Reduction,
[=] (cl::sycl::id<1> index, auto &sum) { [=] (cl::sycl::id<1> index, auto &sum) {
sum +=vec[index]; sum ^=vec[index];
}); });
}); });
theGridAccelerator->wait(); theGridAccelerator->wait();
Double ret = d_sum[0]; Word ret = d_sum[0];
free(d_sum,*theGridAccelerator); free(d_sum,*theGridAccelerator);
std::cout << " svm_reduce finished "<<L<<" sites sum = " << ret <<std::endl;
return ret; return ret;
} }
NAMESPACE_END(Grid);
/*
template <class vobj> template <class vobj>
inline typename vobj::scalar_objectD sumD_gpu_repack(const vobj *lat, Integer osites) inline typename vobj::scalar_objectD sumD_gpu_repack(const vobj *lat, Integer osites)
{ {