EXSWHTEC-302 - Implement tests for memory fence device functions #276
Change-Id: I327527288f90011f09262708dd6372a7c6fc4708
[ROCm/hip-tests commit: 8ef0d724f5]
This commit is contained in:
Mirza Halilčević
committed by
Rakesh Roy
Rakesh Roy
parent
c0f92f2477
commit
70cf4c9dac
@@ -116,6 +116,13 @@ THE SOFTWARE.
|
||||
* @}
|
||||
*/
|
||||
|
||||
/**
|
||||
* @defgroup ThreadfenceTest Memory Fence Functions
|
||||
* @{
|
||||
* This section describes tests for Memory Fence Functions.
|
||||
* @}
|
||||
*/
|
||||
|
||||
/**
|
||||
* @defgroup MemoryTest memory Management APIs
|
||||
* @{
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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)
|
||||
@@ -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 <cmd_options.hh>
|
||||
#include <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
#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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kDevice, true>, 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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kDevice, false>, 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<int> in_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kDevice, false>, 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<int> in_host(LinearAllocs::hipMallocManaged, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kDevice, false>, 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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kDevice, false>, 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));
|
||||
}
|
||||
}
|
||||
@@ -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 <cmd_options.hh>
|
||||
#include <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
#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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kBlock, true>, 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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kBlock, false>, 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<int> in_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kBlock, false>, 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<int> in_host(LinearAllocs::hipMallocManaged, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kBlock, false>, 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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kBlock, false>, 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));
|
||||
}
|
||||
}
|
||||
@@ -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 <cmd_options.hh>
|
||||
#include <hip_test_common.hh>
|
||||
#include <resource_guards.hh>
|
||||
|
||||
#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<ThreadfenceScope::kSystem>(in);
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void ReadKernel(int* out, int* in) {
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
if (tid == 0) {
|
||||
Read<ThreadfenceScope::kSystem>(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<int> in_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> out_dev(LinearAllocs::hipMalloc, 2 * sizeof(int));
|
||||
|
||||
LinearAllocGuard<int> 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<int> in_host(LinearAllocs::hipHostMalloc, 2 * sizeof(int));
|
||||
LinearAllocGuard<int> 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<ThreadfenceScope::kDevice>(out_host.host_ptr(), in_host.host_ptr());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
REQUIRE(!(out_host.host_ptr()[0] == kInitVal1 && out_host.ptr()[1] == kSetVal2));
|
||||
}
|
||||
}
|
||||
@@ -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 <ThreadfenceScope scope> __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 <ThreadfenceScope scope> __host__ __device__ void Write(volatile int* in) {
|
||||
in[0] = kSetVal1;
|
||||
#ifdef __HIP_DEVICE_COMPILE__
|
||||
Threadfence<scope>();
|
||||
#else
|
||||
std::atomic_thread_fence(std::memory_order_seq_cst);
|
||||
#endif
|
||||
in[1] = kSetVal2;
|
||||
}
|
||||
|
||||
template <ThreadfenceScope scope>
|
||||
__host__ __device__ void Read(volatile int* out, volatile int* in) {
|
||||
out[1] = in[1];
|
||||
#ifdef __HIP_DEVICE_COMPILE__
|
||||
Threadfence<scope>();
|
||||
#else
|
||||
std::atomic_thread_fence(std::memory_order_seq_cst);
|
||||
#endif
|
||||
out[0] = in[0];
|
||||
}
|
||||
|
||||
template <ThreadfenceScope scope, bool use_shared_mem>
|
||||
__device__ void ThreadfenceTest(int* out, int* in) {
|
||||
if constexpr (scope == ThreadfenceScope::kBlock || use_shared_mem) {
|
||||
if (threadIdx.x == 0 && blockIdx.x == 0) {
|
||||
Write<scope>(in);
|
||||
} else if (threadIdx.x == 1 && blockIdx.x == 0) {
|
||||
Read<scope>(out, in);
|
||||
}
|
||||
} else if constexpr (scope == ThreadfenceScope::kDevice) {
|
||||
if (threadIdx.x == 0 && blockIdx.x == 0) {
|
||||
Write<scope>(in);
|
||||
} else if (threadIdx.x == 0 && blockIdx.x == 1) {
|
||||
Read<scope>(out, in);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <ThreadfenceScope scope, bool use_shared_mem>
|
||||
__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<scope, use_shared_mem>(out_mem, in_mem);
|
||||
|
||||
if constexpr (use_shared_mem) {
|
||||
__syncthreads();
|
||||
|
||||
if (tid == 0) {
|
||||
out[0] = out_mem[0];
|
||||
out[1] = out_mem[1];
|
||||
}
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user