Merge remote-tracking branch 'origin/intel' into CopyToFixTest

Author: fineg74

SYCL/Assert/assert_in_kernels_ndebug.cpp

@@ -1,5 +1,5 @@
-// FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
-// UNSUPPORTED: cuda || hip
+// FIXME unsupported on HIP until fallback libdevice becomes available
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DNDEBUG %S/assert_in_kernels.cpp -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER

SYCL/BFloat16/bfloat16_type_cuda.cpp

@@ -1,8 +1,6 @@
 // REQUIRES: gpu, cuda
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Xsycl-target-backend --cuda-gpu-arch=sm_80 %s -o %t.out
-// TODO: Currently the CI does not have a sm_80 capable machine. Enable the test
-// execution once it does.
-// RUNx: %t.out
+// RUN: %t.out
 
 //==--------- bfloat16_type_cuda.cpp - SYCL bfloat16 type test -------------==//
 //
@@ -14,4 +12,13 @@
 
 #include "bfloat16_type.hpp"
 
-int main() { return run_tests(); }
+int main() {
+  bool has_bfloat16_aspect = false;
+  for (const auto &plt : sycl::platform::get_platforms()) {
+    if (plt.has(aspect::ext_oneapi_bfloat16))
+      has_bfloat16_aspect = true;
+  }
+
+  if (has_bfloat16_aspect)
+    return run_tests();
+}

SYCL/Basic/accessor/accessor.cpp

@@ -556,6 +556,40 @@ int main() {
       return 1;
     }
   }
