From 00b210e67396f3b36ac23d4e560b880f7298f62f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Alberto=20Cabrera=20P=C3=A9rez?= Date: Thu, 11 Apr 2024 15:55:36 +0100 Subject: [PATCH] [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. --- .../test-e2e/syclcompat/math/math_complex.cpp | 11 +- .../syclcompat/memory/memory_fixt.hpp | 10 +- .../memory/memory_management_shared.cpp | 91 ++++++++++++++ .../memory/memory_management_test2.cpp | 42 ------- .../syclcompat/memory/usm_allocations.cpp | 73 ++--------- .../memory/usm_shared_allocations.cpp | 113 ++++++++++++++++++ 6 files changed, 225 insertions(+), 115 deletions(-) create mode 100644 sycl/test-e2e/syclcompat/memory/memory_management_shared.cpp create mode 100644 sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp diff --git a/sycl/test-e2e/syclcompat/math/math_complex.cpp b/sycl/test-e2e/syclcompat/math/math_complex.cpp index cd81246f07ca3..3b7dba8b24965 100644 --- a/sycl/test-e2e/syclcompat/math/math_complex.cpp +++ b/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 @@ -72,18 +71,20 @@ template 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(1, 1, result_); // Run on device syclcompat::wait(); - assert(*result_ == 1); + syclcompat::memcpy(&h_result_, result_, 1); + assert(h_result_ == 1); assert(cpu_result_ == 1); } }; diff --git a/sycl/test-e2e/syclcompat/memory/memory_fixt.hpp b/sycl/test-e2e/syclcompat/memory/memory_fixt.hpp index 7c613e25a4a4a..e5b8c6ef37972 100644 --- a/sycl/test-e2e/syclcompat/memory/memory_fixt.hpp +++ b/sycl/test-e2e/syclcompat/memory/memory_fixt.hpp @@ -123,7 +123,7 @@ template struct USMTest { skip{should_skip(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 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(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 struct USMTest { sycl::queue q_; syclcompat::dim3 const grid_; syclcompat::dim3 const thread_; - T *d_A; + T *data; size_t size_; bool skip; }; diff --git a/sycl/test-e2e/syclcompat/memory/memory_management_shared.cpp b/sycl/test-e2e/syclcompat/memory/memory_management_shared.cpp new file mode 100644 index 0000000000000..b9641c6864f9f --- /dev/null +++ b/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 + +#include + +#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 s_A(DataW); + syclcompat::shared_memory s_B(DataW); + syclcompat::shared_memory 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; +} diff --git a/sycl/test-e2e/syclcompat/memory/memory_management_test2.cpp b/sycl/test-e2e/syclcompat/memory/memory_management_test2.cpp index 9774faa8e8ba7..afb1f6a5f5a88 100644 --- a/sycl/test-e2e/syclcompat/memory/memory_management_test2.cpp +++ b/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 s_A(DataW); - syclcompat::shared_memory s_B(DataW); - syclcompat::shared_memory 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; } diff --git a/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp b/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp index fb8a8a52da101..78eb465819865 100644 --- a/sycl/test-e2e/syclcompat/memory/usm_allocations.cpp +++ b/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,10 +40,10 @@ template void test_malloc() { if (usm_fixture.skip) return; // Skip unsupported - usm_fixture.d_A = syclcompat::malloc(usm_fixture.size_); + usm_fixture.data = syclcompat::malloc(usm_fixture.size_); usm_fixture.launch_kernel(); usm_fixture.check_result(); - syclcompat::free(usm_fixture.d_A); + syclcompat::free(usm_fixture.data); } template void test_host() { @@ -52,66 +51,37 @@ template void test_host() { USMTest 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(usm_fixture.size_); - usm_fixture.launch_kernel(); - usm_fixture.check_result(); - syclcompat::free(usm_fixture.d_A); -} - -template void test_shared() { - std::cout << __PRETTY_FUNCTION__ << std::endl; - USMTest 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(usm_fixture.size_); + usm_fixture.data = syclcompat::malloc_host(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 usm_fixture; - usm_fixture.d_A = + usm_fixture.data = static_cast(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 usm_fixture; - if (!usm_fixture.q_.get_device().has(sycl::aspect::usm_host_allocations)) return; // Skip unsupported - usm_fixture.d_A = static_cast( + usm_fixture.data = static_cast( 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 usm_fixture; - - if (!usm_fixture.q_.get_device().has(sycl::aspect::usm_shared_allocations)) - return; - - usm_fixture.d_A = static_cast( - 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(); diff --git a/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp b/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp new file mode 100644 index 0000000000000..6e5f7bcfef346 --- /dev/null +++ b/sycl/test-e2e/syclcompat/memory/usm_shared_allocations.cpp @@ -0,0 +1,113 @@ +/*************************************************************************** + * + * 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 + * + * usm_allocations.cpp + * + * Description: + * 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 + +#include +#include + +#include + +#include + +#include "../common.hpp" +#include "memory_common.hpp" +#include "memory_fixt.hpp" + +template void test_shared() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + USMTest usm_fixture; + + if (usm_fixture.skip) + return; // Skip unsupported + + usm_fixture.data = syclcompat::malloc_shared(usm_fixture.size_); + usm_fixture.launch_kernel(); + usm_fixture.check_result(); + syclcompat::free(usm_fixture.data); +} + +void test_non_templated_shared() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + USMTest usm_fixture; + + usm_fixture.data = static_cast( + syclcompat::malloc_shared(usm_fixture.size_ * sizeof(int))); + usm_fixture.launch_kernel(); + usm_fixture.check_result(); + syclcompat::free(usm_fixture.data); +} + +// Test deduce direction +void test_deduce_shared() { + std::cout << __PRETTY_FUNCTION__ << std::endl; + + using memcpy_direction = syclcompat::detail::memcpy_direction; + auto default_queue = syclcompat::get_default_queue(); + + 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, + s_ptr) == + memcpy_direction::device_to_device); + + // * to sys + 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, + 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); + + syclcompat::free(s_ptr); + std::free(sys_ptr); + syclcompat::free(h_ptr); + syclcompat::free(d_ptr); +} + +int main() { + INSTANTIATE_ALL_TYPES(value_type_list, test_shared); + test_non_templated_shared(); + test_deduce_shared(); + + return 0; +}