Merge branch 'intel' of https://github.com/intel/llvm-test-suite into esimd_emulator_ignore_cm_warnings

Author: dongkyunahn-intel

.github/CODEOWNERS

@@ -1,80 +1,35 @@
 * @vladimirlaz @romanovvlad @bader
 
-# AOT compilation
-SYCL/AOT @AGindinson @dm-vodopyanov @AlexeySachkov @romanovvlad
+# Use runtime team as the umbrella for most of the tests
+SYCL @intel/llvm-reviewers-runtime
 
-# Atomics
-SYCL/AtomicRef @AGindinson 
+# AOT compilation
+SYCL/AOT @intel/dpcpp-tools-reviewers
 
-# SYCL RT specific tests
-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
-SYCL/KernelAndProgram @intel/llvm-reviewers-runtime
-SYCL/KernelParams @intel/llvm-reviewers-runtime
-SYCL/Scheduler @intel/llvm-reviewers-runtime
-SYCL/Tracing  @intel/llvm-reviewers-runtime
+#BFloat16 conversion
+SYCL/BFloat16 @AlexeySotkin @intel/dpcpp-tools-reviewers
 
 # Device code split
-SYCL/DeviceCodeSplit @AlexeySachkov  @Fznamznon
-
-# Device library
-SYCL/DeviceLib @vzakhari
-SYCL/DeviceLib/ITTAnnotations @vzakhari @MrSidims @AGindinson
-
-# dot_product API
-SYCL/DotProduct @rdeodhar
+SYCL/DeviceCodeSplit @intel/dpcpp-tools-reviewers
 
 # Explicit SIMD
-SYCL/ESIMD @kbobrovs @v-klochkov @sndmitriev
-
-# Functor
-SYCL/Functor @AlexeySachkov
+SYCL/ESIMD @intel/dpcpp-esimd-reviewers
 
 # Group algorithms
-SYCL/GroupAlgorithm @Pennycook @AlexeySachkov
-SYCL/SubGroup @Pennycook @AlexeySachkov
-SYCL/SubGroupMask @Pennycook @vladimilaz
-
-# Group local memory
-SYCL/GroupLocalMemory @sergey-semenov @Pennycook
-
-# Hierarchical parallelism
-SYCL/HierPar @kbobrovs
+SYCL/GroupAlgorithm @Pennycook @intel/llvm-reviewers-runtime
+SYCL/SubGroup @Pennycook @intel/llvm-reviewers-runtime
+SYCL/SubGroupMask @Pennycook @intel/llvm-reviewers-runtime
+SYCL/GroupLocalMemory @Pennycook @intel/llvm-reviewers-runtime
 
