From 96d4493c3ad62bd332d9424a52d9afa397de71cd Mon Sep 17 00:00:00 2001 From: Rambabu Swargam Date: Fri, 9 Aug 2024 22:22:16 +0530 Subject: [PATCH] SWDEV-468261 - [catch2][dtest] Functional testcases for hipStreamLegacy Change-Id: I37cf900d4cfdc244ba8d1d7e2e8c6d34f535a0c2 --- catch/hipTestMain/config/config_amd_linux | 2 + catch/hipTestMain/config/config_amd_windows | 17 + catch/unit/stream/CMakeLists.txt | 12 + catch/unit/stream/hipStreamLegacy_Ext.cc | 858 ++++++++++++++++++ .../hipStreamLegacy_compiler_options.cc | 152 ++++ catch/unit/stream/hipStreamLegacy_exe.cc | 105 +++ 6 files changed, 1146 insertions(+) create mode 100644 catch/unit/stream/hipStreamLegacy_Ext.cc create mode 100644 catch/unit/stream/hipStreamLegacy_compiler_options.cc create mode 100644 catch/unit/stream/hipStreamLegacy_exe.cc diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 70c6f42dd2..dd8394b552 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -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 ===", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 05491eadba..11879bfa7a 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -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" diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index 751d565765..0e98be542c 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -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) diff --git a/catch/unit/stream/hipStreamLegacy_Ext.cc b/catch/unit/stream/hipStreamLegacy_Ext.cc new file mode 100644 index 0000000000..9b7b1d5542 --- /dev/null +++ b/catch/unit/stream/hipStreamLegacy_Ext.cc @@ -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 +#include +#include +#include + +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 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)); +} diff --git a/catch/unit/stream/hipStreamLegacy_compiler_options.cc b/catch/unit/stream/hipStreamLegacy_compiler_options.cc new file mode 100644 index 0000000000..d324d7c727 --- /dev/null +++ b/catch/unit/stream/hipStreamLegacy_compiler_options.cc @@ -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 +#include +#include + +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)); +} diff --git a/catch/unit/stream/hipStreamLegacy_exe.cc b/catch/unit/stream/hipStreamLegacy_exe.cc new file mode 100644 index 0000000000..6e6d81b791 --- /dev/null +++ b/catch/unit/stream/hipStreamLegacy_exe.cc @@ -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 +#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; +}