[SYCL] Throw an exception when usm_shared_allocations aspect not supported

sycl/source/detail/usm/usm_impl.cpp

@@ -146,11 +146,8 @@ void *alignedAllocInternal(size_t Alignment, size_t Size,
  }
  if (Kind == alloc::shared &&
  !DevImpl->has(sycl::aspect::usm_shared_allocations)) {
- // TODO:: Throw an exception to conform with the specification.
- // Note that many tests will have to be changed to conform with the spec
- // before completing this. That is, the tests will now have to expect
- // exceptions as a result of failed allocations in addition to nullptr
- // being returned depending on the reason why allocation failed.
+ throw sycl::exception(sycl::errc::feature_not_supported,
+ "Device does not support USM shared allocations!");
  }
  void *RetVal = nullptr;
  if (Size == 0)

sycl/test-e2e/Annotated_arg_ptr/annotated_arg.cpp

@@ -75,6 +75,9 @@ MyStruct<T> operator<<(const MyStruct<T> &lhs, const MyStruct<T> &rhs) {
 
 int main() {
  queue Q;
+ if (!Q.get_device().has(aspect::usm_shared_allocations)) {
+ return 0;
+ }
 
  auto *a = malloc_shared<int>(8, Q);
  auto a_ptr = annotated_arg{a};

sycl/test-e2e/Annotated_arg_ptr/annotated_ptr.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 //

sycl/test-e2e/Annotated_usm/annotated_usm_kind.cpp

@@ -125,25 +125,27 @@ template <typename T> void testUsmKind(sycl::queue &q) {
  [&]() { return ATHost(1, q); }, [&]() { return ATHost(1, Ctx); },
  [&]() { return ATAnnotated(1, dev, Ctx, alloc::host); }});
 
- CheckUsmKindAll(
- alloc::shared,
- std::tuple{
- [&]() { return MShared(q); }, [&]() { return MShared(dev, Ctx); },
- [&]() { return MAnnotated(dev, Ctx, alloc::shared); },
- [&]() { return MAnnotated(dev, Ctx, properties{usm_kind_shared}); },
- [&]() { return AShared(1, q); },
- [&]() { return AShared(1, dev, Ctx); },
- [&]() { return AAnnotated(1, dev, Ctx, alloc::shared); },
- [&]() { return TShared(q); }, [&]() { return TShared(dev, Ctx); },
- [&]() { return TAnnotated(dev, Ctx, alloc::shared); },
- [&]() { return TAnnotated(dev, Ctx, properties{usm_kind_shared}); },
- [&]() { return ATShared(1, q); },
- [&]() { return ATShared(1, dev, Ctx); },
- [&]() { return ATAnnotated(1, dev, Ctx, alloc::shared); }});
+ if (q.get_device().has(sycl::aspect::usm_shared_allocations)) {
+ CheckUsmKindAll(
+ alloc::shared,
+ std::tuple{
+ [&]() { return MShared(q); }, [&]() { return MShared(dev, Ctx); },
+ [&]() { return MAnnotated(dev, Ctx, alloc::shared); },
+ [&]() { return MAnnotated(dev, Ctx, properties{usm_kind_shared}); },
+ [&]() { return AShared(1, q); },
+ [&]() { return AShared(1, dev, Ctx); },
+ [&]() { return AAnnotated(1, dev, Ctx, alloc::shared); },
+ [&]() { return TShared(q); }, [&]() { return TShared(dev, Ctx); },
+ [&]() { return TAnnotated(dev, Ctx, alloc::shared); },
+ [&]() { return TAnnotated(dev, Ctx, properties{usm_kind_shared}); },
+ [&]() { return ATShared(1, q); },
+ [&]() { return ATShared(1, dev, Ctx); },
+ [&]() { return ATAnnotated(1, dev, Ctx, alloc::shared); }});
+ }
 }
 
 int main() {
  sycl::queue q;
  testUsmKind<int>(q);
  return 0;
-}
+}

sycl/test-e2e/Basic/group_local_memory.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Basic/large-range.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocation
 // Temporarily add explicit '-O2' to avoid GPU hang issue with O0 optimization.
 // RUN: %{build} -fno-sycl-id-queries-fit-in-int -O2 -o %t.out
 // RUN: env SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE=1 %{run} %t.out

