portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2024-09-19 16:55:37 +01:00
Code Releases Activity

Force initial to identity

Browse Source
This commit is contained in:
Peter Boyle 2024-06-11 17:51:58 +00:00
parent fb9b1d76ca
commit 9563238e9b
1 changed files with 3 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
Grid/lattice/Lattice_reduction_sycl.h
View File

@ -13,6 +13,7 @@ inline typename vobj::scalar_objectD sumD_gpu_tensor(const vobj *lat, Integer os
mysum.resize(1); mysum.resize(1);
sobj *mysum_p = & mysum[0]; sobj *mysum_p = & mysum[0];
sobj identity; zeroit(identity); sobj identity; zeroit(identity);
mysum[0] = identity;
sobj ret ; sobj ret ;
Integer nsimd= vobj::Nsimd(); Integer nsimd= vobj::Nsimd();
@ -80,13 +81,14 @@ template<class Word> Word svm_xor(Word *vec,uint64_t L)
d_sum.resize(1); d_sum.resize(1);
Word *d_sum_p=&d_sum[0]; Word *d_sum_p=&d_sum[0];
Word identity; identity=0; Word identity; identity=0;
d_sum[0] = identity;
const cl::sycl::property_list PropList ({ cl::sycl::property::reduction::initialize_to_identity() }); const cl::sycl::property_list PropList ({ cl::sycl::property::reduction::initialize_to_identity() });
theGridAccelerator->submit([&](cl::sycl::handler &cgh) { theGridAccelerator->submit([&](cl::sycl::handler &cgh) {
auto Reduction = cl::sycl::reduction(d_sum_p,identity,std::bit_xor<>(),PropList); auto Reduction = cl::sycl::reduction(d_sum_p,identity,std::bit_xor<>(),PropList);
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();