Merge remote-tracking branch 'intel_llvm/intel' into simd_view_from_simd

Author: v-klochkov

.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,22 +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
-
-# Function pointers
-SYCL/FunctionPointers @AlexeySachkov
+SYCL/ESIMD @kbobrovs @v-klochkov
 
 # Functor
 SYCL/Functor @AlexeySachkov
 
 # Group algorithms
 SYCL/GroupAlgorithm @Pennycook @AlexeySachkov
 SYCL/SubGroup @Pennycook @AlexeySachkov
+SYCL/SubGroupMask @Pennycook @vladimilaz
 
 # Group local memory
 SYCL/GroupLocalMemory @sergey-semenov @Pennycook
@@ -66,8 +66,14 @@ SYCL/SeparateCompile @AlexeySachkov  @Fznamznon
 # Specialization constant
 SYCL/SpecConstants @kbobrovs
 
-# Specialization constant
-SYCL/USM @jbrodman
+# Unified Shared Memory (USM)
+SYCL/USM @jbrodman @sergey-semenov
 
 # Stream
 SYCL/Basic/stream @againull
+
+#BFloat16 conversion
+SYCL/BFloat16 @AlexeySotkin @MrSidims
+
+# Deprecated features
+SYCL/DeprecatedFeatures @intel/llvm-reviewers-runtime

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 "-device *" %S/Inputs/aot.cpp -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/AOT/multiple-devices.cpp

@@ -12,7 +12,7 @@
 
 // 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: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen "-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
@@ -26,7 +26,7 @@
 // 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
+// RUN: %clangxx -fsycl -fsycl-link-targets=spir64 %t.o -o %t.spv
 // AOT-compile device binary images
 // Neither of AOT tools can compile several files, hence, here is this
 // workaround

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,18 @@
+// 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
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --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.
+
+#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,18 @@
+// REQUIRES: windows
+// 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
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// CHECK-NOT:  One shouldn't see this message
+// FIXME Windows versionprints '(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.
+
+#include "assert_in_kernels.hpp"

SYCL/Assert/assert_in_multiple_tus.cpp

@@ -0,0 +1,17 @@
+// 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
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --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.
+
+#include "assert_in_multiple_tus.hpp"

SYCL/Assert/assert_in_multiple_tus.hpp

@@ -0,0 +1,51 @@
+#include "Inputs/kernels_in_file2.hpp"
+#include <CL/sycl.hpp>
+#include <iostream>
+
+#ifdef DEFINE_NDEBUG_INFILE1
+#define NDEBUG
+#else
+#undef NDEBUG
+#endif
+
+#include <cassert>
+
+using namespace cl::sycl;
+using namespace cl::sycl::access;
+
+static constexpr size_t BUFFER_SIZE = 16;
+
+int checkFunction() {
+  int X = calculus(0);
+  assert(X && "Nil in result");
+  return X;
+}
+
+void enqueueKernel_1_fromFile1(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 Kernel_1>(
+        cl::sycl::nd_range(Buf.get_range(), cl::sycl::range<1>(4)),
+        [=](cl::sycl::id<1> wiID) {
+          int X = 0;
+          if (wiID == 5)
+            X = checkFunction();
+          Acc[wiID] = X;
+        });
+  });
+}
+
+int main(int Argc, const char *Argv[]) {
+
+  queue Q;
+  enqueueKernel_1_fromFile1(&Q);
+  enqueueKernel_2_fromFile2(&Q);
+  Q.wait();
+
+  std::cout << "The test ended." << std::endl;
+  return 0;
+}

SYCL/Assert/assert_in_multiple_tus_one_ndebug.cpp

@@ -0,0 +1,16 @@
+// 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 -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %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
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// CHECK-NOT:  this message from calculus
+// CHECK:      {{.*}}assert_in_multiple_tus.hpp:20: int checkFunction(): global id: [5,0,0],
+// CHECK-SAME: local id: [1,0,0] Assertion `X && "Nil in result"` failed.
+// CHECK-NOT:  this message from file2
+// CHECK-NOT:  The test ended.

SYCL/Assert/assert_in_multiple_tus_one_ndebug_win.cpp

@@ -0,0 +1,16 @@
+// REQUIRES: windows
+// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %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
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// CHECK-NOT:  this message from calculus
+// FIXME Windows versionprints '(null)' instead of '<unknown func>' once in a
+// while for some insane reason.
+// CHECK:      {{.*}}assert_in_multiple_tus.hpp:20: {{<unknown func>|(null)}}: global id: [5,0,0],
+// CHECK-SAME: local id: [1,0,0] Assertion `X && "Nil in result"` failed.
+// CHECK-NOT:  this message from file2
+// CHECK-NOT:  The test ended.