From 5698fed1620570a64dec86c194ec4c8fcfdd0c70 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 14:03:31 +0530 Subject: [PATCH 1/9] SWDEV-327563 - deviceLib build fix for nvidia (#3072) Change-Id: Ic66242caa21f9ba6f41255d84a3fd3e54b481b46 --- tests/catch/unit/deviceLib/CMakeLists.txt | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 7977c7b4de..7355023d09 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -51,10 +51,10 @@ set(AMD_ARCH_SPEC_TEST_SRC # Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906" # having space at the start/end of OFFLOAD_ARCH_STR can cause build failures if(UNIX) - add_custom_target(kerDevAllocMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - add_custom_target(kerDevWriteMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevWriteMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevWriteMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - add_custom_target(kerDevFreeMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevFreeMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevFreeMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - add_custom_target(kerDevAllocSingleKer.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocSingleKer.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocSingleKer.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevAllocMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocMultCO.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevWriteMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevWriteMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevWriteMultCO.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevFreeMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevFreeMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevFreeMultCO.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevAllocSingleKer.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocSingleKer.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocSingleKer.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) endif() if(HIP_PLATFORM MATCHES "amd") From d8efb5fd4d6f1fc3c4d75e1ff4a8bd058b612ec9 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 14:03:51 +0530 Subject: [PATCH 2/9] SWDEV-348684 - Fixed catch-2 test failure (#3073) Change-Id: Ib21c26a3c8e51f0b7f132ecd4ad209dfef1d0689 --- .../streamperthread/hipStreamPerThrdTsts.cc | 20 ++++++++++++------- 1 file changed, 13 insertions(+), 7 deletions(-) diff --git a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index 95f5c27fd5..dec35e100c 100644 --- a/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/tests/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -47,9 +47,7 @@ THE SOFTWARE. #endif #include -#ifdef HT_AMD - #include "hip/hip_cooperative_groups.h" -#endif +#include "hip/hip_cooperative_groups.h" using namespace std::chrono; using namespace cooperative_groups; #if HT_AMD @@ -69,8 +67,12 @@ __global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk, int64_t GpuFrq = (Pk_Clk * 1000); int64_t StrtTck = clock64(); if (index == 0) { - // The following while loop checks the value in ptr for around 4 seconds - while ((clock64() - StrtTck) <= (6 * GpuFrq)) { + // The following while loop holds the execution for ~2 seconds. + // Busy sleep on nvidia + while ((clock64() - StrtTck) <= (2 * GpuFrq)) { + #if HT_AMD + __builtin_amdgcn_s_sleep(10); + #endif } if (WaitEvnt == 1) { *Ad1 = 1; @@ -83,8 +85,12 @@ __global__ void StreamPerThrd(int *Ad, int *Ad1, size_t n, int Pk_Clk, __global__ void StreamPerThrd1(int *A, int Pk_Clk) { int64_t GpuFrq = (Pk_Clk * 1000); int64_t StrtTck = clock64(); - // The following while loop checks the value in ptr for around 3-4 seconds - while ((clock64() - StrtTck) <= (3 * GpuFrq)) { + // The following while loop holds the execution for ~1 second + // Busy sleep on nvidia + while ((clock64() - StrtTck) <= (GpuFrq)) { + #if HT_AMD + __builtin_amdgcn_s_sleep(10); + #endif } *A = 1; } From 5e50c01e44f06ca7effe66a9ab5cc38a839b9775 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 14:04:01 +0530 Subject: [PATCH 3/9] SWDEV-367463 - Add unit test for hipMallocManaged call w/ large size (#3089) Change-Id: I9e9f6687ec837f3aaa1567d8b00f15925b7dd2c2 --- tests/catch/unit/memory/hipMallocManaged.cc | 22 +++++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/tests/catch/unit/memory/hipMallocManaged.cc b/tests/catch/unit/memory/hipMallocManaged.cc index 5921ec122e..1b9766ffa9 100644 --- a/tests/catch/unit/memory/hipMallocManaged.cc +++ b/tests/catch/unit/memory/hipMallocManaged.cc @@ -20,8 +20,10 @@ /* Test Case Description: 1) This testcase verifies the hipMallocManaged basic scenario - supported on all devices - 2) This testcase verifies the hipMallocManaged basic scenario - supported - only on HMM enabled devices + 2) This testcase verifies the hipMallocManaged advanced scenario - supported + only on HMM enabled devices + 3) This testcase verifies that hipMallocManaged returns an OutOfMemory error + for allocations much larger than the available memory - supported on all devices */ #include "hipMallocManagedCommon.hh" @@ -138,3 +140,19 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") { HIP_CHECK(hipFree(B)); REQUIRE(maxError != 0.0f); } + +/* + This testcase verifies that hipMallocManaged returns an OutOfMemory error for allocations much + larger than the available memory - supported on all devices + */ +TEST_CASE("Unit_hipMallocManaged_Large") { + auto managed = HmmAttrPrint(); + if (managed != 1) { + WARN( + "GPU doesn't support hipDeviceAttributeManagedMemory attribute so defaulting to system " + "memory."); + } + + float* A; + HIP_CHECK_ERROR(hipMallocManaged(&A, std::numeric_limits::max()), hipErrorOutOfMemory); +} From 92f705e9bd21868dc39175eeff64305026b438c7 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 14:04:14 +0530 Subject: [PATCH 4/9] SWDEV-368417 - Update FAQ on library support (#3092) Change-Id: I940f840be45e6c45b960ec03ae007472e965cb5e --- docs/markdown/hip_faq.md | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index cdf45fb2c3..e4725190c0 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -105,14 +105,15 @@ However, we can provide a rough summary of the features included in each CUDA SD - Cooperative Launch, Surface Object Management, Version Management ### What libraries does HIP support? -HIP includes growing support for the four key math libraries using hcBlas, hcFft, hcrng and hcsparse, as well as MIOpen for machine intelligence applications. +HIP includes growing support for the four key math libraries using hipBlas, hipFFt, hipRAND and hipSPARSE, as well as MIOpen for machine intelligence applications. These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaced with other HIP applications. The hip interfaces support both ROCm and CUDA paths, with familiar library interfaces. - [hipBlas](https://github.com/ROCmSoftwarePlatform/hipBLAS), which utilizes [rocBlas](https://github.com/ROCmSoftwarePlatform/rocBLAS). -- [hipfft](https://github.com/ROCmSoftwarePlatform/hcFFT) -- [hipsparse](https://github.com/ROCmSoftwarePlatform/hcSPARSE) -- [hiprng](https://github.com/ROCmSoftwarePlatform/hcrng) +- [hipFFt](https://github.com/ROCmSoftwarePlatform/hipfft) +- [hipsSPARSE](https://github.com/ROCmSoftwarePlatform/hipsparse) +- [hipRAND](https://github.com/ROCmSoftwarePlatform/hipRAND) +- [MIOpen](https://github.com/ROCmSoftwarePlatform/MIOpen) Additionally, some of the cublas routines are automatically converted to hipblas equivalents by the HIPIFY tools. These APIs use cublas or hcblas depending on the platform and replace the need to use conditional compilation. From 8c791420357ad7633f43ba04e24fb07dd4d83047 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 14:04:25 +0530 Subject: [PATCH 5/9] SWDEV-357759 - [catch2][dtest] Adding additional functional tests for hipStreamBeginCapture() API (#3094) Change-Id: I7b89dcb6aa24eae0daf2974733ddeb6d3902cafa --- .../config/config_amd_windows_MI2xx.json | 4 +- .../config/config_amd_windows_common.json | 6 +- .../catch/unit/graph/hipStreamBeginCapture.cc | 1068 ++++++++++++++++- 3 files changed, 1074 insertions(+), 4 deletions(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json b/tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json index 2843587cf2..89c484b79b 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json +++ b/tests/catch/hipTestMain/config/config_amd_windows_MI2xx.json @@ -90,6 +90,8 @@ "Unit_hipStreamValue_Wait64_Blocking_NoMask_And", "Unit_hipStreamValue_Wait64_Blocking_NoMask_Nor", "Unit_hipDeviceGetPCIBusId_Negative_PartialFill", - "Unit_hipInit_Negative" + "Unit_hipInit_Negative", + "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime", + "Unit_hipStreamBeginCapture_captureComplexGraph" ] } diff --git a/tests/catch/hipTestMain/config/config_amd_windows_common.json b/tests/catch/hipTestMain/config/config_amd_windows_common.json index ef8466de2f..33da571b15 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -85,10 +85,12 @@ "Unit_hipGraphNodeGetDependentNodes_Functional", "Unit_hipGraphNodeGetDependencies_Functional", "Unit_hipGraphExecChildGraphNodeSetParams_ChildTopology", + "Unit_hipGraphAddEventRecordNode_MultipleRun", + "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime", + "Unit_hipStreamBeginCapture_captureComplexGraph", "Note: needs to be enabled when streamPerThread issues are fixed", "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", - "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", - "Unit_hipGraphAddEventRecordNode_MultipleRun" + "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags" ] } diff --git a/tests/catch/unit/graph/hipStreamBeginCapture.cc b/tests/catch/unit/graph/hipStreamBeginCapture.cc index af550da759..80be624e8a 100644 --- a/tests/catch/unit/graph/hipStreamBeginCapture.cc +++ b/tests/catch/unit/graph/hipStreamBeginCapture.cc @@ -27,15 +27,111 @@ Testcase Scenarios : Functional Capture stream sequence and replay the sequence in multiple iterations. 4) End capture and validate that API returns captured graph for all possible modes on hipStreamPerThread. + 5) Waiting on an event recorded on a captured stream. Initiate capture + on stream1, record an event on stream1, wait for the event on stream2, + end the stream1 capture and Initiate stream capture on stream2 + 5.1) Both streams are created with default flags. + 5.2) Both streams are created with flag = hipStreamCaptureModeGlobal. + 5.3) Both streams are created with different flags. + 5.4) Both streams are created with different priorities. + 5.5) Validate the number of nodes in both the captured graphs. + 6) Colligated Streams capture. Capture operation sequences queued in + 2 streams by overlapping the 2 captures. + 6.1) Both streams are created with default flags. + 6.2) Both streams are created with flag = hipStreamCaptureModeGlobal. + 6.3) Both streams are created with different flags. + 6.4) Both streams are created with different priorities. + 7) Extend the scenario 5.1 for 3 streamsss. + 8) Create 2 streams. Start capturing both stream1 and stream2 at the same + time. On stream1 queue memcpy, kernel and memcpy operations and on stream2 + queue memcpy, kernel and memcpy operations. Execute both the captured + graphs and validate the results. + 9) Capture 2 streams in parallel using threads. Execute the graphs in + sequence in main thread and validate the results. + 9.1) mode = hipStreamCaptureModeGlobal + 9.2) mode = hipStreamCaptureModeThreadLocal + 9.3) mode = hipStreamCaptureModeRelaxed + 10) Queue operations (increment kernels) in 3 streams. Start capturing + the streams after some operations have been queued. This scenario validates + that only operations queued after hipStreamBeginCapture are captured in + the graph. + 11) Detecting invalid capture. Create 2 streams s1 and s2. Start capturing + s1. Create event dependency between s1 and s2 using event record and event + wait. Try capturing s2. hipStreamBeginCapture must return error. + 12) Stream reuse. Capture multiple graphs from the same stream. Validate + graphs are captured correctly. + 13) Test different synchronization during stream capture. + 13.1) Test hipStreamSynchronize. Must return + hipErrorStreamCaptureUnsupported. + 13.2) Test hipDeviceSynchronize. Must return + hipErrorStreamCaptureUnsupported. + 13.3) Test hipDeviceSynchronize. Must return + hipEventSynchronize. + 13.4) Test hipStreamWaitEvent. Must return + hipErrorStreamCaptureIsolation. + 14) End Stream Capture when the stream capture is still in progress. + 14.1) Abruptly end stream capture when stream capture is in progress in + forked stream. hipStreamEndCapture must return + hipErrorStreamCaptureUnjoined. + 14.2) Abruptly end stream capture when operations in forked stream + are still waiting to be captured. hipStreamEndCapture must return + hipErrorStreamCaptureUnjoined. + 15) Testing independent stream capture using multiple GPUs. Capture + a stream in each device context and execute the captured graph in the + context GPU. + 16) Test Nested Stream Capture Functionality: Create 3 streams s1, s2 & s3. + Capture s1, record event e1 on s1, wait for event e1 on s2 and queue + operations in s1. Record event e2 on s2 and wait for it on s3. Queue + operations on both s2 and s3. Record event e4 on s3 and wait for it in s1. + Record event e3 on s2 and wait for it in s1. End stream capture on s1. + Execute the graph and verify the result. + 17) Forked Stream Reuse: In scenario 16, after end capture on s1, queue + operations on both s2 and s3, and capture their graphs. Execute both the + graphs and validate the functionality. + 18) Capture a complex graph containing multiple independent memcpy, kernel + and host nodes. Launch the graph on random input data and validate the + output. + 19) Capture empty streams (parent + forked streams) and validate the + functionality. */ +#include #include #include - +#define INCREMENT_KERNEL_FINALEXP_VAL 7 constexpr size_t N = 1000000; constexpr int LAUNCH_ITERS = 50; +static int gCbackIter = 0; +#define GRIDSIZE 256 +#define BLOCKSIZE 256 +#define CONST_KER1_VAL 3 +#define CONST_KER2_VAL 2 +#define CONST_KER3_VAL 5 +static __global__ void dummyKernel() { + return; +} + +static __global__ void incrementKernel(int *data) { + atomicAdd(data, 1); + return; +} + +static __global__ void myadd(int* A_d, int* B_d) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + A_d[myId] = A_d[myId] + B_d[myId]; +} + +static __global__ void mymul(int* devMem, int value) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + devMem[myId] = devMem[myId] * value; +} + +static void hostNodeCallback(void* data) { + REQUIRE(data == nullptr); + gCbackIter++; +} bool CaptureStreamAndLaunchGraph(float *A_d, float *C_d, float *A_h, float *C_h, hipStreamCaptureMode mode, hipStream_t stream) { @@ -225,4 +321,974 @@ TEST_CASE("Unit_hipStreamBeginCapture_Basic") { HIP_CHECK(hipStreamDestroy(s2)); HIP_CHECK(hipStreamDestroy(s3)); } +/* Local Function + */ +static void interStrmEventSyncCapture(const hipStream_t &stream1, + const hipStream_t &stream2) { + hipGraph_t graph1, graph2; + hipEvent_t event; + hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr}; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(event, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, event, 0)); + dummyKernel<<<1, 1, 0, stream1>>>(); + HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); + HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); + dummyKernel<<<1, 1, 0, stream2>>>(); + dummyKernel<<<1, 1, 0, stream2>>>(); + HIP_CHECK(hipStreamEndCapture(stream2, &graph2)); + // Create Executable Graphs + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + REQUIRE(graphExec1 != nullptr); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + REQUIRE(graphExec2 != nullptr); + size_t numNodes1 = 0, numNodes2 = 0; + HIP_CHECK(hipGraphGetNodes(graph1, nullptr, &numNodes1)); + HIP_CHECK(hipGraphGetNodes(graph2, nullptr, &numNodes2)); + REQUIRE(numNodes1 == 1); + REQUIRE(numNodes2 == 2); + // Execute the Graphs + HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); + HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipStreamSynchronize(stream2)); + // Free + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipEventDestroy(event)); +} +/* Local Function + */ +static void colligatedStrmCapture(const hipStream_t &stream1, + const hipStream_t &stream2) { + hipGraph_t graph1, graph2; + hipEvent_t event; + hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr}; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(event, stream1)); + HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamWaitEvent(stream1, event, 0)); + dummyKernel<<<1, 1, 0, stream1>>>(); + HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); + dummyKernel<<<1, 1, 0, stream2>>>(); + HIP_CHECK(hipStreamEndCapture(stream2, &graph2)); + // Validate end capture is successful + REQUIRE(graph2 != nullptr); + REQUIRE(graph1 != nullptr); + // Create Executable Graphs + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + REQUIRE(graphExec1 != nullptr); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + REQUIRE(graphExec2 != nullptr); + // Execute the Graphs + HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); + HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipStreamSynchronize(stream2)); + // Free + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipEventDestroy(event)); +} +/* Fill input Data + */ +static void init_input(int* a, size_t size) { + unsigned int seed = time(nullptr); + for (size_t i = 0; i < size; i++) { + a[i] = (HipTest::RAND_R(&seed) & 0xFF); + } +} +/* Validate Output + */ +static void validate_output(int* a, int *b, size_t size) { + for (size_t i = 0; i < size; i++) { + REQUIRE(a[i] == (b[i]*b[i])); + } +} +/* Local Function + */ +static void colligatedStrmCaptureFunc(const hipStream_t &stream1, + const hipStream_t &stream2) { + constexpr size_t size = 1024; + constexpr auto blocksPerCU = 6; + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, size); + hipGraph_t graph1, graph2; + int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr}, + *outputVec_d1{nullptr}; + int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr}, + *outputVec_d2{nullptr}; + hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr}; + // host and device allocation + HipTest::initArrays(&inputVec_d1, &outputVec_d1, nullptr, + &inputVec_h1, &outputVec_h1, nullptr, size, false); + HipTest::initArrays(&inputVec_d2, &outputVec_d2, nullptr, + &inputVec_h2, &outputVec_h2, nullptr, size, false); + // Capture 2 streams + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(inputVec_d1, inputVec_h1, sizeof(int) * size, + hipMemcpyDefault, stream1)); + HIP_CHECK(hipMemcpyAsync(inputVec_d2, inputVec_h2, sizeof(int) * size, + hipMemcpyDefault, stream2)); + HipTest::vector_square<<>>( + inputVec_d1, outputVec_d1, size); + HipTest::vector_square<<>>( + inputVec_d2, outputVec_d2, size); + HIP_CHECK(hipMemcpyAsync(outputVec_h1, outputVec_d1, sizeof(int) * size, + hipMemcpyDefault, stream1)); + HIP_CHECK(hipMemcpyAsync(outputVec_h2, outputVec_d2, sizeof(int) * size, + hipMemcpyDefault, stream2)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); + HIP_CHECK(hipStreamEndCapture(stream2, &graph2)); + // Validate end capture is successful + REQUIRE(graph2 != nullptr); + REQUIRE(graph1 != nullptr); + // Create Executable Graphs + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + REQUIRE(graphExec1 != nullptr); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + REQUIRE(graphExec2 != nullptr); + // Execute the Graphs + for (int iter = 0; iter < LAUNCH_ITERS; iter++) { + init_input(inputVec_h1, size); + init_input(inputVec_h2, size); + HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); + HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipStreamSynchronize(stream2)); + validate_output(outputVec_h1, inputVec_h1, size); + validate_output(outputVec_h2, inputVec_h2, size); + } + // Free + HipTest::freeArrays(inputVec_d1, outputVec_d1, nullptr, + inputVec_h1, outputVec_h1, nullptr, false); + HipTest::freeArrays(inputVec_d2, outputVec_d2, nullptr, + inputVec_h2, outputVec_h2, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipGraphDestroy(graph1)); +} +/* Stream Capture thread function + */ +static void threadStrmCaptureFunc(hipStream_t stream, int *inputVec_d, +int *outputVec_d, int *inputVec_h, int *outputVec_h, hipGraph_t *graph, +size_t size, hipStreamCaptureMode mode) { + constexpr auto blocksPerCU = 6; + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, + threadsPerBlock, size); + // Capture stream + HIP_CHECK(hipStreamBeginCapture(stream, mode)); + HIP_CHECK(hipMemcpyAsync(inputVec_d, inputVec_h, sizeof(int) * size, + hipMemcpyDefault, stream)); + HipTest::vector_square<<>>( + inputVec_d, outputVec_d, size); + HIP_CHECK(hipMemcpyAsync(outputVec_h, outputVec_d, sizeof(int) * size, + hipMemcpyDefault, stream)); + HIP_CHECK(hipStreamEndCapture(stream, graph)); +} +/* Local Function for multithreaded tests + */ +static void multithreadedTest(hipStreamCaptureMode mode) { + hipStream_t stream1, stream2; + constexpr size_t size = 1024; + hipGraph_t graph1, graph2; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + int *inputVec_d1{nullptr}, *inputVec_h1{nullptr}, *outputVec_h1{nullptr}, + *outputVec_d1{nullptr}; + int *inputVec_d2{nullptr}, *inputVec_h2{nullptr}, *outputVec_h2{nullptr}, + *outputVec_d2{nullptr}; + hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr}; + // host and device allocation + HipTest::initArrays(&inputVec_d1, &outputVec_d1, nullptr, + &inputVec_h1, &outputVec_h1, nullptr, size, false); + HipTest::initArrays(&inputVec_d2, &outputVec_d2, nullptr, + &inputVec_h2, &outputVec_h2, nullptr, size, false); + // Launch 2 threads to capture the 2 streams into graphs + std::thread t1(threadStrmCaptureFunc, stream1, inputVec_d1, + outputVec_d1, inputVec_h1, outputVec_h1, &graph1, size, mode); + std::thread t2(threadStrmCaptureFunc, stream2, inputVec_d2, + outputVec_d2, inputVec_h2, outputVec_h2, &graph2, size, mode); + t1.join(); + t2.join(); + // Create Executable Graphs + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + // Execute the Graphs + for (int iter = 0; iter < LAUNCH_ITERS; iter++) { + init_input(inputVec_h1, size); + init_input(inputVec_h2, size); + HIP_CHECK(hipGraphLaunch(graphExec1, stream1)); + HIP_CHECK(hipGraphLaunch(graphExec2, stream2)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipStreamSynchronize(stream2)); + validate_output(outputVec_h1, inputVec_h1, size); + validate_output(outputVec_h2, inputVec_h2, size); + } + // Free + HipTest::freeArrays(inputVec_d1, outputVec_d1, nullptr, + inputVec_h1, outputVec_h1, nullptr, false); + HipTest::freeArrays(inputVec_d2, outputVec_d2, nullptr, + inputVec_h2, outputVec_h2, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); +} +/* Test scenario 5.1 + */ +TEST_CASE("Unit_hipStreamBeginCapture_InterStrmEventSync_defaultflag") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + interStrmEventSyncCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 5.2 + */ +TEST_CASE("Unit_hipStreamBeginCapture_InterStrmEventSync_blockingflag") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreateWithFlags(&stream1, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamNonBlocking)); + interStrmEventSyncCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 5.3 + */ +TEST_CASE("Unit_hipStreamBeginCapture_InterStrmEventSync_diffflags") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreateWithFlags(&stream1, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamDefault)); + interStrmEventSyncCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 5.4 + */ +TEST_CASE("Unit_hipStreamBeginCapture_InterStrmEventSync_diffprio") { + hipStream_t stream1, stream2; + int minPriority = 0, maxPriority = 0; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority)); + HIP_CHECK(hipStreamCreateWithPriority(&stream1, hipStreamDefault, + minPriority)); + HIP_CHECK(hipStreamCreateWithPriority(&stream2, hipStreamDefault, + maxPriority)); + interStrmEventSyncCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 6.1 + */ +TEST_CASE("Unit_hipStreamBeginCapture_ColligatedStrmCapture_defaultflag") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + colligatedStrmCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 6.2 + */ +TEST_CASE("Unit_hipStreamBeginCapture_ColligatedStrmCapture_blockingflag") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreateWithFlags(&stream1, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamNonBlocking)); + colligatedStrmCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 6.3 + */ +TEST_CASE("Unit_hipStreamBeginCapture_ColligatedStrmCapture_diffflags") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreateWithFlags(&stream1, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamDefault)); + colligatedStrmCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 6.4 + */ +TEST_CASE("Unit_hipStreamBeginCapture_ColligatedStrmCapture_diffprio") { + hipStream_t stream1, stream2; + int minPriority = 0, maxPriority = 0; + HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority)); + HIP_CHECK(hipStreamCreateWithPriority(&stream1, hipStreamDefault, + minPriority)); + HIP_CHECK(hipStreamCreateWithPriority(&stream2, hipStreamDefault, + maxPriority)); + colligatedStrmCapture(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 7 + */ +TEST_CASE("Unit_hipStreamBeginCapture_multiplestrms") { + hipStream_t stream1, stream2, stream3; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + hipGraph_t graph1, graph2, graph3; + size_t numNodes1 = 0, numNodes2 = 0, numNodes3 = 0; + SECTION("Capture Multiple stream with interdependent events") { + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(event1, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, event1, 0)); + dummyKernel<<<1, 1, 0, stream1>>>(); + HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); + HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(event2, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream3, event2, 0)); + dummyKernel<<<1, 1, 0, stream2>>>(); + HIP_CHECK(hipStreamEndCapture(stream2, &graph2)); + HIP_CHECK(hipStreamBeginCapture(stream3, hipStreamCaptureModeGlobal)); + dummyKernel<<<1, 1, 0, stream3>>>(); + HIP_CHECK(hipStreamEndCapture(stream3, &graph3)); + HIP_CHECK(hipGraphGetNodes(graph1, nullptr, &numNodes1)); + HIP_CHECK(hipGraphGetNodes(graph2, nullptr, &numNodes2)); + HIP_CHECK(hipGraphGetNodes(graph3, nullptr, &numNodes3)); + REQUIRE(numNodes1 == 1); + REQUIRE(numNodes2 == 1); + REQUIRE(numNodes3 == 1); + HIP_CHECK(hipEventDestroy(event2)); + HIP_CHECK(hipEventDestroy(event1)); + } + SECTION("Capture Multiple stream with single event") { + hipEvent_t event1; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(event1, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, event1, 0)); + HIP_CHECK(hipStreamWaitEvent(stream3, event1, 0)); + dummyKernel<<<1, 1, 0, stream1>>>(); + HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); + HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); + dummyKernel<<<1, 1, 0, stream2>>>(); + HIP_CHECK(hipStreamEndCapture(stream2, &graph2)); + HIP_CHECK(hipStreamBeginCapture(stream3, hipStreamCaptureModeGlobal)); + dummyKernel<<<1, 1, 0, stream3>>>(); + HIP_CHECK(hipStreamEndCapture(stream3, &graph3)); + HIP_CHECK(hipGraphGetNodes(graph1, nullptr, &numNodes1)); + HIP_CHECK(hipGraphGetNodes(graph2, nullptr, &numNodes2)); + HIP_CHECK(hipGraphGetNodes(graph3, nullptr, &numNodes3)); + REQUIRE(numNodes1 == 1); + REQUIRE(numNodes2 == 1); + REQUIRE(numNodes3 == 1); + HIP_CHECK(hipEventDestroy(event1)); + } + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 8 + */ +TEST_CASE("Unit_hipStreamBeginCapture_ColligatedStrmCapture_func") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + colligatedStrmCaptureFunc(stream1, stream2); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 9.1 + */ +TEST_CASE("Unit_hipStreamBeginCapture_Multithreaded_Global") { + multithreadedTest(hipStreamCaptureModeGlobal); +} +/* Test scenario 9.2 + */ +TEST_CASE("Unit_hipStreamBeginCapture_Multithreaded_ThreadLocal") { + multithreadedTest(hipStreamCaptureModeThreadLocal); +} +/* Test scenario 9.3 + */ +TEST_CASE("Unit_hipStreamBeginCapture_Multithreaded_Relaxed") { + multithreadedTest(hipStreamCaptureModeRelaxed); +} +/* Test scenario 10 + */ +TEST_CASE("Unit_hipStreamBeginCapture_CapturingFromWithinStrms") { + hipGraph_t graph; + hipStream_t stream1, stream2, stream3; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + hipEvent_t e1, e2, e3; + HIP_CHECK(hipEventCreate(&e1)); + HIP_CHECK(hipEventCreate(&e2)); + HIP_CHECK(hipEventCreate(&e3)); + // Create a device memory of size int and initialize it to 0 + int *devMem{nullptr}, *hostMem{nullptr}; + hostMem = reinterpret_cast(malloc(sizeof(int))); + HIP_CHECK(hipMalloc(&devMem, sizeof(int))); + HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Start Capturing stream1 + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(e1, stream1)); + incrementKernel<<<1, 1, 0, stream2>>>(devMem); + incrementKernel<<<1, 1, 0, stream2>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + HIP_CHECK(hipStreamWaitEvent(stream2, e1, 0)); + HIP_CHECK(hipStreamWaitEvent(stream3, e1, 0)); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream2>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream2>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + HIP_CHECK(hipEventRecord(e2, stream2)); + HIP_CHECK(hipEventRecord(e3, stream3)); + HIP_CHECK(hipStreamWaitEvent(stream1, e2, 0)); + HIP_CHECK(hipStreamWaitEvent(stream1, e3, 0)); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); // End Capture + // Reset device memory + HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Create Executable Graphs + hipGraphExec_t graphExec{nullptr}; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream1)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + REQUIRE((*hostMem) == INCREMENT_KERNEL_FINALEXP_VAL); + HIP_CHECK(hipFree(devMem)); + free(hostMem); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(e3)); + HIP_CHECK(hipEventDestroy(e2)); + HIP_CHECK(hipEventDestroy(e1)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 11 + */ +TEST_CASE("Unit_hipStreamBeginCapture_DetectingInvalidCapture") { + hipStream_t stream1, stream2; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(event, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, event, 0)); + dummyKernel<<<1, 1, 0, stream1>>>(); + // Since stream2 is already in capture mode due to event wait + // hipStreamBeginCapture on stream2 is expected to return error. + REQUIRE(hipSuccess != hipStreamBeginCapture(stream2, + hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 12 + */ +TEST_CASE("Unit_hipStreamBeginCapture_CapturingMultGraphsFrom1Strm") { + hipStream_t stream1; + HIP_CHECK(hipStreamCreate(&stream1)); + hipGraph_t graph[3]; + // Create a device memory of size int and initialize it to 0 + int *devMem{nullptr}, *hostMem{nullptr}; + hostMem = reinterpret_cast(malloc(sizeof(int))); + HIP_CHECK(hipMalloc(&devMem, sizeof(int))); + HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Capture Graph1 + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph[0])); + // Capture Graph2 + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph[1])); + // Capture Graph3 + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph[2])); + // Instantiate and execute all graphs + for (int i = 0; i < 3; i++) { + hipGraphExec_t graphExec{nullptr}; + HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph[i], nullptr, + nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream1)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + REQUIRE((*hostMem) == (i + 1)); + } + HIP_CHECK(hipFree(devMem)); + free(hostMem); + for (int i = 0; i < 3; i++) { + HIP_CHECK(hipGraphDestroy(graph[i])); + } + HIP_CHECK(hipStreamDestroy(stream1)); +} +#if HT_NVIDIA +/* Test scenario 13 + */ +TEST_CASE("Unit_hipStreamBeginCapture_CheckingSyncDuringCapture") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + SECTION("Synchronize stream during capture") { + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + REQUIRE(hipErrorStreamCaptureUnsupported == + hipStreamSynchronize(stream)); + } + SECTION("Synchronize device during capture") { + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + REQUIRE(hipErrorStreamCaptureUnsupported == hipDeviceSynchronize()); + } + SECTION("Synchronize event during capture") { + hipEvent_t e; + HIP_CHECK(hipEventCreate(&e)); + HIP_CHECK(hipEventRecord(e, stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + REQUIRE(hipErrorStreamCaptureUnsupported == hipEventSynchronize(e)); + HIP_CHECK(hipEventDestroy(e)); + } + SECTION("Wait for an event during capture") { + hipEvent_t e; + HIP_CHECK(hipEventCreate(&e)); + HIP_CHECK(hipEventRecord(e, stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + REQUIRE(hipErrorStreamCaptureIsolation == + hipStreamWaitEvent(stream, e, 0)); + HIP_CHECK(hipEventDestroy(e)); + } + SECTION("Query stream during capture") { + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + REQUIRE(hipErrorStreamCaptureUnsupported == hipStreamQuery(stream)); + } + SECTION("Query for an event during capture") { + hipEvent_t e; + HIP_CHECK(hipEventCreate(&e)); + HIP_CHECK(hipEventRecord(e, stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + REQUIRE(hipSuccess != hipEventQuery(e)); + HIP_CHECK(hipEventDestroy(e)); + } + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif +/* Test scenario 14 + */ +TEST_CASE("Unit_hipStreamBeginCapture_EndingCapturewhenCaptureInProgress") { + hipStream_t stream1, stream2; + hipGraph_t graph; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + SECTION("Abruptly end strm capture when in progress in forked strm") { + hipEvent_t e; + HIP_CHECK(hipEventCreate(&e)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + dummyKernel<<<1, 1, 0, stream1>>>(); + HIP_CHECK(hipEventRecord(e, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, e, 0)); + dummyKernel<<<1, 1, 0, stream2>>>(); + REQUIRE(hipErrorStreamCaptureUnjoined == + hipStreamEndCapture(stream1, &graph)); + HIP_CHECK(hipEventDestroy(e)); + } + SECTION("End strm capture when forked strm still has operations") { + hipEvent_t e1, e2; + HIP_CHECK(hipEventCreate(&e1)); + HIP_CHECK(hipEventCreate(&e2)); + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + dummyKernel<<<1, 1, 0, stream1>>>(); + HIP_CHECK(hipEventRecord(e1, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, e1, 0)); + dummyKernel<<<1, 1, 0, stream2>>>(); + HIP_CHECK(hipEventRecord(e2, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream1, e2, 0)); + dummyKernel<<<1, 1, 0, stream2>>>(); + REQUIRE(hipErrorStreamCaptureUnjoined == + hipStreamEndCapture(stream1, &graph)); + HIP_CHECK(hipEventDestroy(e2)); + HIP_CHECK(hipEventDestroy(e1)); + } + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 15 + */ +TEST_CASE("Unit_hipStreamBeginCapture_MultiGPU") { + int devcount = 0; + HIP_CHECK(hipGetDeviceCount(&devcount)); + // If only single GPU is detected then return + if (devcount < 2) { + SUCCEED("skipping the testcases as numDevices < 2"); + return; + } + hipStream_t* stream = reinterpret_cast(malloc( + devcount*sizeof(hipStream_t))); + REQUIRE(stream != nullptr); + hipGraph_t* graph = reinterpret_cast(malloc( + devcount*sizeof(hipGraph_t))); + REQUIRE(graph != nullptr); + int **devMem{nullptr}, **hostMem{nullptr}; + hostMem = reinterpret_cast(malloc(sizeof(int*)*devcount)); + REQUIRE(hostMem != nullptr); + devMem = reinterpret_cast(malloc(sizeof(int*)*devcount)); + REQUIRE(devMem != nullptr); + hipGraphExec_t* graphExec = reinterpret_cast(malloc( + devcount*sizeof(hipGraphExec_t))); + // Capture stream in each device + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + HIP_CHECK(hipStreamCreate(&stream[dev])); + hostMem[dev] = reinterpret_cast(malloc(sizeof(int))); + HIP_CHECK(hipMalloc(&devMem[dev], sizeof(int))); + HIP_CHECK(hipStreamBeginCapture(stream[dev], + hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemsetAsync(devMem[dev], 0, sizeof(int), stream[dev])); + for (int i = 0; i < (dev + 1); i++) { + incrementKernel<<<1, 1, 0, stream[dev]>>>(devMem[dev]); + } + HIP_CHECK(hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int), + hipMemcpyDefault, stream[dev])); + HIP_CHECK(hipStreamEndCapture(stream[dev], &graph[dev])); + } + // Launch the captured graphs in the respective device + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr, + nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec[dev], stream[dev])); + } + // Validate output + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + HIP_CHECK(hipStreamSynchronize(stream[dev])); + REQUIRE((*hostMem[dev]) == (dev + 1)); + } + // Destroy all device resources + for (int dev = 0; dev < devcount; dev++) { + HIP_CHECK(hipSetDevice(dev)); + HIP_CHECK(hipGraphExecDestroy(graphExec[dev])); + HIP_CHECK(hipStreamDestroy(stream[dev])); + } + free(graphExec); + free(hostMem); + free(devMem); + free(stream); + free(graph); +} +/* Test scenario 16 + */ +TEST_CASE("Unit_hipStreamBeginCapture_nestedStreamCapture") { + hipGraph_t graph; + hipStream_t stream1, stream2, stream3; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + hipEvent_t e1, e2, e3, e4; + HIP_CHECK(hipEventCreate(&e1)); + HIP_CHECK(hipEventCreate(&e2)); + HIP_CHECK(hipEventCreate(&e3)); + HIP_CHECK(hipEventCreate(&e4)); + // Create a device memory of size int and initialize it to 0 + int *devMem{nullptr}, *hostMem{nullptr}; + hostMem = reinterpret_cast(malloc(sizeof(int))); + REQUIRE(hostMem != nullptr); + HIP_CHECK(hipMalloc(&devMem, sizeof(int))); + HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Start Capturing stream1 + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(e1, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, e1, 0)); + HIP_CHECK(hipEventRecord(e2, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream3, e2, 0)); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream2>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + incrementKernel<<<1, 1, 0, stream1>>>(devMem); + incrementKernel<<<1, 1, 0, stream2>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + incrementKernel<<<1, 1, 0, stream3>>>(devMem); + HIP_CHECK(hipEventRecord(e3, stream2)); + HIP_CHECK(hipEventRecord(e4, stream3)); + HIP_CHECK(hipStreamWaitEvent(stream1, e4, 0)); + HIP_CHECK(hipStreamWaitEvent(stream1, e3, 0)); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); // End Capture + // Reset device memory + HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Create Executable Graphs + hipGraphExec_t graphExec{nullptr}; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream1)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + REQUIRE((*hostMem) == INCREMENT_KERNEL_FINALEXP_VAL); + HIP_CHECK(hipFree(devMem)); + free(hostMem); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(e4)); + HIP_CHECK(hipEventDestroy(e3)); + HIP_CHECK(hipEventDestroy(e2)); + HIP_CHECK(hipEventDestroy(e1)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 17 + */ +TEST_CASE("Unit_hipStreamBeginCapture_streamReuse") { + hipGraph_t graph1, graph2, graph3; + hipStream_t stream1, stream2, stream3; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + hipEvent_t e1, e2, e3, e4; + HIP_CHECK(hipEventCreate(&e1)); + HIP_CHECK(hipEventCreate(&e2)); + HIP_CHECK(hipEventCreate(&e3)); + HIP_CHECK(hipEventCreate(&e4)); + // Create a device memory of size int and initialize it to 0 + int *devMem1{nullptr}, *hostMem1{nullptr}, *devMem2{nullptr}, + *hostMem2{nullptr}, *devMem3{nullptr}, *hostMem3{nullptr}; + HipTest::initArrays(&devMem1, &devMem2, &devMem3, + &hostMem1, &hostMem2, &hostMem3, 1, false); + HIP_CHECK(hipMemset(devMem1, 0, sizeof(int))); + HIP_CHECK(hipMemset(devMem2, 0, sizeof(int))); + HIP_CHECK(hipMemset(devMem3, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Start Capturing stream1 + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(e1, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, e1, 0)); + HIP_CHECK(hipEventRecord(e2, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream3, e2, 0)); + incrementKernel<<<1, 1, 0, stream1>>>(devMem1); + incrementKernel<<<1, 1, 0, stream2>>>(devMem1); + incrementKernel<<<1, 1, 0, stream3>>>(devMem1); + incrementKernel<<<1, 1, 0, stream1>>>(devMem1); + incrementKernel<<<1, 1, 0, stream2>>>(devMem1); + incrementKernel<<<1, 1, 0, stream3>>>(devMem1); + incrementKernel<<<1, 1, 0, stream3>>>(devMem1); + HIP_CHECK(hipEventRecord(e3, stream2)); + HIP_CHECK(hipEventRecord(e4, stream3)); + HIP_CHECK(hipStreamWaitEvent(stream1, e4, 0)); + HIP_CHECK(hipStreamWaitEvent(stream1, e3, 0)); + HIP_CHECK(hipMemcpyAsync(hostMem1, devMem1, sizeof(int), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); // End Capture + // Start capturing graph2 from stream 2 + HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); + incrementKernel<<<1, 1, 0, stream2>>>(devMem2); + incrementKernel<<<1, 1, 0, stream2>>>(devMem2); + incrementKernel<<<1, 1, 0, stream2>>>(devMem2); + HIP_CHECK(hipMemcpyAsync(hostMem2, devMem2, sizeof(int), + hipMemcpyDefault, stream2)); + HIP_CHECK(hipStreamEndCapture(stream2, &graph2)); // End Capture + // Start capturing graph3 from stream 3 + HIP_CHECK(hipStreamBeginCapture(stream3, hipStreamCaptureModeGlobal)); + incrementKernel<<<1, 1, 0, stream3>>>(devMem3); + incrementKernel<<<1, 1, 0, stream3>>>(devMem3); + incrementKernel<<<1, 1, 0, stream3>>>(devMem3); + incrementKernel<<<1, 1, 0, stream3>>>(devMem3); + incrementKernel<<<1, 1, 0, stream3>>>(devMem3); + HIP_CHECK(hipMemcpyAsync(hostMem3, devMem3, sizeof(int), + hipMemcpyDefault, stream3)); + HIP_CHECK(hipStreamEndCapture(stream3, &graph3)); // End Capture + // Reset device memory + HIP_CHECK(hipMemset(devMem1, 0, sizeof(int))); + HIP_CHECK(hipMemset(devMem2, 0, sizeof(int))); + HIP_CHECK(hipMemset(devMem3, 0, sizeof(int))); + HIP_CHECK(hipDeviceSynchronize()); + // Create Executable Graphs + hipGraphExec_t graphExec{nullptr}; + // Verify graph1 + HIP_CHECK(hipGraphInstantiate(&graphExec, graph1, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream1)); + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + REQUIRE((*hostMem1) == INCREMENT_KERNEL_FINALEXP_VAL); + // Verify graph2 + HIP_CHECK(hipGraphInstantiate(&graphExec, graph2, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream2)); + HIP_CHECK(hipStreamSynchronize(stream2)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + REQUIRE((*hostMem2) == 3); + // Verify graph3 + HIP_CHECK(hipGraphInstantiate(&graphExec, graph3, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream3)); + HIP_CHECK(hipStreamSynchronize(stream3)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + REQUIRE((*hostMem3) == 5); + HipTest::freeArrays(devMem1, devMem2, devMem3, + hostMem1, hostMem2, hostMem3, false); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipGraphDestroy(graph3)); + HIP_CHECK(hipEventDestroy(e4)); + HIP_CHECK(hipEventDestroy(e3)); + HIP_CHECK(hipEventDestroy(e2)); + HIP_CHECK(hipEventDestroy(e1)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} + +/* Test scenario 18 + */ +TEST_CASE("Unit_hipStreamBeginCapture_captureComplexGraph") { + hipGraph_t graph; + hipStream_t stream1, stream2, stream3, stream4, stream5; + // Stream and event create + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + HIP_CHECK(hipStreamCreate(&stream4)); + HIP_CHECK(hipStreamCreate(&stream5)); + hipEvent_t e0, e1, e2, e3, e4, e5, e6; + HIP_CHECK(hipEventCreate(&e0)); + HIP_CHECK(hipEventCreate(&e1)); + HIP_CHECK(hipEventCreate(&e2)); + HIP_CHECK(hipEventCreate(&e3)); + HIP_CHECK(hipEventCreate(&e4)); + HIP_CHECK(hipEventCreate(&e5)); + HIP_CHECK(hipEventCreate(&e6)); + // Allocate Device memory and Host memory + size_t N = GRIDSIZE*BLOCKSIZE; + int *Ah{nullptr}, *Bh{nullptr}, *Ch{nullptr}, *Ad{nullptr}, *Bd{nullptr}; + HipTest::initArrays(&Ad, &Bd, nullptr, &Ah, &Bh, &Ch, N, false); + // Capture streams into graph + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(e0, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream4, e0, 0)); + HIP_CHECK(hipStreamWaitEvent(stream5, e0, 0)); + HIP_CHECK(hipMemcpyAsync(Ad, Ah, (N*sizeof(int)), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipMemcpyAsync(Bd, Bh, (N*sizeof(int)), + hipMemcpyDefault, stream5)); + hipHostFn_t fn = hostNodeCallback; + HIPCHECK(hipLaunchHostFunc(stream4, fn, nullptr)); + HIP_CHECK(hipEventRecord(e1, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, e1, 0)); + int *Ad_2nd_half = Ad + N/2; + int *Ad_1st_half = Ad; + mymul<<>>(Ad_2nd_half, CONST_KER2_VAL); + mymul<<>>(Ad_1st_half, CONST_KER1_VAL); + HIP_CHECK(hipEventRecord(e2, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream3, e2, 0)); + mymul<<>>(Ad_1st_half, CONST_KER3_VAL); + HIPCHECK(hipLaunchHostFunc(stream3, fn, nullptr)); + HIP_CHECK(hipEventRecord(e6, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream1, e6, 0)); + HIP_CHECK(hipEventRecord(e5, stream5)); + HIP_CHECK(hipStreamWaitEvent(stream1, e5, 0)); + myadd<<>>(Ad, Bd); + HIP_CHECK(hipEventRecord(e3, stream3)); + HIP_CHECK(hipStreamWaitEvent(stream1, e3, 0)); + HIP_CHECK(hipEventRecord(e4, stream4)); + HIP_CHECK(hipStreamWaitEvent(stream1, e4, 0)); + HIP_CHECK(hipMemcpyAsync(Ch, Ad, (N*sizeof(int)), + hipMemcpyDefault, stream1)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); // End Capture + // Execute and test the graph + // Create Executable Graphs + hipGraphExec_t graphExec{nullptr}; + // Verify graph1 + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + for (int iter = 0; iter < LAUNCH_ITERS; iter++) { + init_input(Ah, N); + init_input(Bh, N); + HIP_CHECK(hipGraphLaunch(graphExec, stream1)); + HIP_CHECK(hipStreamSynchronize(stream1)); + for (size_t i = 0; i < N; i++) { + if (i > (N/2 - 1)) { + REQUIRE(Ch[i] == (Bh[i] + Ah[i]*CONST_KER2_VAL)); + } else { + REQUIRE(Ch[i] == (Bh[i] + Ah[i]*CONST_KER1_VAL*CONST_KER3_VAL)); + } + } + } + REQUIRE(gCbackIter == (2*LAUNCH_ITERS)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + // Free Device memory and Host memory + HipTest::freeArrays(Ad, Bd, nullptr, Ah, Bh, Ch, false); + // Destroy graph, events and streams + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(e6)); + HIP_CHECK(hipEventDestroy(e5)); + HIP_CHECK(hipEventDestroy(e4)); + HIP_CHECK(hipEventDestroy(e3)); + HIP_CHECK(hipEventDestroy(e2)); + HIP_CHECK(hipEventDestroy(e1)); + HIP_CHECK(hipEventDestroy(e0)); + HIP_CHECK(hipStreamDestroy(stream5)); + HIP_CHECK(hipStreamDestroy(stream4)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} +/* Test scenario 19 + */ +TEST_CASE("Unit_hipStreamBeginCapture_captureEmptyStreams") { + hipGraph_t graph; + hipStream_t stream1, stream2, stream3; + // Stream and event create + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + hipEvent_t e0, e1, e2; + HIP_CHECK(hipEventCreate(&e0)); + HIP_CHECK(hipEventCreate(&e1)); + HIP_CHECK(hipEventCreate(&e2)); + // Capture streams into graph + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(e0, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, e0, 0)); + HIP_CHECK(hipStreamWaitEvent(stream3, e0, 0)); + HIP_CHECK(hipEventRecord(e1, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream1, e1, 0)); + HIP_CHECK(hipEventRecord(e2, stream3)); + HIP_CHECK(hipStreamWaitEvent(stream1, e2, 0)); + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); // End Capture + size_t numNodes = 0; + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + REQUIRE(numNodes == 0); + // Destroy graph, events and streams + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(e2)); + HIP_CHECK(hipEventDestroy(e1)); + HIP_CHECK(hipEventDestroy(e0)); + HIP_CHECK(hipStreamDestroy(stream3)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream1)); +} From 8202097629485c5daeb8091d4a6882e8d11e18dc Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 14:04:36 +0530 Subject: [PATCH 6/9] SWDEV-306410 - Add Unit_hipEventMGpuMThreads* (#3095) Change-Id: Ieaaed1d15b4d258103d8341d4b17d9f83a9e4785 --- tests/catch/unit/event/CMakeLists.txt | 1 + .../unit/event/Unit_hipEventMGpuMThreads.cc | 215 ++++++++++++++++++ 2 files changed, 216 insertions(+) create mode 100644 tests/catch/unit/event/Unit_hipEventMGpuMThreads.cc diff --git a/tests/catch/unit/event/CMakeLists.txt b/tests/catch/unit/event/CMakeLists.txt index 3054e85670..b4d5f2c3b6 100644 --- a/tests/catch/unit/event/CMakeLists.txt +++ b/tests/catch/unit/event/CMakeLists.txt @@ -9,6 +9,7 @@ set(TEST_SRC hipEventCreate.cc hipEventCreateWithFlags.cc hipEventSynchronize.cc + Unit_hipEventMGpuMThreads.cc ) # The test used wait mechanism and doesnt play well with all arch of nvidia diff --git a/tests/catch/unit/event/Unit_hipEventMGpuMThreads.cc b/tests/catch/unit/event/Unit_hipEventMGpuMThreads.cc new file mode 100644 index 0000000000..a6d849065b --- /dev/null +++ b/tests/catch/unit/event/Unit_hipEventMGpuMThreads.cc @@ -0,0 +1,215 @@ +/* + Copyright (c) 2022 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. + */ +#define HIP_TEMPLATE_KERNEL_LAUNCH +#include +#include +#include +#include + +int64_t timeNanos() { + using namespace std::chrono; + static auto t0 = steady_clock::now(); + auto timeSpan = duration_cast(steady_clock::now() - t0); + return timeSpan.count(); +} + +#define WIDTH 1024 + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +__global__ void matrixTranspose(float *out, float *in, const int width) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void matrixTransposeCPUReference(float *output, float *input, + const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; + } + } +} + +void thread_run(const int iThread) { + int i = 0; + int errors = 0; + float eventMs = 1.0f; + float *matrix = nullptr; + float *transposeMatrix = nullptr; + float *cpuTransposeMatrix = nullptr; + float *gpuMatrix = nullptr; + float *gpuTransposeMatrix = nullptr; + hipDeviceProp_t devProp; + memset(&devProp, 0, sizeof(devProp)); + HIP_CHECK(hipGetDeviceProperties(&devProp, iThread)); + fprintf(stderr, "[%d] device name = %s\n", iThread, devProp.name); + + HIP_CHECK(hipSetDevice(iThread)); + hipEvent_t start, stop; + + auto time = timeNanos(); + HIP_CHECK(hipEventCreate(&start)); + fprintf(stderr, "[%d] hipEventCreate(&start) cost cpu time %6.3fms\n", iThread, + (timeNanos() - time)/1000000.0); + + HIP_CHECK(hipEventCreate(&stop)); + + matrix = (float*) malloc(NUM * sizeof(float)); + transposeMatrix = (float*) malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*) malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + matrix[i] = (float) i * 10.0f; + } + + // allocate the memory on the device side + HIP_CHECK(hipMalloc((void**) &gpuMatrix, NUM * sizeof(float))); + HIP_CHECK(hipMalloc((void**) &gpuTransposeMatrix, NUM * sizeof(float))); + + time = timeNanos(); + // Record the start event + // The first call of hipEventRecord will trigger VirtualDevice creation that will trigger building + // of BlitLinearSourceCode, which will cost 200+ ms. + HIP_CHECK(hipEventRecord(start)); + fprintf(stderr, "[%d] hipEventRecord(&start) cost cpu time %6.3fms\n", iThread, + (timeNanos() - time)/1000000.0); + + time = timeNanos(); + // Memory transfer from host to device + HIP_CHECK(hipMemcpy(gpuMatrix, matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); + + // Record the stop event + HIP_CHECK(hipEventRecord(stop)); + HIP_CHECK(hipEventSynchronize(stop)); + + HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + + fprintf(stderr, "[%d] hipMemcpyHostToDevice cost gpu time %6.3fms, cpu time %6.3fms\n", + iThread, eventMs, (timeNanos() - time)/1000000.0); + + // Record the start event + HIP_CHECK(hipEventRecord(start)); + + time = timeNanos(); + // Lauching kernel from host + hipLaunchKernelGGL( + matrixTranspose, + dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, + gpuMatrix, WIDTH); + // Record the stop event + HIP_CHECK(hipEventRecord(stop)); + + fprintf(stderr, "[%d] hipLaunchKernelGGL() cost cpu time %6.3fms\n", iThread, + (timeNanos() - time)/1000000.0); + + HIP_CHECK(hipEventSynchronize(stop)); + HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + + fprintf(stderr, "[%d] kernel Execution cost gpu time %6.3fms, cpu time = %6.3fms\n", + iThread, eventMs, (timeNanos() - time)/1000000.0); + + // Record the start event + HIP_CHECK(hipEventRecord(start)); + + // Memory transfer from device to host + HIP_CHECK(hipMemcpy(transposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), + hipMemcpyDeviceToHost)); + + // Record the stop event + HIP_CHECK(hipEventRecord(stop)); + HIP_CHECK(hipEventSynchronize(stop)); + + HIP_CHECK(hipEventElapsedTime(&eventMs, start, stop)); + + fprintf(stderr, "[%d] hipMemcpyDeviceToHost cost gpu time %6.3fms\n", iThread, eventMs); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, matrix, WIDTH); + + // verify the results + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(transposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + fprintf(stderr, "[%d] FAILED: %d errors\n", iThread, errors); + } else { + fprintf(stderr, "[%d] PASSED\n", iThread); + } + // free the resources on device side + HIP_CHECK(hipFree(gpuMatrix)); + HIP_CHECK(hipFree(gpuTransposeMatrix)); + + // free the resources on host side + free(matrix); + free(transposeMatrix); + free(cpuTransposeMatrix); + REQUIRE(errors == 0); +} + +void testEventMGpuMThreads(int nThreads = 1) { + int iThread = 0; + std::thread *threads = new std::thread[nThreads]; + for (iThread = 0; iThread < nThreads; iThread++) { + threads[iThread] = std::thread(thread_run, iThread); + } + for (iThread = 0; iThread < nThreads; iThread++) { + threads[iThread].join(); + } + delete []threads; +} + +TEST_CASE("Unit_hipEventMGpuMThreads_1") { + testEventMGpuMThreads(1); +} + +TEST_CASE("Unit_hipEventMGpuMThreads_2") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + testEventMGpuMThreads(numDevices); + } else { + SUCCEED("skipped the testcase as number of devices is less than 2"); + } +} + +TEST_CASE("Unit_hipEventMGpuMThreads_3") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + fprintf(stderr, "First round\n"); + testEventMGpuMThreads(numDevices); + fprintf(stderr, "Second round\n"); + testEventMGpuMThreads(numDevices); + } else { + SUCCEED("skipped the testcase as number of devices is less than 2"); + } +} From d76a487187dbe24bdb3eb00e122e64e67241d37a Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 20:15:49 +0530 Subject: [PATCH 7/9] SWDEV-306122 - [catch2][dtest] Added test for hipUserObjectCreate, hipUserObjectRelease, hipUserObjectRetain, hipGraphReleaseUserObject, hipGraphRetainUserObject API (#3061) * SWDEV-306122 - [catch2][dtest] Added test for hipUserObjectCreate, hipUserObjectRelease, hipUserObjectRetain, hipGraphReleaseUserObject, hipGraphRetainUserObject API Change-Id: I7fd25bcc23f0629d92296ab7b20258a6e268a58b --- tests/catch/unit/graph/CMakeLists.txt | 1 + tests/catch/unit/graph/hipUserObjectCreate.cc | 650 ++++++++++++++++++ 2 files changed, 651 insertions(+) create mode 100644 tests/catch/unit/graph/hipUserObjectCreate.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 7f15b7e6ff..e2b66872da 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -80,6 +80,7 @@ set(TEST_SRC hipGraphHostNodeGetParams.cc hipGraphExecChildGraphNodeSetParams.cc hipStreamGetCaptureInfo_v2.cc + hipUserObjectCreate.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipUserObjectCreate.cc b/tests/catch/unit/graph/hipUserObjectCreate.cc new file mode 100644 index 0000000000..d2765d7978 --- /dev/null +++ b/tests/catch/unit/graph/hipUserObjectCreate.cc @@ -0,0 +1,650 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + + +/** + * Functional Test for API - hipUserObjectCreate +1) Call hipUserObjectCreate once and release it by calling hipUserObjectRelease +2) Call hipUserObjectCreate refCount as X and release it by calling + hipUserObjectRelease with same refCount. +3) Call hipUserObjectCreate, retain it by calling hipUserObjectRetain + and release it by calling hipUserObjectRelease twice. +4) Call hipUserObjectCreate with refCount as X, retain it by calling + hipUserObjectRetain with count as Y and release it by calling + hipUserObjectRelease with count as X+Y. + */ + +struct BoxStruct { + int count; + BoxStruct() { + INFO("Constructor called for Struct!\n"); + } + ~BoxStruct() { + INFO("Destructor called for Struct!\n"); + } +}; + +class BoxClass { + public: + BoxClass() { + INFO("Constructor called for Class!\n"); + } + ~BoxClass() { + INFO("Destructor called for Class!\n"); + } +}; + +static void destroyStructObj(void *ptr) { + BoxStruct *ptr1 = reinterpret_cast(ptr); + delete ptr1; +} + +static void destroyClassObj(void *ptr) { + BoxClass *ptr2 = reinterpret_cast(ptr); + delete ptr2; +} + +static void destroyIntObj(void *ptr) { + int *ptr2 = reinterpret_cast(ptr); + delete ptr2; +} + +static void destroyFloatObj(void *ptr) { + float *ptr2 = reinterpret_cast(ptr); + delete ptr2; +} + +/* 1) Call hipUserObjectCreate once and release it by + calling hipUserObjectRelease */ +static void hipUserObjectCreate_Functional_1(void *object, + void destroyObj(void *)) { + hipUserObject_t hObject; + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyObj, 1, + hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipUserObjectRelease(hObject)); +} + +TEST_CASE("Unit_hipUserObjectCreate_Functional_1") { + SECTION("Called with int Object") { + int *object = new int(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_1(object, destroyIntObj); + } + SECTION("Called with float Object") { + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_1(object, destroyFloatObj); + } + SECTION("Called with Class Object") { + BoxClass *object = new BoxClass(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_1(object, destroyClassObj); + } + SECTION("Called with Struct Object") { + BoxStruct *object = new BoxStruct(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_1(object, destroyStructObj); + } +} + +/* 2) Call hipUserObjectCreate refCount as X and release it by + calling hipUserObjectRelease with same refCount. */ +static void hipUserObjectCreate_Functional_2(void *object, + void destroyObj(void *)) { + int refCount = 5; + hipUserObject_t hObject; + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyObj, + refCount, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipUserObjectRelease(hObject, refCount)); +} + +TEST_CASE("Unit_hipUserObjectCreate_Functional_2") { + SECTION("Called with int Object") { + int *object = new int(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_2(object, destroyIntObj); + } + SECTION("Called with float Object") { + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_2(object, destroyFloatObj); + } + SECTION("Called with Class Object") { + BoxClass *object = new BoxClass(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_2(object, destroyClassObj); + } + SECTION("Called with Struct Object") { + BoxStruct *object = new BoxStruct(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_2(object, destroyStructObj); + } +} + +/* 3) Call hipUserObjectCreate, retain it by calling hipUserObjectRetain + and release it by calling hipUserObjectRelease twice. */ +static void hipUserObjectCreate_Functional_3(void *object, + void destroyObj(void *)) { + hipUserObject_t hObject; + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipUserObjectRetain(hObject)); + HIP_CHECK(hipUserObjectRelease(hObject)); + HIP_CHECK(hipUserObjectRelease(hObject)); +} + +TEST_CASE("Unit_hipUserObjectCreate_Functional_3") { + SECTION("Called with int Object") { + int *object = new int(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_3(object, destroyIntObj); + } + SECTION("Called with float Object") { + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_3(object, destroyFloatObj); + } + SECTION("Called with Class Object") { + BoxClass *object = new BoxClass(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_3(object, destroyClassObj); + } + SECTION("Called with Struct Object") { + BoxStruct *object = new BoxStruct(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_3(object, destroyStructObj); + } +} + +/* 4) Call hipUserObjectCreate with refCount as X, retain it by calling + hipUserObjectRetain with count as Y and release it by calling + hipUserObjectRelease with count as X+Y. */ +static void hipUserObjectCreate_Functional_4(void *object, + void destroyObj(void *)) { + int refCount = 5; + int refCountRetain = 8; + hipUserObject_t hObject; + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyObj, + refCount, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipUserObjectRetain(hObject, refCountRetain)); + HIP_CHECK(hipUserObjectRelease(hObject, refCount+refCountRetain)); +} + +TEST_CASE("Unit_hipUserObjectCreate_Functional_4") { + SECTION("Called with int Object") { + int *object = new int(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_4(object, destroyIntObj); + } + SECTION("Called with float Object") { + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_4(object, destroyFloatObj); + } + SECTION("Called with Class Object") { + BoxClass *object = new BoxClass(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_4(object, destroyClassObj); + } + SECTION("Called with Struct Object") { + BoxStruct *object = new BoxStruct(); + REQUIRE(object != nullptr); + hipUserObjectCreate_Functional_4(object, destroyStructObj); + } +} + + +/** + * Negative Test for API - hipUserObjectCreate + 1) Pass User Object as nullptr + 2) Pass object as nullptr + 3) Pass Callback function as nullptr + 4) Pass initialRefcount as 0 + 5) Pass initialRefcount as INT_MAX + 6) Pass flag other than hipUserObjectNoDestructorSync + */ + +TEST_CASE("Unit_hipUserObjectCreate_Negative") { + hipError_t ret; + int *object = new int(); + REQUIRE(object != nullptr); + + hipUserObject_t hObject; + SECTION("Pass User Object as nullptr") { + ret = hipUserObjectCreate(nullptr, object, destroyIntObj, + 1, hipUserObjectNoDestructorSync); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass object as nullptr") { + ret = hipUserObjectCreate(&hObject, nullptr, destroyIntObj, + 1, hipUserObjectNoDestructorSync); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass Callback function as nullptr") { + ret = hipUserObjectCreate(&hObject, object, nullptr, + 1, hipUserObjectNoDestructorSync); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as 0") { + ret = hipUserObjectCreate(&hObject, object, destroyIntObj, + 0, hipUserObjectNoDestructorSync); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as INT_MAX") { + ret = hipUserObjectCreate(&hObject, object, destroyIntObj, + INT_MAX, hipUserObjectNoDestructorSync); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass flag other than hipUserObjectNoDestructorSync") { + ret = hipUserObjectCreate(&hObject, object, destroyIntObj, + 1, hipUserObjectFlags(9)); + REQUIRE(hipErrorInvalidValue == ret); + } +} + +/** + * Negative Test for API - hipUserObjectRelease + 1) Pass User Object as nullptr + 2) Pass initialRefcount as 0 + 3) Pass initialRefcount as INT_MAX + */ + +TEST_CASE("Unit_hipUserObjectRelease_Negative") { + hipError_t ret; + int *object = new int(); + REQUIRE(object != nullptr); + + hipUserObject_t hObject; + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyIntObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + + SECTION("Pass User Object as nullptr") { + ret = hipUserObjectRelease(nullptr, 1); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as 0") { + ret = hipUserObjectRelease(hObject, 0); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as INT_MAX") { + ret = hipUserObjectRelease(hObject, INT_MAX); + REQUIRE(hipSuccess == ret); + } +} + +/** + * Negative Test for API - hipUserObjectRetain + 1) Pass User Object as nullptr + 2) Pass initialRefcount as 0 + 3) Pass initialRefcount as INT_MAX + */ + +TEST_CASE("Unit_hipUserObjectRetain_Negative") { + hipError_t ret; + int *object = new int(); + REQUIRE(object != nullptr); + + hipUserObject_t hObject; + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyIntObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + + SECTION("Pass User Object as nullptr") { + ret = hipUserObjectRetain(nullptr, 1); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as 0") { + ret = hipUserObjectRetain(hObject, 0); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as INT_MAX") { + ret = hipUserObjectRetain(hObject, INT_MAX); + REQUIRE(hipSuccess == ret); + } +} + +TEST_CASE("Unit_hipUserObj_Negative_Test") { + hipError_t ret; + int *object = new int(); + REQUIRE(object != nullptr); + + hipUserObject_t hObject; + // Create a new hObject with 2 reference + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyIntObj, + 2, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + + // Release more than created. + ret = hipUserObjectRelease(hObject, 4); + REQUIRE(hipSuccess == ret); + + // Retain reference to a removed user object + ret = hipUserObjectRetain(hObject, 1); + REQUIRE(hipSuccess == ret); +} + +/** + * Functional Test for API - hipGraphRetainUserObject + */ + +/* 1) Create GraphUserObject and retain it by calling hipGraphRetainUserObject + and release it by calling hipGraphReleaseUserObject. */ + +static void hipGraphRetainUserObject_Functional_1(void *object, + void destroyObj(void *)) { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipUserObject_t hObject; + + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1, + hipGraphUserObjectMove)); + + HIP_CHECK(hipGraphReleaseUserObject(graph, hObject)); + HIP_CHECK(hipUserObjectRelease(hObject)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +TEST_CASE("Unit_hipGraphRetainUserObject_Functional_1") { + SECTION("Called with int Object") { + int *object = new int(); + REQUIRE(object != nullptr); + hipGraphRetainUserObject_Functional_1(object, destroyIntObj); + } + SECTION("Called with float Object") { + float *object = new float(); + REQUIRE(object != nullptr); + hipGraphRetainUserObject_Functional_1(object, destroyFloatObj); + } + SECTION("Called with Class Object") { + BoxClass *object = new BoxClass(); + REQUIRE(object != nullptr); + hipGraphRetainUserObject_Functional_1(object, destroyClassObj); + } + SECTION("Called with Struct Object") { + BoxStruct *object = new BoxStruct(); + REQUIRE(object != nullptr); + hipGraphRetainUserObject_Functional_1(object, destroyStructObj); + } +} + +/* 2) Create UserObject and GraphUserObject and retain using custom reference + count and release it by calling hipGraphReleaseUserObject with count. */ + +TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memcpyNode, kNode; + hipKernelNodeParams kNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kNodeParams.gridDim = dim3(blocks); + kNodeParams.blockDim = dim3(threadsPerBlock); + kNodeParams.sharedMemBytes = 0; + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(), + dependencies.size(), &kNodeParams)); + + dependencies.clear(); + dependencies.push_back(kNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + int refCount = 2; + int refCountRetain = 3; + + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObject_t hObject; + + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyFloatObj, + refCount, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipUserObjectRetain(hObject, refCountRetain)); + HIP_CHECK(hipGraphRetainUserObject(graph, hObject, refCountRetain, + hipGraphUserObjectMove)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HIP_CHECK(hipUserObjectRelease(hObject, refCount + refCountRetain)); + HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, refCountRetain)); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + + +/** + * Negative Test for API - hipGraphRetainUserObject + 1) Pass graph as nullptr + 2) Pass User Object as nullptr + 3) Pass initialRefcount as 0 + 4) Pass initialRefcount as INT_MAX + 5) Pass flag as 0 + 6) Pass flag as INT_MAX + */ + +TEST_CASE("Unit_hipGraphRetainUserObject_Negative") { + hipError_t ret; + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObject_t hObject; + + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyFloatObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + + SECTION("Pass graph as nullptr") { + ret = hipGraphRetainUserObject(nullptr, hObject, 1, + hipGraphUserObjectMove); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass User Object as nullptr") { + ret = hipGraphRetainUserObject(graph, nullptr, 1, + hipGraphUserObjectMove); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as 0") { + ret = hipGraphRetainUserObject(graph, hObject, 0, + hipGraphUserObjectMove); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as INT_MAX") { + ret = hipGraphRetainUserObject(graph, hObject, INT_MAX, + hipGraphUserObjectMove); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass flag as 0") { + ret = hipGraphRetainUserObject(graph, hObject, 1, 0); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass flag as INT_MAX") { + ret = hipGraphRetainUserObject(graph, hObject, 1, INT_MAX); + REQUIRE(hipErrorInvalidValue == ret); + } + + HIP_CHECK(hipUserObjectRelease(hObject, 1)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Negative Test for API - hipGraphReleaseUserObject + 1) Pass graph as nullptr + 2) Pass User Object as nullptr + 3) Pass initialRefcount as 0 + 4) Pass initialRefcount as INT_MAX + */ + +TEST_CASE("Unit_hipGraphReleaseUserObject_Negative") { + hipError_t ret; + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObject_t hObject; + + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyFloatObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1, + hipGraphUserObjectMove)); + + SECTION("Pass graph as nullptr") { + ret = hipGraphReleaseUserObject(nullptr, hObject, 1); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass User Object as nullptr") { + ret = hipGraphReleaseUserObject(graph, nullptr, 1); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as 0") { + ret = hipGraphReleaseUserObject(graph, hObject, 0); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass initialRefcount as INT_MAX") { + ret = hipGraphReleaseUserObject(graph, hObject, INT_MAX); + REQUIRE(hipSuccess == ret); + } + HIP_CHECK(hipUserObjectRelease(hObject, 1)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Basic") { + hipError_t ret; + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + float *object = new float(); + REQUIRE(object != nullptr); + hipUserObject_t hObject; + + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyFloatObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + // Retain graph object with reference count 2 + HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2, + hipGraphUserObjectMove)); + + // Release graph object with reference count more than 2 + ret = hipGraphReleaseUserObject(graph, hObject, 4); + REQUIRE(hipSuccess == ret); + + // Again Retain graph object with reference count 8 + ret = hipGraphRetainUserObject(graph, hObject, 8, + hipGraphUserObjectMove); + REQUIRE(hipSuccess == ret); + // Release graph object with reference count 1 + ret = hipGraphReleaseUserObject(graph, hObject, 1); + REQUIRE(hipSuccess == ret); + + HIP_CHECK(hipUserObjectRelease(hObject, 1)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +TEST_CASE("Unit_hipGraphRetainUserObject_Negative_Null_Object") { + hipError_t ret; + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + float *object = nullptr; // this is used for Null_Object test + hipUserObject_t hObject; + + HIP_CHECK(hipUserObjectCreate(&hObject, object, + destroyFloatObj, + 1, hipUserObjectNoDestructorSync)); + REQUIRE(hObject != nullptr); + // Retain graph object with reference count 2 + HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 2, + hipGraphUserObjectMove)); + + // Release graph object with reference count more than 2 + ret = hipGraphReleaseUserObject(graph, hObject, 4); + REQUIRE(hipSuccess == ret); + + // Again Retain graph object with reference count 8 + ret = hipGraphRetainUserObject(graph, hObject, 8, 0); + REQUIRE(hipSuccess == ret); + // Release graph object with reference count 1 + ret = hipGraphReleaseUserObject(graph, hObject, 1); + REQUIRE(hipSuccess == ret); + + HIP_CHECK(hipUserObjectRelease(hObject, 1)); + HIP_CHECK(hipGraphDestroy(graph)); +} + From 955f1e4979335686e9342fe5842d55a4732bb231 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 20:16:16 +0530 Subject: [PATCH 8/9] SWDEV-346657 - add more hipStreamAcb test (#3064) * SWDEV-346657 - add more hipStreamAcb test Change-Id: If94628ec9c285e9689ce471931b3c40b7d3ece84 --- tests/catch/unit/stream/CMakeLists.txt | 2 + .../stream/hipStreamACb_StrmSyncTiming.cc | 138 ++++++++++++++++++ 2 files changed, 140 insertions(+) create mode 100644 tests/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 0fbb211593..548bb2e52c 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -18,6 +18,7 @@ set(TEST_SRC hipStreamQuery.cc hipStreamWaitEvent.cc hipDeviceGetStreamPriorityRange.cc + hipStreamACb_StrmSyncTiming.cc ) else() set(TEST_SRC @@ -38,6 +39,7 @@ set(TEST_SRC hipStreamSynchronize.cc hipStreamQuery.cc hipDeviceGetStreamPriorityRange.cc + hipStreamACb_StrmSyncTiming.cc ) # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) diff --git a/tests/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc b/tests/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc new file mode 100644 index 0000000000..a607604963 --- /dev/null +++ b/tests/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc @@ -0,0 +1,138 @@ +/* +Copyright (c) 2022 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, INNCLUDING 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 ANNY 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. +*/ + +/** +Testcase Scenario : +Validate behaviour of HIP when multiple hipStreaAddCallback() are called over +multiple Threads. +*/ + +#include +#include +#include +#include +#include + +#ifdef __HIP_PLATFORM_AMD__ +#define HIPRT_CB +#endif + +#define SECONDS_TO_WAIT 2 +#define TO_MICROSECONDS 1000000 + +hipStream_t mystream; +size_t N_elmts = 4096; +bool cbDone = false; +std::atomic Data_mismatch{0}; + +__global__ void vector_square(float* C_d, float* A_d, size_t N_elmts) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < N_elmts; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } + + // Delay the thread 1 + if (offset == 1) { + unsigned long long int wait_t = 3200000000, start = clock64(), cur; + do { + cur = clock64() - start; + } while (cur < wait_t); + } +} + +float *A_h, *C_h; + +static void HIPRT_CB Callback1(hipStream_t stream, hipError_t status, void* userData) { + (void)stream; + (void)status; + (void)userData; + // Validate the data + for (size_t i = 0; i < N_elmts; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + + // Delay the callback completion + std::this_thread::sleep_for(std::chrono::seconds(SECONDS_TO_WAIT)); + cbDone = true; +} + +/** + Test multiple hipStreamAddCallback() called over + multiple Threads. + */ +TEST_CASE("Unit_hipStreamAddCallback_StrmSyncTiming") { + float *A_d, *C_d; + size_t Nbytes = N_elmts * sizeof(float); + + A_h = (float*)malloc(Nbytes); + HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); + C_h = (float*)malloc(Nbytes); + HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); + + // Fill with Phi + i + for (size_t i = 0; i < N_elmts; i++) { + A_h[i] = 1.618f + i; + } + + HIPCHECK(hipMalloc(&A_d, Nbytes)); + HIPCHECK(hipMalloc(&C_d, Nbytes)); + + HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + + HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + + const unsigned threadsPerBlock = 256; + const unsigned blocks = (N_elmts + 255) / threadsPerBlock; + + hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0, mystream, C_d, A_d, + N_elmts); + HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + HIPCHECK(hipStreamAddCallback(mystream, Callback1, NULL, 0)); + + // Wait untill Callback() function changes the cbDone value to true + while (!cbDone) { + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + } + + // Since the callback is supposed to be called only after an implicit stream + // synchronization, and the runtime cannot continue until the callback is done + // hipStreamSynchronize call should not take much time. + auto start = std::chrono::high_resolution_clock::now(); + HIPCHECK(hipStreamSynchronize(mystream)); + auto stop = std::chrono::high_resolution_clock::now(); + auto duration = std::chrono::duration_cast(stop - start); + + HIPCHECK(hipStreamDestroy(mystream)); + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(C_d)); + free(A_h); + free(C_h); + + REQUIRE(Data_mismatch.load() == 0); + // HIP runtime cannot proceed further in the queue until callback completes + // Stream synchronize should not have much task to do after callback + // It should just be an extra empty marker wait + // Therefore the hipStreamSynchronize() in the + // main thread should hardly take any time to complete. + REQUIRE(duration.count() < 100); +} From 0e06303fe2762065c36b62ddbd6fa03367b2ea5b Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 28 Nov 2022 20:17:09 +0530 Subject: [PATCH 9/9] SWDEV-345188 - Windows: fix warnings while building catch (#3071) fixes deprecated warning for 'getenv' on Windows Change-Id: I5ecdaca451a936e843691c4611910ee338b22e46 --- tests/catch/hipTestMain/hip_test_context.cc | 7 +-- tests/catch/include/hip_test_context.hh | 43 ++++++++++++++----- tests/catch/unit/stream/hipStreamGetCUMask.cc | 8 ++-- 3 files changed, 41 insertions(+), 17 deletions(-) diff --git a/tests/catch/hipTestMain/hip_test_context.cc b/tests/catch/hipTestMain/hip_test_context.cc index 27419e4e88..bc38b32f94 100644 --- a/tests/catch/hipTestMain/hip_test_context.cc +++ b/tests/catch/hipTestMain/hip_test_context.cc @@ -4,6 +4,7 @@ #include #include #include +#include "hip_test_context.hh" #include "hip_test_filesystem.hh" void TestContext::detectOS() { @@ -98,11 +99,11 @@ void TestContext::getConfigFiles() { abort(); } - const char* env_config = std::getenv("HIP_CATCH_EXCLUDE_FILE"); + std::string env_config = TestContext::getEnvVar("HIP_CATCH_EXCLUDE_FILE"); LogPrintf("Env Config file: %s", - (env_config != nullptr) ? env_config : "Not found, using common config"); + (!env_config.empty()) ? env_config.c_str() : "Not found, using common config"); // HIP_CATCH_EXCLUDE_FILE is set for custom file path - if (env_config != nullptr) { + if (!env_config.empty()) { if(fs::exists(env_config)) { config_.json_files.push_back(env_config); } diff --git a/tests/catch/include/hip_test_context.hh b/tests/catch/include/hip_test_context.hh index 2ee221330e..931593cc4d 100644 --- a/tests/catch/include/hip_test_context.hh +++ b/tests/catch/include/hip_test_context.hh @@ -32,6 +32,7 @@ THE SOFTWARE. #include #include +// OS Check #if defined(_WIN32) #define HT_WIN 1 #define HT_LINUX 0 @@ -42,6 +43,7 @@ THE SOFTWARE. #error "OS not recognized" #endif +// Platform check #if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) #define HT_AMD 1 #define HT_NVIDIA 0 @@ -52,16 +54,6 @@ THE SOFTWARE. #error "Platform not recognized" #endif -static int _log_enable = (std::getenv("HT_LOG_ENABLE") ? 1 : 0); - -#define LogPrintf(format, ...) \ - { \ - if (_log_enable) { \ - printf(format, __VA_ARGS__); \ - printf("%c", '\n'); \ - } \ - } - typedef struct Config_ { std::vector json_files; // Json files std::string platform; // amd/nvidia @@ -124,6 +116,26 @@ class TestContext { return instance; } + static std::string getEnvVar(std::string var) { + #if defined(_WIN32) + rsize_t MAX_LEN = 4096; + char dstBuf[MAX_LEN]; + size_t dstSize; + if (!::getenv_s(&dstSize, dstBuf, MAX_LEN, var.c_str())) { + return std::string(dstBuf); + } + #elif defined(__linux__) + char* val = std::getenv(var.c_str()); + if (val != NULL) { + return std::string(val); + } + #else + #error "OS not recognized" + #endif + return std::string(""); + } + + bool isWindows() const; bool isLinux() const; bool isNvidia() const; @@ -171,3 +183,14 @@ class TestContext { ~TestContext(); }; + +static bool _log_enable = (!TestContext::getEnvVar("HT_LOG_ENABLE").empty() ? true : false); + +// printing logs +#define LogPrintf(format, ...) \ +{ \ + if(_log_enable) { \ + printf(format, __VA_ARGS__); \ + printf("%c", '\n'); \ + } \ +} \ No newline at end of file diff --git a/tests/catch/unit/stream/hipStreamGetCUMask.cc b/tests/catch/unit/stream/hipStreamGetCUMask.cc index f63acc212a..d7d10dd807 100644 --- a/tests/catch/unit/stream/hipStreamGetCUMask.cc +++ b/tests/catch/unit/stream/hipStreamGetCUMask.cc @@ -37,7 +37,7 @@ TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { std::vector cuMask(maxNum); hipDeviceProp_t props; std::stringstream ss; - char* gCUMask{nullptr}; + std::string gCUMask; std::string globalCUMask(""); std::vector defaultCUMask; @@ -54,8 +54,8 @@ TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { props.name << " with " << props.multiProcessorCount << " CUs"); // Get global CU Mask if exists - gCUMask = getenv("ROC_GLOBAL_CU_MASK"); - if (gCUMask != nullptr && gCUMask[0] != '\0') { + gCUMask = TestContext::getEnvVar("ROC_GLOBAL_CU_MASK"); + if (!gCUMask.empty()) { globalCUMask.assign(gCUMask); for_each(globalCUMask.begin(), globalCUMask.end(), [](char & c) { @@ -155,7 +155,7 @@ TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { INFO("info: reading back CU mask 0x" << ss.str() << " for stream " << stream); - if (!gCUMask) { + if (!gCUMask.empty()) { for (size_t i = 0; i < customMask.size(); i++) { if (customMask[i] != cuMask[i]) { INFO("Error! expected CU mask:" << customMask[i]