From 68f112d576a53c404f55ea34cbad46ce8efbab57 Mon Sep 17 00:00:00 2001 From: Peter Boyle Date: Thu, 10 Oct 2024 22:03:04 +0000 Subject: [PATCH] New software moves cl::sycl --- Grid/lattice/Lattice_reduction_sycl.h | 16 ++++++++-------- Grid/lattice/Lattice_slicesum_core.h | 8 ++++---- 2 files changed, 12 insertions(+), 12 deletions(-) diff --git a/Grid/lattice/Lattice_reduction_sycl.h b/Grid/lattice/Lattice_reduction_sycl.h index 3718e6ea..bc9257b9 100644 --- a/Grid/lattice/Lattice_reduction_sycl.h +++ b/Grid/lattice/Lattice_reduction_sycl.h @@ -16,11 +16,11 @@ inline typename vobj::scalar_objectD sumD_gpu_tensor(const vobj *lat, Integer os Integer nsimd= vobj::Nsimd(); { sycl::buffer abuff(&ret, {1}); - theGridAccelerator->submit([&](cl::sycl::handler &cgh) { - auto Reduction = cl::sycl::reduction(abuff,cgh,identity,std::plus<>()); - cgh.parallel_for(cl::sycl::range<1>{osites}, + theGridAccelerator->submit([&](sycl::handler &cgh) { + auto Reduction = sycl::reduction(abuff,cgh,identity,std::plus<>()); + cgh.parallel_for(sycl::range<1>{osites}, Reduction, - [=] (cl::sycl::id<1> item, auto &sum) { + [=] (sycl::id<1> item, auto &sum) { auto osite = item[0]; sum +=Reduce(lat[osite]); }); @@ -75,11 +75,11 @@ template Word svm_xor(Word *vec,uint64_t L) Word ret = 0; { sycl::buffer abuff(&ret, {1}); - theGridAccelerator->submit([&](cl::sycl::handler &cgh) { - auto Reduction = cl::sycl::reduction(abuff,cgh,identity,std::bit_xor<>()); - cgh.parallel_for(cl::sycl::range<1>{L}, + theGridAccelerator->submit([&](sycl::handler &cgh) { + auto Reduction = sycl::reduction(abuff,cgh,identity,std::bit_xor<>()); + cgh.parallel_for(sycl::range<1>{L}, Reduction, - [=] (cl::sycl::id<1> index, auto &sum) { + [=] (sycl::id<1> index, auto &sum) { sum ^=vec[index]; }); }); diff --git a/Grid/lattice/Lattice_slicesum_core.h b/Grid/lattice/Lattice_slicesum_core.h index f01ba73d..e15055a6 100644 --- a/Grid/lattice/Lattice_slicesum_core.h +++ b/Grid/lattice/Lattice_slicesum_core.h @@ -141,11 +141,11 @@ inline void sliceSumReduction_sycl_small(const vobj *Data, }); for (int r = 0; r < rd; r++) { - theGridAccelerator->submit([&](cl::sycl::handler &cgh) { - auto Reduction = cl::sycl::reduction(&mysum[r],std::plus<>()); - cgh.parallel_for(cl::sycl::range<1>{subvol_size}, + theGridAccelerator->submit([&](sycl::handler &cgh) { + auto Reduction = sycl::reduction(&mysum[r],std::plus<>()); + cgh.parallel_for(sycl::range<1>{subvol_size}, Reduction, - [=](cl::sycl::id<1> item, auto &sum) { + [=](sycl::id<1> item, auto &sum) { auto s = item[0]; sum += rb_p[r*subvol_size+s]; });