sycl/test-e2e/Basic/span.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 //

sycl/test-e2e/Basic/wrapped_usm_pointers.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Complex/sycl_complex_math_test.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%}
 
 // RUN: %{build} -fsycl-device-code-split=per_kernel %{mathflags} -o %t.out

sycl/test-e2e/Complex/sycl_complex_operator_test.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Complex/sycl_complex_pow_test.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%}
 
 // RUN: %{build} -fsycl-device-code-split=per_kernel %{mathflags} -o %t.out

sycl/test-e2e/Complex/sycl_complex_stream_test.cpp

@@ -7,6 +7,8 @@
 
 template <typename T> struct test_sycl_stream_operator {
  bool operator()(sycl::queue &Q, cmplx<T> init) {
+ if (!Q.get_device().has(sycl::aspect::usm_shared_allocations))
+ return true;
  auto *cplx_out = sycl::malloc_shared<experimental::complex<T>>(1, Q);
  cplx_out[0] = experimental::complex<T>(init.re, init.im);
 

sycl/test-e2e/DiscardEvents/discard_events_mixed_calls.cpp

@@ -37,6 +37,8 @@ static constexpr int MAX_ITER_NUM2 = 10;
 void TestHelper(sycl::queue Q,
  const std::function<void(sycl::range<1> Range, int *Harray,
  sycl::buffer<int, 1> Buf)> &Function) {
+ if (!Q.get_device().has(aspect::usm_shared_allocations))
+ return;
  int *Harray = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
  assert(Harray != nullptr);
  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
@@ -113,36 +115,36 @@ void RunTest_USM_Accessor(sycl::queue Q) {
 }
 
 void RunTest_Accessor_USM(sycl::queue Q) {
- TestHelper(
- Q, [&](sycl::range<1> Range, int *Harray, sycl::buffer<int, 1> Buf) {
- {
- sycl::host_accessor HostAcc(Buf);
- for (size_t i = 0; i < BUFFER_SIZE; ++i) {
- HostAcc[i] = 0;
- }
- }
-
- for (int i = 0; i < MAX_ITER_NUM1; ++i)
- IfTrueIncrementBufferAndUSM(Q, Range, Harray, Buf, (i));
-
- for (int i = 0; i < MAX_ITER_NUM2; ++i)
- IfTrueIncrementUSM(Q, Range, Harray, (MAX_ITER_NUM1 + i));
-
- Q.wait();
-
- // check results
- for (size_t i = 0; i < BUFFER_SIZE; ++i) {
- int expected = MAX_ITER_NUM1 + MAX_ITER_NUM2;
- assert(Harray[i] == expected);
- }
- {
- sycl::host_accessor HostAcc(Buf, sycl::read_only);
- for (size_t i = 0; i < BUFFER_SIZE; ++i) {
- int expected = MAX_ITER_NUM1;
- assert(HostAcc[i] == expected);
- }
- }
- });
+ TestHelper(Q,
+   [&](sycl::range<1> Range, int *Harray, sycl::buffer<int, 1> Buf) {
+  {
+  sycl::host_accessor HostAcc(Buf);
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+  HostAcc[i] = 0;
+  }
+  }
+
+  for (int i = 0; i < MAX_ITER_NUM1; ++i)
+  IfTrueIncrementBufferAndUSM(Q, Range, Harray, Buf, (i));
+
+  for (int i = 0; i < MAX_ITER_NUM2; ++i)
+  IfTrueIncrementUSM(Q, Range, Harray, (MAX_ITER_NUM1 + i));
+
+  Q.wait();
+
+  // check results
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+  int expected = MAX_ITER_NUM1 + MAX_ITER_NUM2;
+  assert(Harray[i] == expected);
+  }
+  {
+  sycl::host_accessor HostAcc(Buf, sycl::read_only);
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+  int expected = MAX_ITER_NUM1;
+  assert(HostAcc[i] == expected);
+  }
+  }
+  });
 }
 
 void RunTest_Mixed(sycl::queue Q) {

sycl/test-e2e/DiscardEvents/discard_events_usm.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 
 // RUN: env SYCL_PI_TRACE=2 %{run} %t.out &> %t.txt ; FileCheck %s --input-file %t.txt

sycl/test-e2e/DiscardEvents/discard_events_usm_ooo_queue.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 
 // RUN: env SYCL_PI_TRACE=2 %{run} %t.out &> %t.txt ; FileCheck %s --input-file %t.txt

sycl/test-e2e/DiscardEvents/invalid_event.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // Same hang as on Basic/barrier_order.cpp tracked in
 // https://github.com/intel/llvm/issues/7330.
 // UNSUPPORTED: opencl && gpu

sycl/test-e2e/GroupAlgorithm/exclusive_scan_over_group.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // CPU and ACC not yet supported:
 // Unsupported SPIR-V module SPIRV module requires unsupported capability 6400
 // REQUIRES: gpu

sycl/test-e2e/GroupAlgorithm/root_group.cpp

@@ -33,6 +33,8 @@ void testQueriesAndProperties() {
 
 void testRootGroup() {
  sycl::queue q;
+ if (!q.get_device().has(sycl::aspect::usm_shared_allocations))
+ return;
  const auto bundle =
  sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context());
  const auto kernel = bundle.get_kernel<class RootGroupKernel>();
@@ -66,6 +68,8 @@ void testRootGroup() {
 
 void testRootGroupFunctions() {
  sycl::queue q;
+ if (!q.get_device().has(sycl::aspect::usm_shared_allocations))
+ return;
  const auto bundle =
  sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context());
  const auto kernel = bundle.get_kernel<class RootGroupFunctionsKernel>();

sycl/test-e2e/InOrderEventsExt/get_last_event.cpp

@@ -34,7 +34,8 @@ int main() {
  Failed += Check(Q, "host_task", [&]() {
  return Q.submit([&](sycl::handler &CGH) { CGH.host_task([]() {}); });
  });
-
+ if (!Q.get_device().has(sycl::aspect::usm_shared_allocations))
+ return Failed;
  constexpr size_t N = 64;
  int *Data1 = sycl::malloc_shared<int>(N, Q);
  int *Data2 = sycl::malloc_shared<int>(N, Q);

sycl/test-e2e/InOrderEventsExt/set_external_event.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/KernelAndProgram/disable-caching.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // This test ensures created program/kernels are not retained
 // if and only if caching is disabled.
 

sycl/test-e2e/KernelFusion/sync_two_queues_event_dep.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // For this test, complete_fusion must be supported.
 // RUN: %{build} -o %t.out
 // RUN: env SYCL_RT_WARNING_LEVEL=1 %{run} %t.out 2>&1 | FileCheck %s

sycl/test-e2e/KernelFusion/sync_usm_mem_op.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -fsycl-embed-ir -o %t.out
 // RUN: env SYCL_RT_WARNING_LEVEL=1 %{run} %t.out 2>&1 | FileCheck %s
 

sycl/test-e2e/Reduction/reduction_internal.cpp

@@ -80,7 +80,8 @@ static void test(RedStorage &Storage, RangeTy Range) {
  cgh, Range, ext::oneapi::experimental::empty_properties_t{}, RedSycl,
  [=](auto Item, auto &Red) { Red.combine(T{1}); });
  }).wait();
-
+ if (!q.get_device().has(aspect::usm_shared_allocations))
+ return;
  auto *Result = malloc_shared<T>(1, q);
  q.submit([&](handler &cgh) {
  auto RedAcc = GetRedAcc(cgh);

sycl/test-e2e/Reduction/reduction_range_item.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Reduction/reduction_span.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Reduction/reduction_span_pack.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 //

sycl/test-e2e/Regression/exclusive-scan-char-short.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Regression/group_local_linear_id.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Regression/half_operators.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // REQUIRES: gpu
 // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %{run} %t.out

sycl/test-e2e/Regression/pf-wg-atomic64.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // DISABLED: aspect-atomic64
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

sycl/test-e2e/Regression/range-rounding-this-id.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // This test ensures that this_id returns the correct value
 // even when a kernel is wrapped in a range rounding kernel.
 // RUN: %{build} -o %t.out

sycl/test-e2e/Regression/reduction_64bit_atomic64.cpp

@@ -1,3 +1,4 @@
+// REQUIRES: usm_shared_allocations
 // REQUIRES: aspect-atomic64
 // RUN: %{build} -o %t.out
 //