Merge pull request intel#42 from bb-sycl/xmain

yuxianch · web-flow · commit ebadda3d0636 · 2021-06-25T14:27:11.000+08:00
Auto pulldown and update tc files for xmain branch on 20210625
diff --git a/SYCL/Basic/half_type.cpp b/SYCL/Basic/half_type.cpp
@@ -192,7 +192,52 @@ inline bool bitwise_comparison_fp32(const half val, const uint32_t exp) {
   return reinterpret_cast<const uint32_t &>(fp32) == exp;
 }
 
+constexpr void constexpr_verify_add() {
+  constexpr half a{5.0}, b{2.0}, ref{7.0};
+  constexpr half result = a + b;
+  constexpr half diff = result - ref;
+  constexpr auto sign = diff < 0 ? -1 : 1;
+  static_assert(sign * static_cast<float>(diff) <
+                    std::numeric_limits<cl::sycl::half>::epsilon(),
+                "Constexpr add is wrong");
+}
+
+constexpr void constexpr_verify_sub() {
+  constexpr half a{5.0f}, b{2.0}, ref{3.0};
+  constexpr half result = a - b;
+  constexpr half diff = result - ref;
+  constexpr auto sign = diff < 0 ? -1 : 1;
+  static_assert(sign * static_cast<float>(diff) <
+                    std::numeric_limits<cl::sycl::half>::epsilon(),
+                "Constexpr sub is wrong");
+}
+
+constexpr void constexpr_verify_mul() {
+  constexpr half a{5.0f}, b{2.0}, ref{10.0};
+  constexpr half result = a * b;
+  constexpr half diff = result - ref;
+  constexpr auto sign = diff < 0 ? -1 : 1;
+  static_assert(sign * static_cast<float>(diff) <
+                    std::numeric_limits<cl::sycl::half>::epsilon(),
+                "Constexpr mul is wrong");
+}
+
+constexpr void constexpr_verify_div() {
+  constexpr half a{5.0f}, b{2.0}, ref{2.5};
+  constexpr half result = a / b;
+  constexpr half diff = result - ref;
+  constexpr auto sign = diff < 0 ? -1 : 1;
+  static_assert(sign * static_cast<float>(diff) <
+                    std::numeric_limits<cl::sycl::half>::epsilon(),
+                "Constexpr div is wrong");
+}
+
 int main() {
+  constexpr_verify_add();
+  constexpr_verify_sub();
+  constexpr_verify_mul();
+  constexpr_verify_div();
+
   device dev{default_selector()};
   if (!dev.is_host() && !dev.has_extension("cl_khr_fp16")) {
     std::cout << "This device doesn't support the extension cl_khr_fp16"
diff --git a/SYCL/ESIMD/dp4a.cpp b/SYCL/ESIMD/dp4a.cpp
@@ -6,11 +6,9 @@
 //
 //===----------------------------------------------------------------------===//
 // TODO enable on Windows
-// REQUIRES: linux && gpu
+// REQUIRES: linux && gpu-intel-dg1
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
-// TODO : Enable test for new GPU device
-// XFAIL: *
 
 #include "esimd_test_utils.hpp"
 
diff --git a/SYCL/ESIMD/regression/noinline_bypointers_vadd.cpp b/SYCL/ESIMD/regression/noinline_bypointers_vadd.cpp
@@ -0,0 +1,112 @@
+//==------ noinline_bypointers_vadd.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// UNSUPPORTED: cuda
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// XFAIL: *
+
+// Vector BE bug.
+// Memory access via a pointer type argument is broken if the function was not
+// inlined.
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+#include <iostream>
+
+using namespace cl::sycl;
+using namespace sycl::ext::intel::experimental::esimd;
+
+using ptr = float *;
+static inline constexpr unsigned VL = 32;
+
+SYCL_EXTERNAL ESIMD_NOINLINE void do_add(ptr A, float *B,
+                                         ptr C) SYCL_ESIMD_FUNCTION {
+  simd<float, VL> va;
+  va.copy_from(A);
+  simd<float, VL> vb;
+  vb.copy_from(B);
+  simd<float, VL> vc = va + vb;
+  vc.copy_to(C);
+}
+
+int main(void) {
+  constexpr unsigned Size = 1024;
+  constexpr unsigned GroupSize = 8;
+
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+
+  auto dev = q.get_device();
+  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+  auto ctxt = q.get_context();
+  float *A =
+      static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
+  float *B =
+      static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
+  float *C =
+      static_cast<float *>(malloc_shared(Size * sizeof(float), dev, ctxt));
+
+  for (unsigned i = 0; i < Size; ++i) {
+    A[i] = B[i] = i;
+  }
+
+  // We need that many workitems. Each processes VL elements of data.
+  cl::sycl::range<1> GlobalRange{Size / VL};
+  // Number of workitems in each workgroup.
+  cl::sycl::range<1> LocalRange{GroupSize};
+
+  cl::sycl::nd_range<1> Range(GlobalRange, LocalRange);
+
+  try {
+    auto e = q.submit([&](handler &cgh) {
+      cgh.parallel_for<class Test>(
+          Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+            int i = ndi.get_global_id(0);
+            do_add(ptr{A + i * VL}, B + i * VL, ptr{C + i * VL});
+          });
+    });
+    e.wait();
+  } catch (cl::sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+
+    free(A, ctxt);
+    free(B, ctxt);
+    free(C, ctxt);
+
+    return e.get_cl_code();
+  }
+
+  int err_cnt = 0;
+
+  for (unsigned i = 0; i < Size; ++i) {
+    if (A[i] + B[i] != C[i]) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << C[i] << " != " << A[i]
+                  << " + " << B[i] << "\n";
+      }
+    }
+  }
+
+  free(A, ctxt);
+  free(B, ctxt);
+  free(C, ctxt);
+
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
+              << (Size - err_cnt) << "/" << Size << ")\n";
+    std::cout << "FAILED\n";
+    return 1;
+  }
+
+  std::cout << "Passed\n";
+  return 0;
+}
diff --git a/SYCL/HierPar/hier_par_wgscope_O0.cpp b/SYCL/HierPar/hier_par_wgscope_O0.cpp
@@ -17,9 +17,6 @@
 // code and data in the work group scope, and when the test is compiled with
 // -O0 switch.
 
