Merge branch 'intel' into xpti_tests

Author: Alexander Batashev

.github/CODEOWNERS

@@ -7,6 +7,7 @@ SYCL/AOT @AGindinson @dm-vodopyanov @AlexeySachkov @romanovvlad
 SYCL/AtomicRef @AGindinson 
 
 # SYCL RT specific tests
+SYCL/Assert @intel/llvm-reviewers-runtime
 SYCL/Basic @intel/llvm-reviewers-runtime
 SYCL/Config @intel/llvm-reviewers-runtime
 SYCL/FilterSelector @intel/llvm-reviewers-runtime
@@ -22,20 +23,21 @@ SYCL/DeviceCodeSplit @AlexeySachkov  @Fznamznon
 
 # Device library
 SYCL/DeviceLib @vzakhari
+SYCL/DeviceLib/ITTAnnotations @vzakhari @MrSidims @AGindinson
 
 # dot_product API
 SYCL/DotProduct @rdeodhar
 
 # Explicit SIMD
-SYCL/ESIMD @kbobrovs @DenisBakhvalov
+SYCL/ESIMD @kbobrovs @v-klochkov
 
 # Functor
 SYCL/Functor @AlexeySachkov
 
 # Group algorithms
 SYCL/GroupAlgorithm @Pennycook @AlexeySachkov
 SYCL/SubGroup @Pennycook @AlexeySachkov
-SYCL/GroupMask @Pennycook @vladimilaz
+SYCL/SubGroupMask @Pennycook @vladimilaz
 
 # Group local memory
 SYCL/GroupLocalMemory @sergey-semenov @Pennycook

SYCL/.clang-format

@@ -1,2 +1,2 @@
 BasedOnStyle: LLVM
-CommentPragmas: "(RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK) *:|expected-"
+CommentPragmas: "(RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK[A-Za-z0-9_-]*) *:|expected-"

SYCL/AOT/accelerator.cpp

@@ -8,5 +8,5 @@
 
 // REQUIRES: aoc, accelerator
 
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga-unknown-unknown-sycldevice %S/Inputs/aot.cpp -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga %S/Inputs/aot.cpp -o %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

SYCL/AOT/cpu.cpp

@@ -8,8 +8,8 @@
 
 // REQUIRES: opencl-aot, cpu
 
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice %S/Inputs/aot.cpp -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 %S/Inputs/aot.cpp -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 
 // Test that opencl-aot can handle multiple build options.
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice  %S/Inputs/aot.cpp -Xsycl-target-backend "--bo=-g" -Xsycl-target-backend "--bo=-cl-opt-disable" -o %t2.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64  %S/Inputs/aot.cpp -Xsycl-target-backend "--bo=-g" -Xsycl-target-backend "--bo=-cl-opt-disable" -o %t2.out

SYCL/AOT/gpu.cpp

@@ -10,5 +10,5 @@
 // UNSUPPORTED: cuda
 // CUDA is not compatible with SPIR.
 //
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen-unknown-unknown-sycldevice -Xsycl-target-backend=spir64_gen-unknown-unknown-sycldevice "-device *" %S/Inputs/aot.cpp -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/aot.cpp -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/AOT/multiple-devices.cpp

@@ -6,72 +6,54 @@
 //
 //===----------------------------------------------------------------------===//
 
-// REQUIRES: opencl-aot, ocloc, aoc, cpu, gpu, accelerator, llvm-link, llvm-spirv
+// REQUIRES: opencl-aot, ocloc, aoc, cpu, gpu, accelerator
 // UNSUPPORTED: cuda
 // CUDA is not compatible with SPIR.
 
-// 1-command compilation case
-// Targeting CPU, GPU, FPGA
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice,spir64_gen-unknown-unknown-sycldevice,spir64_fpga-unknown-unknown-sycldevice -Xsycl-target-backend=spir64_gen-unknown-unknown-sycldevice "-device *" %S/Inputs/aot.cpp -o %t_all.out
-// RUN: %HOST_RUN_PLACEHOLDER %t_all.out
-// RUN: %CPU_RUN_PLACEHOLDER %t_all.out
-// RUN: %GPU_RUN_PLACEHOLDER %t_all.out
-// RUN: %ACC_RUN_PLACEHOLDER %t_all.out
+// Produce a fat object for all targets (generic SPIR-V, CPU, GPU, FPGA)
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_gen,spir64_fpga %S/Inputs/aot.cpp -c -o %t.o
 
