From 70cf4c9dac86294bc62814c47218fad205d8a2a4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 18:19:15 +0100 Subject: [PATCH] EXSWHTEC-302 - Implement tests for memory fence device functions #276 Change-Id: I327527288f90011f09262708dd6372a7c6fc4708 [ROCm/hip-tests commit: 8ef0d724f548fb6c25dfe35a8c9e31f00ef19d57] --- .../catch/include/hip_test_defgroups.hh | 7 + projects/hip-tests/catch/unit/CMakeLists.txt | 1 + .../catch/unit/threadfence/CMakeLists.txt | 29 +++ .../catch/unit/threadfence/__threadfence.cc | 201 ++++++++++++++++++ .../unit/threadfence/__threadfence_block.cc | 201 ++++++++++++++++++ .../unit/threadfence/__threadfence_system.cc | 126 +++++++++++ .../unit/threadfence/threadfence_common.hh | 108 ++++++++++ 7 files changed, 673 insertions(+) create mode 100644 projects/hip-tests/catch/unit/threadfence/CMakeLists.txt create mode 100644 projects/hip-tests/catch/unit/threadfence/__threadfence.cc create mode 100644 projects/hip-tests/catch/unit/threadfence/__threadfence_block.cc create mode 100644 projects/hip-tests/catch/unit/threadfence/__threadfence_system.cc create mode 100644 projects/hip-tests/catch/unit/threadfence/threadfence_common.hh diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 7da19422c2..3fb6c774ef 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -116,6 +116,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup ThreadfenceTest Memory Fence Functions + * @{ + * This section describes tests for Memory Fence Functions. + * @} + */ + /** * @defgroup MemoryTest memory Management APIs * @{ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 3074b8b0b5..37f8b73cc5 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -50,6 +50,7 @@ add_subdirectory(complex) add_subdirectory(p2p) add_subdirectory(gcc) add_subdirectory(syncthreads) +add_subdirectory(threadfence) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) diff --git a/projects/hip-tests/catch/unit/threadfence/CMakeLists.txt b/projects/hip-tests/catch/unit/threadfence/CMakeLists.txt new file mode 100644 index 0000000000..51f61f4e2d --- /dev/null +++ b/projects/hip-tests/catch/unit/threadfence/CMakeLists.txt @@ -0,0 +1,29 @@ +# Copyright (c) 2023 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +set(TEST_SRC + __threadfence_block.cc + __threadfence.cc + __threadfence_system.cc +) + +hip_add_exe_to_target(NAME ThreadfenceTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/threadfence/__threadfence.cc b/projects/hip-tests/catch/unit/threadfence/__threadfence.cc new file mode 100644 index 0000000000..781da0bee3 --- /dev/null +++ b/projects/hip-tests/catch/unit/threadfence/__threadfence.cc @@ -0,0 +1,201 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +#include "threadfence_common.hh" + +/** + * @addtogroup __threadfence __threadfence + * @{ + * @ingroup ThreadfenceTest + */ + +/** + * Test Description + * ------------------------ + * - Basic test for a device-wide memory fence on shared memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_Positive_Basic_Shared") { + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HipTest::launchKernel(ThreadfenceTestKernel, 1, 2, + 4 * sizeof(int), nullptr, out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a device-wide memory fence on global memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_Positive_Basic_Global") { + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a device-wide memory fence on page-locked host memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_Positive_Basic_Pinned") { + LinearAllocGuard in_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + in_host.host_ptr()[0] = kInitVal1; + in_host.host_ptr()[1] = kInitVal2; + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_host.host_ptr(), in_host.host_ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(!(out_host.host_ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a device-wide memory fence on managed memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_Positive_Basic_Managed") { + LinearAllocGuard in_host(LinearAllocs::hipMallocManaged, 2 * sizeof(int)); + LinearAllocGuard out_host(LinearAllocs::hipMallocManaged, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + in_host.host_ptr()[0] = kInitVal1; + in_host.host_ptr()[1] = kInitVal2; + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_host.ptr(), in_host.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(!(out_host.host_ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a device-wide memory fence on global peer device memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_Positive_Basic_Peer") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("At least 2 devices are required"); + return; + } + + int can_access_peer; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1)); + REQUIRE(can_access_peer); + + HIP_CHECK(hipSetDevice(0)); + + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HIP_CHECK(hipSetDevice(1)); + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipSetDevice(0)); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/threadfence/__threadfence_block.cc b/projects/hip-tests/catch/unit/threadfence/__threadfence_block.cc new file mode 100644 index 0000000000..43079ea636 --- /dev/null +++ b/projects/hip-tests/catch/unit/threadfence/__threadfence_block.cc @@ -0,0 +1,201 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +#include "threadfence_common.hh" + +/** + * @addtogroup __threadfence_block __threadfence_block + * @{ + * @ingroup ThreadfenceTest + */ + +/** + * Test Description + * ------------------------ + * - Basic test for a block-wide memory fence on shared memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_block.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_block_Positive_Basic_Shared") { + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HipTest::launchKernel(ThreadfenceTestKernel, 1, 2, + 4 * sizeof(int), nullptr, out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a block-wide memory fence on global memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_block.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_block_Positive_Basic_Global") { + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a block-wide memory fence on page-locked host memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_block.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_block_Positive_Basic_Pinned") { + LinearAllocGuard in_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + in_host.host_ptr()[0] = kInitVal1; + in_host.host_ptr()[1] = kInitVal2; + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_host.host_ptr(), in_host.host_ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(!(out_host.host_ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a block-wide memory fence on managed memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_block.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_block_Positive_Basic_Managed") { + LinearAllocGuard in_host(LinearAllocs::hipMallocManaged, 2 * sizeof(int)); + LinearAllocGuard out_host(LinearAllocs::hipMallocManaged, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + in_host.host_ptr()[0] = kInitVal1; + in_host.host_ptr()[1] = kInitVal2; + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_host.ptr(), in_host.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(!(out_host.host_ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a block-wide memory fence on global peer device memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_block.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_block_Positive_Basic_Peer") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("At least 2 devices are required"); + return; + } + + int can_access_peer; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1)); + REQUIRE(can_access_peer); + + HIP_CHECK(hipSetDevice(0)); + + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HIP_CHECK(hipSetDevice(1)); + + HipTest::launchKernel(ThreadfenceTestKernel, 2, 1, 0, nullptr, + out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipSetDevice(0)); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/threadfence/__threadfence_system.cc b/projects/hip-tests/catch/unit/threadfence/__threadfence_system.cc new file mode 100644 index 0000000000..078cf21511 --- /dev/null +++ b/projects/hip-tests/catch/unit/threadfence/__threadfence_system.cc @@ -0,0 +1,126 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +#include "threadfence_common.hh" + +/** + * @addtogroup __threadfence_system __threadfence_system + * @{ + * @ingroup ThreadfenceTest + */ + +__global__ void WriteKernel(int* in) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + + if (tid == 0) { + Write(in); + } +} + +__global__ void ReadKernel(int* out, int* in) { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + + if (tid == 0) { + Read(out, in); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a system-wide memory fence on global peer device memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_system.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_system_Positive_Basic_Peer") { + const auto device_count = HipTest::getDeviceCount(); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("At least 2 devices are required"); + return; + } + + int can_access_peer; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1)); + REQUIRE(can_access_peer); + + HIP_CHECK(hipSetDevice(0)); + + LinearAllocGuard in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + LinearAllocGuard out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int)); + + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[0]), kInitVal1, 1)); + HIP_CHECK(hipMemsetD32(&(in_dev.ptr()[1]), kInitVal2, 1)); + + HipTest::launchKernel(WriteKernel, 1, 1, 0, nullptr, in_dev.ptr()); + + HIP_CHECK(hipSetDevice(1)); + HipTest::launchKernel(ReadKernel, 1, 1, 0, nullptr, out_dev.ptr(), in_dev.ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(out_host.host_ptr(), out_dev.ptr(), 2 * sizeof(int), hipMemcpyDefault)); + + REQUIRE(!(out_host.ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} + +/** + * Test Description + * ------------------------ + * - Basic test for a system-wide memory fence on page-locked host memory. + * + * Test source + * ------------------------ + * - unit/threadfence/__threadfence_system.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit___threadfence_system_Positive_Basic_Host") { + LinearAllocGuard in_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + LinearAllocGuard out_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int)); + + for (int i = 0; i < cmd_options.iterations; ++i) { + in_host.host_ptr()[0] = kInitVal1; + in_host.host_ptr()[1] = kInitVal2; + + HipTest::launchKernel(WriteKernel, 1, 1, 0, nullptr, in_host.host_ptr()); + Read(out_host.host_ptr(), in_host.host_ptr()); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(!(out_host.host_ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2)); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/threadfence/threadfence_common.hh b/projects/hip-tests/catch/unit/threadfence/threadfence_common.hh new file mode 100644 index 0000000000..dc8dca776d --- /dev/null +++ b/projects/hip-tests/catch/unit/threadfence/threadfence_common.hh @@ -0,0 +1,108 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +enum class ThreadfenceScope { kBlock, kDevice, kSystem }; + +template __device__ void Threadfence() { + if constexpr (scope == ThreadfenceScope::kBlock) { + __threadfence_block(); + } else if constexpr (scope == ThreadfenceScope::kDevice) { + __threadfence(); + } else if constexpr (scope == ThreadfenceScope::kSystem) { + __threadfence_system(); + } +} + +static constexpr int kInitVal1 = 1, kInitVal2 = 2; +static constexpr int kSetVal1 = 10, kSetVal2 = 20; + +template __host__ __device__ void Write(volatile int* in) { + in[0] = kSetVal1; +#ifdef __HIP_DEVICE_COMPILE__ + Threadfence(); +#else + std::atomic_thread_fence(std::memory_order_seq_cst); +#endif + in[1] = kSetVal2; +} + +template +__host__ __device__ void Read(volatile int* out, volatile int* in) { + out[1] = in[1]; +#ifdef __HIP_DEVICE_COMPILE__ + Threadfence(); +#else + std::atomic_thread_fence(std::memory_order_seq_cst); +#endif + out[0] = in[0]; +} + +template +__device__ void ThreadfenceTest(int* out, int* in) { + if constexpr (scope == ThreadfenceScope::kBlock || use_shared_mem) { + if (threadIdx.x == 0 && blockIdx.x == 0) { + Write(in); + } else if (threadIdx.x == 1 && blockIdx.x == 0) { + Read(out, in); + } + } else if constexpr (scope == ThreadfenceScope::kDevice) { + if (threadIdx.x == 0 && blockIdx.x == 0) { + Write(in); + } else if (threadIdx.x == 0 && blockIdx.x == 1) { + Read(out, in); + } + } +} + +template +__global__ void ThreadfenceTestKernel(int* out, int* in) { + extern __shared__ int shared_mem[]; + + int tid = blockIdx.x * blockDim.x + threadIdx.x; + + int *out_mem = out, *in_mem = in; + + if constexpr (use_shared_mem) { + if (tid == 0) { + in_mem = &shared_mem[0]; + out_mem = &shared_mem[2]; + + in_mem[0] = in[0]; + in_mem[1] = in[1]; + } + + __syncthreads(); + } + + ThreadfenceTest(out_mem, in_mem); + + if constexpr (use_shared_mem) { + __syncthreads(); + + if (tid == 0) { + out[0] = out_mem[0]; + out[1] = out_mem[1]; + } + } +} \ No newline at end of file