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

Author: rdeodhar

SYCL/Basic/device_event.cpp

@@ -5,8 +5,9 @@
 // TODO: nd_item::barrier() is not implemented on HOST
 // RUNx: %HOST_RUN_PLACEHOLDER %t.run
 //
-// Crashes on AMD
-// XFAIL: hip_amd
+// Crashes on AMD, returns error "Barrier is not supported on the host device
+// yet." with Nvidia.
+// XFAIL: hip_amd || hip_nvidia
 
 //==--------device_event.cpp - SYCL class device_event test ----------------==//
 //

SYCL/Basic/diagnostics/handler.cpp

@@ -1,5 +1,9 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %BE_RUN_PLACEHOLDER %t.out | FileCheck %s
+//
+// Appears to fail on HIP Nvidia because 'no device of requested type available'
+// when constructing a queue with an exception_list.
+// XFAIL: hip_nvidia
 //==------------------- handler.cpp ----------------------------------------==//
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.

SYCL/Basic/group_async_copy.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.run
+// RUN: %clangxx -fsycl -std=c++17 -fsycl-targets=%sycl_triple %s -o %t.run
 // RUN: %GPU_RUN_PLACEHOLDER %t.run
 // RUN: %CPU_RUN_PLACEHOLDER %t.run
 // RUN: %ACC_RUN_PLACEHOLDER %t.run
@@ -13,7 +13,11 @@
 
 using namespace cl::sycl;
 
-template <typename T> class KernelName;
+template <typename T> class TypeHelper;
+
+template <typename T>
+using KernelName = class TypeHelper<typename std::conditional<
+    std::is_same<T, std::byte>::value, unsigned char, T>::type>;
 
 // Define the number of work items to enqueue.
 const size_t NElems = 32;
@@ -24,14 +28,14 @@ template <typename T> void initInputBuffer(buffer<T, 1> &Buf, size_t Stride) {
   auto Acc = Buf.template get_access<access::mode::write>();
   for (size_t I = 0; I < Buf.get_count(); I += WorkGroupSize) {
     for (size_t J = 0; J < WorkGroupSize; J++)
-      Acc[I + J] = I + J + ((J % Stride == 0) ? 100 : 0);
+      Acc[I + J] = static_cast<T>(I + J + ((J % Stride == 0) ? 100 : 0));
   }
 }
 
 template <typename T> void initOutputBuffer(buffer<T, 1> &Buf) {
   auto Acc = Buf.template get_access<access::mode::write>();
   for (size_t I = 0; I < Buf.get_count(); I++)
-    Acc[I] = 0;
+    Acc[I] = static_cast<T>(0);
 }
 
 template <typename T> struct is_vec : std::false_type {};
@@ -48,9 +52,8 @@ template <typename T> bool checkEqual(vec<T, 4> A, size_t B) {
 }
 
 template <typename T>
