amd-srinivas1
6b8a4a23ba
* SWDEV-546345-Added tests for Batch memCpy apis. * Updated tests to use multiple data types * SWDEV-546345-Updated tests for different data types * SWDEV-546345-Updated clang-format * SWDEV-546345-Updated hipMemcpy3DBatchAsync parameter * SWDEV-546345 - Updated hipMemcpy3DBatchAsync tests with hipMemLocationTypeHost --------- Co-authored-by: Rahul Manocha <rmanocha@amd.com> Co-authored-by: Rambabu Swargam <rambabu.swargam@amd.com> Co-authored-by: jainprad <92369414+jainprad@users.noreply.github.com>
394 lines
15 KiB
C++
394 lines
15 KiB
C++
/*
|
|
* 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 <hip_test_common.hh>
|
|
#include <hip_test_defgroups.hh>
|
|
#include <vector>
|
|
/**
|
|
* @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 <typename T>
|
|
void checkArrayContent(hipArray_t array, size_t width, size_t height,
|
|
size_t depth, T expected) {
|
|
std::vector<T> 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<TestType> ? kfloatval1
|
|
: std::is_integral_v<TestType> ? 10
|
|
: 'a';
|
|
const TestType val2 = std::is_floating_point_v<TestType> ? kfloatval2
|
|
: std::is_integral_v<TestType> ? 7
|
|
: 'b';
|
|
const TestType val3 = std::is_floating_point_v<TestType> ? kfloatval3
|
|
: std::is_integral_v<TestType> ? 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<TestType> srcPtr1(elements_3d, val1);
|
|
std::vector<TestType> srcPtr2(elements_3d, val2);
|
|
std::vector<TestType> 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<TestType> ? kfloatval1
|
|
: std::is_integral_v<TestType> ? 10
|
|
: 'a';
|
|
const TestType val2 = std::is_floating_point_v<TestType> ? kfloatval2
|
|
: std::is_integral_v<TestType> ? 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<TestType> srcPtr1(elements_3d, val1);
|
|
std::vector<TestType> srcPtr2(elements_3d, val2);
|
|
|
|
// Device Pointer
|
|
void *dstPtr;
|
|
HIP_CHECK(hipMalloc(&dstPtr, elements_3d * sizeof(TestType)));
|
|
|
|
// Dev Arrays
|
|
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<TestType>();
|
|
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<TestType> 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<TestType>(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.
|
|
* @}
|
|
*/
|