+
+  // placeholder accessor exception  // SYCL2020 4.7.6.9
+  {
+    sycl::queue q;
+    // host device executes kernels via a different method and there
+    // is no good way to throw an exception at this time.
+    if (!q.is_host()) {
+      sycl::range<1> r(4);
+      sycl::buffer<int, 1> b(r);
+      try {
+        sycl::accessor<int, 1, sycl::access::mode::read_write,
+                       sycl::access::target::device,
+                       sycl::access::placeholder::true_t>
+            acc(b);
+
+        q.submit([&](sycl::handler &cgh) {
+          // we do NOT call .require(acc) without which we should throw a
+          // synchronous exception with errc::kernel_argument
+          cgh.parallel_for<class ph>(
+              r, [=](sycl::id<1> index) { acc[index] = 0; });
+        });
+        q.wait_and_throw();
+        assert(false && "we should not be here, missing exception");
+      } catch (sycl::exception &e) {
+        std::cout << "exception received: " << e.what() << std::endl;
+        assert(e.code() == sycl::errc::kernel_argument &&
+               "incorrect error code");
+      } catch (...) {
+        std::cout << "some other exception" << std::endl;
+        return 1;
+      }
+    }
+  }
+
   {
     try {
       int data = -1;

SYCL/Basic/aspects.cpp

@@ -57,6 +57,9 @@ int main() {
     if (plt.has(aspect::fp64)) {
       std::cout << "  fp64" << std::endl;
     }
+    if (plt.has(aspect::ext_oneapi_bfloat16)) {
+      std::cout << " ext_oneapi_bfloat16" << std::endl;
+    }
     if (plt.has(aspect::int64_base_atomics)) {
       std::cout << "  base atomic operations" << std::endl;
     }

SYCL/DeviceLib/cmath_fp64_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include "math_utils.hpp"
 #include <CL/sycl.hpp>
 #include <cmath>

SYCL/DeviceLib/cmath_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include "math_utils.hpp"
 #include <CL/sycl.hpp>
 #include <cmath>

SYCL/DeviceLib/math_fp64_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include "math_utils.hpp"
 #include <CL/sycl.hpp>
 #include <cstdint>

SYCL/DeviceLib/math_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include "math_utils.hpp"
 #include <CL/sycl.hpp>
 #include <cstdint>

SYCL/DeviceLib/std_complex_math_fp64_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include <CL/sycl.hpp>
 #include <array>
 #include <cassert>

SYCL/DeviceLib/std_complex_math_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include <CL/sycl.hpp>
 #include <array>
 #include <cassert>

SYCL/DeviceLib/string_test.cpp

@@ -3,6 +3,11 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
 #include <CL/sycl.hpp>
 #include <cassert>
 #include <cstdint>

SYCL/ESIMD/histogram_raw_send.cpp

@@ -6,7 +6,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu
+// REQUIRES: gpu-intel-gen9
 // UNSUPPORTED: gpu-intel-dg1,cuda,hip
 // UNSUPPORTED: ze_debug-1,ze_debug4
 // RUN: %clangxx -fsycl %s -o %t.out

SYCL/ESIMD/vadd_raw_send.cpp

@@ -6,7 +6,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu
+// REQUIRES: gpu-intel-gen9
 // UNSUPPORTED: gpu-intel-dg1,cuda,hip
 // TODO: esimd_emulator fails due to unimplemented 'raw_send' intrinsic
 // XFAIL: esimd_emulator

SYCL/GroupAlgorithm/exclusive_scan_over_group.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: ze_debug4,ze_debug-1

SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/GroupAlgorithm/reduce_sycl2020.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -I . -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/Plugin/interop-cuda-experimental.cpp

@@ -95,4 +95,62 @@ int main() {
     sycl::queue new_Q(sycl_ctx, sycl::default_selector());
     assert(check_queue(new_Q));
   }
+
+  // Create new event
+  CUevent cu_event;
+
+  CUDA_CHECK(cuCtxSetCurrent(cu_ctx));
+  CUDA_CHECK(cuEventCreate(&cu_event, CU_EVENT_DEFAULT));
+
+  auto sycl_event =
+      sycl::make_event<sycl::backend::ext_oneapi_cuda>(cu_event, sycl_ctx);
+  auto native_event =
+      sycl::get_native<sycl::backend::ext_oneapi_cuda>(sycl_event);
+
+  check_type<sycl::event>(sycl_event);
+  check_type<CUevent>(native_event);
+
+  // Check sycl queue with sycl_ctx still works
+  {
+    sycl::queue new_Q(sycl_ctx, sycl::default_selector());
+    assert(check_queue(new_Q));
+  }
+
+  // Check has_native_event
+  {
+    auto e = Q.submit([&](sycl::handler &cgh) { cgh.single_task([] {}); });
+    assert(sycl::ext::oneapi::cuda::has_native_event(e));
+  }
+
+  {
+    auto e = Q.submit([&](sycl::handler &cgh) { cgh.host_task([] {}); });
+    assert(!sycl::ext::oneapi::cuda::has_native_event(e));
+  }
+
+  // Create new queue
+  CUstream cu_queue;
+  CUDA_CHECK(cuCtxSetCurrent(cu_ctx));
+  CUDA_CHECK(cuStreamCreate(&cu_queue, CU_STREAM_DEFAULT));
+
+  auto sycl_queue =
+      sycl::make_queue<sycl::backend::ext_oneapi_cuda>(cu_queue, sycl_ctx);
+  native_queue = sycl::get_native<sycl::backend::ext_oneapi_cuda>(sycl_queue);
+
+  check_type<sycl::queue>(sycl_queue);
+  check_type<CUstream>(native_queue);
+
+  // Submit some work to new queue
+  assert(check_queue(sycl_queue));
+
+  // Create new queue with Q's native type and submit some work
+  {
+    CUstream Q_native_stream =
+        sycl::get_native<sycl::backend::ext_oneapi_cuda>(Q);
+    sycl::queue new_Q = sycl::make_queue<sycl::backend::ext_oneapi_cuda>(
+        Q_native_stream, Q_sycl_ctx);
+    assert(check_queue(new_Q));
+  }
+
+  // Check Q still works
+  assert(check_queue(Q));
 }

SYCL/README.md

@@ -163,6 +163,9 @@ unavailable.
  * **aot_tool** - Ahead-of-time compilation tools availability;
  * **ocloc**, **opencl-aot** - Specific AOT tool availability;
  * **level_zero_dev_kit** - Level_Zero headers and libraries availability;
+ * **gpu-intel-gen9**  - Intel GPU Gen9  availability;
+ * **gpu-intel-gen11** - Intel GPU Gen11 availability;
+ * **gpu-intel-gen12** - Intel GPU Gen12 availability;
  * **gpu-intel-dg1** - Intel GPU DG1 availability;
  * **gpu-intel-dg2** - Intel GPU DG2 availability;
  * **gpu-intel-pvc** - Intel GPU PVC availability;

SYCL/Sampler/unnormalized-clamp-nearest.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: hip || cuda
+// 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/Sampler/unnormalized-clampedge-nearest.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: hip || cuda
+// 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/Sampler/unnormalized-none-nearest.cpp

@@ -1,4 +1,4 @@
-// UNSUPPORTED: hip || cuda
+// 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/SpecConstants/2020/handler-api.cpp

@@ -13,8 +13,7 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // FIXME: ACC devices use emulation path, which is not yet supported
-// FIXME: CUDA uses emulation path, which is not yet supported
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 
 #include <cstdlib>
 #include <iostream>

SYCL/SpecConstants/2020/host_apis.cpp

@@ -2,7 +2,6 @@
 // RUN:          -fsycl-dead-args-optimization
 // RUN: %BE_RUN_PLACEHOLDER %t.out
 
-// UNSUPPORTED: cuda
 // UNSUPPORTED: hip
 
 #include <sycl/sycl.hpp>

SYCL/SpecConstants/2020/kernel-bundle-api.cpp

@@ -13,8 +13,7 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // FIXME: ACC devices use emulation path, which is not yet supported
-// FIXME: CUDA uses emulation path, which is not yet supported
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 
 #include <cstdlib>
 #include <iostream>

SYCL/SpecConstants/2020/non_native/cuda.cpp

@@ -3,8 +3,5 @@
 // RUN: %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda %S/Inputs/common.cpp -o %t.out
 // RUN: env SYCL_DEVICE_FILTER=cuda %t.out
 
-// TODO: enable this test then compile-time error in sycl-post-link is fixed
-// UNSUPPORTED: cuda
-
 // This test checks correctness of SYCL2020 non-native specialization constants
 // on CUDA device

SYCL/SpecConstants/2020/vector-convolution-demo.cpp

@@ -2,7 +2,7 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
-// UNSUPPORTED: cuda || hip
+// UNSUPPORTED: hip
 
 // This test checks the spenario of using specialization constants with an
 // 'array of array' as well as a 'stuct with an array of array' types for

SYCL/lit.cfg.py

@@ -100,12 +100,19 @@
     config.substitutions.append( ('%obj_ext', '.o') )
 config.substitutions.append( ('%sycl_include',  config.sycl_include ) )
 
+# Intel GPU FAMILY availability
+if lit_config.params.get('gpu-intel-gen9', False):
+    config.available_features.add('gpu-intel-gen9')
+if lit_config.params.get('gpu-intel-gen11', False):
+    config.available_features.add('gpu-intel-gen11')
+if lit_config.params.get('gpu-intel-gen12', False):
+    config.available_features.add('gpu-intel-gen12')
+
+# Intel GPU DEVICE availability
 if lit_config.params.get('gpu-intel-dg1', False):
     config.available_features.add('gpu-intel-dg1')
-
 if lit_config.params.get('gpu-intel-dg2', False):
     config.available_features.add('gpu-intel-dg2')
-
 if lit_config.params.get('gpu-intel-pvc', False):
     config.available_features.add('gpu-intel-pvc')