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

Author: v-klochkov

.github/CODEOWNERS

@@ -10,6 +10,7 @@ SYCL/AtomicRef @AGindinson
 SYCL/Assert @intel/llvm-reviewers-runtime
 SYCL/Basic @intel/llvm-reviewers-runtime
 SYCL/Config @intel/llvm-reviewers-runtime
+SYCL/DiscardEvents @intel/llvm-reviewers-runtime
 SYCL/FilterSelector @intel/llvm-reviewers-runtime
 SYCL/HostInteropTask @intel/llvm-reviewers-runtime
 SYCL/InorderQueue @intel/llvm-reviewers-runtime
@@ -29,7 +30,7 @@ SYCL/DeviceLib/ITTAnnotations @vzakhari @MrSidims @AGindinson
 SYCL/DotProduct @rdeodhar
 
 # Explicit SIMD
-SYCL/ESIMD @kbobrovs @v-klochkov
+SYCL/ESIMD @kbobrovs @v-klochkov @sndmitriev
 
 # Functor
 SYCL/Functor @AlexeySachkov

SYCL/Basic/barrier_order.cpp

@@ -1,3 +1,4 @@
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/Basic/enqueue_barrier.cpp

@@ -5,7 +5,7 @@
 
 // The test is failing sporadically on Windows OpenCL RTs
 // Disabling on windows until fixed
-// UNSUPPORTED: windows
+// UNSUPPORTED: hip_amd, windows
 
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/fpga_device_selector.hpp>

SYCL/Basic/parallel_for_user_types.cpp

@@ -0,0 +1,85 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test performs basic check of supporting user defined class that are
+// implicitly converted from sycl::item/sycl::nd_item in parallel_for.
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+template <int Dimensions> class item_wrapper {
+public:
+  item_wrapper(sycl::item<Dimensions> it) : m_item(it) {}
+
+  size_t get() { return m_item; }
+
+private:
+  sycl::item<Dimensions> m_item;
+};
+
+template <int Dimensions> class nd_item_wrapper {
+public:
+  nd_item_wrapper(sycl::nd_item<Dimensions> it) : m_item(it) {}
+
+  size_t get() { return m_item.get_global_linear_id(); }
+
+private:
+  sycl::nd_item<Dimensions> m_item;
+};
+
+int main() {
+  sycl::queue q;
+
+  // Initialize data array
+  const int sz = 16;
+  int data[sz] = {0};
+  for (int i = 0; i < sz; ++i) {
+    data[i] = i;
+  }
+
+  // Check user defined sycl::item wrapper
+  sycl::buffer<int> data_buf(data, sz);
+  q.submit([&](sycl::handler &h) {
+    auto buf_acc = data_buf.get_access<sycl::access::mode::read_write>(h);
+    h.parallel_for(sycl::range<1>{sz},
+                   [=](item_wrapper<1> item) { buf_acc[item.get()] += 1; });
+  });
+  q.wait();
+  bool failed = false;
+
+  {
+    auto buf_acc = data_buf.get_access<sycl::access::mode::read>();
+    for (int i = 0; i < sz; ++i) {
+      failed |= (buf_acc[i] != i + 1);
+    }
+    if (failed) {
+      std::cout << "item_wrapper check failed" << std::endl;
+      return 1;
+    }
+  }
+
+  // Check user defined sycl::nd_item wrapper
+  q.submit([&](sycl::handler &h) {
+    auto buf_acc = data_buf.get_access<sycl::access::mode::read_write>(h);
+    h.parallel_for(sycl::nd_range<1>{sz, 2},
+                   [=](nd_item_wrapper<1> item) { buf_acc[item.get()] += 1; });
+  });
+  q.wait();
+
+  {
+    auto buf_acc = data_buf.get_access<sycl::access::mode::read>();
+    for (int i = 0; i < sz; ++i) {
+      failed |= (buf_acc[i] != i + 2);
+    }
+    if (failed) {
+      std::cout << "nd_item_wrapper check failed" << std::endl;
+      return 1;
+    }
+  }
+
+  std::cout << "Test passed" << std::endl;
+  return 0;
+}

SYCL/Basic/stream/stream.cpp

@@ -4,9 +4,6 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
 // RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
-//
-// Missing built-ins on AMD
-// XFAIL: hip_amd
 
 //==------------------ stream.cpp - SYCL stream basic test -----------------==//
 //

SYCL/Basic/submit_barrier.cpp

@@ -1,3 +1,4 @@
+// UNSUPPORTED: hip_amd
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/DeviceLib/built-ins/nan.cpp

@@ -5,8 +5,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t_gpu.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// XFAIL: cuda
-
 #include <CL/sycl.hpp>
 
 #include <cassert>

SYCL/DeviceLib/built-ins/scalar_math.cpp

@@ -3,9 +3,6 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
-//
-// Incorrect results with hip on AMD
-// XFAIL: hip_amd
 
 #include <CL/sycl.hpp>
 

SYCL/DeviceLib/built-ins/vector_relational.cpp

@@ -1,3 +1,5 @@
+// FIXME unsupported on windows (opencl and level-zero) until fix of libdevice
+// UNSUPPORTED: windows && (opencl || level_zero)
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/DeviceLib/cmath_fp64_test.cpp

@@ -1,3 +1,5 @@
+// FIXME unsupported on windows (opencl) until fix of libdevice fails
+// UNSUPPORTED: windows && opencl
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/DeviceLib/cmath_test.cpp

