From fa51cf8b3eb53168ac67f2be761bd3ea8c5ca366 Mon Sep 17 00:00:00 2001 From: SrinivasaRao Date: Thu, 28 Nov 2024 16:47:50 +0530 Subject: [PATCH] SWDEV-493289-[catch2][dtest]-Negative Tests for the hipGraphExecBatchMemOpNodeSetParams API Change-Id: I52fde7f8bf7ab8275d8484aa52685830689df174 [ROCm/hip-tests commit: 87d771b1555a8832d2f6ec5fdc14c9e42e75091c] --- .../hip-tests/catch/unit/graph/CMakeLists.txt | 3 +- .../hipGraphExecBatchMemOpNodeSetParams.cc | 156 ++++++++++++++++++ 2 files changed, 158 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/graph/hipGraphExecBatchMemOpNodeSetParams.cc diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index 6a69414168..0168652b89 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -82,7 +82,8 @@ set(TEST_SRC hipStreamIsCapturing.cc hipStreamIsCapturing_old.cc hipGraphMemFreeNodeGetParams.cc - hipGraphAsyncUserObj.cc) + hipGraphAsyncUserObj.cc + hipGraphExecBatchMemOpNodeSetParams.cc) if(HIP_PLATFORM MATCHES "amd") set(AMD_SRC diff --git a/projects/hip-tests/catch/unit/graph/hipGraphExecBatchMemOpNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphExecBatchMemOpNodeSetParams.cc new file mode 100644 index 0000000000..ef7186ea48 --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipGraphExecBatchMemOpNodeSetParams.cc @@ -0,0 +1,156 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +/** + * @addtogroup hipGraphExecBatchMemOpNodeSetParams + hipGraphExecBatchMemOpNodeSetParams + * @{ + * @ingroup GraphTest + * `hipError_t hipGraphExecBatchMemOpNodeSetParams(hipGraphExec_t hGraphExec, + hipGraphNode_t hNode, const hipBatchMemOpNodeParams* nodeParams)` + * - Sets the parameters for a batch mem op node in the given graphExec + */ + +/** + * Test Description + * ------------------------ + * - Verify the Negative cases of the hipGraphExecBatchMemOpNodeSetParams API. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecBatchMemOpNodeSetParams.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGraphExecBatchMemOpNodeSetParams_NegativeTsts") { + HIP_CHECK(hipInit(0)); + hipGraph_t graph, graph1; + hipGraphExec_t graphExec; + hipCtx_t ctx; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&ctx, 0, device)); + // Create a HIP graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&graph1, 0)); + INFO("Graph created."); + + static hipStreamBatchMemOpParams paramArray[2], newParamArray[2]; + std::vector opsArray(1); + HIP_CHECK(hipMalloc((void **)&opsArray[0], sizeof(uint32_t))); + + paramArray[0].operation = hipStreamMemOpWriteValue32; + paramArray[0].writeValue.address = opsArray[0]; + paramArray[0].writeValue.value = 1000; + paramArray[0].writeValue.flags = 0x0; + // paramArray[0].writeValue.alias = 0; + + paramArray[1].operation = hipStreamMemOpWaitValue32; + paramArray[1].waitValue.address = opsArray[0]; + paramArray[1].waitValue.value = 1000; + paramArray[1].waitValue.flags = hipStreamWaitValueEq; + // paramArray[i].waitValue.alias = 0; + + int totalOps = 2; + // Setup the batch memory operation node parameters + hipBatchMemOpNodeParams batchNodeParams; + batchNodeParams.ctx = ctx; // Use the current HIP context + batchNodeParams.count = totalOps; // Total number of memory operations + batchNodeParams.paramArray = + paramArray; // Pointer to the array of memory operations + batchNodeParams.flags = 0; // No special flags + + // Add a batch memory operation node to the graph + hipGraphNode_t batchMemOpNode, batchMemOpNode_1; + HIP_CHECK(hipGraphAddBatchMemOpNode(&batchMemOpNode, graph, nullptr, 0, + &batchNodeParams)); + HIP_CHECK(hipGraphAddBatchMemOpNode(&batchMemOpNode_1, graph1, nullptr, 0, + &batchNodeParams)); + INFO("hipGraphAddBatchMemOpNode added successfully."); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + INFO("Graph instantiated."); + for (int i = 0; i < totalOps; i++) { + newParamArray[i] = paramArray[i]; + } + newParamArray[0].writeValue.value = 2000; + newParamArray[1].waitValue.flags = hipStreamWaitValueGte; + + hipBatchMemOpNodeParams newBatchNodeParams; + newBatchNodeParams.ctx = ctx; + newBatchNodeParams.count = totalOps; + newBatchNodeParams.paramArray = newParamArray; + newBatchNodeParams.flags = 0; + + hipBatchMemOpNodeParams invalidNewBatchNodeParams; + invalidNewBatchNodeParams.ctx = ctx; + invalidNewBatchNodeParams.count = 400; + invalidNewBatchNodeParams.paramArray = newParamArray; + invalidNewBatchNodeParams.flags = -4; + + SECTION("Graph Executable as nullptr") { + HIP_CHECK_ERROR(hipGraphExecBatchMemOpNodeSetParams(nullptr, batchMemOpNode, + &newBatchNodeParams), + hipErrorInvalidValue); + } + + SECTION("Batch Memory Node as nullptr") { + HIP_CHECK_ERROR(hipGraphExecBatchMemOpNodeSetParams(graphExec, nullptr, + &newBatchNodeParams), + hipErrorInvalidValue); + } + // Disabled for NVIDIA due to the defect SWDEV-502247 + #if HT_AMD + SECTION("Batch node Parameters as nullptr") { + HIP_CHECK_ERROR( + hipGraphExecBatchMemOpNodeSetParams(graphExec, batchMemOpNode, nullptr), + hipErrorInvalidValue); + } + #endif + SECTION("Irrelevant Batch Node") { + HIP_CHECK_ERROR(hipGraphExecBatchMemOpNodeSetParams( + graphExec, batchMemOpNode_1, &newBatchNodeParams), + hipErrorInvalidValue); + } + // Disabled due to defect SWDEV-502219 + #if 0 + SECTION("Invalid Batch Node Params") { + HIP_CHECK_ERROR(hipGraphExecBatchMemOpNodeSetParams( + graphExec, batchMemOpNode, &invalidNewBatchNodeParams), + hipErrorInvalidValue); + } + #endif + SECTION("Unchanged Batch node Parameters") { + HIP_CHECK_ERROR(hipGraphExecBatchMemOpNodeSetParams(graphExec, + batchMemOpNode, &batchNodeParams), hipSuccess); + } + HIP_CHECK(hipFree((void *)opsArray[0])); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipCtxPopCurrent(&ctx)); + HIP_CHECK(hipCtxDestroy(ctx)); +} +/** + * End doxygen group GraphTest. + * @} + */ +