-typename std::enable_if<!is_vec<T>::value, bool>::type checkEqual(T A,
-                                                                  size_t B) {
-  T TB = B;
+typename std::enable_if_t<!is_vec<T>::value, bool> checkEqual(T A, size_t B) {
+  T TB = static_cast<T>(B);
   return A == TB;
 }
 
@@ -67,7 +70,16 @@ template <typename T> std::string toString(vec<T, 4> A) {
 }
 
 template <typename T = void>
-typename std::enable_if<!is_vec<T>::value, std::string>::type toString(T A) {
+typename std::enable_if_t<
+    !is_vec<T>::value && std::is_same<T, std::byte>::value, std::string>
+toString(T A) {
+  return std::to_string((unsigned char)A);
+}
+
+template <typename T = void>
+typename std::enable_if_t<
+    !is_vec<T>::value && !std::is_same<T, std::byte>::value, std::string>
+toString(T A) {
   return std::to_string(A);
 }
 
@@ -156,6 +168,8 @@ int main() {
       return 1;
     if (test<cl::sycl::cl_bool>(Stride))
       return 1;
+    if (test<std::byte>(Stride))
+      return 1;
   }
 
   std::cout << "Test passed.\n";

SYCL/Basic/handler/interop_task.cpp

@@ -11,7 +11,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "CL/sycl/access/access.hpp"
 #include <CL/sycl.hpp>
 #include <CL/sycl/backend/opencl.hpp>
 

SYCL/Basic/handler/run_on_host_intel.cpp

@@ -10,7 +10,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "CL/sycl/access/access.hpp"
 #include <CL/sycl.hpp>
 
 #include "../../helpers.hpp"

SYCL/Basic/host-task-dependency.cpp

@@ -5,7 +5,7 @@
 //
 // TODO: Behaviour is unstable for level zero on Windows. Enable when fixed.
 // TODO: The test is sporadically fails on CUDA. Enable when fixed.
-// UNSUPPORTED: (windows && level_zero) || cuda
+// UNSUPPORTED: (windows && level_zero) || cuda || hip_nvidia
 
 #define SYCL2020_DISABLE_DEPRECATION_WARNINGS
 

SYCL/Basic/image/image_max_size.cpp

@@ -3,8 +3,9 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: cuda || hip || (windows && opencl && gpu)
 // CUDA does not support info::device::image3d_max_width query.
+// TODO: Irregular runtime fails on Windows/opencl:gpu require analysis.
 
 // The test checks that 'image' with max allowed sizes is handled correctly.
 

SYCL/Basic/kernel_bundle/kernel_bundle_api.cpp

@@ -330,8 +330,7 @@ int main() {
           sycl::link({KernelBundleObject1, KernelBundleObject2},
                      std::vector<sycl::device>{});
         },
-        "Not all devices are in the set of associated devices for input "
-        "bundles or vector of devices is empty");
+        "Vector of devices is empty");
 
     std::cerr << "Mismatched contexts for join" << std::endl;
     checkException(

SYCL/Basic/kernel_info.cpp

@@ -6,8 +6,8 @@
 // Fail is flaky for level_zero, enable when fixed.
 // UNSUPPORTED: level_zero
 //
-// Failing on HIP AMD
-// XFAIL: hip_amd
+// Failing on HIP AMD and HIP Nvidia
+// XFAIL: hip_amd || hip_nvidia
 
 //==--- kernel_info.cpp - SYCL kernel info test ----------------------------==//
 //
@@ -32,7 +32,6 @@ int main() {
   kernel krn = KB.get_kernel(KernelID);
 
   q.submit([&](handler &cgh) {
-    cgh.use_kernel_bundle(KB);
     auto acc = buf.get_access<access::mode::read_write>(cgh);
     cgh.single_task<class SingleTask>(krn, [=]() { acc[0] = acc[0] + 1; });
   });

SYCL/Basic/memory-consumption.cpp

@@ -1,6 +1,7 @@
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
+// XFAIL: hip_nvidia
 // Issue #106: The test failing sporadically on OpenCL platform due to
 // processing OCL_ICD_FILENAMES debug environment variable which causes
 // extra memory allocation on device creation.

SYCL/Basic/partition_supported.cpp

@@ -2,7 +2,9 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
-
+//
+// Nvidia should not allow sub_devices but does not throw corresponding error.
+// XFAIL: hip_nvidia
 /* Check that:
 1) [info::device::partition_properties]: returns the partition properties
 supported by this SYCL device; a vector of info::partition_property. If this

SYCL/Basic/queue/queue.cpp

@@ -1,5 +1,7 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %BE_RUN_PLACEHOLDER %t.out
+//
+// XFAIL: hip_nvidia
 //==--------------- queue.cpp - SYCL queue test ----------------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

SYCL/Basic/queue/release.cpp

@@ -2,6 +2,8 @@
 // RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
 // RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+//
+// XFAIL: hip_nvidia
 
 #include <CL/sycl.hpp>
 int main() {

SYCL/Basic/stream/auto_flush.cpp

@@ -3,6 +3,8 @@
 // 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
+//
+// XFAIL: hip_nvidia
 //==-------------- copy.cpp - SYCL stream obect auto flushing test ---------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

SYCL/Config/kernel_from_file.cpp

@@ -3,8 +3,9 @@
 
 // FIXME Disabled fallback assert as it'll require either online linking or
 // explicit offline linking step here
-// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT %cxx_std_optionc++17 -fsycl-device-only -fno-sycl-use-bitcode -Xclang -fsycl-int-header=%t.h -c %s -o %t.spv -Xclang -verify-ignore-unexpected=note,warning -Wno-sycl-strict
-// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT %cxx_std_optionc++17 %include_option %t.h %s -o %t.out %sycl_options -Xclang -verify-ignore-unexpected=note,warning
+// FIXME separate compilation requires -fno-sycl-dead-args-optimization
+// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT %cxx_std_optionc++17 -fsycl-device-only -fno-sycl-use-bitcode -fno-sycl-dead-args-optimization -Xclang -fsycl-int-header=%t.h -c %s -o %t.spv -Xclang -verify-ignore-unexpected=note,warning -Wno-sycl-strict
+// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT %cxx_std_optionc++17 %include_option %t.h %s -o %t.out %sycl_options -fno-sycl-dead-args-optimization -Xclang -verify-ignore-unexpected=note,warning
 // RUN: %BE_RUN_PLACEHOLDER env SYCL_USE_KERNEL_SPV=%t.spv %t.out | FileCheck %s
 // CHECK: Passed