[SYCL][COMPAT] shared memory test split from main tests (#13063)

Some tests were being completely skipped entirely because a part of them
required `aspect::usm_shared_allocations`.
This PR splits the tests in various files to disable only the usages of
shared memory.

sycl/test-e2e/syclcompat/math/math_complex.cpp

@@ -30,8 +30,7 @@
 //===---------------------------------------------------------------===//
 
 // REQUIRES: aspect-fp64
-// REQUIRES: usm_shared_allocations
-// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
+// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
 // RUN: %{run} %t.out
 
 #include <complex>
@@ -72,18 +71,20 @@ template <auto F> class ComplexLauncher {
 protected:
  int *result_;
  int cpu_result_{0};
+ int h_result_;
 
 public:
  ComplexLauncher() {
- result_ = (int *)syclcompat::malloc_shared(sizeof(int));
- *result_ = 0;
+ result_ = (int *)syclcompat::malloc(sizeof(int));
+ syclcompat::memset(result_, 0, sizeof(int));
  };
  ~ComplexLauncher() { syclcompat::free(result_); }
  void launch() {
  F(&cpu_result_); // Run on host
  syclcompat::launch<F>(1, 1, result_); // Run on device
  syclcompat::wait();
- assert(*result_ == 1);
+ syclcompat::memcpy<int>(&h_result_, result_, 1);
+ assert(h_result_ == 1);
  assert(cpu_result_ == 1);
  }
 };

sycl/test-e2e/syclcompat/memory/memory_fixt.hpp

@@ -123,7 +123,7 @@ template <typename T> struct USMTest {
  skip{should_skip<T>(syclcompat::get_current_device())} {}
 
  void launch_kernel() {
- auto &dd_A = d_A;
+ auto &dd_A = data;
  return q_
  .submit([&](sycl::handler &cgh) {
  cgh.parallel_for(
@@ -135,15 +135,15 @@ template <typename T> struct USMTest {
  // Check result is identity vector
  // Handles memcpy for USM device alloc
  void check_result() {
- sycl::usm::alloc ptr_type = sycl::get_pointer_type(d_A, q_.get_context());
+ sycl::usm::alloc ptr_type = sycl::get_pointer_type(data, q_.get_context());
  assert(ptr_type != sycl::usm::alloc::unknown);
 
  T *result;
  if (ptr_type == sycl::usm::alloc::device) {
  result = static_cast<T *>(std::malloc(sizeof(T) * size_));
- syclcompat::memcpy(result, d_A, sizeof(T) * size_);
+ syclcompat::memcpy(result, data, sizeof(T) * size_);
  } else {
- result = d_A;
+ result = data;
  }
 
  for (size_t i = 0; i < size_; i++) {
@@ -157,7 +157,7 @@ template <typename T> struct USMTest {
  sycl::queue q_;
  syclcompat::dim3 const grid_;
  syclcompat::dim3 const thread_;
- T *d_A;
+ T *data;
  size_t size_;
  bool skip;
 };

sycl/test-e2e/syclcompat/memory/memory_management_shared.cpp

@@ -0,0 +1,91 @@
+/***************************************************************************
+ *
+ * Copyright (C) Codeplay Software Ltd.
+ *
+ * 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
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ * SYCLcompat API
+ *
+ * memory_management_shared.cpp
+ *
+ * Description:
+ * memory operations tests with shared memory
+ **************************************************************************/
+
+// The original source was under the license below:
+// ====------ memory_management_test2.cpp---------- -*- C++ -* ----===////
+//
+// 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: usm_shared_allocations
+// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <syclcompat/memory.hpp>
+
+#include "../common.hpp"
+#include "memory_common.hpp"
+
+constexpr size_t DataW = 100;
+constexpr size_t DataH = 100;
+
+void test_shared_memory() {
+ std::cout << __PRETTY_FUNCTION__ << std::endl;
+
+ syclcompat::shared_memory<float, 1> s_A(DataW);
+ syclcompat::shared_memory<float, 1> s_B(DataW);
+ syclcompat::shared_memory<float, 1> s_C(DataW);
+
+ s_A.init();
+ s_B.init();
+ s_C.init();
+
+ for (int i = 0; i < DataW; i++) {
+ s_A[i] = 1.0f;
+ s_B[i] = 2.0f;
+ }
+
+ {
+ syclcompat::get_default_queue().submit([&](sycl::handler &cgh) {
+ float *d_A = s_A.get_ptr();
+ float *d_B = s_B.get_ptr();
+ float *d_C = s_C.get_ptr();
+ cgh.parallel_for(sycl::range<1>(DataW), [=](sycl::id<1> id) {
+ int i = id[0];
+ float *A = d_A;
+ float *B = d_B;
+ float *C = d_C;
+ C[i] = A[i] + B[i];
+ });
+ });
+ syclcompat::get_default_queue().wait_and_throw();
+ }
+
+ // verify hostD
+ for (int i = 0; i < DataW; i++) {
+ for (int j = 0; j < DataH; j++) {
+ assert(fabs(s_C[i] - s_A[i] - s_B[i]) <= 1e-5);
+ }
+ }
+}
+
+int main() {
+ test_shared_memory();
+
+ return 0;
+}

sycl/test-e2e/syclcompat/memory/memory_management_test2.cpp

@@ -30,7 +30,6 @@
 //
 // ===----------------------------------------------------------------------===//
 
-// REQUIRES: usm_shared_allocations
 // RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
 // RUN: %{run} %t.out
 
@@ -217,46 +216,6 @@ void test_global_memory() {
  }
 }
 
-void test_shared_memory() {
- std::cout << __PRETTY_FUNCTION__ << std::endl;
-
- syclcompat::shared_memory<float, 1> s_A(DataW);
- syclcompat::shared_memory<float, 1> s_B(DataW);
- syclcompat::shared_memory<float, 1> s_C(DataW);
-
- s_A.init();
- s_B.init();
- s_C.init();
-
- for (int i = 0; i < DataW; i++) {
- s_A[i] = 1.0f;
- s_B[i] = 2.0f;
- }
-
- {
- syclcompat::get_default_queue().submit([&](sycl::handler &cgh) {
- float *d_A = s_A.get_ptr();
- float *d_B = s_B.get_ptr();
- float *d_C = s_C.get_ptr();
- cgh.parallel_for(sycl::range<1>(DataW), [=](sycl::id<1> id) {
- int i = id[0];
- float *A = d_A;
- float *B = d_B;
- float *C = d_C;
- C[i] = A[i] + B[i];
- });
- });
- syclcompat::get_default_queue().wait_and_throw();
- }
-
- // verify hostD
- for (int i = 0; i < DataW; i++) {
- for (int j = 0; j < DataH; j++) {
- assert(fabs(s_C[i] - s_A[i] - s_B[i]) <= 1e-5);
- }
- }
-}
-
 void test_constant_memory() {
  std::cout << __PRETTY_FUNCTION__ << std::endl;
 
@@ -366,7 +325,6 @@ int main() {
  test_memcpy_pitched_q();
 
  test_global_memory();
- test_shared_memory();
  test_constant_memory();
  return 0;
 }

sycl/test-e2e/syclcompat/memory/usm_allocations.cpp

@@ -20,7 +20,6 @@
  * USM allocation tests
  **************************************************************************/
 
-// REQUIRES: usm_shared_allocations
 // RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
 // RUN: %{run} %t.out
 
@@ -41,77 +40,48 @@ template <typename T> void test_malloc() {
  if (usm_fixture.skip)
  return; // Skip unsupported
 
- usm_fixture.d_A = syclcompat::malloc<T>(usm_fixture.size_);
+ usm_fixture.data = syclcompat::malloc<T>(usm_fixture.size_);
  usm_fixture.launch_kernel();
  usm_fixture.check_result();
- syclcompat::free(usm_fixture.d_A);
+ syclcompat::free(usm_fixture.data);
 }
 
 template <typename T> void test_host() {
  std::cout << __PRETTY_FUNCTION__ << std::endl;
  USMTest<T> usm_fixture;
  if (usm_fixture.skip)
  return; // Skip unsupported
-
  if (!usm_fixture.q_.get_device().has(sycl::aspect::usm_host_allocations))
  return; // Skip unsupported
- usm_fixture.d_A = syclcompat::malloc_host<T>(usm_fixture.size_);
- usm_fixture.launch_kernel();
- usm_fixture.check_result();
- syclcompat::free(usm_fixture.d_A);
-}
-
-template <typename T> void test_shared() {
- std::cout << __PRETTY_FUNCTION__ << std::endl;
- USMTest<T> usm_fixture;
- if (usm_fixture.skip)
- return; // Skip unsupported
 
- if (!usm_fixture.q_.get_device().has(sycl::aspect::usm_shared_allocations))
- return; // Skip unsupported
- usm_fixture.d_A = syclcompat::malloc_shared<T>(usm_fixture.size_);
+ usm_fixture.data = syclcompat::malloc_host<T>(usm_fixture.size_);
  usm_fixture.launch_kernel();
  usm_fixture.check_result();
- syclcompat::free(usm_fixture.d_A);
+ syclcompat::free(usm_fixture.data);
 }
 
 void test_non_templated_malloc() {
  std::cout << __PRETTY_FUNCTION__ << std::endl;
  USMTest<int> usm_fixture;
 
- usm_fixture.d_A =
+ usm_fixture.data =
  static_cast<int *>(syclcompat::malloc(usm_fixture.size_ * sizeof(int)));
  usm_fixture.launch_kernel();
  usm_fixture.check_result();
- syclcompat::free(usm_fixture.d_A);
+ syclcompat::free(usm_fixture.data);
 }
 
 void test_non_templated_host() {
  std::cout << __PRETTY_FUNCTION__ << std::endl;
  USMTest<int> usm_fixture;
-
  if (!usm_fixture.q_.get_device().has(sycl::aspect::usm_host_allocations))
  return; // Skip unsupported
 
- usm_fixture.d_A = static_cast<int *>(
+ usm_fixture.data = static_cast<int *>(
  syclcompat::malloc_host(usm_fixture.size_ * sizeof(int)));
  usm_fixture.launch_kernel();
  usm_fixture.check_result();
- syclcompat::free(usm_fixture.d_A);
-}
-
-void test_non_templated_shared() {
- std::cout << __PRETTY_FUNCTION__ << std::endl;
- USMTest<int> usm_fixture;
-
- if (!usm_fixture.q_.get_device().has(sycl::aspect::usm_shared_allocations))
- return;
-
- usm_fixture.d_A = static_cast<int *>(
- syclcompat::malloc_shared(usm_fixture.size_ * sizeof(int)));
- usm_fixture.launch_kernel();
- usm_fixture.check_result();
- syclcompat::free(usm_fixture.d_A);
+ syclcompat::free(usm_fixture.data);
 }
 
 // Test deduce direction
@@ -120,11 +90,12 @@ void test_deduce() {
 
  using memcpy_direction = syclcompat::detail::memcpy_direction;
  auto default_queue = syclcompat::get_default_queue();
+ if (!default_queue.get_device().has(sycl::aspect::usm_host_allocations))
+ return; // Skip unsupported
 
  int *h_ptr = (int *)syclcompat::malloc_host(sizeof(int));
  int *sys_ptr = (int *)std::malloc(sizeof(int));
  int *d_ptr = (int *)syclcompat::malloc(sizeof(int));
- int *s_ptr = (int *)syclcompat::malloc_shared(sizeof(int));
 
  // * to host
  assert(syclcompat::detail::deduce_memcpy_direction(default_queue, h_ptr,
@@ -135,9 +106,6 @@ void test_deduce() {
  assert(syclcompat::detail::deduce_memcpy_direction(default_queue, h_ptr,
  d_ptr) ==
  memcpy_direction::device_to_device);
- assert(syclcompat::detail::deduce_memcpy_direction(default_queue, h_ptr,
- s_ptr) ==
- memcpy_direction::device_to_device);
 
  // * to sys
  assert(syclcompat::detail::deduce_memcpy_direction(
@@ -148,8 +116,6 @@ void test_deduce() {
  assert(syclcompat::detail::deduce_memcpy_direction(default_queue, sys_ptr,
  d_ptr) ==
  memcpy_direction::device_to_host);
- assert(syclcompat::detail::deduce_memcpy_direction(
- default_queue, sys_ptr, s_ptr) == memcpy_direction::host_to_host);
 
  // * to dev
  assert(syclcompat::detail::deduce_memcpy_direction(default_queue, d_ptr,
@@ -161,39 +127,20 @@ void test_deduce() {
  assert(syclcompat::detail::deduce_memcpy_direction(default_queue, d_ptr,
  d_ptr) ==
  memcpy_direction::device_to_device);
- assert(syclcompat::detail::deduce_memcpy_direction(default_queue, d_ptr,
- s_ptr) ==
- memcpy_direction::device_to_device);
-
- // * to shared
- assert(syclcompat::detail::deduce_memcpy_direction(default_queue, s_ptr,
- h_ptr) ==
- memcpy_direction::device_to_device);
- assert(syclcompat::detail::deduce_memcpy_direction(
- default_queue, s_ptr, sys_ptr) == memcpy_direction::host_to_host);
- assert(syclcompat::detail::deduce_memcpy_direction(default_queue, s_ptr,
- d_ptr) ==
- memcpy_direction::device_to_device);
- assert(syclcompat::detail::deduce_memcpy_direction(default_queue, s_ptr,
- s_ptr) ==
- memcpy_direction::device_to_device);
 
  std::free(sys_ptr);
  syclcompat::free(h_ptr);
  syclcompat::free(d_ptr);
- syclcompat::free(s_ptr);
 }
 
 int main() {
  INSTANTIATE_ALL_TYPES(value_type_list, test_malloc);
  INSTANTIATE_ALL_TYPES(value_type_list, test_host);
- INSTANTIATE_ALL_TYPES(value_type_list, test_shared);
 
  // Avoid combinatorial explosion by only testing non-templated
  // syclcompat::malloc with int type
  test_non_templated_malloc();
  test_non_templated_host();
- test_non_templated_shared();
 
  test_deduce();