@@ -1,3 +1,5 @@
+// FIXME unsupported on windows (opencl) until fix of libdevice fails
+// UNSUPPORTED: windows && opencl
 // RUN: %clangxx -fsycl -fno-builtin %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/DeviceLib/math_fp64_test.cpp

@@ -1,3 +1,5 @@
+// FIXME unsupported on windows (opencl) until fix of libdevice fails
+// UNSUPPORTED: windows && opencl
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/DeviceLib/math_test.cpp

@@ -1,3 +1,5 @@
+// FIXME unsupported on windows (opencl) until fix of libdevice fails
+// UNSUPPORTED: windows && opencl
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/DiscardEvents/discard_events_accessors.cpp

@@ -0,0 +1,106 @@
+// FIXME unsupported on level_zero until L0 Plugin support becomes available for
+// discard_queue_events
+// UNSUPPORTED: level_zero
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is `nullptr` for
+// piEnqueueKernelLaunch for USM kernel using local accessor, but
+// is not `nullptr` for kernel using buffer accessor.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: The test passed.
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+static constexpr int MAGIC_NUM = -1;
+static constexpr size_t BUFFER_SIZE = 16;
+
+void RunKernelHelper(sycl::queue Q,
+                     const std::function<void(int *Harray)> &TestFunction) {
+  int *Harray = sycl::malloc_host<int>(BUFFER_SIZE, Q);
+  assert(Harray != nullptr);
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    Harray[i] = MAGIC_NUM;
+  }
+
+  TestFunction(Harray);
+
+  // Checks result
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    size_t expected = i + 10;
+    assert(Harray[i] == expected);
+  }
+  free(Harray, Q);
+}
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(props);
+  sycl::range<1> Range(BUFFER_SIZE);
+
+  RunKernelHelper(Q, [&](int *Harray) {
+    Q.submit([&](sycl::handler &CGH) {
+      const size_t LocalMemSize = BUFFER_SIZE;
+      using LocalAccessor =
+          sycl::accessor<int, 1, sycl::access::mode::read_write,
+                         sycl::access::target::local>;
+      LocalAccessor LocalAcc(LocalMemSize, CGH);
+
+      CGH.parallel_for<class kernel_using_local_memory>(
+          Range, [=](sycl::item<1> itemID) {
+            size_t i = itemID.get_id(0);
+            int *Ptr = LocalAcc.get_pointer();
+            Ptr[i] = i + 5;
+            Harray[i] = Ptr[i] + 5;
+          });
+    });
+    Q.wait();
+  });
+
+  RunKernelHelper(Q, [&](int *Harray) {
+    sycl::buffer<int, 1> Buf(Range);
+    Q.submit([&](sycl::handler &CGH) {
+      auto Acc = Buf.get_access<sycl::access::mode::read_write>(CGH);
+      CGH.parallel_for<class kernel_using_buffer_accessor>(
+          Range, [=](sycl::item<1> itemID) {
+            size_t i = itemID.get_id(0);
+            Harray[i] = i + 10;
+            Acc[i] = i + 20;
+          });
+    });
+    Q.wait();
+
+    // Checks result
+    auto HostAcc = Buf.get_access<sycl::access::mode::read>();
+    for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+      size_t expected = i + 20;
+      assert(HostAcc[i] == expected);
+    }
+  });
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}

SYCL/DiscardEvents/discard_events_host_task.cpp

@@ -0,0 +1,57 @@
+// If necessary, the test can be removed as run_on_host_intel() is deprecated
+// and host_task() which should be used instead does not use the PI call
+// piEnqueueNativeKernel
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is `nullptr` for
+// piEnqueueNativeKernel.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// CHECK: ---> piEnqueueNativeKernel(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: The test passed.
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+
+void CheckArray(sycl::queue Q, int *x, size_t buffer_size, int expected) {
+  Q.wait();
+  for (size_t i = 0; i < buffer_size; ++i)
+    assert(x[i] == expected);
+}
+
+static constexpr size_t BUFFER_SIZE = 16;
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list Props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(Props);
+
+  int *x = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(x != nullptr);
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.run_on_host_intel([=]() {
+      for (size_t i = 0; i < BUFFER_SIZE; ++i)
+        x[i] = 8;
+    });
+  });
+  CheckArray(Q, x, BUFFER_SIZE, 8);
+
+  Q.wait();
+  free(x, Q);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}

SYCL/DiscardEvents/discard_events_kernel_using_assert.hpp

@@ -0,0 +1,45 @@
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+static constexpr int MAGIC_NUM = -1;
+static constexpr size_t BUFFER_SIZE = 16;
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list Props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(Props);
+
+  sycl::range<1> Range(BUFFER_SIZE);
+  int *Harray = sycl::malloc_host<int>(BUFFER_SIZE, Q);
+  if (Harray == nullptr) {
+    return -1;
+  }
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    Harray[i] = MAGIC_NUM;
+  }
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.parallel_for<class kernel_using_assert>(
+        Range, [=](sycl::item<1> itemID) {
+          size_t i = itemID.get_id(0);
+          Harray[i] = i + 10;
+          assert(Harray[i] == i + 10 && "assert message");
+        });
+  });
+  Q.wait();
+
+  // Checks result
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    size_t expected = i + 10;
+    if (Harray[i] != expected)
+      return -1;
+  }
+  free(Harray, Q);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}