From e38cf4dc8604634daddcc6a6b05ab1187ec83efc Mon Sep 17 00:00:00 2001 From: Tuowen Zhao Date: Mon, 4 Oct 2021 16:23:27 +0000 Subject: Update formatting --- subgroup.cpp | 40 ++++++++++++++++++++-------------------- 1 file changed, 20 insertions(+), 20 deletions(-) (limited to 'subgroup.cpp') diff --git a/subgroup.cpp b/subgroup.cpp index 59e873d..f55a1f8 100644 --- a/subgroup.cpp +++ b/subgroup.cpp @@ -9,17 +9,18 @@ using namespace cl::sycl; -template -inline void dev_shl(ONEAPI::sub_group &SG, T &res, T l, T r, unsigned kn, unsigned cw, unsigned cid) { +template +inline void dev_shl(ONEAPI::sub_group &SG, T &res, T l, T r, unsigned kn, unsigned cw, + unsigned cid) { T l_tmp = SG.shuffle_down(l, cw - (kn)); T r_tmp = SG.shuffle_up(r, kn); - res = (cid) < kn? l_tmp : r_tmp; + res = (cid) < kn ? l_tmp : r_tmp; } void printInfo(device &Device) { - std::cout << "Using OpenCL " << (Device.is_cpu() ? "CPU" : "GPU") - << " device {" << Device.get_info() - << "} from {" << Device.get_info() << "}" << std::endl; + std::cout << "Using OpenCL " << (Device.is_cpu() ? "CPU" : "GPU") << " device {" + << Device.get_info() << "} from {" + << Device.get_info() << "}" << std::endl; auto dot_num_groups = Device.get_info(); auto dot_wgsize = Device.get_info(); @@ -67,18 +68,17 @@ int main() { auto b = B.get_access(cgh); auto c = C.get_access(cgh); - cgh.parallel_for( - nworkitem, [=](nd_item<1> WIid) { - ONEAPI::sub_group SG = WIid.get_sub_group(); - int sglid = SG.get_local_id().get(0); - uint32_t i = WIid.get_global_id(0); - int* cp = &c[0]; - const int* ap = &a[0]; - const int* bp = &b[0]; -// c[i] = a[i] + b[i]; - dev_shl(SG, cp[i], ap[i], bp[i], 2, 8, sglid & 7); - c[i] = WIid.get_group(0); - }); + cgh.parallel_for(nworkitem, [=](nd_item<1> WIid) { + ONEAPI::sub_group SG = WIid.get_sub_group(); + int sglid = SG.get_local_id().get(0); + uint32_t i = WIid.get_global_id(0); + int *cp = &c[0]; + const int *ap = &a[0]; + const int *bp = &b[0]; + // c[i] = a[i] + b[i]; + dev_shl(SG, cp[i], ap[i], bp[i], 2, 8, sglid & 7); + c[i] = WIid.get_group(0); + }); }; std::cout << "submitting" << std::endl; auto e = Queue.submit(kernel); @@ -87,8 +87,8 @@ int main() { std::cout << "waited" << std::endl; auto c_r = C.get_access(); -// for (int i = 0; i < N; ++i) -// std::cout << i << " + " << N - i << " = " << c[i] << std::endl; + // for (int i = 0; i < N; ++i) + // std::cout << i << " + " << N - i << " = " << c[i] << std::endl; return 0; } -- cgit v1.2.3-70-g09d2