diff --git a/catch/unit/stream/CMakeLists.txt b/catch/unit/stream/CMakeLists.txt index 7ca31d0a9b..0795d86b50 100644 --- a/catch/unit/stream/CMakeLists.txt +++ b/catch/unit/stream/CMakeLists.txt @@ -19,8 +19,7 @@ set(TEST_SRC hipLaunchHostFunc.cc hipStreamGetDevice.cc hipStreamLegacy_Ext.cc - hipStreamLegacy_compiler_options.cc -) + hipStreamLegacy_compiler_options.cc) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} hipStreamGetCUMask.cc hipStreamWithCUMask.cc @@ -28,7 +27,8 @@ if(HIP_PLATFORM MATCHES "amd") hipStreamGetFlags_spt.cc hipStreamGetPriority_spt.cc hipStreamQuery_spt.cc - hipStreamSynchronize_spt.cc) + hipStreamSynchronize_spt.cc + hipStreamBatchMemOp.cc) else() set(TEST_SRC ${TEST_SRC} diff --git a/catch/unit/stream/hipStreamBatchMemOp.cc b/catch/unit/stream/hipStreamBatchMemOp.cc new file mode 100644 index 0000000000..6d0d88c9b7 --- /dev/null +++ b/catch/unit/stream/hipStreamBatchMemOp.cc @@ -0,0 +1,120 @@ +/* +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 hipStreamBatchMemOp hipStreamBatchMemOp + * @{ + * @ingroup StreamTest + * `hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count, + hipStreamBatchMemOpParams* paramArray, unsigned + int flags);` - + * Enqueues an array of stream memory operations in the stream. + */ +/** + * Test Description + * ------------------------ + * - Verify the Negative cases of the hipStreamBatchMemOp API. + * Test source + * ------------------------ + * - unit/stream/hipStreamBatchMemOp.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipStreamBatchMemOp_Negative_Tests") { + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + int totalOps = 2; + static hipStreamBatchMemOpParams paramArray[2], invalidParamArray[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; + + invalidParamArray[0].operation = hipStreamMemOpBarrier; + invalidParamArray[0].writeValue.address = opsArray[0]; + invalidParamArray[0].writeValue.value = 1000; + invalidParamArray[0].writeValue.flags = 32; + invalidParamArray[0].writeValue.alias = 0; + + invalidParamArray[1].operation = hipStreamMemOpBarrier; + invalidParamArray[1].waitValue.address = opsArray[0]; + invalidParamArray[1].waitValue.value = 1000; + invalidParamArray[1].waitValue.flags = hipStreamWaitValueEq; + + SECTION("Stream as a nullptr") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(nullptr, totalOps, paramArray, 0), + hipSuccess); + } + + SECTION("Invalid Stream") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(reinterpret_cast(-1), + totalOps, paramArray, 0), + hipErrorContextIsDestroyed); + } + + SECTION("Parameter Array as a nullptr") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(stream, totalOps, nullptr, 0), + hipErrorInvalidValue); + } + + SECTION("More than 256 Total Operations") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(stream, 1000, paramArray, 0), + hipErrorInvalidValue); + } + + SECTION("Total Operations less than 0") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(stream, -4, paramArray, 0), + hipErrorInvalidValue); + } + SECTION("Total Operations Zero") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(stream, 0, paramArray, 0), + hipSuccess); + } + + SECTION("Flag value not Zero") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(stream, totalOps, paramArray, -6), + hipErrorInvalidValue); + } + + // Disabled due to defect SWDEV-502219 + #if 0 + SECTION("InValid Parameter Array") { + HIP_CHECK_ERROR(hipStreamBatchMemOp(stream, totalOps, invalidParamArray, 0), + hipErrorInvalidValue); + } + #endif + HIP_CHECK(hipFree((void *)opsArray[0])); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group StreamTest. + * @} + */