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

Author: Alexander Batashev

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/diagnostics/non-uniform-wk-gp-test.cpp

@@ -13,12 +13,13 @@
 using namespace cl::sycl;
 
 int test() {
-  try {
-    queue q = queue();
-    auto device = q.get_device();
-    auto deviceName = device.get_info<cl::sycl::info::device::name>();
-    std::cout << " Device Name: " << deviceName << std::endl;
+  queue q = queue();
+  auto device = q.get_device();
+  auto deviceName = device.get_info<cl::sycl::info::device::name>();
+  std::cout << " Device Name: " << deviceName << std::endl;
 
+  int res = 1;
+  try {
     const int N = 1;
     q.submit([&](handler &cgh) {
       cl::sycl::stream kernelout(108 * 64 + 128, 64, cgh);
@@ -29,23 +30,23 @@ int test() {
                       << cl::sycl::endl;
           });
     });
-
+    std::cout << "Test failed: no exception thrown." << std::endl;
   } catch (sycl::runtime_error &E) {
     if (std::string(E.what()).find(
             "Non-uniform work-groups are not supported by the target device") !=
         std::string::npos) {
       std::cout << E.what() << std::endl;
       std::cout << "Test passed: caught the expected error." << std::endl;
-      return 0;
+      res = 0;
     } else {
       std::cout << E.what() << std::endl;
       std::cout << "Test failed: received error is incorrect." << std::endl;
-      return 1;
     }
   }
+  q.wait();
 
   std::cout << "Test passed: results are correct." << std::endl;
-  return 0;
+  return res;
 }
 
 int main() {

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/fpga_tests/fpga_latency_control_lsu.cpp

@@ -0,0 +1,63 @@
+// REQUIRES: aoc, accelerator
+// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==- fpga_latency_control_lsu.cpp - SYCL FPGA latency control on LSU test -==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+using namespace sycl;
+
+using PrefetchingLSU = ext::intel::experimental::lsu<
+    ext::intel::experimental::prefetch<true>,
+    ext::intel::experimental::statically_coalesce<false>>;
+
+using BurstCoalescedLSU = ext::intel::experimental::lsu<
+    ext::intel::experimental::burst_coalesce<true>,
+    ext::intel::experimental::statically_coalesce<false>>;
+
+int test_latency_control(queue Queue) {
+  std::vector<float> input_data = {1.23f};
+  std::vector<float> output_data = {.0f};
+
+  {
+    buffer input_buffer(input_data);
+    buffer output_buffer(output_data);
+
+    Queue.submit([&](handler &cgh) {
+      auto input_accessor = input_buffer.get_access<access::mode::read>(cgh);
+
+      auto output_accessor = output_buffer.get_access<access::mode::write>(cgh);
+
+      cgh.single_task<class kernel>([=] {
+        auto in_ptr = input_accessor.get_pointer();
+        auto out_ptr = output_accessor.get_pointer();
+
+        float value = PrefetchingLSU::load<
+            ext::intel::experimental::latency_anchor_id<0>>(in_ptr);
+
+        BurstCoalescedLSU::store<ext::intel::experimental::latency_constraint<
+            0, ext::intel::experimental::type::exact, 5>>(out_ptr, value);
+      });
+    });
+  }
+
+  if (output_data[0] != input_data[0]) {
+    std::cout << "Unexpected read from output_data: " << output_data[0]
+              << ", v.s. expected " << input_data[0] << std::endl;
+
+    return -1;
+  }
+  return 0;
+}
+
+int main() {
+  queue Queue{ext::intel::fpga_emulator_selector{}};
+
+  return test_latency_control(Queue);
+}

SYCL/Basic/fpga_tests/fpga_latency_control_pipe.cpp

@@ -0,0 +1,60 @@
+// REQUIRES: aoc, accelerator
+// RUN: %clangxx -fsycl -fintelfpga %s -o %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//== fpga_latency_control_pipe.cpp - SYCL FPGA latency control on pipe test ==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+using namespace sycl;
+
+using Pipe1 = ext::intel::experimental::pipe<class PipeClass1, int, 8>;
+using Pipe2 = ext::intel::experimental::pipe<class PipeClass2, int, 8>;
+
+int test_latency_control(queue Queue) {
+  std::vector<int> input_data = {1};
+  std::vector<int> output_data = {0};
+
+  {
+    buffer input_buffer(input_data);
+    buffer output_buffer(output_data);
+
+    Queue.submit([&](handler &cgh) {
+      auto input_accessor = input_buffer.get_access<access::mode::read>(cgh);
+
+      auto output_accessor = output_buffer.get_access<access::mode::write>(cgh);
+
+      cgh.single_task<class kernel>([=] {
+        Pipe1::write(input_accessor[0]);
+
+        int value =
+            Pipe1::read<ext::intel::experimental::latency_anchor_id<0>>();
+
+        Pipe2::write<ext::intel::experimental::latency_anchor_id<1>,
+                     ext::intel::experimental::latency_constraint<
+                         0, ext::intel::experimental::type::exact, 2>>(value);
+
+        output_accessor[0] = Pipe2::read();
+      });
+    });
+  }
+
+  if (output_data[0] != input_data[0]) {
+    std::cout << "Unexpected read from output_data: " << output_data[0]
+              << ", v.s. expected " << input_data[0] << std::endl;
+
+    return -1;
+  }
+  return 0;
+}
+
+int main() {
+  queue Queue{ext::intel::fpga_emulator_selector{}};
+
+  return test_latency_control(Queue);
+}

SYCL/Basic/intel-ext-device.cpp

@@ -40,6 +40,7 @@ int main(int argc, char **argv) {
       int numSlices = 0;
       int numSubslices = 0;
       int numEUsPerSubslice = 0;
+      int numHWThreadsPerEU = 0;
       for (const auto &dev : plt.get_devices()) {
         std::cout << "Platform #" << pltCount++ << ":" << std::endl;
         if (dev.has(aspect::gpu)) {
@@ -90,6 +91,12 @@ int main(int argc, char **argv) {
               std::cout << "Number of EUs per subslice = " << numEUsPerSubslice
                         << std::endl;
             }
+            if (dev.has(aspect::ext_intel_gpu_hw_threads_per_eu)) {
+              numHWThreadsPerEU =
+                  dev.get_info<info::device::ext_intel_gpu_hw_threads_per_eu>();
+              std::cout << "Number of HW threads per EU = " << numHWThreadsPerEU
+                        << std::endl;
+            }
             if (dev.has(aspect::ext_intel_max_mem_bandwidth)) {
               // not supported yet
               long m =

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/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/ESIMD/api/functional/ctors/ctor_copy.cpp

@@ -43,7 +43,7 @@ struct initializer {
 
 // Descriptor class for the case of calling constructor in variable declaration
 // context.
-struct var_declaration {
+struct var_decl {
   static std::string get_description() { return "variable declaration"; }
 
   template <typename DataT, int NumElems>
@@ -57,7 +57,7 @@ struct var_declaration {
 
 // Descriptor class for the case of calling constructor in rvalue in an
 // expression context.
-struct rval_in_expression {
+struct rval_in_expr {
   static std::string get_description() { return "rvalue in an expression"; }
 
   template <typename DataT, int NumElems>
@@ -92,42 +92,19 @@ class const_ref {
   }
 };
 
-template <typename DataT, typename TestT>
-using run_test_with_one_elem = test<DataT, 1, TestT>;
-
-template <typename DataT, typename TestT>
-using run_test_with_eight_elems = test<DataT, 8, TestT>;
-
-template <typename DataT, typename TestT>
-using run_test_with_sixteen_elems = test<DataT, 16, TestT>;
-
-template <typename DataT, typename TestT>
-using run_test_with_thirty_two_elems = test<DataT, 32, TestT>;
-
-template <typename TestT, typename... T>
-bool run_verification_for_type(sycl::queue &queue,
-                               const named_type_pack<T...> &types) {
-  bool passed{true};
-
-  passed &= for_all_types<run_test_with_one_elem, TestT>(types, queue);
-  passed &= for_all_types<run_test_with_eight_elems, TestT>(types, queue);
-  passed &= for_all_types<run_test_with_sixteen_elems, TestT>(types, queue);
-  passed &= for_all_types<run_test_with_thirty_two_elems, TestT>(types, queue);
-  return passed;
-}
-
 int main(int argc, char **argv) {
   sycl::queue queue{esimd_test::ESIMDSelector{},
                     esimd_test::createExceptionHandler()};
 
   bool passed{true};
 
-  auto types{get_tested_types<tested_types::all>()};
+  const auto types{get_tested_types<tested_types::all>()};
+  const auto dims{get_all_dimensions()};
 
-  passed &= run_verification_for_type<initializer>(queue, types);
-  passed &= run_verification_for_type<var_declaration>(queue, types);
-  passed &= run_verification_for_type<rval_in_expression>(queue, types);
-  passed &= run_verification_for_type<const_ref>(queue, types);
+  passed &= for_all_types_and_dims<test, initializer>(types, dims, queue);
+  passed &= for_all_types_and_dims<test, var_decl>(types, dims, queue);
+  passed &= for_all_types_and_dims<test, rval_in_expr>(types, dims, queue);
+  passed &= for_all_types_and_dims<test, const_ref>(types, dims, queue);
 
   std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
   return passed ? 0 : 1;

SYCL/ESIMD/api/functional/ctors/ctor_default.cpp

@@ -42,7 +42,7 @@ struct initializer {
 
 // Descriptor class for the case of calling constructor in variable declaration
 // context
-struct var_declaration {
+struct var_decl {
   static std::string get_description() { return "variable declaration"; }
 
   template <typename DataT, int NumElems>
@@ -54,7 +54,7 @@ struct var_declaration {
 
 // Descriptor class for the case of calling constructor in rvalue in an
 // expression context
-struct rval_in_expression {
+struct rval_in_expr {
   static std::string get_description() { return "rvalue in an expression"; }
 
   template <typename DataT, int NumElems>
@@ -116,44 +116,19 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
   }
 };
 
-template <typename DataT, typename TestT>
-using run_test_with_one_elem = test<DataT, 1, TestT>;
-
-template <typename DataT, typename TestT>
-using run_test_with_eight_elems = test<DataT, 8, TestT>;
-
-template <typename DataT, typename TestT>
-using run_test_with_sixteen_elems = test<DataT, 16, TestT>;
-
-template <typename DataT, typename TestT>
-using run_test_with_thirty_two_elems = test<DataT, 32, TestT>;
-
-template <typename TestT, typename... T>
-bool run_verification_with_chosen_test_type(
-    sycl::queue &queue, const named_type_pack<T...> &types) {
-  bool passed{true};
-
-  passed &= for_all_types<run_test_with_one_elem, TestT>(types, queue);
-  passed &= for_all_types<run_test_with_eight_elems, TestT>(types, queue);
-  passed &= for_all_types<run_test_with_sixteen_elems, TestT>(types, queue);
-  passed &= for_all_types<run_test_with_thirty_two_elems, TestT>(types, queue);
-  return passed;
-}
-
 int main(int argc, char **argv) {
   sycl::queue queue{esimd_test::ESIMDSelector{},
                     esimd_test::createExceptionHandler()};
 
   bool passed{true};
 
   const auto types{get_tested_types<tested_types::all>()};
+  const auto dims{get_all_dimensions()};
 
-  passed &= run_verification_with_chosen_test_type<initializer>(queue, types);
-  passed &=
-      run_verification_with_chosen_test_type<var_declaration>(queue, types);
-  passed &=
-      run_verification_with_chosen_test_type<rval_in_expression>(queue, types);
-  passed &= run_verification_with_chosen_test_type<const_ref>(queue, types);
+  passed &= for_all_types_and_dims<test, initializer>(types, dims, queue);
+  passed &= for_all_types_and_dims<test, var_decl>(types, dims, queue);
+  passed &= for_all_types_and_dims<test, rval_in_expr>(types, dims, queue);
+  passed &= for_all_types_and_dims<test, const_ref>(types, dims, queue);
 
   std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
   return passed ? 0 : 1;