-// The test is failing on CUDA after intel/llvm#3779
-// XFAIL: cuda
-
 #include "Inputs/hier_par_wgscope_impl.hpp"
 
 int main() { return run(); }
diff --git a/SYCL/Regression/2020-spec-constants-debug-info.cpp b/SYCL/Regression/2020-spec-constants-debug-info.cpp
@@ -0,0 +1,29 @@
+// RUN: %clangxx -fsycl -g %s -o %t.out
+// RUN: %clangxx -fsycl -g -O0 %s -o %t.out
+// RUN: %clangxx -fsycl -g -O2 %s -o %t.out
+//
+// The idea of this test is to make sure that we can compile the following
+// simple example without crashes/assertions firing at llvm-spirv step due to
+// debug info corrupted by sycl-post-link
+
+#include <sycl/sycl.hpp>
+
+constexpr sycl::specialization_id<int> test_id_1{42};
+
+int main() {
+
+  sycl::queue q;
+  {
+    sycl::buffer<double, 1> Buf{sycl::range{1}};
+    q.submit([&](sycl::handler &cgh) {
+      auto Acc = Buf.get_access<sycl::access::mode::read_write>(cgh);
+      cgh.set_specialization_constant<test_id_1>(1);
+      cgh.single_task<class Kernel1>([=](sycl::kernel_handler kh) {
+        Acc[0] = kh.get_specialization_constant<test_id_1>();
+      });
+    });
+    auto Acc = Buf.get_access<sycl::access::mode::read>();
+    assert(Acc[0] == 1);
+  }
+  return 0;
+}
diff --git a/SYCL/Sampler/normalized-clampedge-linear-float.cpp b/SYCL/Sampler/normalized-clampedge-linear-float.cpp
@@ -3,13 +3,10 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // XFAIL: cuda
-// UNSUPPORTED: level_zero && windows
 
 // CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
 // type (such as unorm_int8)
 
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
-
 /*
     This file sets up an image, initializes it with data,
     and verifies that the data is sampled correctly with a
diff --git a/SYCL/Sampler/normalized-clampedge-nearest.cpp b/SYCL/Sampler/normalized-clampedge-nearest.cpp
@@ -2,9 +2,7 @@
 // RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
-// UNSUPPORTED: level_zero && windows
 
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/Sampler/normalized-mirror-linear-float.cpp b/SYCL/Sampler/normalized-mirror-linear-float.cpp
@@ -4,12 +4,10 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 
 // XFAIL: cuda
-// UNSUPPORTED: level_zero && windows
 
 // CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
 // type (such as unorm_int8)
 
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/Sampler/normalized-mirror-nearest.cpp b/SYCL/Sampler/normalized-mirror-nearest.cpp
@@ -3,11 +3,9 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
 // XFAIL: cuda
-// UNSUPPORTED: level_zero && windows
 
 // CUDA is not handling repeat or mirror correctly with normalized coordinates.
 // Waiting on a fix.
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/Sampler/normalized-none-linear-float.cpp b/SYCL/Sampler/normalized-none-linear-float.cpp
@@ -4,12 +4,10 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 
 // XFAIL: cuda
-// UNSUPPORTED: level_zero && windows
 
 // CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
 // type (such as unorm_int8)
 
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/Sampler/normalized-none-nearest.cpp b/SYCL/Sampler/normalized-none-nearest.cpp
@@ -2,9 +2,7 @@
 // RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
-// UNSUPPORTED: level_zero && windows
 
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/Sampler/normalized-repeat-linear-float.cpp b/SYCL/Sampler/normalized-repeat-linear-float.cpp
@@ -4,12 +4,10 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 
 // XFAIL: cuda
-// UNSUPPORTED: level_zero && windows
 
 // CUDA works with image_channel_type::fp32, but not with any 8-bit per channel
 // type (such as unorm_int8)
 
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/Sampler/normalized-repeat-nearest.cpp b/SYCL/Sampler/normalized-repeat-nearest.cpp
@@ -3,11 +3,9 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
 // XFAIL: cuda
-// UNSUPPORTED: level_zero && windows
 
 // CUDA is not handling repeat or mirror correctly with normalized coordinates.
 // Waiting on a fix.
-// LevelZero on Windows hangs with normalized coordinates. Waiting on fix.
 
 /*
     This file sets up an image, initializes it with data,
diff --git a/SYCL/SpecConstants/2020/common.hpp b/SYCL/SpecConstants/2020/common.hpp
@@ -29,8 +29,7 @@ inline bool operator!=(const custom_type_nested &lhs,
 
 inline std::ostream &operator<<(std::ostream &out,
                                 const custom_type_nested &v) {
-  return out << "custom_type_nested { .c = " << v.c << ", .f = " << v.f << "}"
-             << std::endl;
+  return out << "custom_type_nested { .c = " << v.c << ", .f = " << v.f << "}";
 }
 
 struct custom_type {
@@ -54,14 +53,15 @@ inline bool operator!=(const custom_type &lhs, const custom_type &rhs) {
 
 inline std::ostream &operator<<(std::ostream &out, const custom_type &v) {
   return out << "custom_type { .n = \n\t" << v.n << ",\n .ull = " << v.ull
-             << "}" << std::endl;
+             << "}";
 }
 
 template <typename T>
-bool check_value(const T &got, const T &ref, const std::string &variable_name) {
+bool check_value(const T &ref, const T &got, const std::string &variable_name) {
   if (got != ref) {
     std::cout << "Unexpected value of " << variable_name << ": " << got
               << " (got) vs " << ref << " (expected)" << std::endl;
+    return false;
   }
 
   return true;
diff --git a/SYCL/SpecConstants/2020/kernel-bundle-api.cpp b/SYCL/SpecConstants/2020/kernel-bundle-api.cpp
diff --git a/config_sycl/esimd_regression_noinline_bypointers_vadd.info b/config_sycl/esimd_regression_noinline_bypointers_vadd.info
diff --git a/config_sycl/regression_2020_spec_constants_debug_info.info b/config_sycl/regression_2020_spec_constants_debug_info.info
diff --git a/config_sycl/specconstants_2020_kernel_bundle_api.info b/config_sycl/specconstants_2020_kernel_bundle_api.info
diff --git a/llvm_test_suite_sycl.xml b/llvm_test_suite_sycl.xml
diff --git a/llvm_test_suite_sycl_valgrind.xml b/llvm_test_suite_sycl_valgrind.xml