-// FIXME: Change the behavior when proper automation for assert support is
-// introduced. For the time being, AOT flow can't detect if specific extension
-// is available for this or that device. The automation to be introduced is to
-// query native binary generator on specific features.
-//
-// Produce object file, spirv, device images to combine these differently
-// at link-time, thus testing various AOT-compiled images configurations
-// RUN: %clangxx -fsycl %S/Inputs/aot.cpp -c -o %t.o
-// RUN: %clangxx -fsycl -fsycl-link-targets=spir64-unknown-unknown-sycldevice %t.o -o %t.spv
-// AOT-compile device binary images
-// Neither of AOT tools can compile several files, hence, here is this
-// workaround
-// RUN: %llvm_spirv -r %sycl_libs_dir/libsycl-fallback-cassert.spv -o=%T/fallback-cassert.bc
-// RUN: %llvm_spirv -r %t.spv -o=%t.bc
-// RUN: %llvm_link %t.bc %T/fallback-cassert.bc -o=%t2.bc
-// RUN: %llvm_spirv %t2.bc -o=%t.spv
-// RUN: opencl-aot %t.spv -o=%t_cpu.ir --device=cpu
-// RUN: ocloc -file %t.spv -spirv_input -output %t_gen.out -output_no_suffix -device cfl
-// RUN: aoc %t.spv -o %t_fpga.aocx -sycl -dep-files=%t.d
+// CPU, GPU, FPGA
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_all_aot.out
+// RUN: %HOST_RUN_PLACEHOLDER %t_all_aot.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_all_aot.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_all_aot.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_all_aot.out
 
 // CPU, GPU
-// RUN: %clangxx -fsycl -fsycl-add-targets=spir64_x86_64:%t_cpu.ir,spir64_gen:%t_gen.out %t.o -o %t_cpu_gpu.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_cpu_gpu.out
 // RUN: %HOST_RUN_PLACEHOLDER %t_cpu_gpu.out
 // RUN: %CPU_RUN_PLACEHOLDER %t_cpu_gpu.out
 // RUN: %GPU_RUN_PLACEHOLDER %t_cpu_gpu.out
 
 // CPU, FPGA
-// RUN: %clangxx -fsycl -fsycl-add-targets=spir64_x86_64:%t_cpu.ir,spir64_fpga:%t_fpga.aocx %t.o -o %t_cpu_fpga.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_fpga %t.o -o %t_cpu_fpga.out
 // RUN: %HOST_RUN_PLACEHOLDER %t_cpu_fpga.out
 // RUN: %CPU_RUN_PLACEHOLDER %t_cpu_fpga.out
 // RUN: %ACC_RUN_PLACEHOLDER %t_cpu_fpga.out
 
 // GPU, FPGA
-// RUN: %clangxx -fsycl -fsycl-add-targets=spir64_gen:%t_gen.out,spir64_fpga:%t_fpga.aocx %t.o -o %t_gpu_fpga.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_gpu_fpga.out
 // RUN: %HOST_RUN_PLACEHOLDER %t_gpu_fpga.out
 // RUN: %GPU_RUN_PLACEHOLDER %t_gpu_fpga.out
 // RUN: %ACC_RUN_PLACEHOLDER %t_gpu_fpga.out
 
 // No AOT-compiled image for CPU
-// RUN: %clangxx -fsycl -fsycl-add-targets=spir64:%t.spv,spir64_gen:%t_gen.out,spir64_fpga:%t_fpga.aocx %t.o -o %t_spv_gpu_fpga.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_spv_gpu_fpga.out
 // RUN: %CPU_RUN_PLACEHOLDER %t_spv_gpu_fpga.out
 // Check that execution on AOT-compatible devices is unaffected
 // RUN: %GPU_RUN_PLACEHOLDER %t_spv_gpu_fpga.out
 // RUN: %ACC_RUN_PLACEHOLDER %t_spv_gpu_fpga.out
 
 // No AOT-compiled image for GPU
-// RUN: %clangxx -fsycl -fsycl-add-targets=spir64:%t.spv,spir64_x86_64:%t_cpu.ir,spir64_fpga:%t_fpga.aocx %t.o -o %t_spv_cpu_fpga.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_fpga %t.o -o %t_spv_cpu_fpga.out
 // RUN: %GPU_RUN_PLACEHOLDER %t_spv_cpu_fpga.out
 // Check that execution on AOT-compatible devices is unaffected
 // RUN: %CPU_RUN_PLACEHOLDER %t_spv_cpu_fpga.out
 // RUN: %ACC_RUN_PLACEHOLDER %t_spv_cpu_fpga.out
 
 // No AOT-compiled image for FPGA
