diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index ef001acd62..9001ab394b 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -87,6 +87,8 @@ set(TEST_SRC hipMemsetD32Async.cc hipMemsetD8Async.cc hipMemcpy2DArrayToArray.cc + hipMemcpy3DBatchAsync.cc + hipMemcpyBatchAsync.cc ) if(UNIX) diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy3DBatchAsync.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy3DBatchAsync.cc new file mode 100644 index 0000000000..ca5be81574 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy3DBatchAsync.cc @@ -0,0 +1,393 @@ +/* + * Copyright (c) 2025 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 +/** + * @addtogroup hipMemcpy3DBatchAsync hipMemcpy3DBatchAsync + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy3DBatchAsync(size_t numOps, struct hipMemcpy3DBatchOp* + opList, size_t* failIdx, unsigned long long flags, hipStream_t stream + __dparm(0))` - + * Perform Batch of 3D copies. + */ +// Helper to check array content +template +void checkArrayContent(hipArray_t array, size_t width, size_t height, + size_t depth, T expected) { + std::vector hostBuf(width * height * depth, 0); + hipMemcpy3DParms copyParms{}; + copyParms.srcArray = array; + copyParms.dstPtr = + make_hipPitchedPtr(hostBuf.data(), width * sizeof(T), width, height); + copyParms.extent = make_hipExtent(width, height, depth); + copyParms.kind = hipMemcpyDeviceToHost; + copyParms.srcPos = make_hipPos(0, 0, 0); + copyParms.dstPos = make_hipPos(0, 0, 0); + HIP_CHECK(hipMemcpy3D(©Parms)); + for (size_t i = 0; i < width * height * depth; ++i) { + INFO("Array FAILURE at Index: " << i << "\nval : " << hostBuf[i] + << " expected:" << expected); + REQUIRE(hostBuf[i] == expected); + } +} +/** + * Test Description + * ------------------------ + * - Test case to verify the Asynchronus 3D batch memory copy. + * 1. Test case verifies below batch pointer to pointer mem copy operations. + * 2. Op1: Host -> Device Copy + * 3. Op2: Device -> Device Copy + * 4. Op3: Device -> Host Copy + * 5. Op4: Host -> Host Copy + * 6. Prepare hipMemcpy3DBatchOp Array with appropriate data for ptr-ptr copy. + * 7. Create Stream. + * 8. Launch the hipMemcpy3DBatchAsync with appropriate fields. + * 9. Validate the data. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpy3DBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpy3DBatchAsync_Ptr2PtrBatchOps", "", char, int, + float) { + constexpr auto kfloatval1 = -1.5f; + constexpr auto kfloatval2 = 2.25f; + constexpr auto kfloatval3 = -0.75f; + const TestType val1 = std::is_floating_point_v ? kfloatval1 + : std::is_integral_v ? 10 + : 'a'; + const TestType val2 = std::is_floating_point_v ? kfloatval2 + : std::is_integral_v ? 7 + : 'b'; + const TestType val3 = std::is_floating_point_v ? kfloatval3 + : std::is_integral_v ? 3 + : 'c'; + + constexpr int numOps = 4; + constexpr int numW = 16; + constexpr int numH = 16; + constexpr int depth = 10; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + hipExtent extent = make_hipExtent(numW * sizeof(TestType), numH, depth); + size_t elements_3d = numW * numH * depth; + + // Source Pointers + std::vector srcPtr1(elements_3d, val1); + std::vector srcPtr2(elements_3d, val2); + std::vector srcPtr3(elements_3d, val3); + + // Device Pointers + void *dstPtr1, *dstPtr2; + HIP_CHECK(hipMalloc(&dstPtr1, elements_3d * sizeof(TestType))); + HIP_CHECK(hipMalloc(&dstPtr2, elements_3d * sizeof(TestType))); + + // Prepare batch ops array + hipMemcpy3DBatchOp ops[numOps]; + + // Op 1: Host pointer -> Device pointer + ops[0].src.type = hipMemcpyOperandTypePointer; + ops[0].src.op.ptr.ptr = srcPtr1.data(); + ops[0].src.op.ptr.rowLength = extent.width; + ops[0].src.op.ptr.layerHeight = extent.height; + ops[0].src.op.ptr.locHint.type = hipMemLocationTypeHost; + ops[0].src.op.ptr.locHint.id = 0; + ops[0].dst.type = hipMemcpyOperandTypePointer; + ops[0].dst.op.ptr.ptr = dstPtr1; + ops[0].dst.op.ptr.rowLength = extent.width; + ops[0].dst.op.ptr.layerHeight = extent.height; + ops[0].dst.op.ptr.locHint.type = hipMemLocationTypeDevice; + ops[0].dst.op.ptr.locHint.id = 0; + ops[0].extent = extent; + ops[0].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[0].flags = hipMemcpyFlagDefault; + + // Op 2: device pointer -> device pointer + ops[1].src.type = hipMemcpyOperandTypePointer; + ops[1].src.op.ptr.ptr = dstPtr1; + ops[1].src.op.ptr.rowLength = extent.width; + ops[1].src.op.ptr.layerHeight = extent.height; + ops[1].src.op.ptr.locHint.type = hipMemLocationTypeDevice; + ops[1].src.op.ptr.locHint.id = 0; + ops[1].dst.type = hipMemcpyOperandTypePointer; + ops[1].dst.op.ptr.ptr = dstPtr2; + ops[1].dst.op.ptr.rowLength = extent.width; + ops[1].dst.op.ptr.layerHeight = extent.height; + ops[1].dst.op.ptr.locHint.type = hipMemLocationTypeDevice; + ops[1].dst.op.ptr.locHint.id = 0; + ops[1].extent = extent; + ops[1].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[1].flags = hipMemcpyFlagDefault; + + // Op 3: Device pointer -> Host pointer + ops[2].src.type = hipMemcpyOperandTypePointer; + ops[2].src.op.ptr.ptr = dstPtr2; + ops[2].src.op.ptr.rowLength = extent.width; + ops[2].src.op.ptr.layerHeight = extent.height; + ops[2].src.op.ptr.locHint.type = hipMemLocationTypeDevice; + ops[2].src.op.ptr.locHint.id = 0; + ops[2].dst.type = hipMemcpyOperandTypePointer; + ops[2].dst.op.ptr.ptr = srcPtr2.data(); + ops[2].dst.op.ptr.rowLength = extent.width; + ops[2].dst.op.ptr.layerHeight = extent.height; + ops[2].dst.op.ptr.locHint.type = hipMemLocationTypeHost; + ops[2].dst.op.ptr.locHint.id = 0; + ops[2].extent = extent; + ops[2].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[2].flags = hipMemcpyFlagDefault; + + // Op 4: Host pointer -> Host pointer + ops[3].src.type = hipMemcpyOperandTypePointer; + ops[3].src.op.ptr.ptr = srcPtr2.data(); + ops[3].src.op.ptr.rowLength = extent.width; + ops[3].src.op.ptr.layerHeight = extent.height; + ops[3].src.op.ptr.locHint.type = hipMemLocationTypeHost; + ops[3].src.op.ptr.locHint.id = 0; + ops[3].dst.type = hipMemcpyOperandTypePointer; + ops[3].dst.op.ptr.ptr = srcPtr3.data(); + ops[3].dst.op.ptr.rowLength = extent.width; + ops[3].dst.op.ptr.layerHeight = extent.height; + ops[3].dst.op.ptr.locHint.type = hipMemLocationTypeHost; + ops[3].dst.op.ptr.locHint.id = 0; + ops[3].extent = extent; + ops[3].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[3].flags = hipMemcpyFlagDefault; + + // Launch the batch + size_t failIdx; + unsigned long long flags = 0; + HIP_CHECK(hipMemcpy3DBatchAsync(numOps, ops, &failIdx, flags, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validation + for (size_t i = 0; i < elements_3d; ++i) { + INFO("Array FAILURE at Index: " << i << "\nval : " << srcPtr3[i]); + REQUIRE(srcPtr3[i] == val1); + } + + // Cleanup + HIP_CHECK(hipFree(dstPtr1)); + HIP_CHECK(hipFree(dstPtr2)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * Test Description + * ------------------------ + * - Test case to verify the Asynchronus 3D batch memory copy. + * 1. Test case verifies below batch mem copy operations. + * 2. Op1: Host -> Array + * 3. Op2: Array -> Device ptr + * 4. Op3: Device ptr -> Array + * 5. Op4: Array -> Array + * 6. Op5: Array -> Host + * 7. Prepare hipMemcpy3DBatchOp Array with appropriate data for ptr-ptr copy. + * 8. Create Stream. + * 9. Launch the hipMemcpy3DBatchAsync with appropriate fields. + * 10. Vaidate the data. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpy3DBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpy3DBatchAsync_ArrayMemCpyBatchOps", "", char, + int, float) { + CHECK_IMAGE_SUPPORT + constexpr auto kfloatval1 = -1.5f; + constexpr auto kfloatval2 = 2.25f; + const TestType val1 = std::is_floating_point_v ? kfloatval1 + : std::is_integral_v ? 10 + : 'a'; + const TestType val2 = std::is_floating_point_v ? kfloatval2 + : std::is_integral_v ? 7 + : 'b'; + constexpr int numOps = 5; + constexpr int numW = 16; + constexpr int numH = 16; + constexpr int depth = 10; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + hipExtent extent = make_hipExtent(numW, numH, depth); + size_t elements_3d = extent.width * extent.height * extent.depth; + + // Host Pointers + std::vector srcPtr1(elements_3d, val1); + std::vector srcPtr2(elements_3d, val2); + + // Device Pointer + void *dstPtr; + HIP_CHECK(hipMalloc(&dstPtr, elements_3d * sizeof(TestType))); + + // Dev Arrays + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + hipArray_t array1, array2, array3; + HIP_CHECK(hipMalloc3DArray(&array1, &channelDesc, extent, 0)); + HIP_CHECK(hipMalloc3DArray(&array2, &channelDesc, extent, 0)); + HIP_CHECK(hipMalloc3DArray(&array3, &channelDesc, extent, 0)); + + // Fill dev Array with val2 + std::vector tmpHost(elements_3d, val2); + hipMemcpy3DParms fillParms{}; + fillParms.srcPtr = + make_hipPitchedPtr(tmpHost.data(), extent.width * sizeof(TestType), + extent.width, extent.height); + fillParms.dstArray = array1; + fillParms.extent = extent; + fillParms.kind = hipMemcpyHostToDevice; + HIP_CHECK(hipMemcpy3D(&fillParms)); + + // Prepare batch ops array + hipMemcpy3DBatchOp ops[numOps]; + + // Op 1: host ptr -> device array + ops[0].src.type = hipMemcpyOperandTypePointer; + ops[0].src.op.ptr.ptr = srcPtr1.data(); + ops[0].src.op.ptr.rowLength = extent.width; + ops[0].src.op.ptr.layerHeight = extent.height; + ops[0].src.op.ptr.locHint.type = hipMemLocationTypeHost; + ops[0].src.op.ptr.locHint.id = 0; + ops[0].dst.type = hipMemcpyOperandTypeArray; + ops[0].dst.op.array.array = array1; + ops[0].dst.op.array.offset = {0, 0, 0}; + ops[0].extent = extent; + ops[0].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[0].flags = hipMemcpyFlagDefault; + + // Op 2: device array -> dev ptr + ops[1].src.type = hipMemcpyOperandTypeArray; + ops[1].src.op.array.array = array1; + ops[1].src.op.array.offset = {0, 0, 0}; + ops[1].dst.type = hipMemcpyOperandTypePointer; + ops[1].dst.op.ptr.ptr = dstPtr; + ops[1].dst.op.ptr.rowLength = extent.width; + ops[1].dst.op.ptr.layerHeight = extent.height; + ops[1].dst.op.ptr.locHint.type = hipMemLocationTypeDevice; + ops[1].dst.op.ptr.locHint.id = 0; + ops[1].extent = extent; + ops[1].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[1].flags = hipMemcpyFlagDefault; + + // Op 3: dev ptr -> device array + ops[2].src.type = hipMemcpyOperandTypePointer; + ops[2].src.op.ptr.ptr = dstPtr; + ops[2].src.op.ptr.rowLength = extent.width; + ops[2].src.op.ptr.layerHeight = extent.height; + ops[2].src.op.ptr.locHint.type = hipMemLocationTypeDevice; + ops[2].src.op.ptr.locHint.id = 0; + ops[2].dst.type = hipMemcpyOperandTypeArray; + ops[2].dst.op.array.array = array2; + ops[2].dst.op.array.offset = {0, 0, 0}; + ops[2].extent = extent; + ops[2].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[2].flags = hipMemcpyFlagDefault; + + // Op 4: hip array -> hip array + ops[3].src.type = hipMemcpyOperandTypeArray; + ops[3].src.op.array.array = array2; + ops[3].src.op.array.offset = {0, 0, 0}; + ops[3].dst.type = hipMemcpyOperandTypeArray; + ops[3].dst.op.array.array = array3; + ops[3].dst.op.array.offset = {0, 0, 0}; + ops[3].extent = extent; + ops[3].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[3].flags = hipMemcpyFlagDefault; + + // Op 5: device array -> host ptr + ops[4].src.type = hipMemcpyOperandTypeArray; + ops[4].src.op.array.array = array3; + ops[4].src.op.array.offset = {0, 0, 0}; + ops[4].dst.type = hipMemcpyOperandTypePointer; + ops[4].dst.op.ptr.ptr = srcPtr2.data(); + ops[4].dst.op.ptr.rowLength = extent.width; + ops[4].dst.op.ptr.layerHeight = extent.height; + ops[4].dst.op.ptr.locHint.type = hipMemLocationTypeHost; + ops[4].dst.op.ptr.locHint.id = 0; + ops[4].extent = extent; + ops[4].srcAccessOrder = hipMemcpySrcAccessOrderStream; + ops[4].flags = hipMemcpyFlagDefault; + + // Launch the batch + size_t failIdx; + unsigned long long flags = 0; + HIP_CHECK(hipMemcpy3DBatchAsync(numOps, ops, &failIdx, flags, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Check Random Array data + checkArrayContent(array2, extent.width, extent.height, extent.depth, + val1); + // Check Final data. + for (size_t i = 0; i < elements_3d; ++i) { + INFO("Pointer Copy Failure at Index: " << i << "\nval : " << srcPtr2[i]); + REQUIRE(srcPtr2[i] == val1); + } + + // Cleanup + HIP_CHECK(hipFree(dstPtr)); + HIP_CHECK(hipFreeArray(array1)); + HIP_CHECK(hipFreeArray(array2)); + HIP_CHECK(hipFreeArray(array3)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * Test Description + * ------------------------ + * - Test case to verify the negative cases of hipMemcpy3DBatchAsync. + * 1. Num of Operations as 0. + * 2. Non Zero flag. + * 3. Ops array as nullptr + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpy3DBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemcpy3DBatchAsync_NegativeTests") { + CHECK_IMAGE_SUPPORT + const int numOps = 2; + hipStream_t stream = NULL; + HIP_CHECK(hipStreamCreate(&stream)); + size_t failIdx; + unsigned long long flags = 0; + hipMemcpy3DBatchOp ops[numOps]; + SECTION("Zero Operations") { + HIP_CHECK_ERROR(hipMemcpy3DBatchAsync(0, ops, &failIdx, flags, stream), + hipErrorInvalidValue); + } + SECTION("Non Zero flag") { + HIP_CHECK_ERROR(hipMemcpy3DBatchAsync(numOps, ops, &failIdx, 2, stream), + hipErrorInvalidValue); + } + SECTION("Ops array as nullptr") { + HIP_CHECK_ERROR( + hipMemcpy3DBatchAsync(numOps, nullptr, &failIdx, flags, stream), + hipErrorInvalidValue); + } + // Cleanup + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyBatchAsync.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyBatchAsync.cc new file mode 100644 index 0000000000..d9ba495c12 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyBatchAsync.cc @@ -0,0 +1,388 @@ +/* + * Copyright (c) 2025 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 +/** + * @addtogroup hipMemcpyBatchAsync hipMemcpyBatchAsync + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpyBatchAsync(void** dsts, void** srcs, size_t* sizes, + size_t count, hipMemcpyAttributes* attrs, size_t* attrsIdxs, size_t numAttrs, + size_t* failIdx, hipStream_t stream __dparm(0))` + - + * Perform Batch of 1D copies. + */ +/** + * Test Description + * ------------------------ + * - Test case to verify the 1D batch memory copy. + * 1. Create Array of device pointers(Src, Dst). + * 2. Set the MemcpyBatch params. As of now no support for memcpy Attributes. + * 3. Perform batch memcpy operation from deviceptr to deviceptr. + * 4. Validate data on host. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpyBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +#if HT_AMD +TEMPLATE_TEST_CASE("Unit_hipMemcpyBatchAsync_D2D_Functional", "", char, int, + float) { + const size_t count = 2; + size_t numAttrs = 0; + const size_t arrSize = 4096; + const size_t size = 4096 * sizeof(TestType); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + constexpr auto kfloatval1 = 2.25f; + constexpr auto kfloatval2 = 0.25f; + const TestType val1 = std::is_floating_point_v ? kfloatval1 + : std::is_integral_v ? 10 + : 'a'; + const TestType val2 = std::is_floating_point_v ? kfloatval2 + : std::is_integral_v ? 4 + : 'b'; + + // Allocate buffers for pointer-ptr copy + void *srcPtr[count], *dstPtr[count]; + std::vector> hostPtr1( + count, std::vector(arrSize, val1)); + std::vector> hostPtr2( + count, std::vector(arrSize, val2)); + size_t sizes[2]; + size_t attrsIdxs[1]; + for (int i = 0; i < count; i++) { + HIP_CHECK(hipMalloc(&srcPtr[i], size)); + HIP_CHECK(hipMalloc(&dstPtr[i], size)); + HIP_CHECK( + hipMemcpy(srcPtr[i], hostPtr2[i].data(), size, hipMemcpyHostToDevice)); + sizes[i] = size; + } + attrsIdxs[0] = 0; + size_t failIdx; + + HIP_CHECK(hipMemcpyBatchAsync(dstPtr, srcPtr, sizes, count, nullptr, + attrsIdxs, numAttrs, &failIdx, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + // validation + for (int i = 0; i < count; i++) { + HIP_CHECK( + hipMemcpy(hostPtr1[i].data(), dstPtr[i], size, hipMemcpyDeviceToHost)); + for (int j = 0; j < arrSize; j++) { + INFO("Array FAILURE at Index: " << i << " " << j + << "\nval : " << hostPtr1[i][j]); + REQUIRE(hostPtr1[i][j] == val2); + } + } + // Clean up + for (int i = 0; i < count; i++) { + HIP_CHECK(hipFree(srcPtr[i])); + HIP_CHECK(hipFree(dstPtr[i])); + } + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify the 1D batch memory copy. + * 1. Create Array of device pointers(Src, Dst). + * 2. Set the MemcpyBatch params. As of now no support for memcpy Attributes. + * 3. Perform batch memcpy operation From Host to Device. + * 4. Validate data on host. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpyBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyBatchAsync_H2D_Functional", "", char, int, + float) { + const size_t count = 2; + size_t numAttrs = 0; + const size_t arrSize = 4096; + const size_t size = 4096 * sizeof(TestType); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Allocate buffers for pointer-ptr copy + void *hostSrcPtr[count], *dstPtr[count]; + constexpr auto kfloatval1 = 2.25f; + constexpr auto kfloatval2 = 0.25f; + const TestType val1 = std::is_floating_point_v ? kfloatval1 + : std::is_integral_v ? 10 + : 'a'; + const TestType val2 = std::is_floating_point_v ? kfloatval2 + : std::is_integral_v ? 4 + : 'b'; + std::vector> hostPtr( + count, std::vector(arrSize, val2)); + std::array arr; + arr.fill(val1); + size_t sizes[2]; + size_t attrsIdxs[1]; + for (int i = 0; i < count; i++) { + hostSrcPtr[i] = arr.data(); + HIP_CHECK(hipMalloc(&dstPtr[i], size)); + sizes[i] = size; + } + attrsIdxs[0] = 0; + size_t failIdx; + + HIP_CHECK(hipMemcpyBatchAsync(dstPtr, hostSrcPtr, sizes, count, nullptr, + attrsIdxs, numAttrs, &failIdx, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + // validation + for (int i = 0; i < count; i++) { + HIP_CHECK( + hipMemcpy(hostPtr[i].data(), dstPtr[i], size, hipMemcpyDeviceToHost)); + for (int j = 0; j < arrSize; j++) { + INFO("Array FAILURE at Index: " << i << " " << j + << "\nval : " << hostPtr[i][j]); + REQUIRE(hostPtr[i][j] == val1); + } + } + // Clean up + for (int i = 0; i < count; i++) { + HIP_CHECK(hipFree(dstPtr[i])); + } + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify the 1D batch memory copy. + * 1. Create Array of device pointers(Src, Dst). + * 2. Set the MemcpyBatch params. As of now no support for memcpy Attributes. + * 3. Perform batch memcpy operation From Device to Host. + * 4. Validate data on host. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpyBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyBatchAsync_D2H_Functional", "", char, int, + float) { + const size_t count = 2; + size_t numAttrs = 0; + const size_t arrSize = 4096; + const size_t size = 4096 * sizeof(TestType); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + constexpr auto kfloatval1 = 2.25f; + constexpr auto kfloatval2 = 0.25f; + const TestType val1 = std::is_floating_point_v ? kfloatval1 + : std::is_integral_v ? 10 + : 'a'; + const TestType val2 = std::is_floating_point_v ? kfloatval2 + : std::is_integral_v ? 4 + : 'b'; + // Allocate buffers for pointer-ptr copy + TestType *hostDstPtr[count]; + void *deviceSrcPtr[count]; + std::vector> hostPtr( + count, std::vector(arrSize, val1)); + std::array arr; + arr.fill(val2); + size_t sizes[2]; + size_t attrsIdxs[1]; + for (int i = 0; i < count; i++) { + hostDstPtr[i] = arr.data(); + HIP_CHECK(hipMalloc(&deviceSrcPtr[i], size)); + HIP_CHECK(hipMemcpy(deviceSrcPtr[i], hostPtr[i].data(), size, + hipMemcpyHostToDevice)); + sizes[i] = size; + } + attrsIdxs[0] = 0; + size_t failIdx; + + HIP_CHECK(hipMemcpyBatchAsync(reinterpret_cast(hostDstPtr), + deviceSrcPtr, sizes, count, nullptr, attrsIdxs, + numAttrs, &failIdx, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + // validation + for (int i = 0; i < count; i++) { + for (int j = 0; j < arrSize; j++) { + INFO("Array FAILURE at Index: " << i << " " << j + << "\nval : " << hostDstPtr[i][j]); + REQUIRE(hostDstPtr[i][j] == val1); + } + } + // Clean up + for (int i = 0; i < count; i++) { + HIP_CHECK(hipFree(deviceSrcPtr[i])); + } + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * Test Description + * ------------------------ + * - Test case to verify the 1D batch memory copy. + * 1. Create Array of device pointers(Src, Dst). + * 2. Set the MemcpyBatch params. As of now no support for memcpy Attributes. + * 3. Perform batch memcpy operation From Host to Host. + * 4. Validate data on host. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpyBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyBatchAsync_H2H_Functional", "", char, int, + float) { + const size_t count = 2; + size_t numAttrs = 0; + const size_t arrSize = 4096; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + constexpr auto kfloatval1 = 2.25; + const TestType val1 = std::is_floating_point_v ? kfloatval1 + : std::is_integral_v ? 10 + : 'a'; + constexpr auto kfloatval2 = 0.25f; + const TestType val2 = std::is_floating_point_v ? kfloatval2 + : std::is_integral_v ? 4 + : 'b'; + + // Allocate buffers for pointer-ptr copy + TestType *hostDstPtr[count], *hostSrcPtr[count]; + std::array arr1, arr2; + arr1.fill(val1); + arr2.fill(val2); + size_t sizes[2]; + size_t attrsIdxs[1]; + for (int i = 0; i < count; i++) { + hostDstPtr[i] = arr1.data(); + hostSrcPtr[i] = arr2.data(); + sizes[i] = arrSize * sizeof(TestType); + } + attrsIdxs[0] = 0; + size_t failIdx; + + HIP_CHECK(hipMemcpyBatchAsync(reinterpret_cast(hostDstPtr), + reinterpret_cast(hostSrcPtr), sizes, + count, nullptr, attrsIdxs, numAttrs, &failIdx, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + // validation + for (int i = 0; i < count; i++) { + for (int j = 0; j < arrSize; j++) { + INFO("Array FAILURE at Index: " << i << " " << j + << "\nval : " << hostDstPtr[i][j]); + REQUIRE(hostDstPtr[i][j] == val2); + } + } + // Clean up + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif +/** + * Test Description + * ------------------------ + * - Test case to verify the negative cases of hipMemcpyBatchAsync. + * 1. Dst Array as nullptr. + * 2. Src Array as nullptr. + * 3. Operations Count as 0. + * 4. Num of attributes as 0. + * 5. Sizes Array as nullptr. + * 6. Attr Array as nullptr. + * 7. AttrsIdxs Array as nullptr. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpyBatchAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemcpyBatchAsync_NegativeTsts") { + const size_t count = 2; + size_t numAttrs = 0; + size_t sizes[2]; + size_t attrsIdxs[1]; + const size_t size = 4096 * sizeof(char); + hipStream_t stream = NULL; + HIP_CHECK(hipStreamCreate(&stream)); + void *srcPtr[count], *dstPtr[count]; + for (int i = 0; i < count; i++) { + HIP_CHECK(hipMalloc(&srcPtr[i], size)); + HIP_CHECK(hipMalloc(&dstPtr[i], size)); + sizes[i] = size; + } + + attrsIdxs[0] = 0; + size_t failIdx; + SECTION("Dst Array as nullptr") { + HIP_CHECK_ERROR(hipMemcpyBatchAsync(nullptr, srcPtr, sizes, count, nullptr, + attrsIdxs, numAttrs, &failIdx, stream), + hipErrorInvalidValue); + } + SECTION("Src Array as nullptr") { + HIP_CHECK_ERROR(hipMemcpyBatchAsync(dstPtr, nullptr, sizes, count, nullptr, + attrsIdxs, numAttrs, &failIdx, stream), + hipErrorInvalidValue); + } + SECTION("Count as zero") { + HIP_CHECK_ERROR(hipMemcpyBatchAsync(dstPtr, srcPtr, sizes, 0, nullptr, + attrsIdxs, numAttrs, &failIdx, stream), + hipErrorInvalidValue); + } + SECTION("sizes Array as nullptr") { + HIP_CHECK_ERROR(hipMemcpyBatchAsync(dstPtr, srcPtr, nullptr, count, nullptr, + attrsIdxs, numAttrs, &failIdx, stream), + hipErrorInvalidValue); + } +#if 0 // Enable these tests when support for memcpy attributes is enabled. + SECTION("Number of Attributes as zero") { + HIP_CHECK_ERROR( + hipMemcpyBatchAsync(dstPtr, srcPtr, sizes, count, attr, attrsIdxs, 0, &failIdx, stream), + hipErrorInvalidValue); + } + SECTION("Attr Array as nullptr") { + HIP_CHECK_ERROR(hipMemcpyBatchAsync(dstPtr, srcPtr, sizes, count, nullptr, attrsIdxs, numAttrs, + &failIdx, stream), + hipErrorInvalidValue); + } + + SECTION("attrsIdxs Array as nullptr") { + HIP_CHECK_ERROR(hipMemcpyBatchAsync(dstPtr, srcPtr, sizes, count, attr, nullptr, numAttrs, + &failIdx, stream), + hipErrorInvalidValue); + } +#endif + // Clean up + for (int i = 0; i < count; i++) { + HIP_CHECK(hipFree(srcPtr[i])); + HIP_CHECK(hipFree(dstPtr[i])); + } + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */