From d676a0143101cefdb14cae7edb25c5fb932b80a0 Mon Sep 17 00:00:00 2001 From: DURGESH KROTTAPALLI Date: Wed, 16 Jun 2021 03:03:11 +0530 Subject: [PATCH] SWDEV-292393 - [catch2][dtest] Tests for hipMemcpy3D and hipMemcpyParam2D APIs Added functional and negative scenarios for hipMemcpy3D and hipMemcpyParam2D APIs in catch2 framework Change-Id: I3473eb952c23aba6bc500d78ca93d52a9f7a5d6f [ROCm/hip commit: aa6b01f793c1697a6039ad86472956f4f8afa6de] --- .../tests/catch/unit/memory/CMakeLists.txt | 4 + .../tests/catch/unit/memory/hipMemcpy3D.cc | 622 +++++++++++++++ .../catch/unit/memory/hipMemcpy3DAsync.cc | 736 ++++++++++++++++++ .../catch/unit/memory/hipMemcpyParam2D.cc | 380 +++++++++ .../unit/memory/hipMemcpyParam2DAsync.cc | 487 ++++++++++++ 5 files changed, 2229 insertions(+) create mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpy3D.cc create mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpy3DAsync.cc create mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyParam2D.cc create mode 100644 projects/hip/tests/catch/unit/memory/hipMemcpyParam2DAsync.cc diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index 5e8c1f975b..df180f4b87 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -6,6 +6,10 @@ set(TEST_SRC hipMemcpy2DToArrayAsync.cc hipMemcpyPeer.cc hipMemcpyPeerAsync.cc + hipMemcpy3D.cc + hipMemcpy3DAsync.cc + hipMemcpyParam2D.cc + hipMemcpyParam2DAsync.cc ) # Create shared lib of all tests diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy3D.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy3D.cc new file mode 100644 index 0000000000..af51e115e1 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy3D.cc @@ -0,0 +1,622 @@ +/* +Copyright (c) 2021 - present 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. +*/ + +/* + * This testfile verifies the following scenarios of hipMemcpy3D API + * + * 1. Verifying hipMemcpy3D API for H2D,D2D and D2H scenarios for + different datatypes and sizes. + * 2. Verifying Negative Scenarios + * 3. Verifying Extent validation scenarios by passing 0 + * 4. Verifying hipMemcpy3D API by allocating Memory in + * one GPU and trigger hipMemcpy3D from peer GPU + * + */ + +#include +#include + +static constexpr auto width{10}; +static constexpr auto height{10}; +static constexpr auto depth{10}; + +template +class Memcpy3D { + int width, height, depth; + unsigned int size; + hipArray *arr, *arr1; + hipChannelFormatKind formatKind; + hipMemcpy3DParms myparms; + T* hData; + public: + Memcpy3D(int l_width, int l_height, int l_depth, + hipChannelFormatKind l_format); + void simple_Memcpy3D(); + void Extent_Validation(); + void NegativeTests(); + void AllocateMemory(); + void DeAllocateMemory(); + void SetDefaultData(); + void D2D_DeviceMem_OnDiffDevice(); + void D2H_H2D_DeviceMem_OnDiffDevice(); +}; + +/* + * This API sets the default values of hipMemcpy3DParms structure + */ +template +void Memcpy3D::SetDefaultData() { + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); +} + +/* + * Constructor initalized width,depth and height + */ +template +Memcpy3D::Memcpy3D(int l_width, int l_height, int l_depth, + hipChannelFormatKind l_format) { + width = l_width; + height = l_height; + depth = l_depth; + formatKind = l_format; +} + +/* + * Allocating Memory and initalizing data for both + * device and host variables + */ +template +void Memcpy3D::AllocateMemory() { + size = width * height * depth * sizeof(T); + hData = reinterpret_cast(malloc(size)); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, + depth), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&arr1, &channelDesc, make_hipExtent(width, height, + depth), hipArrayDefault)); +} + +/* + * DeAllocates the Memory of device and host variables + */ +template +void Memcpy3D::DeAllocateMemory() { + HIP_CHECK(hipFreeArray(arr)); + HIP_CHECK(hipFreeArray(arr1)); + free(hData); +} + +/* + * This API verifies both H2D & D2H functionalities of hipMemcpy3D API + * by allocating memory in one GPU and calling the hipMemcpy3D API + * from another GPU. + * H2D case: + * Input : "hData" is initialized with the respective offset value + * Output: Destination array "arr" variable. + * + * D2H case: + * Input: "arr" array variable from the above output + * Output: "hOutputData" variable data is copied from "arr" variable + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ +template +void Memcpy3D::D2H_H2D_DeviceMem_OnDiffDevice() { + HIP_CHECK(hipSetDevice(0)); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (peerAccess) { + AllocateMemory(); + // Memory is allocated on device 0 and Memcpy3DAsync triggered from device 1 + HIP_CHECK(hipSetDevice(1)); + + // H2D Scenario + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + HIP_CHECK(hipDeviceSynchronize()); + + // Device to host + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + SetDefaultData(); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), + width, height); + myparms.srcArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + HIP_CHECK(hipDeviceSynchronize()); + + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); + free(hOutputData); + DeAllocateMemory(); + } else { + SUCCEED("Skipped the test as there is no peer access\n"); + } +} +/* + * This API verifies both D2D functionalities of hipMemcpy3D API + * by allocating memory in one GPU and calling the hipMemcpy3D API + * from another GPU. + * + * D2D case: + * Input : "arr" variable is initialized with the "hData" variable in GPU-0 + * Output: "arr2" variable in GPU-0 + * + * hipMemcpy3D API is triggered from GPU-1 + * The "arr2" variable is then copied to "hOutputData" for validating + * the result + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ +template +void Memcpy3D::D2D_DeviceMem_OnDiffDevice() { + HIP_CHECK(hipSetDevice(0)); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1)); + if (peerAccess) { + AllocateMemory(); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + + // Host to device copy + myparms.srcPtr = make_hipPitchedPtr(hData, + width * sizeof(T), + width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + hipArray *arr2; + hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(T)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&arr2, &channelDesc1, + make_hipExtent(width, height, + depth), hipArrayDefault)); + + // Allocating Mem on GPU device 0 and trigger hipMemcpy3D from GPU 1 + HIP_CHECK(hipSetDevice(1)); + + // D2D Scenario + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcArray = arr; + myparms.dstArray = arr2; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + HIP_CHECK(hipDeviceSynchronize()); + + // For validating the D2D copy copying it again to hOutputData and + // verifying it with iniital data hData + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + SetDefaultData(); + + // Device to host + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), + width, height); + myparms.srcArray = arr2; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkArray(hData, hOutputData, width, height, depth); + + // DeAllocating Memory + free(hOutputData); + DeAllocateMemory(); + } else { + SUCCEED("Skipped the test as there is no peer access\n"); + } +} +/* + * This API verifies all the negative scenarios of hipMemcpy3D API + */ +template +void Memcpy3D::NegativeTests() { + HIP_CHECK(hipSetDevice(0)); + AllocateMemory(); + + // Initialization of data + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + + SECTION("Nullptr to destination array") { + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = nullptr; + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Nullptr to source array") { + myparms.srcArray = nullptr; + myparms.dstPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing both Source ptr and array") { + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.srcArray = arr; + myparms.dstArray = arr1; + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing both destination ptr and array") { + myparms.dstPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; + myparms.srcArray = arr1; + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing Max value to extent") { + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(std::numeric_limits::max(), + std::numeric_limits::max(), + std::numeric_limits::max()); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing Source pitchedPtr as nullptr") { + myparms.srcPtr = make_hipPitchedPtr(nullptr, width * sizeof(T), + width, height); + myparms.dstArray = arr; + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing Dst pitchedPtr as nullptr") { + myparms.dstPtr = make_hipPitchedPtr(nullptr, width * sizeof(T), + width, height); + myparms.srcArray = arr; + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing width > max width size in extent") { + myparms.extent = make_hipExtent(width+1 , height, depth); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing hgt > max width size in extent") { + myparms.extent = make_hipExtent(width , height+1, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing depth > max width size in extent") { + myparms.extent = make_hipExtent(width , height, depth+1); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing dst width pos > max allocated width") { + myparms.dstPos = make_hipPos(width+1, 0, 0); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing dst height pos > max allocated hgt") { + myparms.dstPos = make_hipPos(0, height+1, 0); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing dst depth pos > max allocated depth") { + myparms.dstPos = make_hipPos(0, 0, depth+1); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing src width pos > max allocated width") { + myparms.srcPos = make_hipPos(width+1, 0, 0); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing src height pos > max allocated hgt") { + myparms.srcPos = make_hipPos(0, height+1, 0); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing src height pos > max allocated hgt") { + myparms.srcPos = make_hipPos(0, 0, depth+1); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + SECTION("Passing src array size > dst array size") { + // Passing src array size greater than destination array size + hipArray *arr2; + hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(T)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&arr2, &channelDesc1, + make_hipExtent(3, 3 + , 3), hipArrayDefault)); + myparms.srcArray = arr; + myparms.dstArray = arr2; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) != hipSuccess); + } + + // DeAllocation of memory + DeAllocateMemory(); +} + +/* + * This API verifies the Extent validation Scenarios + */ +template +void Memcpy3D::Extent_Validation() { + HIP_CHECK(hipSetDevice(0)); + AllocateMemory(); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + SECTION("Passing Extent as 0") { + myparms.extent = make_hipExtent(0 , 0, 0); + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + } + SECTION("Passing Width 0 in Extent") { + myparms.extent = make_hipExtent(0 , height, depth); + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + } + SECTION("Passing Height 0 in Extent") { + myparms.extent = make_hipExtent(width , 0, depth); + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + } + SECTION("Passing Depth 0 in Extent") { + myparms.extent = make_hipExtent(width , height, 0); + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + } + SECTION("Passing Depth 0 in Extent") { + REQUIRE(hipMemcpy3D(nullptr) != hipSuccess); + } + DeAllocateMemory(); +} + +/* + * This API verifies H2H-D2D-D2H functionalities of hipMemcpy3D API + * + * Input : "arr" variable is initialized with the "hData" variable in GPU-0 + * Output: "arr1" variable in GPU-0 + * + * The "arr1" variable is then copied to "hOutputData" for validating + * the result + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ + +template +void Memcpy3D::simple_Memcpy3D() { + HIP_CHECK(hipSetDevice(0)); + AllocateMemory(); + + // Host to Device + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + + // Array to Array + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + + // Device to host + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), width, height); + myparms.srcArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3D(&myparms) == hipSuccess); + + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); + + // DeAllocating the Memory + free(hOutputData); + DeAllocateMemory(); +} +/* + This testcase performs hipMemcpy3D API validation for + different datatypes and different sizes +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpy3D_Basic", "[hipMemcpy3D]", + int, unsigned int, float) { + int numDevices = 0; + auto i = GENERATE(10, 100, 1024, 10*1024); + auto j = GENERATE(10, 100); + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + if (std::is_same::value) { + Memcpy3D memcpy3d_obj(i, j, j, hipChannelFormatKindFloat); + memcpy3d_obj.simple_Memcpy3D(); + } else if (std::is_same::value) { + Memcpy3D memcpy3d_obj(i, j, j, hipChannelFormatKindUnsigned); + memcpy3d_obj.simple_Memcpy3D(); + } else if (std::is_same::value) { + Memcpy3D memcpy3d_obj(i, j, j, hipChannelFormatKindSigned); + memcpy3d_obj.simple_Memcpy3D(); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* +This testcase performs the extent validation scenarios of +hipMemcpy3D API +*/ +TEST_CASE("Unit_hipMemcpy3D_ExtentValidation") { + Memcpy3D memcpy3d(width, height, depth, + hipChannelFormatKindSigned); + memcpy3d.Extent_Validation(); +} + +/* +This testcase performs the negative scenarios of +hipMemcpy3D API +*/ +TEST_CASE("Unit_hipMemcpy3D_multiDevice-Negative") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + Memcpy3D memcpy3d(width, height, depth, + hipChannelFormatKindSigned); + memcpy3d.NegativeTests(); + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* +This testcase performs the D2H,H2D and D2D on peer +GPU device +*/ +TEST_CASE("Unit_hipMemcpy3D_multiDevice-OnPeerDevice") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + SECTION("D2H & H2D On DiffDevice") { + Memcpy3D memcpy3d_d2h_obj(width, height, depth, + hipChannelFormatKindFloat); + memcpy3d_d2h_obj.D2H_H2D_DeviceMem_OnDiffDevice(); + } + + SECTION("D2D On DiffDevice") { + Memcpy3D memcpy3d_d2d_obj(width, height, depth, + hipChannelFormatKindFloat); + memcpy3d_d2d_obj.D2D_DeviceMem_OnDiffDevice(); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy3DAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy3DAsync.cc new file mode 100644 index 0000000000..2327958e70 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy3DAsync.cc @@ -0,0 +1,736 @@ +/* +Copyright (c) 2021 - present 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. +*/ + +/* + * This testfile verifies the following Scenarios of hipMemcpy3DAsync API + + * 1. Verifying hipMemcpy3DAsync API for H2D,D2D and D2H scenarios + * 2. Verifying Negative Scenarios + * 3. Verifying Extent validation scenarios by passing 0 + * 4. Verifying hipMemcpy3DAsync API by allocating Memory in + * one GPU and trigger hipMemcpy3D from peer GPU + * 5. D2D where src and dst memory on GPU-0 and stream on GPU-1 +*/ + +#include +#include + +static constexpr auto width{10}; +static constexpr auto height{10}; +static constexpr auto depth{10}; + +template +class Memcpy3DAsync { + int width, height, depth; + unsigned int size; + hipArray *arr, *arr1; + hipChannelFormatKind formatKind; + hipMemcpy3DParms myparms; + T* hData; + hipStream_t stream; + public: + Memcpy3DAsync(int l_width, int l_height, int l_depth, + hipChannelFormatKind l_format); + void simple_Memcpy3DAsync(); + void Extent_Validation(); + void NegativeTests(); + void AllocateMemory(); + void DeAllocateMemory(); + void SetDefaultData(); + void D2D_SameDeviceMem_StreamDiffDevice(); + void D2D_DeviceMem_OnDiffDevice(); + void D2H_H2D_DeviceMem_OnDiffDevice(); +}; + +/* + * This API sets the default values of hipMemcpy3DParms structure + */ +template +void Memcpy3DAsync::SetDefaultData() { + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); +} + +/* + * Constructor initalized width,depth and height + */ +template +Memcpy3DAsync::Memcpy3DAsync(int l_width, int l_height, int l_depth, + hipChannelFormatKind l_format) { + width = l_width; + height = l_height; + depth = l_depth; + formatKind = l_format; +} + +/* + * Allocating Memory and initalizing data for both + * device and host variables + */ +template +void Memcpy3DAsync::AllocateMemory() { + size = width * height * depth * sizeof(T); + hData = reinterpret_cast(malloc(size)); + memset(hData, 0, size); + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, + depth), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&arr1, &channelDesc, make_hipExtent(width, height, + depth), hipArrayDefault)); +} + +/* + * DeAllocates the Memory of device and host variables + */ +template +void Memcpy3DAsync::DeAllocateMemory() { + hipFreeArray(arr); + hipFreeArray(arr1); + free(hData); + hipStreamDestroy(stream); +} + +/* + * This API verifies both H2D & D2H functionalities of hipMemcpy3DAsync API + * by allocating memory in one GPU and calling the hipMemcpy3DAsync API + * from another GPU. + * H2D case: + * Input : "hData" is initialized with the respective offset value + * Output: Destination array "arr" variable. + * + * D2H case: + * Input: "arr" array variable from the above output + * Output: "hOutputData" variable data is copied from "arr" variable + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ +template +void Memcpy3DAsync::D2H_H2D_DeviceMem_OnDiffDevice() { + HIP_CHECK(hipSetDevice(0)); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (peerAccess) { + AllocateMemory(); + + // Memory is allocated on device 0 and Memcpy3DAsyncAsync + // triggered from device 1 + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + + // Host to Device + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + SetDefaultData(); + + // Device to host + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), + width, height); + myparms.srcArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); + free(hOutputData); + + // DeAllocating the Memory + DeAllocateMemory(); + } else { + SUCCEED("Skipped the test as there is no peer access"); + } +} + +/* + * This API verifies both D2D functionalities of hipMemcpy3DAsync API + * by allocating memory in one GPU and calling the hipMemcpy3DAsync API + * from another GPU. + * + * D2D case: + * Input : "arr" variable is initialized with the "hData" variable in GPU-0 + * Output: "arr2" variable in GPU-0 + * + * hipMemcpy3DAsync API is triggered from GPU-1 + * The "arr2" variable is then copied to "hOutputData" for validating + * the result + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ +template +void Memcpy3DAsync::D2D_DeviceMem_OnDiffDevice() { + HIP_CHECK(hipSetDevice(0)); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 0, 1)); + if (peerAccess) { + // Allocating Memory and setting default data + AllocateMemory(); + hipStream_t stream1; + HIP_CHECK(hipStreamCreate(&stream1)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + + // Host to Device Scenario + myparms.srcPtr = make_hipPitchedPtr(hData, + width * sizeof(T), + width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream1) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream1)); + + // Allocating Mem on GPU device 0 and trigger hipMemcpy3DAsync from GPU 1 + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + hipArray *arr2; + hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(T)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&arr2, &channelDesc1, + make_hipExtent(width, height, + depth), hipArrayDefault)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + + // Device to Device + myparms.srcArray = arr; + myparms.dstArray = arr2; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // For validating the D2D copy copying it again to hOutputData and + // verifying it with iniital data hData + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + SetDefaultData(); + + // Device to host + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), + width, height); + myparms.srcArray = arr2; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream)== hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); + + // Deleting the memory + free(hOutputData); + DeAllocateMemory(); + } else { + SUCCEED("Skipped the test as there is no peer access"); + } +} + +/* + * This API verifies all the negative scenarios of hipMemcpy3D API +*/ +template +void Memcpy3DAsync::NegativeTests() { + HIP_CHECK(hipSetDevice(0)); + AllocateMemory(); + HIP_CHECK(hipStreamCreate(&stream)); + + // Initialization of data + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.extent = make_hipExtent(width , height, depth); +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + + SECTION("Nullptr to destination array") { + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = nullptr; + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Nullptr to source array") { + myparms.srcArray = nullptr; + myparms.dstPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing both Source ptr and array") { + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.srcArray = arr; + myparms.dstArray = arr1; + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing both destination ptr and array") { + myparms.dstPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; + myparms.srcArray = arr1; + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing Max value to extent") { + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(std::numeric_limits::max(), + std::numeric_limits::max(), + std::numeric_limits::max()); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing Source pitchedPtr as nullptr") { + myparms.srcPtr = make_hipPitchedPtr(nullptr, width * sizeof(T), + width, height); + myparms.dstArray = arr; + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing Dst pitchedPtr as nullptr") { + myparms.dstPtr = make_hipPitchedPtr(nullptr, width * sizeof(T), + width, height); + myparms.srcArray = arr; + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing width > max width size in extent") { + myparms.extent = make_hipExtent(width+1 , height, depth); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing hgt > max width size in extent") { + myparms.extent = make_hipExtent(width , height+1, depth); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + myparms.dstArray = arr; + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing depth > max width size in extent") { + myparms.extent = make_hipExtent(width , height, depth+1); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing dst width pos > max allocated width") { + myparms.dstPos = make_hipPos(width+1, 0, 0); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing dst height pos > max allocated hgt") { + myparms.dstPos = make_hipPos(0, height+1, 0); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing dst depth pos > max allocated depth") { + myparms.dstPos = make_hipPos(0, 0, depth+1); + myparms.dstArray = arr; + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), + width, height); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing src width pos > max allocated width") { + myparms.srcPos = make_hipPos(width+1, 0, 0); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing src height pos > max allocated hgt") { + myparms.srcPos = make_hipPos(0, height+1, 0); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing src height pos > max allocated hgt") { + myparms.srcPos = make_hipPos(0, 0, depth+1); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing src array size > dst array size") { + hipArray *arr2; + hipChannelFormatDesc channelDesc1 = hipCreateChannelDesc(sizeof(T)*8, + 0, 0, 0, formatKind); + HIP_CHECK(hipMalloc3DArray(&arr2, &channelDesc1, + make_hipExtent(3, 3 + , 3), hipArrayDefault)); + myparms.srcArray = arr; + myparms.dstArray = arr2; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) != hipSuccess); + } + + SECTION("Passing nullptr to hipMemcpy3DAsync") { + REQUIRE(hipMemcpy3DAsync(nullptr, stream) != hipSuccess); + } + // DeAllocating of memory + DeAllocateMemory(); +} + +/* + * This API verifies both D2D functionalities of hipMemcpy3DAsync API + * by allocating memory in one GPU and calling the hipMemcpy3DAsync API + * from another GPU. + * + * D2D case: + * Input : "arr" variable is initialized with the "hData" variable in GPU-0 + * Output: "arr1" variable in GPU-0 + * + * Stream on GPU-1 + * The "arr2" variable is then copied to "hOutputData" for validating + * the result + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ +template +void Memcpy3DAsync::D2D_SameDeviceMem_StreamDiffDevice() { + HIP_CHECK(hipSetDevice(0)); + AllocateMemory(); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&stream)); + SetDefaultData(); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + + // Host to Device + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Array to Array + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + SetDefaultData(); + + // Device to host + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), width, height); + myparms.srcArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); + + // Deallocating the memory + free(hOutputData); + DeAllocateMemory(); +} + +/* + * This API verifies the Extent validation Scenarios + */ +template +void Memcpy3DAsync::Extent_Validation() { + HIP_CHECK(hipSetDevice(0)); + AllocateMemory(); + HIP_CHECK(hipStreamCreate(&stream)); + // Passing extent as 0 + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + SECTION("Passing Extent as 0") { + myparms.extent = make_hipExtent(0 , 0, 0); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + } + SECTION("Passing Width 0 in Extent") { + myparms.extent = make_hipExtent(0 , height, depth); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + } + SECTION("Passing Height 0 in Extent") { + myparms.extent = make_hipExtent(width , 0, depth); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + } + SECTION("Passing Depth 0 in Extent") { + myparms.extent = make_hipExtent(width , height, 0); + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + } + DeAllocateMemory(); +} + +/* + * This API verifies H2H-D2D-D2H functionalities of hipMemcpy3DAsync API + * + * Input : "arr" variable is initialized with the "hData" variable in GPU-0 + * Output: "arr1" variable in GPU-0 + * + * The "arr1" variable is then copied to "hOutputData" for validating + * the result + * + * Validating the result by comparing "hData" and "hOutputData" variables + */ +template +void Memcpy3DAsync::simple_Memcpy3DAsync() { + HIP_CHECK(hipSetDevice(0)); + + // Allocating the Memory + AllocateMemory(); + HIP_CHECK(hipStreamCreate(&stream)); + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + + // Host to Device + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyHostToDevice; +#else + myparms.kind = hipMemcpyHostToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Array to Array + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.srcArray = arr; + myparms.dstArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToDevice; +#else + myparms.kind = hipMemcpyDeviceToDevice; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + T *hOutputData = reinterpret_cast(malloc(size)); + memset(hOutputData, 0, size); + + // Device to host + memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); + SetDefaultData(); + myparms.dstPtr = make_hipPitchedPtr(hOutputData, + width * sizeof(T), width, height); + myparms.srcArray = arr1; +#ifdef __HIP_PLATFORM_NVCC__ + myparms.kind = cudaMemcpyDeviceToHost; +#else + myparms.kind = hipMemcpyDeviceToHost; +#endif + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validating the result + HipTest::checkArray(hData, hOutputData, width, height, depth); + + // DeAllocating the memory + free(hOutputData); + DeAllocateMemory(); +} +/* +This testcase verifies hipMemcpyAsync for different datatypes +and different sizes +*/ +TEMPLATE_TEST_CASE("Unit_hipMemcpy3DAsync_Basic", + "[hipMemcpy3DAsync]", + int, unsigned int, float) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + auto i = GENERATE(10, 100, 1024, 10*1024); + auto j = GENERATE(10, 100); + if (numDevices > 1) { + if (std::is_same::value) { + Memcpy3DAsync memcpy3d_obj(i, j, j, + hipChannelFormatKindSigned); + memcpy3d_obj.simple_Memcpy3DAsync(); + } else if (std::is_same::value) { + Memcpy3DAsync memcpy3d_obj(i, j, j, + hipChannelFormatKindUnsigned); + memcpy3d_obj.simple_Memcpy3DAsync(); + } else if (std::is_same::value) { + Memcpy3DAsync memcpy3d_obj(i, j, j, + hipChannelFormatKindFloat); + memcpy3d_obj.simple_Memcpy3DAsync(); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* +This testcase performs the extent validation scenarios of +hipMemcpy3D API +*/ +TEST_CASE("Unit_hipMemcpy3DAsync_ExtentValidation") { + Memcpy3DAsync memcpy3d(width, height, depth, + hipChannelFormatKindSigned); + memcpy3d.Extent_Validation(); +} + +/* +This testcase performs the negative scenarios of +hipMemcpy3DAsync API +*/ +TEST_CASE("Unit_hipMemcpy3DAsync_multiDevice-Negative") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + Memcpy3DAsync memcpy3d(width, height, depth, + hipChannelFormatKindSigned); + memcpy3d.NegativeTests(); + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* +This testcase performs the D2H,H2D and D2D on peer +GPU device +*/ +TEST_CASE("Unit_hipMemcpy3DAsync_multiDevice-D2D") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + SECTION("D2D on different Device") { + Memcpy3DAsync memcpy3d_d2d_obj(width, height, depth, + hipChannelFormatKindFloat); + memcpy3d_d2d_obj.D2D_DeviceMem_OnDiffDevice(); + } + + SECTION("D2H and H2D on different device") { + Memcpy3DAsync memcpy3d_d2h_obj(width, height, depth, + hipChannelFormatKindFloat); + memcpy3d_d2h_obj.D2H_H2D_DeviceMem_OnDiffDevice(); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* +This testcase checks hipMemcpy3DAsync API by +allocating memory in one GPU and creating stream +in another GPU +*/ +TEST_CASE("Unit_hipMemcpy3DAsync_multiDevice-DiffStream") { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + Memcpy3DAsync memcpy3dAsync(width, height, depth, + hipChannelFormatKindSigned); + memcpy3dAsync.D2D_SameDeviceMem_StreamDiffDevice(); + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyParam2D.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyParam2D.cc new file mode 100644 index 0000000000..7420954b9f --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyParam2D.cc @@ -0,0 +1,380 @@ +/* +Copyright (c) 2021-present 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. +*/ + +/* +This testfile verifies the following scenarios of hipMemcpyParam2D API +1. Negative Scenarios +2. Extent Validation Scenarios +3. D2D copy for different datatypes +4. H2D and D2H copy for different datatypes +*/ + +#include +#include + +static constexpr size_t NUM_W{10}; +static constexpr size_t NUM_H{10}; +/* + * This testcase verifies D2D functionality of hipMemcpyParam2D API + * Input: Intializing "A_d" device variable with "C_h" host variable + * Output: "A_d" device variable to "E_d" device variable + * + * Validating the result by copying "E_d" to "A_h" and checking + * it with the initalized data "C_h". + * + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyParam2D_multiDevice-D2D", + "[hipMemcpyParam2D]", + char, float, int, + double, long double) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + // Initialize and Allocating Memory + HIP_CHECK(hipSetDevice(0)); + TestType* A_h{nullptr}, *C_h{nullptr}, *A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(TestType)}; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + HIP_CHECK(hipSetDevice(1)); + char *E_d; + size_t pitch_E; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&E_d), + &pitch_E, width, NUM_H)); + + // Initalizing A_d with C_h + HIP_CHECK(hipMemcpy2D(A_d, pitch_A, C_h, width, + NUM_W * sizeof(TestType), NUM_H, hipMemcpyHostToDevice)); + + // Device to Device + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = E_d; + desc.dstDevice = hipDeviceptr_t(E_d); + desc.dstPitch = pitch_E; + desc.WidthInBytes = NUM_W * sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + + // Copying E_d to A_h + HIP_CHECK(hipMemcpy2D(A_h, width, E_d, pitch_E, + NUM_W * sizeof(TestType), NUM_H, + hipMemcpyDeviceToHost)); + + // Validating the result + REQUIRE(HipTest::checkArray(A_h, C_h, NUM_W, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, false); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* + * This testcase verifies H2D & D2H functionality of hipMemcpyParam2D API + * H2D case: + * Input: "C_h" host variable initialized with default data + * Output: "A_d" device variable + * + * D2H case: + * Input: "A_d" device variable from the previous output + * OutPut: "A_h" variable + * + * Validating the result by comparing "A_h" to "C_h" + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyParam2D_multiDevice-H2D-D2H", + "[hipMemcpyParam2D]", char, float, + int, double, long double) { + // 1 refers to pinned host memory and 0 refers + // to unpinned memory + auto memory_type = GENERATE(0, 1); + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipSetDevice(0)); + + // Initialize and Allocating Memory + TestType* A_h{nullptr}, *C_h{nullptr}, + *A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(TestType)}; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + + // Based on memory type (pinned/unpinned) allocating memory + if (memory_type) { + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, true); + } else { + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + } + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + // Host to Device + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.srcMemoryType = hipMemoryTypeHost; +#endif + desc.srcHost = C_h; + desc.srcDevice = hipDeviceptr_t(C_h); + desc.srcPitch = width; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = A_d; + desc.dstDevice = hipDeviceptr_t(A_d); + desc.dstPitch = pitch_A; + desc.WidthInBytes = NUM_W*sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + + // Device to Host + memset(&desc, 0x0, sizeof(hip_Memcpy2D)); +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.dstMemoryType = hipMemoryTypeHost; +#endif + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = NUM_W*sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + + // Validating the result + REQUIRE(HipTest::checkArray(A_h, C_h, NUM_W, NUM_H) == true); + + // DeAllocating the Memory + HIP_CHECK(hipFree(A_d)); + if (memory_type) { + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, true); + } else { + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, false); + } + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} +/* + * This testcase verifies the extent validation scenarios + */ +TEST_CASE("Unit_hipMemcpyParam2D_ExtentValidation") { + // Allocating memory and Initializing the data + HIP_CHECK(hipSetDevice(0)); + char* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, + * A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(char)}; + constexpr auto memsetval{100}; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + HipTest::initArrays(nullptr, nullptr, nullptr, + &B_h, nullptr, nullptr, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + HipTest::setDefaultData(NUM_W*NUM_H, B_h, nullptr, nullptr); + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, NUM_W, NUM_H)); + + // Device to Host + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.dstMemoryType = hipMemoryTypeHost; +#endif + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = NUM_W; + desc.Height = NUM_H; + + SECTION("Destination Pitch is 0") { + desc.dstPitch = 0; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + } + + SECTION("Source Pitch is 0") { + desc.srcPitch = 0; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + } + + SECTION("Height is 0") { + desc.Height = 0; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + REQUIRE(HipTest::checkArray(A_h, B_h, NUM_W, NUM_H) == true); + } + + SECTION("Width is 0") { + desc.WidthInBytes = 0; + REQUIRE(hipMemcpyParam2D(&desc) == hipSuccess); + REQUIRE(HipTest::checkArray(A_h, B_h, NUM_W, NUM_H) == true); + } + + // DeAllocating the Memory + HIP_CHECK(hipFree(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +} + +/* + * This testcase verifies the negative scenarios + */ +TEST_CASE("Unit_hipMemcpyParam2D_Negative") { + HIP_CHECK(hipSetDevice(0)); + + // Allocating and Initializing the data + float* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, + * A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(float)}; + constexpr auto memsetval{100}; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, &C_h, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, B_h, C_h); + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, NUM_W, NUM_H)); + + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.dstMemoryType = hipMemoryTypeHost; +#endif + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = NUM_W; + desc.Height = NUM_H; + + SECTION("Null Pointer to Source Device Pointer") { + desc.srcDevice = hipDeviceptr_t(nullptr); + REQUIRE(hipMemcpyParam2D(&desc) != hipSuccess); + } + + SECTION("Null Pointer to Destination Device Pointer") { + memset(&desc, 0x0, sizeof(hip_Memcpy2D)); +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.srcMemoryType = hipMemoryTypeHost; +#endif + desc.srcHost = A_h; + desc.srcDevice = hipDeviceptr_t(A_h); + desc.srcPitch = width; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = A_d; + desc.dstDevice = hipDeviceptr_t(nullptr); + desc.dstPitch = pitch_A; + desc.WidthInBytes = NUM_W; + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2D(&desc) != hipSuccess); + } + + SECTION("Null Pointer to both Src & Dst Device Pointer") { + desc.srcDevice = hipDeviceptr_t(nullptr); + desc.dstDevice = hipDeviceptr_t(nullptr); + REQUIRE(hipMemcpyParam2D(&desc) != hipSuccess); + } + + SECTION("Width > src/dest pitches") { + desc.WidthInBytes = pitch_A+1; + REQUIRE(hipMemcpyParam2D(&desc) != hipSuccess); + } + + // DeAllocating the Memory + HIP_CHECK(hipFree(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +} diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyParam2DAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyParam2DAsync.cc new file mode 100644 index 0000000000..dac78feab4 --- /dev/null +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyParam2DAsync.cc @@ -0,0 +1,487 @@ +/* +Copyright (c) 2021-present 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. +*/ + +/* +This testfile verifies the following scenarios of hipMemcpyParam2DAsync API +1. Negative Scenarios +2. Extent Validation Scenarios +3. D2D copy for different datatypes +4. H2D and D2H copy for different datatypes +5. Device context change scenario where memory allocated in one GPU + stream created in another GPU +*/ + +#include +#include + +static constexpr size_t NUM_W{10}; +static constexpr size_t NUM_H{10}; +/* + * This testcase verifies D2D functionality of hipMemcpyParam2DAsync API + * Where Memory is allocated in GPU-0 and stream is created in GPU-1 + * + * Input: Intializing "A_d" device variable with "C_h" host variable + * Output: "A_d" device variable to "E_d" device variable + * + * Validating the result by copying "E_d" to "A_h" and checking + * it with the initalized data "C_h". + * + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice", + "[hipMemcpyParam2DAsync]", char, float, int, + double, long double) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + // Allocating and Initializing the data + HIP_CHECK(hipSetDevice(0)); + TestType* A_h{nullptr}, *C_h{nullptr}, *A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(TestType)}; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + TestType *E_d{nullptr}; + size_t pitch_E; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&E_d), + &pitch_E, width, NUM_H)); + + // Initalizing A_d with C_h + HIP_CHECK(hipMemcpy2D(A_d, pitch_A, C_h, width, + NUM_W*sizeof(TestType), NUM_H, + hipMemcpyHostToDevice)); + HIP_CHECK(hipSetDevice(1)); + hipStream_t stream; + hipStreamCreate(&stream); + + // Device to Device + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = E_d; + desc.dstDevice = hipDeviceptr_t(E_d); + desc.dstPitch = pitch_E; + desc.WidthInBytes = NUM_W*sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Copying the result E_d to A_h host variable + HIP_CHECK(hipMemcpy2D(A_h, width, E_d, pitch_E, + NUM_W*sizeof(TestType), NUM_H, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + // Validating the result + REQUIRE(HipTest::checkArray(A_h, C_h, NUM_W, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFree(E_d)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, false); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* + * This testcase verifies D2D functionality of hipMemcpyParam2DAsync API + * Input: Intializing "A_d" device variable with "C_h" host variable + * Output: "A_d" device variable to "E_d" device variable + * + * Validating the result by copying "E_d" to "A_h" and checking + * it with the initalized data "C_h". + * + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyParam2DAsync_multiDevice-D2D", + "[hipMemcpyParam2DAsync]", char, + int, float, double, long double) { + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + // Allocating and Initializing the data + HIP_CHECK(hipSetDevice(0)); + TestType* A_h{nullptr}, *C_h{nullptr}, *A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(TestType)}; + hipStream_t stream; + hipStreamCreate(&stream); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + HIP_CHECK(hipSetDevice(1)); + TestType *E_d; + size_t pitch_E; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&E_d), + &pitch_E, width, NUM_H)); + + // Initializing A_d with C_h + HIP_CHECK(hipMemcpy2D(A_d, pitch_A, C_h, width, + NUM_W*sizeof(TestType), NUM_H, hipMemcpyHostToDevice)); + + // Device to Device + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = E_d; + desc.dstDevice = hipDeviceptr_t(E_d); + desc.dstPitch = pitch_E; + desc.WidthInBytes = NUM_W*sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Copying the result E_d to A_h host variable + HIP_CHECK(hipMemcpy2D(A_h, width, E_d, pitch_E, + NUM_W*sizeof(TestType), NUM_H, hipMemcpyDeviceToHost)); + + // Validating the result + REQUIRE(HipTest::checkArray(A_h, C_h, NUM_W, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, false); + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} + +/* + * This testcase verifies H2D & D2H functionality of hipMemcpyParam2DAsync API + * H2D case: + * Input: "C_h" host variable initialized with default data + * Output: "A_d" device variable + * + * D2H case: + * Input: "A_d" device variable from the previous output + * OutPut: "A_h" variable + * + * Validating the result by comparing "A_h" to "C_h" + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpyParam2DAsync_multiDevice-H2D-D2H", + "[hipMemcpyParam2DAsync]", char, + int, float, double, long double) { + // 1 refers to pinned host memory and 0 refers + // to unpinned memory + auto memory_type = GENERATE(0, 1); + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + // Allocating and Initializing the data + HIP_CHECK(hipSetDevice(0)); + TestType* A_h{nullptr}, *C_h{nullptr}, + *A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(TestType)}; + hipStream_t stream; + + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + + // Based on memory type (pinned/unpinned) allocating memory + if (memory_type) { + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, true); + } else { + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + } + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + int peerAccess = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + if (!peerAccess) { + SUCCEED("Skipped the test as there is no peer access"); + } else { + // Host to Device + hip_Memcpy2D desc = {}; + hipStreamCreate(&stream); +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.srcMemoryType = hipMemoryTypeHost; +#endif + desc.srcHost = C_h; + desc.srcDevice = hipDeviceptr_t(C_h); + desc.srcPitch = width; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = A_d; + desc.dstDevice = hipDeviceptr_t(A_d); + desc.dstPitch = pitch_A; + desc.WidthInBytes = NUM_W*sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Device to Host + memset(&desc, 0x0, sizeof(hip_Memcpy2D)); +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.dstMemoryType = hipMemoryTypeHost; +#endif + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = NUM_W*sizeof(TestType); + desc.Height = NUM_H; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Validating the result + REQUIRE(HipTest::checkArray(A_h, C_h, NUM_W, NUM_H) == true); + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); + if (memory_type) { + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, true); + } else { + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, nullptr, C_h, false); + } + } + } else { + SUCCEED("skipping the testcases as numDevices < 2"); + } +} +/* + * This testcase verifies the extent validation scenarios + */ +TEST_CASE("Unit_hipMemcpyParam2DAsync_ExtentValidation") { + HIP_CHECK(hipSetDevice(0)); + char* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, + * A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(char)}; + constexpr auto memsetval{100}; + hipStream_t stream; + hipStreamCreate(&stream); + + // Allocating and Initializing the data + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, nullptr, &C_h, + width*NUM_H, false); + HipTest::initArrays(nullptr, nullptr, nullptr, + &B_h, nullptr, nullptr, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, nullptr, C_h); + HipTest::setDefaultData(NUM_W*NUM_H, B_h, nullptr, nullptr); + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, NUM_W, NUM_H)); + + // Device to Host + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.dstMemoryType = hipMemoryTypeHost; +#endif + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = NUM_W; + desc.Height = NUM_H; + + SECTION("Destination Pitch is 0") { + desc.dstPitch = 0; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + } + + SECTION("Source Pitch is 0") { + desc.srcPitch = 0; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + } + + SECTION("Height is 0") { + desc.Height = 0; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, B_h, NUM_W, NUM_H) == true); + } + + SECTION("Width is 0") { + desc.Height = 0; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(HipTest::checkArray(A_h, B_h, NUM_W, NUM_H) == true); + } + + // DeAllocating the Memory + HIP_CHECK(hipFree(A_d)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +} + +/* + * This testcase verifies the negative scenarios + */ +TEST_CASE("Unit_hipMemcpyParam2DAsync_Negative") { + HIP_CHECK(hipSetDevice(0)); + float* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, + * A_d{nullptr}; + size_t pitch_A; + size_t width{NUM_W * sizeof(float)}; + constexpr auto memsetval{100}; + hipStream_t stream; + hipStreamCreate(&stream); + + // Allocating and Initializing the data + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), + &pitch_A, width, NUM_H)); + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, &C_h, + width*NUM_H, false); + HipTest::setDefaultData(NUM_W*NUM_H, A_h, B_h, C_h); + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, NUM_W, NUM_H)); + + // Device to Host + hip_Memcpy2D desc = {}; +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.srcMemoryType = hipMemoryTypeDevice; +#endif + desc.srcHost = A_d; + desc.srcDevice = hipDeviceptr_t(A_d); + desc.srcPitch = pitch_A; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.dstMemoryType = hipMemoryTypeHost; +#endif + desc.dstHost = A_h; + desc.dstDevice = hipDeviceptr_t(A_h); + desc.dstPitch = width; + desc.WidthInBytes = NUM_W; + desc.Height = NUM_H; + + SECTION("Null Pointer to Source Device Pointer") { + desc.srcDevice = hipDeviceptr_t(nullptr); + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) != hipSuccess); + } + + SECTION("Null Pointer to Destination Device Pointer") { + memset(&desc, 0x0, sizeof(hip_Memcpy2D)); +#ifdef __HIP_PLATFORM_NVCC__ + desc.srcMemoryType = CU_MEMORYTYPE_HOST; +#else + desc.srcMemoryType = hipMemoryTypeHost; +#endif + desc.srcHost = A_h; + desc.srcDevice = hipDeviceptr_t(A_h); + desc.srcPitch = width; +#ifdef __HIP_PLATFORM_NVCC__ + desc.dstMemoryType = CU_MEMORYTYPE_DEVICE; +#else + desc.dstMemoryType = hipMemoryTypeDevice; +#endif + desc.dstHost = A_d; + desc.dstDevice = hipDeviceptr_t(nullptr); + desc.dstPitch = pitch_A; + desc.WidthInBytes = NUM_W; + desc.Height = NUM_H; + + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) != hipSuccess); + } + + SECTION("Null Pointer to both Src & Dst Device Pointer") { + desc.srcDevice = hipDeviceptr_t(nullptr); + desc.dstDevice = hipDeviceptr_t(nullptr); + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) != hipSuccess); + } + + SECTION("Width > src/dest pitches") { + desc.WidthInBytes = pitch_A+1; + REQUIRE(hipMemcpyParam2DAsync(&desc, stream) != hipSuccess); + } + + // DeAllocating the memory + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamSynchronize(stream)); + HipTest::freeArrays(nullptr, nullptr, nullptr, + A_h, B_h, C_h, false); +}