-// RUN: %clangxx -fsycl -fsycl-add-targets=spir64:%t.spv,spir64_x86_64:%t_cpu.ir,spir64_gen:%t_gen.out %t.o -o %t_spv_cpu_gpu.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_spv_cpu_gpu.out
 // RUN: %ACC_RUN_PLACEHOLDER %t_spv_cpu_gpu.out
 // Check that execution on AOT-compatible devices is unaffected
 // RUN: %CPU_RUN_PLACEHOLDER %t_spv_cpu_gpu.out

SYCL/AOT/with-llvm-bc.cpp

@@ -8,8 +8,8 @@
 
 // REQUIRES: cpu, dump_ir
 
-// RUN: %clangxx -fsycl -fsycl-targets=spir64-unknown-unknown-sycldevice -c %S/Inputs/aot.cpp -o %t.o
-// RUN: %clangxx -fsycl -fsycl-link-targets=spir64-unknown-unknown-sycldevice %t.o -o %t.spv
+// RUN: %clangxx -fsycl -fsycl-targets=spir64 -c %S/Inputs/aot.cpp -o %t.o
+// RUN: %clangxx -fsycl -fsycl-link-targets=spir64 %t.o -o %t.spv
 // RUN: llvm-spirv -r %t.spv -o %t.bc
 // RUN: %clangxx -fsycl -fsycl-add-targets=spir64:%t.bc %t.o -o %t.out
 //

SYCL/Assert/Inputs/kernels_in_file2.cpp

@@ -0,0 +1,45 @@
+#include "kernels_in_file2.hpp"
+
+#ifdef DEFINE_NDEBUG_INFILE2
+#define NDEBUG
+#else
+#undef NDEBUG
+#endif
+
+#include <cassert>
+
+using namespace cl::sycl;
+using namespace cl::sycl::access;
+
+int calculus(int X) {
+  assert(X && "this message from calculus");
+  return X * 2;
+}
+
+void check_nil(int value) { assert(value && "this message from file2"); }
+
+static constexpr size_t BUFFER_SIZE = 4;
+
+void enqueueKernel_1_fromFile2(queue *Q) {
+  cl::sycl::range<1> numOfItems{BUFFER_SIZE};
+  cl::sycl::buffer<int, 1> Buf(numOfItems);
+
+  Q->submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class kernel1_from_separate_file>(
+        numOfItems, [=](cl::sycl::id<1> wiID) { check_nil(Acc[wiID]); });
+  });
+}
+
+void enqueueKernel_2_fromFile2(queue *Q) {
+  cl::sycl::range<1> numOfItems{BUFFER_SIZE};
+  cl::sycl::buffer<int, 1> Buf(numOfItems);
+
+  Q->submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class kernel2_from_separate_file>(
+        numOfItems, [=](cl::sycl::id<1> wiID) { check_nil(Acc[wiID]); });
+  });
+}

SYCL/Assert/Inputs/kernels_in_file2.hpp

@@ -0,0 +1,7 @@
+#include <CL/sycl.hpp>
+
+SYCL_EXTERNAL int calculus(int X);
+
+void enqueueKernel_1_fromFile2(sycl::queue *Q);
+
+void enqueueKernel_2_fromFile2(sycl::queue *Q);

SYCL/Assert/assert_in_kernels.cpp

@@ -0,0 +1,22 @@
+// REQUIRES: linux
+// FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// Shouldn't fail on ACC as fallback assert isn't enqueued there
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt
+//
+// CHECK-NOT:  One shouldn't see this message
+// CHECK:      {{.*}}assert_in_kernels.hpp:26: void kernelFunc2(int *, int): global id: [{{[0,2]}},0,0], local id: [0,0,0]
+// CHECK-SAME: Assertion `Buf[wiID] == 0 && "from assert statement"` failed.
+// CHECK-NOT:  test aborts earlier, one shouldn't see this message
+// CHECK-NOT:  The test ended.
+//
+// CHECK-ACC-NOT: {{.*}}assert_in_kernels.hpp:26: void kernelFunc2(int *, int): global id: [{{[0,2]}},0,0], local id: [0,0,0]
+// CHECK-ACC: The test ended.
+
+#include "assert_in_kernels.hpp"