-# Inline assembler
-SYCL/InlineAsm @AlexeySachkov
+# Plugin interface for Level Zero
+SYCL/Plugin/*level-zero* @intel/dpcpp-l0-pi-reviewers
+SYCL/Plugin/*level_zero* @intel/dpcpp-l0-pi-reviewers
 
-# Online compiler
-SYCL/OnlineCompiler @v-klochkov
-
-# Plugin interface
-SYCL/Plugin @smaslov-intel
-
-# Reduction algorithms
-SYCL/Reduction @v-klochkov
-
-# Image sampler
-SYCL/Sampler @cperkinsintel
+# Printf
+SYCL/Printf @intel/dpcpp-tools-reviewers
 
 # Separate compilation
-SYCL/SeparateCompile @AlexeySachkov  @Fznamznon
+SYCL/SeparateCompile @intel/dpcpp-tools-reviewers
 
 # Specialization constant
-SYCL/SpecConstants @kbobrovs
-
-# 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/SpecConstants @intel/dpcpp-tools-reviewers

SYCL/AtomicRef/add.cpp

@@ -12,10 +12,14 @@ using namespace sycl;
 
 // Floating-point types do not support pre- or post-increment
 template <> void add_test<float>(queue q, size_t N) {
-  add_fetch_test<::sycl::ext::oneapi::atomic_ref, float>(q, N);
-  add_fetch_test<::sycl::atomic_ref, float>(q, N);
-  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, float>(q, N);
-  add_plus_equal_test<::sycl::atomic_ref, float>(q, N);
+  add_fetch_test<::sycl::ext::oneapi::atomic_ref,
+                 access::address_space::global_space, float>(q, N);
+  add_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
+                 float>(q, N);
+  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
+                      access::address_space::global_space, float>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
+                      float>(q, N);
 }
 
 int main() {

SYCL/AtomicRef/add.h

@@ -11,7 +11,8 @@ using namespace sycl::ext::oneapi;
 
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
-          typename T, typename Difference = T>
+          access::address_space address_space, typename T,
+          typename Difference = T>
 void add_fetch_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -27,7 +28,7 @@ void add_fetch_test(queue q, size_t N) {
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
         auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
-                             access::address_space::global_space>(sum[0]);
+                             address_space>(sum[0]);
         out[gid] = atm.fetch_add(Difference(1));
       });
     });
@@ -48,7 +49,8 @@ void add_fetch_test(queue q, size_t N) {
 
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
-          typename T, typename Difference = T>
+          access::address_space address_space, typename T,
+          typename Difference = T>
 void add_plus_equal_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -64,7 +66,7 @@ void add_plus_equal_test(queue q, size_t N) {
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
         auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
-                             access::address_space::global_space>(sum[0]);
+                             address_space>(sum[0]);
         out[gid] = atm += Difference(1);
       });
     });
@@ -85,7 +87,8 @@ void add_plus_equal_test(queue q, size_t N) {
 
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
-          typename T, typename Difference = T>
+          access::address_space address_space, typename T,
+          typename Difference = T>
 void add_pre_inc_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -101,7 +104,7 @@ void add_pre_inc_test(queue q, size_t N) {
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
         auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
-                             access::address_space::global_space>(sum[0]);
+                             address_space>(sum[0]);
         out[gid] = ++atm;
       });
     });
@@ -122,7 +125,8 @@ void add_pre_inc_test(queue q, size_t N) {
 
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
-          typename T, typename Difference = T>
+          access::address_space address_space, typename T,
+          typename Difference = T>
 void add_post_inc_test(queue q, size_t N) {
   T sum = 0;
   std::vector<T> output(N);
@@ -138,7 +142,7 @@ void add_post_inc_test(queue q, size_t N) {
       cgh.parallel_for(range<1>(N), [=](item<1> it) {
         int gid = it.get_id(0);
         auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
-                             access::address_space::global_space>(sum[0]);
+                             address_space>(sum[0]);
         out[gid] = atm++;
       });
     });
@@ -159,12 +163,34 @@ void add_post_inc_test(queue q, size_t N) {
 
 template <typename T, typename Difference = T>
 void add_test(queue q, size_t N) {
-  add_fetch_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
-  add_fetch_test<::sycl::atomic_ref, T, Difference>(q, N);
-  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
-  add_plus_equal_test<::sycl::atomic_ref, T, Difference>(q, N);
-  add_pre_inc_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
-  add_pre_inc_test<::sycl::atomic_ref, T, Difference>(q, N);
-  add_post_inc_test<::sycl::ext::oneapi::atomic_ref, T, Difference>(q, N);
-  add_post_inc_test<::sycl::atomic_ref, T, Difference>(q, N);
+  add_fetch_test<::sycl::ext::oneapi::atomic_ref,
+                 access::address_space::global_space, T, Difference>(q, N);
+  add_fetch_test<::sycl::atomic_ref, access::address_space::global_space, T,
+                 Difference>(q, N);
+  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
+                      access::address_space::global_space, T, Difference>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
+                      T, Difference>(q, N);
+  add_pre_inc_test<::sycl::ext::oneapi::atomic_ref,
+                   access::address_space::global_space, T, Difference>(q, N);
+  add_pre_inc_test<::sycl::atomic_ref, access::address_space::global_space, T,
+                   Difference>(q, N);
+  add_post_inc_test<::sycl::ext::oneapi::atomic_ref,
+                    access::address_space::global_space, T, Difference>(q, N);
+  add_post_inc_test<::sycl::atomic_ref, access::address_space::global_space, T,
+                    Difference>(q, N);
+}
+
+template <typename T, typename Difference = T>
+void add_generic_test(queue q, size_t N) {
+  add_fetch_test<::sycl::atomic_ref, access::address_space::generic_space, T,
+                 Difference>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::generic_space,
+                      T, Difference>(q, N);
+  add_pre_inc_test<::sycl::atomic_ref, access::address_space::generic_space, T,
+                   Difference>(q, N);
+  add_post_inc_test<::sycl::atomic_ref, access::address_space::generic_space, T,
+                    Difference>(q, N);
+  add_post_inc_test<::sycl::atomic_ref, access::address_space::global_space, T,
+                    Difference>(q, N);
 }

SYCL/AtomicRef/add_atomic64.cpp

@@ -11,10 +11,14 @@ using namespace sycl;
 
 // Floating-point types do not support pre- or post-increment
 template <> void add_test<double>(queue q, size_t N) {
-  add_fetch_test<::sycl::ext::oneapi::atomic_ref, double>(q, N);
-  add_fetch_test<::sycl::atomic_ref, double>(q, N);
-  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref, double>(q, N);
-  add_plus_equal_test<::sycl::atomic_ref, double>(q, N);
+  add_fetch_test<::sycl::ext::oneapi::atomic_ref,
+                 access::address_space::global_space, double>(q, N);
+  add_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
+                 double>(q, N);
+  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
+                      access::address_space::global_space, double>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
+                      double>(q, N);
 }
 
 int main() {

SYCL/AtomicRef/add_atomic64_generic.cpp

@@ -0,0 +1,52 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
+// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// CUDA backend has had no support for the generic address space yet
+// XFAIL: cuda || hip
+
+#include "add.h"
+#include <iostream>
+using namespace sycl;
+
+// Floating-point types do not support pre- or post-increment
+template <> void add_generic_test<double>(queue q, size_t N) {
+  add_fetch_test<::sycl::atomic_ref, access::address_space::generic_space,
+                 double>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::generic_space,
+                      double>(q, N);
+}
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  add_generic_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    add_generic_test<long>(q, N);
+    add_generic_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    add_generic_test<long long>(q, N);
+    add_generic_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    add_generic_test<char *, ptrdiff_t>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/add_generic.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
+// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// CUDA backend has had no support for the generic address space yet
+// XFAIL: cuda || hip
+
+#include "add.h"
+#include <iostream>
+using namespace sycl;
+
+// Floating-point types do not support pre- or post-increment
+template <> void add_generic_test<float>(queue q, size_t N) {
+  add_fetch_test<::sycl::atomic_ref, access::address_space::generic_space,
+                 float>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::generic_space,
+                      float>(q, N);
+}
+
+int main() {
+  queue q;
+
+  constexpr int N = 32;
+  add_generic_test<int>(q, N);
+  add_generic_test<unsigned int>(q, N);
+  add_generic_test<float>(q, N);
+
+  // Include long tests if they are 32 bits wide
+  if constexpr (sizeof(long) == 4) {
+    add_generic_test<long>(q, N);
+    add_generic_test<unsigned long>(q, N);
+  }
+
+  // Include pointer tests if they are 32 bits wide
+  if constexpr (sizeof(char *) == 4) {
+    add_generic_test<char *, ptrdiff_t>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/assignment.h

@@ -11,12 +11,12 @@ using namespace sycl::ext::oneapi;
 
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
-          typename T>
+          access::address_space address_space, typename T>
 class assignment_kernel;
 
 template <template <typename, memory_order, memory_scope, access::address_space>
           class AtomicRef,
-          typename T>
+          access::address_space address_space, typename T>
 void assignment_test(queue q, size_t N) {
   T initial = T(N);
   T assignment = initial;
@@ -25,11 +25,11 @@ void assignment_test(queue q, size_t N) {
     q.submit([&](handler &cgh) {
       auto st =
           assignment_buf.template get_access<access::mode::read_write>(cgh);
-      cgh.parallel_for<assignment_kernel<AtomicRef, T>>(
+      cgh.parallel_for<assignment_kernel<AtomicRef, address_space, T>>(
           range<1>(N), [=](item<1> it) {
             size_t gid = it.get_id(0);
             auto atm = AtomicRef<T, memory_order::relaxed, memory_scope::device,
-                                 access::address_space::global_space>(st[0]);
+                                 address_space>(st[0]);
             atm = T(gid);
           });
     });
@@ -42,6 +42,13 @@ void assignment_test(queue q, size_t N) {
 }
 
 template <typename T> void assignment_test(queue q, size_t N) {
-  assignment_test<::sycl::ext::oneapi::atomic_ref, T>(q, N);
-  assignment_test<::sycl::atomic_ref, T>(q, N);
+  assignment_test<::sycl::ext::oneapi::atomic_ref,
+                  access::address_space::global_space, T>(q, N);
+  assignment_test<::sycl::atomic_ref, access::address_space::global_space, T>(
+      q, N);
+}
+
+template <typename T> void assignment_generic_test(queue q, size_t N) {
+  assignment_test<::sycl::atomic_ref, access::address_space::generic_space, T>(
+      q, N);
 }