Merge remote-tracking branch 'upstream/intel' into filter

Author: bso-intel

.github/CODEOWNERS

@@ -29,15 +29,13 @@ SYCL/DotProduct @rdeodhar
 # Explicit SIMD
 SYCL/ESIMD @kbobrovs @DenisBakhvalov
 
-# Function pointers
-SYCL/FunctionPointers @AlexeySachkov
-
 # Functor
 SYCL/Functor @AlexeySachkov
 
 # Group algorithms
 SYCL/GroupAlgorithm @Pennycook @AlexeySachkov
 SYCL/SubGroup @Pennycook @AlexeySachkov
+SYCL/GroupMask @Pennycook @vladimilaz
 
 # Group local memory
 SYCL/GroupLocalMemory @sergey-semenov @Pennycook
@@ -66,11 +64,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/spec_const_aot.cpp

@@ -36,8 +36,8 @@ int main(int argc, char **argv) {
             << "\n";
   cl::sycl::program prog(q.get_context());
 
-  cl::sycl::ONEAPI::experimental::spec_constant<int32_t, MyInt32Const> i32 =
-      prog.set_spec_constant<MyInt32Const>(10);
+  cl::sycl::ext::oneapi::experimental::spec_constant<int32_t, MyInt32Const>
+      i32 = prog.set_spec_constant<MyInt32Const>(10);
 
   prog.build_with_kernel_type<Kernel>();
 

SYCL/AtomicRef/accessor.cpp

@@ -10,7 +10,7 @@
 #include <numeric>
 #include <vector>
 using namespace sycl;
-using namespace sycl::ONEAPI;
+using namespace sycl::ext::oneapi;
 
 // Equivalent to add_test from add.cpp
 // Uses atomic_accessor instead of atomic_ref

SYCL/AtomicRef/add.cpp

@@ -3,188 +3,36 @@
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-#include <CL/sycl.hpp>
-#include <algorithm>
-#include <cassert>
+#include "add.h"
 #include <iostream>
-#include <numeric>
-#include <vector>
 using namespace sycl;
-using namespace sycl::ONEAPI;
-
-template <typename T, typename Difference = T>
-void add_fetch_test(queue q, size_t N) {
-  T sum = 0;
-  std::vector<T> output(N);
-  std::fill(output.begin(), output.end(), T(0));
-  {
-    buffer<T> sum_buf(&sum, 1);
-    buffer<T> output_buf(output.data(), output.size());
-
-    q.submit([&](handler &cgh) {
-      auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
-      auto out =
-          output_buf.template get_access<access::mode::discard_write>(cgh);
-      cgh.parallel_for(range<1>(N), [=](item<1> it) {
-        int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
-        out[gid] = atm.fetch_add(Difference(1));
-      });
-    });
-  }
-
-  // All work-items increment by 1, so final value should be equal to N
-  assert(sum == T(N));
-
-  // Fetch returns original value: will be in [0, N-1]
-  auto min_e = std::min_element(output.begin(), output.end());
-  auto max_e = std::max_element(output.begin(), output.end());
-  assert(*min_e == T(0) && *max_e == T(N - 1));
-
-  // Intermediate values should be unique
-  std::sort(output.begin(), output.end());
-  assert(std::unique(output.begin(), output.end()) == output.end());
-}
-
-template <typename T, typename Difference = T>
-void add_plus_equal_test(queue q, size_t N) {
-  T sum = 0;
-  std::vector<T> output(N);
-  std::fill(output.begin(), output.end(), T(0));
-  {
-    buffer<T> sum_buf(&sum, 1);
-    buffer<T> output_buf(output.data(), output.size());
-
-    q.submit([&](handler &cgh) {
-      auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
-      auto out =
-          output_buf.template get_access<access::mode::discard_write>(cgh);
-      cgh.parallel_for(range<1>(N), [=](item<1> it) {
-        int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
-        out[gid] = atm += Difference(1);
-      });
-    });
-  }
-
-  // All work-items increment by 1, so final value should be equal to N
-  assert(sum == T(N));
-
-  // += returns updated value: will be in [1, N]
-  auto min_e = std::min_element(output.begin(), output.end());
-  auto max_e = std::max_element(output.begin(), output.end());
-  assert(*min_e == T(1) && *max_e == T(N));
-
-  // Intermediate values should be unique
-  std::sort(output.begin(), output.end());
-  assert(std::unique(output.begin(), output.end()) == output.end());
-}
-
-template <typename T, typename Difference = T>
-void add_pre_inc_test(queue q, size_t N) {
-  T sum = 0;
-  std::vector<T> output(N);
-  std::fill(output.begin(), output.end(), T(0));
-  {
-    buffer<T> sum_buf(&sum, 1);
-    buffer<T> output_buf(output.data(), output.size());
-
-    q.submit([&](handler &cgh) {
-      auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
-      auto out =
-          output_buf.template get_access<access::mode::discard_write>(cgh);
-      cgh.parallel_for(range<1>(N), [=](item<1> it) {
-        int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
-        out[gid] = ++atm;
-      });
-    });
-  }
-
-  // All work-items increment by 1, so final value should be equal to N
-  assert(sum == T(N));
-
-  // Pre-increment returns updated value: will be in [1, N]
-  auto min_e = std::min_element(output.begin(), output.end());
-  auto max_e = std::max_element(output.begin(), output.end());
-  assert(*min_e == T(1) && *max_e == T(N));
-
-  // Intermediate values should be unique
-  std::sort(output.begin(), output.end());
-  assert(std::unique(output.begin(), output.end()) == output.end());
-}
-
-template <typename T, typename Difference = T>
-void add_post_inc_test(queue q, size_t N) {
-  T sum = 0;
-  std::vector<T> output(N);
-  std::fill(output.begin(), output.end(), T(0));
-  {
-    buffer<T> sum_buf(&sum, 1);
-    buffer<T> output_buf(output.data(), output.size());
-
-    q.submit([&](handler &cgh) {
-      auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
-      auto out =
-          output_buf.template get_access<access::mode::discard_write>(cgh);
-      cgh.parallel_for(range<1>(N), [=](item<1> it) {
-        int gid = it.get_id(0);
-        auto atm = atomic_ref<T, memory_order::relaxed, memory_scope::device,
-                              access::address_space::global_space>(sum[0]);
-        out[gid] = atm++;
-      });
-    });
-  }
-
-  // All work-items increment by 1, so final value should be equal to N
-  assert(sum == T(N));
-
-  // Post-increment returns original value: will be in [0, N-1]
-  auto min_e = std::min_element(output.begin(), output.end());
-  auto max_e = std::max_element(output.begin(), output.end());
-  assert(*min_e == T(0) && *max_e == T(N - 1));
-
-  // Intermediate values should be unique
-  std::sort(output.begin(), output.end());
-  assert(std::unique(output.begin(), output.end()) == output.end());
-}
-
-template <typename T, typename Difference = T>
-void add_test(queue q, size_t N) {
-  add_fetch_test<T, Difference>(q, N);
-  add_plus_equal_test<T, Difference>(q, N);
-  add_pre_inc_test<T, Difference>(q, N);
-  add_post_inc_test<T, Difference>(q, N);
-}
 
 // Floating-point types do not support pre- or post-increment
 template <> void add_test<float>(queue q, size_t N) {
   add_fetch_test<float>(q, N);
   add_plus_equal_test<float>(q, N);
 }
-template <> void add_test<double>(queue q, size_t N) {
-  add_fetch_test<double>(q, N);
-  add_plus_equal_test<double>(q, N);
-}
 
 int main() {
   queue q;
-  std::string version = q.get_device().get_info<info::device::version>();
 
   constexpr int N = 32;
   add_test<int>(q, N);
   add_test<unsigned int>(q, N);
-  add_test<long>(q, N);
-  add_test<unsigned long>(q, N);
-  add_test<long long>(q, N);
-  add_test<unsigned long long>(q, N);
   add_test<float>(q, N);
-  add_test<double>(q, N);
-  add_test<char *, ptrdiff_t>(q, N);
+
+  // Include long tests if they are 32 bits wide
+  if constexpr (sizeof(long) == 4) {
+    add_test<long>(q, N);
+    add_test<unsigned long>(q, N);
+  }
+
+  // Include pointer tests if they are 32 bits wide
+  if constexpr (sizeof(char *) == 4) {
+    add_test<char *, ptrdiff_t>(q, N);
+  }
 
   std::cout << "Test passed." << std::endl;
 }