From 5368e3ecd8f17dc9d7f5dff87790c6490ecbd1a3 Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Thu, 30 Jul 2020 14:40:01 +0300 Subject: [PATCH 1/2] [NFC][SYCL] fix LIT tests - Remove sub_group/common_ocl.cpp because it duplicates sub_group/common.cpp and directly use OpenCL API that causes instability on some configurations. - Fix sub_group/shuffle*.cpp tests to align with shuffle_xor restrictions mentioned in spec: "If the result of the XOR is greater than max_sub_group_size then it is considered out-of-range" --- sycl/test/sub_group/common_ocl.cpp | 106 ----------------------------- sycl/test/sub_group/shuffle.hpp | 42 ++++++++---- 2 files changed, 30 insertions(+), 118 deletions(-) delete mode 100644 sycl/test/sub_group/common_ocl.cpp diff --git a/sycl/test/sub_group/common_ocl.cpp b/sycl/test/sub_group/common_ocl.cpp deleted file mode 100644 index 232e6c6c11acc..0000000000000 --- a/sycl/test/sub_group/common_ocl.cpp +++ /dev/null @@ -1,106 +0,0 @@ -// REQUIRES: opencl - -// RUN: %clang_cc1 -x cl -cl-std=CL2.0 %S/sg.cl -triple spir64-unknown-unknown -emit-llvm-bc -o %T/kernel_ocl.bc -include opencl-c.h -// RUN: llvm-spirv %T/kernel_ocl.bc -o %T/kernel_ocl.spv -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -L %opencl_libs_dir -lOpenCL -// RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out %T/kernel_ocl.spv -// RUN: %GPU_RUN_PLACEHOLDER %t.out %T/kernel_ocl.spv -// RUN: %ACC_RUN_PLACEHOLDER %t.out %T/kernel_ocl.spv - -//==--- common_ocl.cpp - basic SG methods in SYCL vs OpenCL ---*- C++ -*---==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include "helper.hpp" -#include -#include -#include -#include - -using namespace cl::sycl; -struct Data { - unsigned int local_id; - unsigned int local_range; - unsigned int max_local_range; - unsigned int group_id; - unsigned int group_range; -}; - -void check(queue &Queue, const int G, const int L, const char *SpvFile) { - try { - nd_range<1> NdRange(G, L); - buffer oclbuf(G); - buffer syclbuf(G); - - std::ifstream File(SpvFile, std::ios::binary); - if (!File.is_open()) { - std::cerr << std::strerror(errno); - throw compile_program_error("Cannot open SPIRV file\n", PI_INVALID_VALUE); - } - File.seekg(0, std::ios::end); - vector_class Spv(File.tellg()); - File.seekg(0); - File.read(Spv.data(), Spv.size()); - File.close(); - int Err; - cl_program ClProgram = clCreateProgramWithIL(Queue.get_context().get(), - Spv.data(), Spv.size(), &Err); - CHECK_OCL_CODE(Err); - CHECK_OCL_CODE( - clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr)); - program Prog(Queue.get_context(), ClProgram); - Queue.submit([&](handler &cgh) { - auto oclacc = oclbuf.get_access(cgh); - cgh.set_args(oclacc); - cgh.parallel_for(NdRange, Prog.get_kernel("ocl_subgr")); - }); - auto oclacc = oclbuf.get_access(); - - Queue.submit([&](handler &cgh) { - auto syclacc = syclbuf.get_access(cgh); - cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { - intel::sub_group SG = NdItem.get_sub_group(); - syclacc[NdItem.get_global_id()].local_id = SG.get_local_id().get(0); - syclacc[NdItem.get_global_id()].local_range = - SG.get_local_range().get(0); - syclacc[NdItem.get_global_id()].max_local_range = - SG.get_max_local_range().get(0); - syclacc[NdItem.get_global_id()].group_id = SG.get_group_id().get(0); - syclacc[NdItem.get_global_id()].group_range = SG.get_group_range().get(0); - }); - }); - auto syclacc = syclbuf.get_access(); - for (int j = 0; j < G; j++) { - exit_if_not_equal(syclacc[j].local_id, oclacc[j].local_id, "local_id"); - exit_if_not_equal(syclacc[j].local_range, oclacc[j].local_range, - "local_range"); - exit_if_not_equal(syclacc[j].max_local_range, oclacc[j].max_local_range, - "max_local_range"); - exit_if_not_equal(syclacc[j].group_id, oclacc[j].group_id, "group_id"); - exit_if_not_equal(syclacc[j].group_range, oclacc[j].group_range, - "group_range"); - } - } catch (exception e) { - std::cout << "SYCL exception caught: " << e.what(); - exit(1); - } -} -int main(int argc, char **argv) { - queue Queue; - if (!core_sg_supported(Queue.get_device()) || argc != 2) { - std::cout << "Skipping test\n"; - return 0; - } - - check(Queue, 240, 80, argv[1]); - check(Queue, 8, 4, argv[1]); - check(Queue, 24, 12, argv[1]); - check(Queue, 1024, 256, argv[1]); - std::cout << "Test passed." << std::endl; - return 0; -} diff --git a/sycl/test/sub_group/shuffle.hpp b/sycl/test/sub_group/shuffle.hpp index 94c82ab99c2d1..7c16121febc0f 100644 --- a/sycl/test/sub_group/shuffle.hpp +++ b/sycl/test/sub_group/shuffle.hpp @@ -8,8 +8,7 @@ #include "helper.hpp" #include -template -class sycl_subgr; +template class sycl_subgr; using namespace cl::sycl; @@ -66,8 +65,9 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { acc_up[NdItem.get_global_id()] = SG.shuffle_up(vwggid, sgid); /* Save GID+SGID */ acc_down[NdItem.get_global_id()] = SG.shuffle_down(vwggid, sgid); - /* Save GID XOR SGID */ - acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(vwggid, sgid); + /* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + acc_xor[NdItem.get_global_id()] = + SG.shuffle_xor(vwggid, sgid % SG.get_max_local_range()[0]); }); }); auto acc = buf.template get_access(); @@ -81,12 +81,18 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { size_t sg_size = sgsizeacc[0]; int SGid = 0; + int SGLid = 0; + int SGBeginGid = 0; for (int j = 0; j < G; j++) { if (j % L % sg_size == 0) { SGid++; + SGLid = 0; + SGBeginGid = j; } if (j % L == 0) { SGid = 0; + SGLid = 0; + SGBeginGid = j; } /*GID of middle element in every subgroup*/ exit_if_not_equal_vec( @@ -115,8 +121,11 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { exit_if_not_equal_vec(acc2_up[j], vec(j - SGid + sg_size), "shuffle2_up"); } - /* GID XOR SGID */ - exit_if_not_equal_vec(acc_xor[j], vec(j ^ SGid), "shuffle_xor"); + /* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + exit_if_not_equal_vec(acc_xor[j], + vec(SGBeginGid + (SGLid ^ (SGid % sg_size))), + "shuffle_xor"); + SGLid++; } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); @@ -124,8 +133,7 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { } } -template -void check(queue &Queue, size_t G = 240, size_t L = 60) { +template void check(queue &Queue, size_t G = 240, size_t L = 60) { try { nd_range<1> NdRange(G, L); buffer buf2(G); @@ -171,8 +179,9 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { acc_up[NdItem.get_global_id()] = SG.shuffle_up(wggid, sgid); /* Save GID+SGID */ acc_down[NdItem.get_global_id()] = SG.shuffle_down(wggid, sgid); - /* Save GID XOR SGID */ - acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(wggid, sgid); + /* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + acc_xor[NdItem.get_global_id()] = + SG.shuffle_xor(wggid, sgid % SG.get_max_local_range()[0]); }); }); auto acc = buf.template get_access(); @@ -186,13 +195,20 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { size_t sg_size = sgsizeacc[0]; int SGid = 0; + int SGLid = 0; + int SGBeginGid = 0; for (int j = 0; j < G; j++) { if (j % L % sg_size == 0) { SGid++; + SGLid = 0; + SGBeginGid = j; } if (j % L == 0) { SGid = 0; + SGLid = 0; + SGBeginGid = j; } + /*GID of middle element in every subgroup*/ exit_if_not_equal(acc[j], j / L * L + SGid * sg_size + sg_size / 2, "shuffle"); @@ -215,8 +231,10 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) { if (j % L - SGid + sg_size < L) /* Do not go out LG*/ exit_if_not_equal(acc2_up[j], j - SGid + sg_size, "shuffle2_up"); } - /* GID XOR SGID */ - exit_if_not_equal(acc_xor[j], j ^ SGid, "shuffle_xor"); + /* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + exit_if_not_equal(acc_xor[j], SGBeginGid + (SGLid ^ (SGid % sg_size)), + "shuffle_xor"); + SGLid++; } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); From 85cd6133991a201dcadbad3855a1fdf5c421c140 Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Thu, 30 Jul 2020 17:59:55 +0300 Subject: [PATCH 2/2] Fix additional test --- sycl/test/sub_group/generic-shuffle.cpp | 54 ++++++++++++++++++------- 1 file changed, 39 insertions(+), 15 deletions(-) diff --git a/sycl/test/sub_group/generic-shuffle.cpp b/sycl/test/sub_group/generic-shuffle.cpp index d2d7e191dfa32..d4dbe84906537 100644 --- a/sycl/test/sub_group/generic-shuffle.cpp +++ b/sycl/test/sub_group/generic-shuffle.cpp @@ -18,8 +18,7 @@ #include "helper.hpp" #include #include -template -class pointer_kernel; +template class pointer_kernel; using namespace cl::sycl; @@ -59,8 +58,9 @@ void check_pointer(queue &Queue, size_t G = 240, size_t L = 60) { /* Save GID+SGID */ acc_down[NdItem.get_global_id()] = SG.shuffle_down(ptr, sgid); - /* Save GID XOR SGID */ - acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(ptr, sgid); + /* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + acc_xor[NdItem.get_global_id()] = + SG.shuffle_xor(ptr, sgid % SG.get_max_local_range()[0]); }); }); auto acc = buf.template get_access(); @@ -71,30 +71,44 @@ void check_pointer(queue &Queue, size_t G = 240, size_t L = 60) { size_t sg_size = sgsizeacc[0]; int SGid = 0; + int SGLid = 0; + int SGBeginGid = 0; for (int j = 0; j < G; j++) { if (j % L % sg_size == 0) { SGid++; + SGLid = 0; + SGBeginGid = j; } if (j % L == 0) { SGid = 0; + SGLid = 0; + SGBeginGid = j; } /*GID of middle element in every subgroup*/ - exit_if_not_equal(acc[j], static_cast(0x0) + (j / L * L + SGid * sg_size + sg_size / 2), + exit_if_not_equal(acc[j], + static_cast(0x0) + + (j / L * L + SGid * sg_size + sg_size / 2), "shuffle"); /* Value GID+SGID for all element except last SGID in SG*/ if (j % L % sg_size + SGid < sg_size && j % L + SGid < L) { - exit_if_not_equal(acc_down[j], static_cast(0x0) + (j + SGid), "shuffle_down"); + exit_if_not_equal(acc_down[j], static_cast(0x0) + (j + SGid), + "shuffle_down"); } /* Value GID-SGID for all element except first SGID in SG*/ if (j % L % sg_size >= SGid) { - exit_if_not_equal(acc_up[j], static_cast(0x0) + (j - SGid), "shuffle_up"); + exit_if_not_equal(acc_up[j], static_cast(0x0) + (j - SGid), + "shuffle_up"); } - /* GID XOR SGID */ - exit_if_not_equal(acc_xor[j], static_cast(0x0) + (j ^ SGid), "shuffle_xor"); + /* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + exit_if_not_equal(acc_xor[j], + static_cast(0x0) + + (SGBeginGid + (SGLid ^ (SGid % sg_size))), + "shuffle_xor"); + SGLid++; } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); @@ -145,8 +159,9 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) { /* Save GID+SGID */ acc_down[NdItem.get_global_id()] = SG.shuffle_down(val, sgid); - /* Save GID XOR SGID */ - acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(val, sgid); + /* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + acc_xor[NdItem.get_global_id()] = + SG.shuffle_xor(val, sgid % SG.get_max_local_range()[0]); }); }); auto acc = buf.template get_access(); @@ -157,17 +172,23 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) { size_t sg_size = sgsizeacc[0]; int SGid = 0; + int SGLid = 0; + int SGBeginGid = 0; for (int j = 0; j < G; j++) { if (j % L % sg_size == 0) { SGid++; + SGLid = 0; + SGBeginGid = j; } if (j % L == 0) { SGid = 0; + SGLid = 0; + SGBeginGid = j; } /*GID of middle element in every subgroup*/ - exit_if_not_equal(acc[j], values[j / L * L + SGid * sg_size + sg_size / 2], - "shuffle"); + exit_if_not_equal( + acc[j], values[j / L * L + SGid * sg_size + sg_size / 2], "shuffle"); /* Value GID+SGID for all element except last SGID in SG*/ if (j % L % sg_size + SGid < sg_size && j % L + SGid < L) { @@ -179,8 +200,11 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) { exit_if_not_equal(acc_up[j], values[j - SGid], "shuffle_up"); } - /* GID XOR SGID */ - exit_if_not_equal(acc_xor[j], values[j ^ SGid], "shuffle_xor"); + /* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */ + exit_if_not_equal(acc_xor[j], + values[SGBeginGid + (SGLid ^ (SGid % sg_size))], + "shuffle_xor"); + SGLid++; } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what();