From decab3841a72afee59a7a9bc740bae934dd56d29 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 25 Jul 2022 13:25:25 +0530 Subject: [PATCH 01/10] Update Jenkinsfile Add backendLabel to stage name for better readability in pipeline views --- .jenkins/Jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/.jenkins/Jenkinsfile b/.jenkins/Jenkinsfile index a4c767be38..4f29347af6 100644 --- a/.jenkins/Jenkinsfile +++ b/.jenkins/Jenkinsfile @@ -28,7 +28,7 @@ def hipBuildTest(String backendLabel) { } } } - stage("Build - HIT framework") { + stage("Build - HIT framework ${backendLabel}") { // Running the build on hipamd workspace dir("${WORKSPACE}/hipamd") { sh """#!/usr/bin/env bash @@ -52,7 +52,7 @@ def hipBuildTest(String backendLabel) { } } timeout(time: 1, unit: 'HOURS') { - stage('HIP Unit Tests - HIT framework') { + stage('HIP Unit Tests - HIT framework ${backendLabel}') { dir("${WORKSPACE}/hipamd/build") { sh """#!/usr/bin/env bash set -x @@ -67,7 +67,7 @@ def hipBuildTest(String backendLabel) { } } } - stage("Build - Catch2 framework") { + stage("Build - Catch2 framework ${backendLabel}") { // Running the build on hipamd workspace dir("${WORKSPACE}/hipamd") { sh """#!/usr/bin/env bash @@ -94,7 +94,7 @@ def hipBuildTest(String backendLabel) { } } timeout(time: 1, unit: 'HOURS') { - stage('HIP Unit Tests - Catch2 framework') { + stage('HIP Unit Tests - Catch2 framework ${backendLabel}') { dir("${WORKSPACE}/hipamd/build") { sh """#!/usr/bin/env bash set -x From 15575466ac0ed7cd3d8aa746633874418a0f5cab Mon Sep 17 00:00:00 2001 From: searlmc1 Date: Mon, 25 Jul 2022 22:17:06 -0700 Subject: [PATCH 02/10] SWDEV-347383 - Remove code that sets -std and -fms-* flags (#2820) --- bin/hipcc.pl | 5 ----- 1 file changed, 5 deletions(-) diff --git a/bin/hipcc.pl b/bin/hipcc.pl index 645ae62d03..fa833c023d 100755 --- a/bin/hipcc.pl +++ b/bin/hipcc.pl @@ -187,11 +187,6 @@ if ($HIP_PLATFORM eq "amd") { print ("HIP_CLANG_TARGET=$HIP_CLANG_TARGET\n"); } - if ($isWindows) { - $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility"; - } else { - $HIPCXXFLAGS .= " -std=c++11"; - } $HIPCXXFLAGS .= " -isystem \"$HIP_CLANG_INCLUDE_PATH/..\""; $HIPCFLAGS .= " -isystem \"$HIP_CLANG_INCLUDE_PATH/..\""; $HIPLDFLAGS .= " -L\"$HIP_LIB_PATH\""; From e68136227a16bf436ff241d6f6cf459ae20c047f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 10:47:49 +0530 Subject: [PATCH 03/10] SWDEV-306122 - [catch2][dtest] hipGraphExecMemcpyNodeSetParamsToSymbol API (#2824) Change-Id: I3f533b88f9fadcb17bde6577212cf1561fa49273 Co-authored-by: mbhiutra --- .../config/config_amd_windows.json | 3 +- tests/catch/unit/graph/CMakeLists.txt | 1 + ...hipGraphExecMemcpyNodeSetParamsToSymbol.cc | 316 ++++++++++++++++++ 3 files changed, 319 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index c755d65649..71c8656333 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -44,6 +44,7 @@ "Unit_hipGraphKernelNodeSetParams_Functional", "Unit_hipGraphExecKernelNodeSetParams_Negative", "Unit_hipGraphExecKernelNodeSetParams_Functional", - "Unit_hipGraphMemcpyNodeSetParams1D_Functional" + "Unit_hipGraphMemcpyNodeSetParams1D_Functional", + "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional" ] } diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index b07f15891b..dddb5af1ba 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -73,6 +73,7 @@ set(TEST_SRC hipGraphExecKernelNodeSetParams.cc hipGraphLaunch.cc hipGraphMemcpyNodeSetParams1D.cc + hipGraphExecMemcpyNodeSetParamsToSymbol.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc b/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc new file mode 100644 index 0000000000..e83bfe9cea --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsToSymbol.cc @@ -0,0 +1,316 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios of hipGraphExecMemcpyNodeSetParamsToSymbol API: +Functional : +1) Allocate global symbol memory, Instantiate a graph with memcpy node, + obtain executable graph and update the node params with set exec api call. + Make sure they are taking effect. +2) Allocate const symbol memory, Instantiate a graph with memcpy node, + obtain executable graph and update the node params with set exec api call. + Make sure they are taking effect. +Negative : +1) Pass hGraphExec as nullptr and check if api returns error. +2) Pass GraphNode as nullptr and check if api returns error. +3) Pass symbol ptr as nullptr, api expected to return error code. +4) Pass source ptr as nullptr, api expected to return error code. +5) Pass count as zero, api expected to return error code. +6) Pass offset+count greater than allocated size, api expected to return error code. +7) Pass same symbol pointer as source ptr and destination ptr, api expected to return error code. +8) Pass Pass both dstn ptr and source ptr as 2 different symbol ptr, api expected to return error code. +9) Copy from device ptr to host ptr but pass kind as different, api expected to return error code. +10) Check with other graph node but pass same graphExec, api expected to return error code. +*/ + +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE], globalOut[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + +__global__ void MemcpyToSymbolExecKernel(int* B_d) { + for (int i = 0 ; i < SIZE; i++) { + B_d[i] = globalIn[i]; + } +} + +__global__ void MemcpyToConstSymbolExecKernel(int* B_d) { + for (int i = 0 ; i < SIZE; i++) { + B_d[i] = globalConst[i]; + } +} + +/* This testcase verifies negative scenarios of + hipGraphExecMemcpyNodeSetParamsToSymbol API */ +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipError_t ret; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + C_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + HIP_CHECK(hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), A_d, + Nbytes, 0, + hipMemcpyDeviceToDevice)); + + // Instantiate the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + SECTION("Pass hGraphExec as nullptr") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(nullptr, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass GraphNode as nullptr") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + nullptr, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass symbol ptr as nullptr") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + nullptr, + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidSymbol == ret); + } + SECTION("Pass source ptr as nullptr") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count as zero") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, 0, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass offset+count greater than allocated size") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 10, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass same symbol pointer as source ptr and destination ptr") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass both dstn ptr and source ptr as 2 different symbol ptr") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + HIP_SYMBOL(globalOut), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Copy from device ptr to host ptr but pass kind as different") { + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_h, + Nbytes, 0, + hipMemcpyDeviceToHost); + REQUIRE(hipSuccess != ret); + } + SECTION("Check with other graph node but pass same graphExec") { + hipGraph_t graph1; + hipGraphNode_t memcpyToSymbolNode1{}; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode1, graph1, + nullptr, + 0, + HIP_SYMBOL(globalOut), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + ret = hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode1, + HIP_SYMBOL(globalOut), + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + + HipTest::freeArrays(A_d, B_d, C_d, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +static +void hipGraphExecMemcpyNodeSetParamsToSymbol_GlobalMem(bool useConstVar) { + constexpr size_t Nbytes = SIZE * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE); + hipGraphNode_t memcpytosymbolkernel, memcpyD2H_B; + hipKernelNodeParams kernelNodeParams{}; + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + if (useConstVar) { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalConst), + C_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + C_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + // Adding Kernel node + void* kernelArgs1[] = {&B_d}; + if (useConstVar) { + kernelNodeParams.func = + reinterpret_cast(MemcpyToConstSymbolExecKernel); + } else { + kernelNodeParams.func = reinterpret_cast(MemcpyToSymbolExecKernel); + } + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memcpytosymbolkernel, graph, + dependencies.data(), dependencies.size(), + &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(memcpytosymbolkernel); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(), + dependencies.size(), B_h, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + // Update the node with source pointer from C_d to A_d + if (useConstVar) { + HIP_CHECK(hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalConst), A_d, + Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphExecMemcpyNodeSetParamsToSymbol(graphExec, + memcpyToSymbolNode, + HIP_SYMBOL(globalIn), A_d, + Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, C_d, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphExecMemcpyNodeSetParamsToSymbol Functional scenario. +1) Allocate global symbol memory, Instantiate a graph with memcpy node, + obtain executable graph and update the node params with set exec api call. + Make sure they are taking effect. +2) Allocate const symbol memory, Instantiate a graph with memcpy node, + obtain executable graph and update the node params with set exec api call. + Make sure they are taking effect. + */ +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional") { + SECTION("Check and update with Global Device Symbol Memory") { + hipGraphExecMemcpyNodeSetParamsToSymbol_GlobalMem(false); + } + SECTION("Check and update with Constant Global Device Symbol Memory") { + hipGraphExecMemcpyNodeSetParamsToSymbol_GlobalMem(true); + } +} + From c5d856348e4c0132696692ae9adcfcd84e7aafac Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 11:46:32 +0530 Subject: [PATCH 04/10] SWDEV-344211 - Fix hipHostMalloc tests (#2823) - Increase allocation size to 30MB to make sure new page table allocations happen - Change validation criteria to check for increase in VRAM usage for visible HIP devices Change-Id: Ie5bf1a89ceda0ff50140f426235047adcca0fdea --- tests/catch/multiproc/hipHostMallocTestsMproc.cc | 8 ++++++-- tests/src/runtimeApi/memory/hipHostMallocTests.cpp | 8 ++++++-- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/tests/catch/multiproc/hipHostMallocTestsMproc.cc b/tests/catch/multiproc/hipHostMallocTestsMproc.cc index cca3c396aa..2e9a186187 100644 --- a/tests/catch/multiproc/hipHostMallocTestsMproc.cc +++ b/tests/catch/multiproc/hipHostMallocTestsMproc.cc @@ -41,7 +41,7 @@ Testcase Scenarios : * Defines */ #define LIB_ROCMSMI "librocm_smi64.so" -#define ALLOC_SIZE (1*1024*1024) +#define ALLOC_SIZE (30*1024*1024) /** * Global variables @@ -221,7 +221,11 @@ static bool validatePageTableAllocations(const char *devList, int visDevCnt) { } for (indx = 0; indx < numdev; indx++) { - if (current[indx] - prev[indx]) + // For visible hip devices, there should be increase in VRAM usage + // due to page table allocations + // For NON visible hip devices, there can be reduction in VRAM usage + // due to removal of page tables from them + if (current[indx] > prev[indx]) changeCnt++; } diff --git a/tests/src/runtimeApi/memory/hipHostMallocTests.cpp b/tests/src/runtimeApi/memory/hipHostMallocTests.cpp index 7243813ef3..013b533529 100644 --- a/tests/src/runtimeApi/memory/hipHostMallocTests.cpp +++ b/tests/src/runtimeApi/memory/hipHostMallocTests.cpp @@ -61,7 +61,7 @@ Testcase Scenarios : */ #define LIB_ROCMSMI "librocm_smi64.so" #define NUM_BYTES 1000 -#define ALLOC_SIZE (1*1024*1024) +#define ALLOC_SIZE (30*1024*1024) /** @@ -246,7 +246,11 @@ bool validatePageTableAllocations(const char *devList, int devCnt) { } for (indx = 0; indx < numdev; indx++) { - if (current[indx] - prev[indx]) + // For visible hip devices, there should be increase in VRAM usage + // due to page table allocations + // For NON visible hip devices, there can be reduction in VRAM usage + // due to removal of page tables from them + if (current[indx] > prev[indx]) changeCnt++; } From 21ec83b73442c4b6688772f4ae87b5579138297c Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 11:47:03 +0530 Subject: [PATCH 05/10] SWDEV-327563 - Enable fixed tests linux and windows v2 (#2822) Change-Id: I094b5c8bda458917a074af9b8638074362c1aa56 --- .../catch/hipTestMain/config/config_amd_windows.json | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 71c8656333..ac7afac6c5 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -5,19 +5,11 @@ "Unit_printf_flags", "Unit_printf_specifier", "Unit_hipTextureMipmapObj2D_Check", - "Unit_hipManagedKeyword_MultiGpu", "Unit_hipGraphAddHostNode_ClonedGraphwithHostNode", "Unit_hipEventIpc", "Unit_hipMalloc3D_Negative", "Unit_hipPointerGetAttribute_MappedMem", - "Unit_hipStreamBeginCapture_BasicFunctional", - "Unit_hipStreamBeginCapture_hipStreamPerThread", - "Unit_hiprtc_functional", "Unit_hipStreamValue_Write", - "Unit_hipStreamGetCaptureInfo_BasicFunctional", - "Unit_hipStreamGetCaptureInfo_hipStreamPerThread", - "Unit_hipStreamGetCaptureInfo_UniqueID", - "Unit_hipStreamGetCaptureInfo_ArgValidation", "Unit_hipMemPoolApi_Basic", "Unit_hipMemPoolApi_BasicAlloc", "Unit_hipMemPoolApi_BasicTrim", @@ -41,10 +33,6 @@ "Unit_hipArrayCreate_happy - float2", "Unit_hipArrayCreate_happy - float4", "Unit_hipMemVmm_Basic", - "Unit_hipGraphKernelNodeSetParams_Functional", - "Unit_hipGraphExecKernelNodeSetParams_Negative", - "Unit_hipGraphExecKernelNodeSetParams_Functional", - "Unit_hipGraphMemcpyNodeSetParams1D_Functional", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional" ] } From d029b1e039b8023ec9cc863185370b77c08b119b Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 11:47:18 +0530 Subject: [PATCH 06/10] SWDEV-338376 - Driver Alias for Resourcetype (#2821) Change-Id: I621fe95020a3b2d82034cf5979e834b6e13f2981 --- include/hip/driver_types.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/driver_types.h b/include/hip/driver_types.h index fb02516354..30b3977a12 100644 --- a/include/hip/driver_types.h +++ b/include/hip/driver_types.h @@ -142,7 +142,7 @@ typedef enum HIPresourcetype_enum { HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY = 0x01, /**< Mipmapped array resource */ HIP_RESOURCE_TYPE_LINEAR = 0x02, /**< Linear resource */ HIP_RESOURCE_TYPE_PITCH2D = 0x03 /**< Pitch 2D resource */ -} HIPresourcetype; +} HIPresourcetype, hipResourcetype; /** * hip address modes */ From 7db8bde590d4465d68f884109cf317a5e6fd194a Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 12:14:26 +0530 Subject: [PATCH 07/10] SWDEV-344691 - hipMalloc3DArray fails if width is more than capacity. (#2819) Change-Id: I3eeaa9ef2b3197b3e5c5f03ccc1ba069dfb64fb7 --- tests/catch/hipTestMain/config/config_amd_windows.json | 8 ++++++++ tests/catch/unit/memory/hipMemcpy3D.cc | 8 ++++++-- tests/catch/unit/memory/hipMemcpy3DAsync.cc | 6 +++++- 3 files changed, 19 insertions(+), 3 deletions(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index ac7afac6c5..6c129f8778 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -34,5 +34,13 @@ "Unit_hipArrayCreate_happy - float4", "Unit_hipMemVmm_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional" + "Unit_hipMallocManaged_MultiChunkMultiDevice", + "Unit_hipMallocManaged_TwoPointers - int", + "Unit_hipMallocManaged_TwoPointers - float", + "Unit_hipMallocManaged_TwoPointers - double", + "Unit_hipMallocManaged_DeviceContextChange - unsigned char", + "Unit_hipMallocManaged_DeviceContextChange - int", + "Unit_hipMallocManaged_DeviceContextChange - float", + "Unit_hipMallocManaged_DeviceContextChange - double" ] } diff --git a/tests/catch/unit/memory/hipMemcpy3D.cc b/tests/catch/unit/memory/hipMemcpy3D.cc index 75fc8efe0c..4af0883639 100644 --- a/tests/catch/unit/memory/hipMemcpy3D.cc +++ b/tests/catch/unit/memory/hipMemcpy3D.cc @@ -552,9 +552,13 @@ void Memcpy3D::simple_Memcpy3D() { */ TEMPLATE_TEST_CASE("Unit_hipMemcpy3D_Basic", "[hipMemcpy3D]", int, unsigned int, float) { - int numDevices = 0; - auto i = GENERATE(10, 100, 1024, 10*1024); + int device = -1; + HIP_CHECK(hipGetDevice(&device)); + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop,device)); + auto i = GENERATE_COPY(10, 100, 1024, prop.maxTexture3D[0]); auto j = GENERATE(10, 100); + int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { if (std::is_same::value) { diff --git a/tests/catch/unit/memory/hipMemcpy3DAsync.cc b/tests/catch/unit/memory/hipMemcpy3DAsync.cc index 5918a287c9..1f690fa410 100644 --- a/tests/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy3DAsync.cc @@ -662,7 +662,11 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy3DAsync_Basic", int, unsigned int, float) { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); - auto i = GENERATE(10, 100, 1024, 10*1024); + int device = -1; + HIP_CHECK(hipGetDevice(&device)); + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop,device)); + auto i = GENERATE_COPY(10, 100, 1024, prop.maxTexture3D[0]); auto j = GENERATE(10, 100); if (numDevices > 1) { if (std::is_same::value) { From e411ebc467afd5503bb829ad2da6c71c43db420f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 15:07:42 +0530 Subject: [PATCH 08/10] SWDEV-343380 - Added hipDrvGetErrorName/String (#2825) Change-Id: I1f6df9e801d02e56635c6e9546b71e08b31ea7ce --- include/hip/hip_runtime_api.h | 24 +++++++++++++++++++++--- 1 file changed, 21 insertions(+), 3 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 53d88f65aa..55b430aada 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1842,7 +1842,7 @@ hipError_t hipGetLastError(void); */ hipError_t hipPeekAtLastError(void); /** - * @brief Return name of the specified error code in text form. + * @brief Return hip error as text string form. * * @param hip_error Error code to convert to name. * @return const char pointer to the NULL-terminated error name @@ -1856,11 +1856,29 @@ const char* hipGetErrorName(hipError_t hip_error); * @param hipError Error code to convert to string. * @return const char pointer to the NULL-terminated error string * - * @warning : on HCC, this function returns the name of the error (same as hipGetErrorName) - * * @see hipGetErrorName, hipGetLastError, hipPeakAtLastError, hipError_t */ const char* hipGetErrorString(hipError_t hipError); +/** + * @brief Return hip error as text string form. + * + * @param [in] hipError Error code to convert to string. + * @param [out] const char pointer to the NULL-terminated error string + * @return #hipSuccess, #hipErrorInvalidValue + * + * @see hipGetErrorName, hipGetLastError, hipPeakAtLastError, hipError_t + */ +hipError_t hipDrvGetErrorName(hipError_t hipError, const char** errorString); +/** + * @brief Return handy text string message to explain the error which occurred + * + * @param [in] hipError Error code to convert to string. + * @param [out] const char pointer to the NULL-terminated error string + * @return #hipSuccess, #hipErrorInvalidValue + * + * @see hipGetErrorName, hipGetLastError, hipPeakAtLastError, hipError_t + */ +hipError_t hipDrvGetErrorString(hipError_t hipError, const char** errorString); // end doxygen Error /** * @} From a2e3642e5b4fd3bfd39063100e81f935451e54b6 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 15:08:04 +0530 Subject: [PATCH 09/10] SWDEV-306122 - [catch2][dtest] hipGraph tests for hipGraphNodeGetDependentNodes api (#2826) Change-Id: Ie9a81fa8829d48ace19ac3fa522f957504c3ccc0 --- .../config/config_amd_windows.json | 4 +- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphNodeGetDependentNodes.cc | 366 ++++++++++++++++++ 3 files changed, 370 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphNodeGetDependentNodes.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 6c129f8778..1393a73ae7 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -41,6 +41,8 @@ "Unit_hipMallocManaged_DeviceContextChange - unsigned char", "Unit_hipMallocManaged_DeviceContextChange - int", "Unit_hipMallocManaged_DeviceContextChange - float", - "Unit_hipMallocManaged_DeviceContextChange - double" + "Unit_hipMallocManaged_DeviceContextChange - double", + "Unit_hipGraphNodeGetDependentNodes_Functional", + "Unit_hipGraphNodeGetDependentNodes_ParamValidation" ] } diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index dddb5af1ba..159a475c8c 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -74,6 +74,7 @@ set(TEST_SRC hipGraphLaunch.cc hipGraphMemcpyNodeSetParams1D.cc hipGraphExecMemcpyNodeSetParamsToSymbol.cc + hipGraphNodeGetDependentNodes.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphNodeGetDependentNodes.cc b/tests/catch/unit/graph/hipGraphNodeGetDependentNodes.cc new file mode 100644 index 0000000000..d71ee72489 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphNodeGetDependentNodes.cc @@ -0,0 +1,366 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + +Testcase Scenarios +------------------ +Functional: +1) Create a graph and add nodes with dependencies. Query for dependent nodes of the node passed and verify the result with dependencies defined. +2) When pDependentNodes is passed as nullptr, verify pNumDependentNodes returns the number of dependent nodes. +3) When pNumDependentNodes is higher than the actual number of dependent nodes, the remaining entries in pDependentNodes will be set to NULL, + and the number of nodes actually obtained will be returned in pNumDependentNodes. +4) When pNumDependentNodes is lesser than the actual number of dependent nodes, api should return the requested number of nodes in pDependentNodes. + +Argument Validation: +1) Add a single node in graph and pass the node to api. Verify the api returns dependent nodes as 0. +2) Pass node as nullptr and verify api doesn’t crash, returns error code. +3) Pass pNumDependentNodes as nullptr and verify api doesn’t crash, returns error code. +4) Pass node as un-initialized/invalid parameter and verify api returns error code. + +*/ + +#include +#include +#include + +static __global__ void updateResult(int* C_d, int* Res_d, int val, + int64_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (int64_t i = NELEM - stride + offset; i >= 0; i -= stride) { + Res_d[i] = C_d[i] + val; + } +} + +static __global__ void vectorSum(const int* A_d, const int* B_d, + const int* C_d, int* Res_d, size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < NELEM; i += stride) { + Res_d[i] = A_d[i] + B_d[i] + C_d[i]; + } +} + +/** + * Verify api when GetDependent nodes is requested + * for actual number of nodes. + */ +static void queryActualNumOfDepNodes(const std::vector &Nlist, + hipGraphNode_t kernel_vecSqr, size_t numDeps) { + hipGraphNode_t* depnodes; + int numBytes = sizeof(hipGraphNode_t) * numDeps; + depnodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(depnodes != nullptr); + HIP_CHECK(hipGraphNodeGetDependentNodes(kernel_vecSqr, depnodes, &numDeps)); + REQUIRE(numDeps == Nlist.size()); + + // Verify all dependent nodes are present in the node entries returned + for (auto Node : Nlist) { + bool found = false; + for (size_t i = 0; i < numDeps; i++) { + if (Node == depnodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Dependent node " << Node << " not present in returned list"); + REQUIRE(false); + } + } + free(depnodes); +} + +/** + * Verify api when GetDependent nodes queried + * for greater number than actual number of nodes. + */ +static void queryGreaterNumOfDepNodes(const std::vector &Nlist, + hipGraphNode_t kernel_vecSqr, size_t numDeps) { + constexpr auto addlEntries = 4; + hipGraphNode_t* depnodes; + size_t totDeps = numDeps + addlEntries; + int numBytes = sizeof(hipGraphNode_t) * totDeps; + depnodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(depnodes != nullptr); + HIP_CHECK(hipGraphNodeGetDependentNodes(kernel_vecSqr, depnodes, &totDeps)); + REQUIRE(totDeps == Nlist.size()); + + for (auto i = numDeps; i < numDeps + addlEntries; i++) { + REQUIRE(depnodes[i] == nullptr); + } + + // Verify all dependent nodes are present in the node entries returned + for (auto Node : Nlist) { + bool found = false; + for (size_t i = 0; i < numDeps; i++) { + if (Node == depnodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Dependent node " << Node << " not present in returned list"); + REQUIRE(false); + } + } + free(depnodes); +} + +/** + * Verify api when GetDependent nodes queried + * for lesser number than actual number of nodes. + */ +static void queryLesserNumOfDepNodes(const std::vector &Nlist, + hipGraphNode_t kernel_vecSqr, size_t numDeps) { + size_t totDeps = numDeps - 1; + hipGraphNode_t* depnodes; + int numBytes = sizeof(hipGraphNode_t) * totDeps; + size_t count{}; + depnodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(depnodes != nullptr); + HIP_CHECK(hipGraphNodeGetDependentNodes(kernel_vecSqr, depnodes, &totDeps)); + REQUIRE(totDeps == Nlist.size() - 1); + + // Verify all dependent nodes are present in the node entries returned + for (auto Node : Nlist) { + for (size_t i = 0; i < totDeps; i++) { + if (Node == depnodes[i]) { + count++; + break; + } + } + } + REQUIRE(count == totDeps); + free(depnodes); +} + +/** + * Functional Test for getting dependent nodes in graph and verifying execution + */ +TEST_CASE("Unit_hipGraphNodeGetDependentNodes_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraphNode_t kernel_vecSqr{}, kernel_vecAdd{}; + hipGraphNode_t kernelmod1{}, kernelmod2{}, kernelmod3{}; + hipGraphNode_t memcpyD2H{}, memcpyH2D_A{}; + hipKernelNodeParams kernelNodeParams{}; + hipGraph_t graph{}; + size_t numDeps{}; + hipStream_t streamForGraph; + int *A_d, *C_d; + int *A_h, *C_h; + int *Res1_d, *Res2_d, *Res3_d; + int *Sum_d, *Sum_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HipTest::initArrays(&A_d, &C_d, &Sum_d, &A_h, &C_h, &Sum_h, N); + HipTest::initArrays(&Res1_d, &Res2_d, &Res3_d, + nullptr, nullptr, nullptr, N); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + // Initialize input buffer and vecsqr result + for (size_t i = 0; i < N; ++i) { + A_h[i] = i + 1; + C_h[i] = A_h[i] * A_h[i]; + } + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgsVS[] = {&A_d, &C_d, reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(HipTest::vector_square); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgsVS); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSqr, graph, &memcpyH2D_A, 1, + &kernelNodeParams)); + + // Create multiple nodes dependent on vecSqr node. + // Dependent nodes takes vecSqr input and computes output independently. + std::vector nodelist; + int incValue1{1}; + void* kernelArgs1[] = {&C_d, &Res1_d, &incValue1, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(updateResult); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelmod1, graph, &kernel_vecSqr, 1, + &kernelNodeParams)); + nodelist.push_back(kernelmod1); + + int incValue2{2}; + void* kernelArgs2[] = {&C_d, &Res2_d, &incValue2, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(updateResult); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelmod2, graph, &kernel_vecSqr, 1, + &kernelNodeParams)); + nodelist.push_back(kernelmod2); + + int incValue3{3}; + void* kernelArgs3[] = {&C_d, &Res3_d, &incValue3, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(updateResult); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs3); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelmod3, graph, &kernel_vecSqr, 1, + &kernelNodeParams)); + nodelist.push_back(kernelmod3); + + HIP_CHECK(hipGraphNodeGetDependentNodes(kernel_vecSqr, nullptr, &numDeps)); + REQUIRE(numDeps == nodelist.size()); + + // Verify api When Dependent nodes are requested for actual number of nodes. + queryActualNumOfDepNodes(nodelist, kernel_vecSqr, numDeps); + + // Verify api When Dependent nodes are requested for more than + // actual number of nodes. + queryGreaterNumOfDepNodes(nodelist, kernel_vecSqr, numDeps); + + // Verify api When Dependent nodes are requested for less than + // actual number of nodes. + queryLesserNumOfDepNodes(nodelist, kernel_vecSqr, numDeps); + + // Compute sum from all dependent nodes + void* kernelArgsAdd[] = {&Res1_d, &Res2_d, &Res3_d, &Sum_d, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(vectorSum); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgsAdd); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, + nodelist.data(), nodelist.size(), + &kernelNodeParams)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, &kernel_vecAdd, 1, + Sum_h, Sum_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if ( Sum_h[i] != ( (C_h[i] + incValue1) + + (C_h[i] + incValue2) + + (C_h[i] + incValue3) ) ) { + INFO("Sum not matching at " << i << " Sum_h[i] " << Sum_h[i] + << " C_h[i] " << C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, C_d, Sum_d, A_h, C_h, Sum_h, false); + HipTest::freeArrays(Res1_d, Res2_d, Res3_d, + nullptr, nullptr, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Test performs api parameter validation by passing various values + * as input and output parameters and validates the behavior. + * Test will include both negative and positive scenarios. + */ +TEST_CASE("Unit_hipGraphNodeGetDependentNodes_ParamValidation") { + hipGraph_t graph{}; + const int numBytes = 100; + size_t numDeps{1}; + hipGraphNode_t memsetNode{}, depnodes{}; + hipError_t ret{}; + char *A_d; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipMalloc(&A_d, numBytes)); + hipMemsetParams memsetParams{}; + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 1; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numBytes * sizeof(char); + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, + 0, &memsetParams)); + + SECTION("single node in graph") { + ret = hipGraphNodeGetDependentNodes(memsetNode, &depnodes, &numDeps); + + // Api expected to return success and no dependent nodes. + REQUIRE(ret == hipSuccess); + REQUIRE(numDeps == 0); + } + + SECTION("node as nullptr") { + ret = hipGraphNodeGetDependentNodes(nullptr, &depnodes, &numDeps); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("NumDependentNodes as nullptr") { + ret = hipGraphNodeGetDependentNodes(memsetNode, &depnodes, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("node as un-initialized/invalid parameter") { + hipGraphNode_t uninit_node{}; + ret = hipGraphNodeGetDependentNodes(uninit_node, &depnodes, &numDeps); + REQUIRE(ret == hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(A_d)); +} From 1c4822f24a99448c14f375ac9dd54cde2c598a54 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 26 Jul 2022 15:09:35 +0530 Subject: [PATCH 10/10] SWDEV-306122 - [catch2][dtest] hipGraph tests for hipGraphNodeGetDependencies api (#2827) Change-Id: I017ba7891c84adb4d24d857b31800723bd43ed59 --- .../config/config_amd_windows.json | 6 +- tests/catch/unit/graph/CMakeLists.txt | 1 + .../unit/graph/hipGraphNodeGetDependencies.cc | 368 ++++++++++++++++++ 3 files changed, 373 insertions(+), 2 deletions(-) create mode 100644 tests/catch/unit/graph/hipGraphNodeGetDependencies.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 1393a73ae7..057a8a97f6 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -33,7 +33,7 @@ "Unit_hipArrayCreate_happy - float2", "Unit_hipArrayCreate_happy - float4", "Unit_hipMemVmm_Basic", - "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional" + "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional", "Unit_hipMallocManaged_MultiChunkMultiDevice", "Unit_hipMallocManaged_TwoPointers - int", "Unit_hipMallocManaged_TwoPointers - float", @@ -43,6 +43,8 @@ "Unit_hipMallocManaged_DeviceContextChange - float", "Unit_hipMallocManaged_DeviceContextChange - double", "Unit_hipGraphNodeGetDependentNodes_Functional", - "Unit_hipGraphNodeGetDependentNodes_ParamValidation" + "Unit_hipGraphNodeGetDependentNodes_ParamValidation", + "Unit_hipGraphNodeGetDependencies_Functional", + "Unit_hipGraphNodeGetDependencies_ParamValidation" ] } diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 159a475c8c..07e62fe261 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -75,6 +75,7 @@ set(TEST_SRC hipGraphMemcpyNodeSetParams1D.cc hipGraphExecMemcpyNodeSetParamsToSymbol.cc hipGraphNodeGetDependentNodes.cc + hipGraphNodeGetDependencies.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphNodeGetDependencies.cc b/tests/catch/unit/graph/hipGraphNodeGetDependencies.cc new file mode 100644 index 0000000000..572e1fba0b --- /dev/null +++ b/tests/catch/unit/graph/hipGraphNodeGetDependencies.cc @@ -0,0 +1,368 @@ +/* +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, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + +Testcase Scenarios +------------------ +Functional: +1) Create a graph and add nodes with dependencies. Try to fetch dependencies of a node added and verify the api returns the node's + dependencies that were defined. +2) When pDependencies is passed as nullptr, verify pNumDependencies returns actual number of dependencies of node. +3) When pNumDependencies is higher than the actual number of dependencies, the remaining entries in pDependencies will be set to NULL, + and the number of nodes actually obtained will be returned in pNumDependencies. +4) When pNumDependencies is lesser than the actual number of dependencies, api should return the requested number of dependencies. + +Argument Validation: +1) Verify the api returns pNumDependencies(0) when node passed is a root node. +2) Pass node as nullptr and verify api doesn’t crash, returns error code. +3) Pass pNumDependencies as nullptr and verify api doesn’t crash, returns error code. +4) Pass node as un-initialized/invalid parameter and verify api returns error code. + +*/ + +#include +#include +#include + +static __global__ void updateResult(int* C_d, int* Res_d, int val, + int64_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (int64_t i = NELEM - stride + offset; i >= 0; i -= stride) { + Res_d[i] = C_d[i] + val; + } +} + +static __global__ void vectorSum(const int* A_d, const int* B_d, + const int* C_d, int* Res_d, size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < NELEM; i += stride) { + Res_d[i] = A_d[i] + B_d[i] + C_d[i]; + } +} + +/** + * Verify api when GetDependencies is requested + * for actual number of nodes. + */ +static void queryActualNumOfDeps(const std::vector &Nlist, + hipGraphNode_t kernel_vecAdd, size_t numDeps) { + hipGraphNode_t* depnodes; + int numBytes = sizeof(hipGraphNode_t) * numDeps; + depnodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(depnodes != nullptr); + HIP_CHECK(hipGraphNodeGetDependencies(kernel_vecAdd, depnodes, &numDeps)); + REQUIRE(numDeps == Nlist.size()); + + // Verify all dependencies are present in the node entries returned + for (auto Node : Nlist) { + bool found = false; + for (size_t i = 0; i < numDeps; i++) { + if (Node == depnodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Dependency node " << Node << " not present in returned list"); + REQUIRE(false); + } + } + free(depnodes); +} + +/** + * Verify api when GetDependencies queried + * for greater number than actual number of nodes. + */ +static void queryGreaterNumOfDeps(const std::vector &Nlist, + hipGraphNode_t kernel_vecAdd, size_t numDeps) { + constexpr auto addlEntries = 4; + hipGraphNode_t* depnodes; + size_t totDeps = numDeps + addlEntries; + int numBytes = sizeof(hipGraphNode_t) * totDeps; + depnodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(depnodes != nullptr); + HIP_CHECK(hipGraphNodeGetDependencies(kernel_vecAdd, depnodes, &totDeps)); + REQUIRE(totDeps == Nlist.size()); + + for (auto i = numDeps; i < numDeps + addlEntries; i++) { + REQUIRE(depnodes[i] == nullptr); + } + + // Verify all dependencies are present in the node entries returned + for (auto Node : Nlist) { + bool found = false; + for (size_t i = 0; i < numDeps; i++) { + if (Node == depnodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Dependency node " << Node << " not present in returned list"); + REQUIRE(false); + } + } + free(depnodes); +} + +/** + * Verify api when GetDependencies queried + * for lesser number than actual number of nodes. + */ +static void queryLesserNumOfDeps(const std::vector &Nlist, + hipGraphNode_t kernel_vecAdd, size_t numDeps) { + size_t totDeps = numDeps - 1; + hipGraphNode_t* depnodes; + int numBytes = sizeof(hipGraphNode_t) * totDeps; + size_t count{}; + depnodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(depnodes != nullptr); + HIP_CHECK(hipGraphNodeGetDependencies(kernel_vecAdd, depnodes, &totDeps)); + REQUIRE(totDeps == Nlist.size() - 1); + + // Verify all dependencies are present in the node entries returned + for (auto Node : Nlist) { + for (size_t i = 0; i < totDeps; i++) { + if (Node == depnodes[i]) { + count++; + break; + } + } + } + REQUIRE(count == totDeps); + free(depnodes); +} + +/** + * Functional Test for getting dependencies of node in graph and verifying execution + */ +TEST_CASE("Unit_hipGraphNodeGetDependencies_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraphNode_t kernel_vecSqr{}, kernel_vecAdd{}; + hipGraphNode_t kernelmod1{}, kernelmod2{}, kernelmod3{}; + hipGraphNode_t memcpyD2H{}, memcpyH2D_A{}; + hipKernelNodeParams kernelNodeParams{}; + hipGraph_t graph{}; + size_t numDeps{}; + hipStream_t streamForGraph; + int *A_d, *C_d; + int *A_h, *C_h; + int *Res1_d, *Res2_d, *Res3_d; + int *Sum_d, *Sum_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HipTest::initArrays(&A_d, &C_d, &Sum_d, &A_h, &C_h, &Sum_h, N); + HipTest::initArrays(&Res1_d, &Res2_d, &Res3_d, + nullptr, nullptr, nullptr, N); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + // Initialize input buffer and vecsqr result + for (size_t i = 0; i < N; ++i) { + A_h[i] = i + 1; + C_h[i] = A_h[i] * A_h[i]; + } + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgsVS[] = {&A_d, &C_d, reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(HipTest::vector_square); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgsVS); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSqr, graph, &memcpyH2D_A, 1, + &kernelNodeParams)); + + // Create multiple nodes dependent on vecSqr node. + // Dependent nodes takes vecSqr input and computes output independently. + std::vector nodelist; + int incValue1{1}; + void* kernelArgs1[] = {&C_d, &Res1_d, &incValue1, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(updateResult); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelmod1, graph, &kernel_vecSqr, 1, + &kernelNodeParams)); + nodelist.push_back(kernelmod1); + + int incValue2{2}; + void* kernelArgs2[] = {&C_d, &Res2_d, &incValue2, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(updateResult); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelmod2, graph, &kernel_vecSqr, 1, + &kernelNodeParams)); + nodelist.push_back(kernelmod2); + + int incValue3{3}; + void* kernelArgs3[] = {&C_d, &Res3_d, &incValue3, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(updateResult); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs3); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelmod3, graph, &kernel_vecSqr, 1, + &kernelNodeParams)); + nodelist.push_back(kernelmod3); + + // Compute sum from all dependent nodes + void* kernelArgsAdd[] = {&Res1_d, &Res2_d, &Res3_d, &Sum_d, + reinterpret_cast(&NElem)}; + memset(&kernelNodeParams, 0, sizeof(kernelNodeParams)); + kernelNodeParams.func = + reinterpret_cast(vectorSum); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgsAdd); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, + nodelist.data(), nodelist.size(), + &kernelNodeParams)); + + HIP_CHECK(hipGraphNodeGetDependencies(kernel_vecAdd, nullptr, &numDeps)); + REQUIRE(numDeps == nodelist.size()); + + // Verify api When Dependencies are requested for actual number of nodes. + queryActualNumOfDeps(nodelist, kernel_vecAdd, numDeps); + + // Verify api When Dependencies are requested for more than + // actual number of nodes. + queryGreaterNumOfDeps(nodelist, kernel_vecAdd, numDeps); + + // Verify api When Dependencies are requested for less than + // actual number of nodes. + queryLesserNumOfDeps(nodelist, kernel_vecAdd, numDeps); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, &kernel_vecAdd, 1, + Sum_h, Sum_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if ( Sum_h[i] != ( (C_h[i] + incValue1) + + (C_h[i] + incValue2) + + (C_h[i] + incValue3) ) ) { + INFO("Sum not matching at " << i << " Sum_h[i] " << Sum_h[i] + << " C_h[i] " << C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, C_d, Sum_d, A_h, C_h, Sum_h, false); + HipTest::freeArrays(Res1_d, Res2_d, Res3_d, + nullptr, nullptr, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/** + * Test performs api parameter validation by passing various values + * as input and output parameters and validates the behavior. + * Test will include both negative and positive scenarios. + */ +TEST_CASE("Unit_hipGraphNodeGetDependencies_ParamValidation") { + hipGraph_t graph{}; + const int numBytes = 100; + size_t numDeps{1}; + hipGraphNode_t memsetNode{}, depnodes{}; + hipError_t ret{}; + char *A_d; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipMalloc(&A_d, numBytes)); + hipMemsetParams memsetParams{}; + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 1; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numBytes * sizeof(char); + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, + 0, &memsetParams)); + + SECTION("node passed is a root node") { + ret = hipGraphNodeGetDependencies(memsetNode, &depnodes, &numDeps); + + // Api expected to return success and no dependencies. + REQUIRE(ret == hipSuccess); + REQUIRE(numDeps == 0); + } + + SECTION("node as nullptr") { + ret = hipGraphNodeGetDependencies(nullptr, &depnodes, &numDeps); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("NumDependencies as nullptr") { + ret = hipGraphNodeGetDependencies(memsetNode, &depnodes, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("node as un-initialized/invalid parameter") { + hipGraphNode_t uninit_node{}; + ret = hipGraphNodeGetDependencies(uninit_node, &depnodes, &numDeps); + REQUIRE(ret == hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(A_d)); +}