AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

SWDEV-468261 - [catch2][dtest] Functional testcases for hipStreamLegacy

Sfoglia il codice sorgente 
Change-Id: I37cf900d4cfdc244ba8d1d7e2e8c6d34f535a0c2
This commit is contained in:
Rambabu Swargam
2024-08-09 22:22:16 +05:30
parent fb5e1d33d9
commit 96d4493c3a
6 ha cambiato i file con 1146 aggiunte e 0 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
catch/hipTestMain/config/config_amd_linux
+2
Vedi File
@@ -957,6 +957,8 @@
"Unit_hipDeviceGetGraphMemAttribute_Functional_Multi_Device",
"=== SWDEV-457316 Below tests are disabled temporarily to avoid combined PSDB ===",
"Unit_hipGraphAddMemFreeNode_Negative_NotSupported",
"=== SWDEV-478776 Below test is disabled due to fix is not avaialable in mainline ===",
"Unit_hipStreamLegacy_hipStreamSynchronize",
#endif
#if defined gfx908
"=== Below test soft hang in stress test on 29/08/23 ===",
catch/hipTestMain/config/config_amd_windows
+17
Vedi File
@@ -1150,6 +1150,23 @@
"=== SWDEV-475482 - Disable tests to merge clr change",
"Unit_hipCreateTextureObject_LinearResource",
"Unit_hipCreateTextureObject_Pitch2DResource",
"=== SWDEV-478745 Below tests are disabled due latest code not available in windows PSDB ===",
"Unit_hipStreamLegacy_WithBlockingStream",
"Unit_hipStreamLegacy_MultipleThreads",
"Unit_hipStreamLegacy_NegetiveCase",
"Unit_hipStreamLegacy_WithNonBlockingStream",
"Unit_hipStreamLegacy_WithStreamPerThread",
"Unit_hipStreamLegacy_MultiDevice",
"Unit_hipStreamLegacy_H2H_H2D_D2D_D2H_Default",
"Unit_hipStreamLegacy_MultiDeviceMultiOperation",
"Unit_hipStreamLegacy_TwoThreadsEachOneDiffOperation",
"Unit_hipStreamLegacy_TwoDevicesEachOneDiffOperation",
"Unit_hipStreamLegacy_TwoThreadsInTwoDevicesEachOneDiffOperation",
"Unit_hipStreamLegacy_InChildProcess",
"Unit_hipStreamLegacy_WithKernel",
"Unit_hipStreamLegacy_hipStreamSynchronize",
"Unit_hipStreamLegacy_WithSptCompilerOption",
"Unit_hipStreamLegacy_TwoThreadsDiffOperationWithSptCompOption",
"========================================================================================",
#endif
"End of json"
catch/unit/stream/CMakeLists.txt
+12
Vedi File
@@ -18,6 +18,8 @@ set(TEST_SRC
hipStreamACb_StrmSyncTiming.cc
hipLaunchHostFunc.cc
hipStreamGetDevice.cc
hipStreamLegacy_Ext.cc
hipStreamLegacy_compiler_options.cc
)
if(HIP_PLATFORM MATCHES "amd")
@@ -35,7 +37,17 @@ else()
# COMPILE_FLAGS -std=c++17)
endif()
if(HIP_PLATFORM MATCHES "amd")
set_source_files_properties(hipStreamLegacy_compiler_options.cc PROPERTIES COMPILE_OPTIONS "-fgpu-default-stream=per-thread")
endif()
if(HIP_PLATFORM MATCHES "nvidia")
set_source_files_properties(hipStreamLegacy_compiler_options.cc PROPERTIES COMPILE_OPTIONS "--default-stream=per-thread")
endif()
hip_add_exe_to_target(
NAME StreamTest
TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests
COMPILE_OPTIONS -std=c++17 COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
add_executable(hipStreamLegacy_exe EXCLUDE_FROM_ALL hipStreamLegacy_exe.cc)
add_dependencies(StreamTest hipStreamLegacy_exe)
catch/unit/stream/hipStreamLegacy_Ext.cc
+858
Vedi File
@@ -0,0 +1,858 @@
/*
Copyright (c) 2024 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_helper.hh>
#include <utils.hh>
#include <hip_test_process.hh>
static constexpr int N = 2 * 1024 * 1024;
static constexpr size_t NBYTES = N * sizeof(int);
/**
* Local Function to fill the array with given value
*/
static void fillHostArray(int *arr, int size, int value) {
for ( int i = 0; i < size; i++ ) {
arr[i] = value;
}
}
/**
* Kernel to fill the array with given value
*/
static __global__ void fillArray(int *arr, int size, int value) {
for ( int i = 0; i < size; i++ ) {
arr[i] = value;
}
}
/**
* Local Function to fill the device array with given value
*/
static void fillDeviceArray(int *arr, int size, int value) {
fillArray<<<1, 1>>>(arr, size, value);
}
/**
* In addOneKernel function, all elements of the array a increased by 1
*/
static __global__ void addOneKernel(int *a, int size) {
int offset = blockDim.x * blockIdx.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for ( int i = offset; i < size; i+=stride ) {
a[i] += 1;
}
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario:-
* - Do Memory copy Asynchronously H2D in user defined blocking stream,
* - do Memory copy Asynchronously D2H in Legacy stream.
* - The task 2 which is in legacy stream should wait till the
* - task 1 completes its execution which is in userdefined stream.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_WithBlockingStream") {
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 1);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
fillDeviceArray(devArr, N, 2);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 3);
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 1 \n");
REQUIRE(hostArrDst[i] == 1);
}
HIP_CHECK(hipStreamDestroy(stream));
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
/**
* Local Function to perform the below operations:-
* Do Memory copy Asynchronously H2D in Legacy stream
* and then do Memory copy Asynchronously D2H in Legacy stream.
* Task 2 should wait till the execution of task 1.
*/
void launchFunction(hipStream_t stream) {
int *hostArrSrc = new int[N];
REQUIRE_THREAD(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 5);
int *devArr = nullptr;
HIP_CHECK_THREAD(hipMalloc(&devArr, NBYTES));
REQUIRE_THREAD(devArr != nullptr);
fillDeviceArray(devArr, N, 6);
int *hostArrDst = new int[N];
REQUIRE_THREAD(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 7);
HIP_CHECK_THREAD(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, stream));
HIP_CHECK_THREAD(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, stream));
for ( int i = 0; i < N; i++ ) {
if ( hostArrDst[i] != 5 ) {
std::cout << "At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 5 \n" << std::endl;
REQUIRE_THREAD(false);
}
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK_THREAD(hipFree(devArr));
}
/**
* Test Description
* ------------------------
* - This test case tests hipStreamLegacy in multi threaded scenario:-
* - All threads should launch successfully and run independently
* - and uses the same legacy stream.
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_MultipleThreads") {
const unsigned int threadsSupported = std::thread::hardware_concurrency();
const int numberOfThreads = (threadsSupported >= 10) ? 10 : threadsSupported;
std::vector<std::thread> threads;
for (int t = 0; t < numberOfThreads; t++) {
threads.push_back(std::thread(launchFunction, hipStreamLegacy));
}
for (int t = 0; (t < numberOfThreads) && (t < threads.size()); t++) {
threads[t].join();
}
HIP_CHECK_THREAD_FINALIZE();
}
/**
* Test Description
* ------------------------
* - Pass the hipStreamLegacy to hipStreamBeginCapture() api
* - and the api should return hipErrorStreamCaptureUnsupported.
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_NegetiveCase") {
hipStream_t stream = hipStreamLegacy;
REQUIRE(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)
== hipErrorStreamCaptureUnsupported);
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario:-
* - Do Memory copy Asynchronously H2D in user defined Non-blocking stream
* - and do Memory copy Asynchronously D2H in Legacy stream.
* - The task 2 which is in legacy stream should not wait till the
* - task 1 completes its execution, two streams should run concurrently.
* - And the host thread should wait for the task 2 which is in legacy stream.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_WithNonBlockingStream") {
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 10);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
fillDeviceArray(devArr, N, 11);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 12);
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 10 or 11 \n");
REQUIRE(((hostArrDst[i] == 10) || (hostArrDst[i] == 11)));
}
HIP_CHECK(hipStreamDestroy(stream));
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario:-
* - Do Memory copy Asynchronously H2D using hipStreamPerThread
* - and do Memory copy Asynchronously D2H using hipStreamLegacy.
* - The task 2 which is in legacy stream should wait till the
* - task 1 completes its execution which is in hipStreamPerThread.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_WithStreamPerThread") {
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 15);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
fillDeviceArray(devArr, N, 16);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, hipStreamPerThread));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 15 \n");
REQUIRE(hostArrDst[i] == 15);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario with all available devices :-
* - Do Memory copy Asynchronously H2D in Legacy stream
* - and do Memory copy Asynchronously D2H in Legacy stream.
* - The task 2 which is in legacy stream should wait till the
* - task 1 completes its execution, in all the devices.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_MultiDevice") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
HipTest::HIP_SKIP_TEST("Skipping because this machine has total GPUs < 2");
return;
}
for (int deviceId = 0; deviceId < deviceCount; deviceId++) {
HIP_CHECK(hipSetDevice(deviceId));
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 20);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
fillDeviceArray(devArr, N, 21);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 22);
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 20 "
" For deviceId : " << deviceId << "\n");
REQUIRE(hostArrDst[i] == 20);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
}
/**
* Test Description
* ------------------------
* - This test case testing the hipStreamLegacy with hipMemcpyAsync
* - in all the ways like H2H, H2D, D2D, D2H, also with the hipMemcpyDefault.
* - All the operations should success.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_H2H_H2D_D2D_D2H_Default") {
int *hostArr1 = new int[N];
REQUIRE(hostArr1 != nullptr);
fillHostArray(hostArr1, N, 30);
int *hostArr2 = new int[N];
REQUIRE(hostArr2 != nullptr);
fillHostArray(hostArr2, N, 31);
int *devArr1 = nullptr;
HIP_CHECK(hipMalloc(&devArr1, NBYTES));
REQUIRE(devArr1 != nullptr);
fillDeviceArray(devArr1, N, 32);
int *devArr2 = nullptr;
HIP_CHECK(hipMalloc(&devArr2, NBYTES));
REQUIRE(devArr2 != nullptr);
fillDeviceArray(devArr2, N, 33);
int *hostArr3 = new int[N];
REQUIRE(hostArr3 != nullptr);
fillHostArray(hostArr3, N, 34);
int *hostArr4 = new int[N];
REQUIRE(hostArr4 != nullptr);
fillHostArray(hostArr4, N, 35);
HIP_CHECK(hipMemcpyAsync(hostArr2, hostArr1, NBYTES,
hipMemcpyHostToHost, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(devArr1, hostArr2, NBYTES,
hipMemcpyHostToDevice, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(devArr2, devArr1, NBYTES,
hipMemcpyDeviceToDevice, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(hostArr3, devArr2, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(hostArr4, hostArr3, NBYTES,
hipMemcpyDefault, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArr4[i] <<
" Expected value : 30 \n");
REQUIRE(hostArr4[i] == 30);
}
delete[] hostArr1;
delete[] hostArr2;
HIP_CHECK(hipFree(devArr1));
HIP_CHECK(hipFree(devArr2));
delete[] hostArr3;
delete[] hostArr4;
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario with two devices:-
* - h1Dev0 -> d1Dev0 : memcpy from host to device 0
* - d1Dev0 -> d1Dev1 : memcpy from device 0 to device 1
* - d1Dev1 -> d2Dev1 : memcpy from device to device with in device 1
* - d2Dev1 -> d2Dev0 : memcpy from device 1 to device 0
* - d2Dev0 -> h2Dev0 : memcpy from device 0 to host
* - The opeations in device 0 and device 1 should run concurrently.
* - Multiple opeations in multiple devices should run without
* - any conflicts.
* - And the final host array should have value other than 45,
* - since host thread waits.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_MultiDeviceMultiOperation") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
HipTest::HIP_SKIP_TEST("Skipping because this machine has total GPUs < 2");
return;
}
int currentDevice = 0;
int peerDevice = 1;
// Set arrays in device 0
HIP_CHECK(hipSetDevice(currentDevice));
int *h1Dev0 = new int[N];
REQUIRE(h1Dev0 != nullptr);
fillHostArray(h1Dev0, N, 40);
int *d1Dev0 = nullptr;
HIP_CHECK(hipMalloc(&d1Dev0, NBYTES));
REQUIRE(d1Dev0 != nullptr);
fillDeviceArray(d1Dev0, N, 41);
// Set arrays in device 1
HIP_CHECK(hipSetDevice(peerDevice));
int *d1Dev1 = nullptr;
HIP_CHECK(hipMalloc(&d1Dev1, NBYTES));
REQUIRE(d1Dev1 != nullptr);
fillDeviceArray(d1Dev1, N, 42);
int *d2Dev1 = nullptr;
HIP_CHECK(hipMalloc(&d2Dev1, NBYTES));
REQUIRE(d2Dev1 != nullptr);
fillDeviceArray(d2Dev1, N, 43);
// Set destination arrays in device 0
HIP_CHECK(hipSetDevice(currentDevice));
int *d2Dev0 = nullptr;
HIP_CHECK(hipMalloc(&d2Dev0, NBYTES));
REQUIRE(d2Dev0 != nullptr);
fillDeviceArray(d2Dev0, N, 44);
int *h2Dev0 = new int[N];
REQUIRE(h2Dev0 != nullptr);
fillHostArray(h2Dev0, N, 45);
// Do operations in current device
HIP_CHECK(hipSetDevice(currentDevice));
HIP_CHECK(hipMemcpyAsync(d1Dev0, h1Dev0, NBYTES,
hipMemcpyHostToHost, hipStreamLegacy));
// Copy from current device to peer device
HIP_CHECK(hipMemcpyPeerAsync(d1Dev1, peerDevice, // des
d1Dev0, currentDevice, // src
NBYTES, hipStreamLegacy));
// Do operations in peer device
HIP_CHECK(hipSetDevice(peerDevice));
HIP_CHECK(hipMemcpyAsync(d2Dev1, d1Dev1, NBYTES,
hipMemcpyDeviceToDevice, hipStreamLegacy));
// Copy from peer device to current device
HIP_CHECK(hipMemcpyPeerAsync(d2Dev0, currentDevice, // des
d2Dev1, peerDevice, // src
NBYTES, hipStreamLegacy));
// Finally copy daat to hostArr4
HIP_CHECK(hipSetDevice(currentDevice));
HIP_CHECK(hipMemcpyAsync(h2Dev0, d2Dev0, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << h2Dev0[i] <<
" Expected value : 40/41/42/43/44 \n");
REQUIRE(h2Dev0[i] != 45);
}
HIP_CHECK(hipSetDevice(currentDevice));
delete[] h1Dev0;
delete[] h2Dev0;
HIP_CHECK(hipFree(d1Dev0));
HIP_CHECK(hipFree(d2Dev0));
HIP_CHECK(hipSetDevice(peerDevice));
HIP_CHECK(hipFree(d1Dev1));
HIP_CHECK(hipFree(d2Dev1));
HIP_CHECK(hipSetDevice(currentDevice));
}
/*
* Local helper function to copy data from host to device
*/
static void copyFromHostToDevice(int* hostArr, int* devArr) {
HIP_CHECK(hipMemcpyAsync(devArr, hostArr, NBYTES,
hipMemcpyHostToDevice,
hipStreamLegacy));
}
/*
* Local helper function to copy data from device to host
*/
static void copyFromDeviceToHost(int* devArr, int* hostArr) {
HIP_CHECK(hipMemcpyAsync(hostArr, devArr, NBYTES,
hipMemcpyDeviceToHost,
hipStreamLegacy));
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario:-
* - Launch two threads,
* - In thread 1 : H -> D Copy
* - In thread 2 : D -> H Copy
* - These two thredas should run sequentially as they are using one stream.
* - Note : Joined first thread before launching second thread just to avoid,
* - the scenario of second thread can be launched before first
* - thread while scheduling.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_TwoThreadsEachOneDiffOperation") {
const unsigned int threadsSupported = std::thread::hardware_concurrency();
if (threadsSupported < 2) {
HipTest::HIP_SKIP_TEST("Skipping due to machine does't " \
"support two concurrent threads");
return;
}
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 50);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
fillDeviceArray(devArr, N, 51);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 52);
std::thread H2D_Thread(copyFromHostToDevice, hostArrSrc, devArr);
H2D_Thread.join();
std::thread D2H_Thread(copyFromDeviceToHost, devArr, hostArrDst);
D2H_Thread.join();
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 50 \n");
REQUIRE(hostArrDst[i] == 50);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
/**
* Test Description
* ------------------------
* - This test case tests the following scenario with two devices:-
* - devArrDev0 -> devArrDev1 : memcpy from device 0 to device 1
* - devArrDev1 -> hostArrDst : memcpy from device 1 to host
* - The opeations in device 0 and device 1 should run concurrently.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_TwoDevicesEachOneDiffOperation") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
HipTest::HIP_SKIP_TEST("Skipping because this machine has total GPUs < 2");
return;
}
// Set arrays in device 0
HIP_CHECK(hipSetDevice(0));
int *devArrDev0 = nullptr;
HIP_CHECK(hipMalloc(&devArrDev0, NBYTES));
REQUIRE(devArrDev0 != nullptr);
fillDeviceArray(devArrDev0, N, 500);
// Set arrays in device 1
HIP_CHECK(hipSetDevice(1));
int *devArrDev1 = nullptr;
HIP_CHECK(hipMalloc(&devArrDev1, NBYTES));
REQUIRE(devArrDev1 != nullptr);
fillDeviceArray(devArrDev1, N, 501);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 502);
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipMemcpyPeerAsync(devArrDev1, 1, // des
devArrDev0, 0, // src
NBYTES,
hipStreamLegacy));
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArrDev1, NBYTES,
hipMemcpyDeviceToHost,
hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 500 or 501 \n");
REQUIRE(((hostArrDst[i] == 500) || (hostArrDst[i] == 501)));
}
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipFree(devArrDev1));
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipFree(devArrDev0));
delete[] hostArrDst;
}
/*
* Local helper function to copy data from device 0 to device 1
*/
static void operationsInDev0(int* devArrDev0, int* devArrDev1) {
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipMemcpyPeerAsync(devArrDev1, 1, // des
devArrDev0, 0, // src
NBYTES,
hipStreamLegacy));
}
/*
* Local helper function to copy data from device to host
*/
static void operationsInDev1(int* devArrDev1, int* hostArrDst) {
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArrDev1, NBYTES,
hipMemcpyDeviceToHost,
hipStreamLegacy));
}
/**
* Test Description
* ------------------------
* - This test case tests the below scenario with two devices in two threads:-
* - In thread 1 Dev 0 : devArrDev0 -> devArrDev1 : memcpy from dev 0 to dev 1
* - In thread 2 Dev 1 : devArrDev1 -> hostArrDst : memcpy from dev 1 to host
* - The opeations in device 0 and device 1, thread 1 and thread 2
* - should run concurrently.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_TwoThreadsInTwoDevicesEachOneDiffOperation") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
HipTest::HIP_SKIP_TEST("Skipping because this machine has total GPUs < 2");
return;
}
// Set arrays in device 0
HIP_CHECK(hipSetDevice(0));
int *devArrDev0 = nullptr;
HIP_CHECK(hipMalloc(&devArrDev0, NBYTES));
REQUIRE(devArrDev0 != nullptr);
fillDeviceArray(devArrDev0, N, 999);
// Set arrays in device 1
HIP_CHECK(hipSetDevice(1));
int *devArrDev1 = nullptr;
HIP_CHECK(hipMalloc(&devArrDev1, NBYTES));
REQUIRE(devArrDev1 != nullptr);
fillDeviceArray(devArrDev1, N, 888);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 777);
HIP_CHECK(hipSetDevice(0));
std::thread dev0Thread(operationsInDev0, devArrDev0, devArrDev1);
std::thread dev1Thread(operationsInDev1, devArrDev1, hostArrDst);
dev0Thread.join();
dev1Thread.join();
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 999 or 888 \n");
REQUIRE(((hostArrDst[i] == 999) || (hostArrDst[i] == 888)));
}
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipFree(devArrDev1));
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipFree(devArrDev0));
delete[] hostArrDst;
}
/**
* Test Description
* ------------------------
* - This test creates the child process and use the hipStreamLegacy flag
* - in the child process. The hipStreamLegacy flag should work properly
* - in child process.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_InChildProcess") {
hip::SpawnProc proc("hipStreamLegacy_exe", true);
REQUIRE(proc.run() == 0);
}
/**
* Test Description
* ------------------------
* - This test case, tests the hipStreamLegacy with Kernel.
* - Kernel launch should success and should give proper result.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_WithKernel") {
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 1);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 6);
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, hipStreamLegacy));
addOneKernel<<<1, 1, 0, hipStreamLegacy>>>(devArr, N);
addOneKernel<<<1, 1, 0, hipStreamLegacy>>>(devArr, N);
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 3 \n");
REQUIRE(hostArrDst[i] == 3);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
/**
* Test Description
* ------------------------
* - This test case, tests the hipStreamSynchronize with
* - the hipStreamLegacy. It should work without any conflicts.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_Ext.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_hipStreamSynchronize") {
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillHostArray(hostArrSrc, N, 1);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillHostArray(hostArrDst, N, 3);
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice,
hipStreamLegacy));
HIP_CHECK(hipStreamSynchronize(hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost,
hipStreamLegacy));
HIP_CHECK(hipStreamSynchronize(hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 1 \n");
REQUIRE(hostArrDst[i] == 1);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
catch/unit/stream/hipStreamLegacy_compiler_options.cc
+152
Vedi File
@@ -0,0 +1,152 @@
/*
Copyright (c) 2024 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_helper.hh>
#include <utils.hh>
static constexpr unsigned int N = 2 * 1024 * 1024;
static constexpr size_t NBYTES = N * sizeof(int);
/**
* Local Function to fill the array with given value
*/
static void fillArr(int *arr, unsigned int size, int value) {
for ( int i = 0; i < size; i++ ) {
arr[i] = value;
}
}
/**
* Test Description
* ------------------------
* - This test, tests the following scenario with per thread compiler option:-
* - Do Memory copy Asynchronously H2D in Legacy stream
* - and do Memory copy Asynchronously D2H in Legacy stream.
* - The task 2 should wait till the execution of task 1.
* - Even it is compiled with stream per thread, since one stream for both
* - these should be in sync.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_compiler_options.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_WithSptCompilerOption") {
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillArr(hostArrSrc, N, 1);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillArr(hostArrDst, N, 3);
HIP_CHECK(hipMemcpyAsync(devArr, hostArrSrc, NBYTES,
hipMemcpyHostToDevice, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(hostArrDst, devArr, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 1 \n");
REQUIRE(hostArrDst[i] == 1);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
/*
* Local helper function to copy data from host to device
*/
static void copyHostToDevice(int* hostArr, int* devArr) {
HIP_CHECK(hipMemcpyAsync(devArr, hostArr, NBYTES,
hipMemcpyHostToDevice,
hipStreamLegacy));
}
/*
* Local helper function to copy data from device to host
*/
static void copyDeviceToHost(int* devArr, int* hostArr) {
HIP_CHECK(hipMemcpyAsync(hostArr, devArr, NBYTES,
hipMemcpyDeviceToHost,
hipStreamLegacy));
}
/**
* Test Description
* ------------------------
* - This test, tests the following scenario with per thread compiler option:-
* - Launch two threads,
* - In thread 1 : H -> D Copy
* - In thread 2 : D -> H Copy
* - These two threads should run successfully and update the data.
* - Note : Joined first thread before launching second thread to avoid,
* - the scenario of second thread can be launched before first
* - thread while scheduling.
* Test source
* ------------------------
* - unit/stream/hipStreamLegacy_compiler_options.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.3
*/
TEST_CASE("Unit_hipStreamLegacy_TwoThreadsDiffOperationWithSptCompOption") {
const unsigned int threadsSupported = std::thread::hardware_concurrency();
if (threadsSupported < 2) {
HipTest::HIP_SKIP_TEST("Skipping due to machine does't " \
"support two concurrent threads");
return;
}
int *hostArrSrc = new int[N];
REQUIRE(hostArrSrc != nullptr);
fillArr(hostArrSrc, N, 50);
int *devArr = nullptr;
HIP_CHECK(hipMalloc(&devArr, NBYTES));
REQUIRE(devArr != nullptr);
int *hostArrDst = new int[N];
REQUIRE(hostArrDst != nullptr);
fillArr(hostArrDst, N, 52);
std::thread H2D_Thread(copyHostToDevice, hostArrSrc, devArr);
H2D_Thread.join();
std::thread D2H_Thread(copyDeviceToHost, devArr, hostArrDst);
D2H_Thread.join();
for ( int i = 0; i < N; i++ ) {
INFO("At index : " << i << " Got value : " << hostArrDst[i] <<
" Expected value : 50 \n");
REQUIRE(hostArrDst[i] == 50);
}
delete[] hostArrSrc;
delete[] hostArrDst;
HIP_CHECK(hipFree(devArr));
}
catch/unit/stream/hipStreamLegacy_exe.cc
+105
Vedi File
@@ -0,0 +1,105 @@
/*
Copyright (c) 2024 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <iostream>
#include "hip/hip_runtime.h"
static constexpr int N = 2 * 1024 * 1024;
static constexpr size_t NBYTES = N * sizeof(int);
#define HIP_CHECK(error) \
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && \
(localError != hipErrorPeerAccessAlreadyEnabled)) { \
printf("error: '%s'(%d) from %s at %s:%d\n", \
hipGetErrorString(localError), \
localError, #error, __FUNCTION__, __LINE__);\
return -1; \
} \
}
/**
* Used hipStreamLegacy in child process and copies H2H, H2D, D2D, D2H
*/
int main() {
constexpr int elementVal = 3456;
int *hostArr1 = new int[N];
if (hostArr1 == nullptr) {
std::cout << "Invalid hostArr1" << std::endl;
return -1;
}
for ( int i = 0; i < N; i++ ) {
hostArr1[i] = elementVal;
}
int *hostArr2 = new int[N];
if (hostArr2 == nullptr) {
std::cout << "Invalid hostArr2" << std::endl;
return -1;
}
int *devArr1 = nullptr;
HIP_CHECK(hipMalloc(&devArr1, NBYTES));
if (devArr1 == nullptr) {
std::cout << "Invalid devArr1" << std::endl;
return -1;
}
int *devArr2 = nullptr;
HIP_CHECK(hipMalloc(&devArr2, NBYTES));
if (devArr2 == nullptr) {
std::cout << "Invalid devArr2" << std::endl;
return -1;
}
int *hostArr3 = new int[N];
if (hostArr3 == nullptr) {
std::cout << "Invalid hostArr3" << std::endl;
return -1;
}
HIP_CHECK(hipMemcpyAsync(hostArr2, hostArr1, NBYTES,
hipMemcpyHostToHost, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(devArr1, hostArr2, NBYTES,
hipMemcpyHostToDevice, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(devArr2, devArr1, NBYTES,
hipMemcpyDeviceToDevice, hipStreamLegacy));
HIP_CHECK(hipMemcpyAsync(hostArr3, devArr2, NBYTES,
hipMemcpyDeviceToHost, hipStreamLegacy));
for (int i = 0; i < N; i++) {
if (hostArr3[i] != elementVal) {
std::cout << "At index : = " << i <<
" Got value : " << hostArr3[i] <<
" Expected value : " << elementVal << std::endl;
return -1;
}
}
delete[] hostArr1;
delete[] hostArr2;
HIP_CHECK(hipFree(devArr1));
HIP_CHECK(hipFree(devArr2));
delete[] hostArr3;
return 0;
}