SYCL/Assert/assert_in_kernels.hpp

@@ -0,0 +1,72 @@
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+using namespace cl::sycl::access;
+
+void kernelFunc1(int *Buf, int wiID) {
+  Buf[wiID] = 9;
+  assert(Buf[wiID] != 0 && "One shouldn't see this message");
+}
+
+void assertTest1(queue &Q, buffer<int, 1> &Buf) {
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class Kernel_1>(
+        Buf.get_range(),
+        [=](cl::sycl::id<1> wiID) { kernelFunc1(&Acc[0], wiID); });
+  });
+}
+
+void kernelFunc2(int *Buf, int wiID) {
+  if (wiID % 2 != 0)
+    Buf[wiID] = 0;
+  assert(Buf[wiID] == 0 && "from assert statement");
+}
+
+void assertTest2(queue &Q, buffer<int, 1> &Buf) {
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class Kernel_2>(
+        Buf.get_range(),
+        [=](cl::sycl::id<1> wiID) { kernelFunc2(&Acc[0], wiID); });
+  });
+}
+
+void kernelFunc3(int *Buf, int wiID) {
+  if (wiID == 0)
+    assert(false && "test aborts earlier, one shouldn't see this message");
+  Buf[wiID] = 9;
+}
+
+void assertTest3(queue &Q, buffer<int, 1> &Buf) {
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class Kernel_3>(
+        Buf.get_range(),
+        [=](cl::sycl::id<1> wiID) { kernelFunc3(&Acc[0], wiID); });
+  });
+}
+
+int main(int Argc, const char *Argv[]) {
+  std::array<int, 4> Vec = {1, 2, 3, 4};
+  cl::sycl::range<1> numOfItems{Vec.size()};
+  cl::sycl::buffer<int, 1> Buf(Vec.data(), numOfItems);
+
+  queue Q;
+  assertTest1(Q, Buf);
+  Q.wait();
+
+  assertTest2(Q, Buf);
+  Q.wait();
+
+  assertTest3(Q, Buf);
+  Q.wait();
+
+  std::cout << "The test ended." << std::endl;
+  return 0;
+}

SYCL/Assert/assert_in_kernels_ndebug.cpp

@@ -0,0 +1,11 @@
+// FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DNDEBUG %S/assert_in_kernels.cpp -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+//
+// CHECK-NOT: One shouldn't see this message
+// CHECK-NOT: from assert statement
+// CHECK-NOT: test aborts earlier, one shouldn't see this message
+// CHECK: The test ended.

SYCL/Assert/assert_in_kernels_win.cpp

@@ -0,0 +1,23 @@
+// REQUIRES: windows
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// Shouldn't fail on ACC as fallback assert isn't enqueued there
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt
+//
+// CHECK-NOT:  One shouldn't see this message
+// FIXME Windows version prints '(null)' instead of '<unknown func>' once in a
+// while for some insane reason.
+// CHECK:      {{.*}}assert_in_kernels.hpp:26: {{<unknown func>|(null)}}: global id: [{{[0,2]}},0,0], local id: [0,0,0]
+// CHECK-SAME: Assertion `Buf[wiID] == 0 && "from assert statement"` failed.
+// CHECK-NOT:  test aborts earlier, one shouldn't see this message
+// CHECK-NOT:  The test ended.
+//
+// CHECK-ACC-NOT: {{.*}}assert_in_kernels.hpp:26: {{<unknown func>|(null)}}: global id: [{{[0,2]}},0,0], local id: [0,0,0]
+// CHECK-ACC: The test ended.
+
+#include "assert_in_kernels.hpp"

SYCL/Assert/assert_in_multiple_tus.cpp

@@ -0,0 +1,21 @@
+// REQUIRES: linux
+// FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// Shouldn't fail on ACC as fallback assert isn't enqueued there
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt
+//
+// CHECK:      {{.*}}kernels_in_file2.cpp:15: int calculus(int): global id: [5,0,0], local id: [1,0,0]
+// CHECK-SAME: Assertion `X && "this message from calculus"` failed.
+// CHECK-NOT:  this message from file2
+// CHECK-NOT:  The test ended.
+//
+// CHECK-ACC-NOT: {{.*}}kernels_in_file2.cpp:15: int calculus(int): global id: [5,0,0], local id: [1,0,0]
+// CHECK-ACC: The test ended.
+
+#include "assert_in_multiple_tus.hpp"