Rewrite tests using malloc_device and fix casing

sycl/test-e2e/DiscardEvents/invalid_event.cpp

@@ -2,7 +2,6 @@
 // https://github.com/intel/llvm/issues/7330.
 // UNSUPPORTED: opencl && gpu
 // RUN: %{build} -o %t.out
-// REQUIRES: aspect-usm_shared_allocations
 // RUN: %{run} %t.out
 
 // The test checks that each PI call to the queue returns a discarded event
@@ -19,9 +18,9 @@ void QueueAPIsReturnDiscardedEvent(sycl::queue Q) {
  sycl::range<1> range(BUFFER_SIZE);
 
  auto Dev = Q.get_device();
- int *x = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+ int *x = sycl::malloc_device<int>(BUFFER_SIZE, Q);
  assert(x != nullptr);
- int *y = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+ int *y = sycl::malloc_device<int>(BUFFER_SIZE, Q);
  assert(y != nullptr);
 
  sycl::event DiscardedEvent;

sycl/test-e2e/GroupAlgorithm/root_group.cpp

@@ -41,10 +41,10 @@ void testRootGroup() {
  max_num_work_group_sync>(q);
  const auto props = sycl::ext::oneapi::experimental::properties{
  sycl::ext::oneapi::experimental::use_root_sync};
- sycl::buffer<int> data_buf{sycl::range{maxWGs * WorkGroupSize}};
+ sycl::buffer<int> dataBuf{sycl::range{maxWGs * WorkGroupSize}};
  const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
  q.submit([&](sycl::handler &h) {
- sycl::accessor data{data_buf, h};
+ sycl::accessor data{dataBuf, h};
  h.parallel_for<class RootGroupKernel>(
  range, props, [=](sycl::nd_item<1> it) {
  auto root = it.ext_oneapi_get_root_group();
@@ -59,7 +59,7 @@ void testRootGroup() {
  data[root.get_local_id()] = sum;
  });
  });
- sycl::host_accessor data{data_buf};
+ sycl::host_accessor data{dataBuf};
  const int workItemCount = static_cast<int>(range.get_global_range().size());
  for (int i = 0; i < workItemCount; i++) {
  assert(data[i] == (workItemCount - 1));
@@ -78,10 +78,10 @@ void testRootGroupFunctions() {
  sycl::ext::oneapi::experimental::use_root_sync};
 
  constexpr int testCount = 10;
- sycl::buffer<bool> testResults_buf{sycl::range{testCount}};
+ sycl::buffer<bool> testResultsBuf{sycl::range{testCount}};
  const auto range = sycl::nd_range<1>{maxWGs * WorkGroupSize, WorkGroupSize};
  q.submit([&](sycl::handler &h) {
- sycl::accessor testResults{testResults_buf, h};
+ sycl::accessor testResults{testResultsBuf, h};
  h.parallel_for<class RootGroupFunctionsKernel>(
  range, props, [=](sycl::nd_item<1> it) {
  const auto root = it.ext_oneapi_get_root_group();
@@ -113,7 +113,7 @@ void testRootGroupFunctions() {
  }
  });
  });
- sycl::host_accessor testResults{testResults_buf};
+ sycl::host_accessor testResults{testResultsBuf};
  for (int i = 0; i < testCount; i++) {
  assert(testResults[i]);
  }

sycl/test-e2e/InOrderEventsExt/set_external_event.cpp

@@ -15,8 +15,8 @@ int main() {
  sycl::queue Q1{Ctx, Dev, {sycl::property::queue::in_order{}}};
  sycl::queue Q2{Ctx, Dev, {sycl::property::queue::in_order{}}};
 
- sycl::buffer<int> DevDatabuf{sycl::range{N}};
- sycl::accessor DevData{DevDatabuf};
+ sycl::buffer<int> DevDataBuf{sycl::range{N}};
+ sycl::accessor DevData{DevDataBuf};
  int *HostData = (int *)malloc(N * sizeof(int) * 10);
 
  for (size_t I = 0; I < 10; ++I) {
@@ -52,6 +52,6 @@ int main() {
  }
  }
  }
-
+ free(HostData);
  return Failures;
 }

sycl/test-e2e/Reduction/reduction_internal.cpp

@@ -80,13 +80,13 @@ 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();
- sycl::buffer<T> Result_buf{sycl::range{1}};
+ sycl::buffer<T> ResultBuf{sycl::range{1}};
  q.submit([&](handler &cgh) {
- sycl::accessor Result{Result_buf, cgh};
+ sycl::accessor Result{ResultBuf, cgh};
  auto RedAcc = GetRedAcc(cgh);
  cgh.single_task([=]() { Result[0] = RedAcc[0]; });
  });
- sycl::host_accessor Result{Result_buf};
+ sycl::host_accessor Result{ResultBuf};
  auto N = get_global_range(Range).size();
  int Expected = InitToIdentity ? N : Init + N;
 #if defined(__PRETTY_FUNCTION__)

sycl/test-e2e/Reduction/reduction_range_item.cpp

@@ -3,15 +3,14 @@
 
 // Windows doesn't yet have full shutdown().
 // UNSUPPORTED: ze_debug && windows
-// REQUIRES: aspect-usm_shared_allocations
 #include <sycl/sycl.hpp>
 
 using namespace sycl;
 
 int main() {
  queue q;
- auto *RedMem = malloc_shared<int>(1, q);
- auto *Success = malloc_shared<bool>(1, q);
+ auto *RedMem = malloc_device<int>(1, q);
+ auto *Success = malloc_device<bool>(1, q);
  *Success = true;
 
  *RedMem = 0;
@@ -24,11 +23,15 @@ int main() {
  *Success = false;
  })
  .wait();
-
- assert(*RedMem == 7);
- assert(*Success);
-
- *RedMem = 0;
+ int RedMemHost;
+ bool SuccessHost;
+ q.memcpy(&RedMemHost, RedMem, sizeof(int)).wait();
+ q.memcpy(&SuccessHost, Success, sizeof(bool)).wait();
+ assert(RedMemHost == 7);
+ assert(SuccessHost);
+
+ RedMemHost = 0;
+ q.memcpy(RedMem, &RedMemHost, sizeof(int)).wait();
  q.parallel_for(range<2>{1030, 7}, reduction(RedMem, std::plus<int>{}),
  [=](item<2> Item, auto &Red) {
  Red += 1;
@@ -44,8 +47,10 @@ int main() {
  })
  .wait();
 
- assert(*RedMem == 1030 * 7);
- assert(*Success);
+ q.memcpy(&RedMemHost, RedMem, sizeof(int)).wait();
+ q.memcpy(&SuccessHost, Success, sizeof(bool)).wait();
+ assert(RedMemHost == 1030 * 7);
+ assert(SuccessHost);
 
  free(RedMem, q);
  free(Success, q);

sycl/test-e2e/Reduction/reduction_span.cpp

@@ -3,7 +3,6 @@
 
 // Windows doesn't yet have full shutdown().
 // UNSUPPORTED: ze_debug && windows
-// REQUIRES: aspect-usm_shared_allocations
 // This test performs basic checks of reductions initialized with a sycl::span
 
 #include <sycl/sycl.hpp>
@@ -44,7 +43,7 @@ template <size_t N, typename T, typename BinaryOperation, typename Range,
 void test(queue Q, Range Rng, T Identity, T Value) {
 
  // Initialize output to identity value
- T *Output = malloc_shared<T>(N, Q);
+ T *Output = malloc_device<T>(N, Q);
  Q.parallel_for(range<1>{N}, [=](id<1> I) { Output[I] = Identity; }).wait();
 
  // Perform generalized "histogram" with N bins
@@ -70,11 +69,13 @@ void test(queue Q, Range Rng, T Identity, T Value) {
  }
 
  bool Passed = true;
+ T OutputHost[N];
+ Q.memcpy(OutputHost, Output, N * sizeof(T)).wait();
  for (size_t I = 0; I < N; ++I) {
  if (I < Size % N) {
- Passed &= (Output[I] == Expected);
+ Passed &= (OutputHost[I] == Expected);
  } else {
- Passed &= (Output[I] == ExpectedRemainder);
+ Passed &= (OutputHost[I] == ExpectedRemainder);
  }
  }
 

sycl/test-e2e/Reduction/reduction_span_pack.cpp

@@ -6,7 +6,6 @@
 
 // Windows doesn't yet have full shutdown().
 // UNSUPPORTED: ze_debug && windows
-// REQUIRES: aspect-usm_shared_allocations
 // This test performs basic checks of reductions initialized with a pack
 // containing at least one sycl::span
 
@@ -49,9 +48,9 @@ template <size_t N, typename T, typename BinaryOperation, typename Range,
 void test1(queue Q, Range Rng, T Identity, T Value) {
 
  // Initialize output to identity value
- int *Sum = malloc_shared<int>(1, Q);
+ int *Sum = malloc_device<int>(1, Q);
  Q.single_task([=]() { *Sum = 0; }).wait();
- T *Output = malloc_shared<T>(N, Q);
+ T *Output = malloc_device<T>(N, Q);
  Q.parallel_for(range<1>{N}, [=](id<1> I) { Output[I] = Identity; }).wait();
 
  // Perform generalized "histogram" with N bins
@@ -81,11 +80,13 @@ void test1(queue Q, Range Rng, T Identity, T Value) {
  }
 
  bool Passed = true;
+ T OutputHost[N];
+ Q.memcpy(OutputHost, Output, N * sizeof(T)).wait();
  for (size_t I = 0; I < N; ++I) {
  if (I < Size % N) {
- Passed &= (Output[I] == Expected);
+ Passed &= (OutputHost[I] == Expected);
  } else {
- Passed &= (Output[I] == ExpectedRemainder);
+ Passed &= (OutputHost[I] == ExpectedRemainder);
  }
  }
  Passed &= (*Sum == Size);
@@ -101,9 +102,9 @@ template <size_t N, typename T, typename BinaryOperation, typename Range,
 void test2(queue Q, Range Rng, T Identity, T Value) {
 
  // Initialize output to identity value
- int *Output1 = malloc_shared<int>(N, Q);
+ int *Output1 = malloc_device<int>(N, Q);
  Q.parallel_for(range<1>{N}, [=](id<1> I) { Output1[I] = 0; }).wait();
- T *Output2 = malloc_shared<T>(N, Q);
+ T *Output2 = malloc_device<T>(N, Q);
  Q.parallel_for(range<1>{N}, [=](id<1> I) { Output2[I] = Identity; }).wait();
 
  // Perform generalized "histogram" with N bins
@@ -121,7 +122,10 @@ void test2(queue Q, Range Rng, T Identity, T Value) {
  } else /*if (SubmissionMode == submission_mode::queue) */ {
  Q.parallel_for(Rng, Redu1, Redu2, Kern).wait();
  }
-
+ int Output1Host[N];
+ T Output2Host[N];
+ Q.memcpy(Output1Host, Output1, N * sizeof(int)).wait();
+ Q.memcpy(Output2Host, Output2, N * sizeof(T)).wait();
  size_t Size = getLinearSize(Rng);
  bool Passed = true;
  // Span1
@@ -132,12 +136,11 @@ void test2(queue Q, Range Rng, T Identity, T Value) {
  ExpectedRemainder = Expected;
  Expected += 1;
  }
-
  for (size_t I = 0; I < N; ++I) {
  if (I < Size % N) {
- Passed &= (Output1[I] == Expected);
+ Passed &= (Output1Host[I] == Expected);
  } else {
- Passed &= (Output1[I] == ExpectedRemainder);
+ Passed &= (Output1Host[I] == ExpectedRemainder);
  }
  }
  }
@@ -153,9 +156,9 @@ void test2(queue Q, Range Rng, T Identity, T Value) {
 
  for (size_t I = 0; I < N; ++I) {
  if (I < Size % N) {
- Passed &= (Output2[I] == Expected);
+ Passed &= (Output2Host[I] == Expected);
  } else {
- Passed &= (Output2[I] == ExpectedRemainder);
+ Passed &= (Output2Host[I] == ExpectedRemainder);
  }
  }
  }

sycl/test-e2e/Regression/reduction_64bit_atomic64.cpp

@@ -2,7 +2,6 @@
 // RUN: %{build} -o %t.out
 //
 // RUN: %{run} %t.out
-// REQUIRES: aspect-usm_shared_allocations
 // Tests that a previously known case for reduction doesn't cause a requirement
 // for atomic64.
 // TODO: When aspect requirements are added to testing, this test could be set
@@ -19,7 +18,7 @@ using namespace sycl;
 int main() {
  queue Q;
 
- long long *Out = malloc_shared<long long>(1, Q);
+ long long *Out = malloc_device<long long>(1, Q);
 
  // Case 1: nd_range reduction with 64-bit integer and either sycl::plus,
  // sycl::minimum or sycl::maximum. group_reduce_and_atomic_cross_wg strategy