/* * 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 hipMemcpy3DPeer hipMemcpy3DPeer* @{ *@ingroup MemoryTest* `hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms* p)` - *Performs 3D memory copies between devices. */ /** * Test Description * ------------------------ * - Test case to verify the 3D peer to peer device copy. * 1. Allocate device memory for two Arrays (array_1, array_2). * 2. Fill the source array with some data with hipMemcpy3D. * 3. Transfer data between two peer devices with hipMemcpy3DPeer. * 4. Verify the data on host. * Test source * ------------------------ * - catch/unit/memory/hipMemcpy3DPeer.cc * Test requirements * ------------------------ * - HIP_VERSION >= 7.1 */ TEST_CASE("Unit_hipMemcpy3DPeer_BasicFunctional") { CHECK_IMAGE_SUPPORT constexpr int numW = 16; constexpr int numH = 16; constexpr int depth = 4; size_t volume = numW * numH * depth; hipExtent extent = make_hipExtent(numW, numH, depth); const auto device_count = HipTest::getDeviceCount(); if (device_count <= 1) { std::string msg = "Invalid Device Count. Hence Skipping the test.. "; HipTest::HIP_SKIP_TEST(msg.c_str()); } const auto src_device = GENERATE_COPY(range(0, device_count)); const auto dst_device = GENERATE_COPY(range(0, device_count)); if (src_device == dst_device) { std::string msg = "Both Source and Destination device ids are same."; INFO("Src device: " << src_device << ", Dst device: " << dst_device); HipTest::HIP_SKIP_TEST(msg.c_str()); } HIP_CHECK(hipSetDevice(src_device)); int can_access_peer = 0; HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device)); if (!can_access_peer) { std::string msg = "Skipped as peer access cannot be enabled between devices " + std::to_string(src_device) + " " + std::to_string(dst_device); HipTest::HIP_SKIP_TEST(msg.c_str()); } // Array-1 Memory allocation hipChannelFormatDesc channelDesc_1 = hipCreateChannelDesc(); hipArray_t array_1; HIP_CHECK(hipMalloc3DArray(&array_1, &channelDesc_1, extent, 0)); // Set the array memory std::vectortmpHost(volume, 0xb); hipMemcpy3DParms fillParms{}; fillParms.srcPos = make_hipPos(0, 0, 0); fillParms.dstPos = make_hipPos(0, 0, 0); fillParms.srcPtr = make_hipPitchedPtr(tmpHost.data(), numW, numW, numH); fillParms.dstArray = array_1; fillParms.extent = extent; fillParms.kind = hipMemcpyHostToDevice; HIP_CHECK(hipMemcpy3D(&fillParms)); // Array-2 Memory allocation HIP_CHECK(hipSetDevice(dst_device)); hipChannelFormatDesc channelDesc_2 = hipCreateChannelDesc(); hipArray_t array_2; HIP_CHECK(hipMalloc3DArray(&array_2, &channelDesc_2, extent, 0)); HIP_CHECK(hipSetDevice(src_device)); // Copy data to peer-peer device hipMemcpy3DPeerParms peerCopyParams{}; peerCopyParams.dstArray = array_2; peerCopyParams.dstDevice = dst_device; peerCopyParams.dstPos = make_hipPos(0, 0, 0); peerCopyParams.dstPtr = make_hipPitchedPtr(0, 0, 0, 0); peerCopyParams.srcArray = array_1; peerCopyParams.srcDevice = src_device; peerCopyParams.srcPos = make_hipPos(0, 0, 0); peerCopyParams.srcPtr = make_hipPitchedPtr(0, 0, 0, 0); peerCopyParams.extent = extent; HIP_CHECK(hipMemcpy3DPeer(&peerCopyParams)); // Copy data from Device Array- host ptr std::vectorhostBuf(volume, 0xa); hipMemcpy3DParms copyParms{}; copyParms.srcArray = array_2; copyParms.dstPtr = make_hipPitchedPtr(hostBuf.data(), numW, numW, numH); copyParms.extent = extent; copyParms.kind = hipMemcpyDeviceToHost; HIP_CHECK(hipMemcpy3D(©Parms)); // Validation for (size_t i = 0; i < volume; i++) { INFO("Array FAILURE at Index: "<< i << "\nval : " < max width size in extent * 7. Passing height > max height size in extent * 8. Passing depth > max depth size in extent * Test source * ------------------------ * - catch/unit/memory/hipMemcpy3DPeer.cc * Test requirements * ------------------------ * - HIP_VERSION >= 7.1 */ TEST_CASE("Unit_hipMemcpy3DPeer_NegativeTsts") { CHECK_IMAGE_SUPPORT constexpr int numW = 16; constexpr int numH = 16; constexpr int depth = 4; hipExtent extent = make_hipExtent(numW, numH, depth); const auto device_count = HipTest::getDeviceCount(); const auto src_device = GENERATE_COPY(range(0, device_count)); const auto dst_device = GENERATE_COPY(range(0, device_count)); HIP_CHECK(hipSetDevice(src_device)); // Array-2 Memory allocation hipArray_t array_1; hipChannelFormatDesc channelDesc_1 = hipCreateChannelDesc(); HIP_CHECK(hipMalloc3DArray(&array_1, &channelDesc_1, extent, 0)); // Array-2 Memory allocation HIP_CHECK(hipSetDevice(dst_device)); hipArray_t array_2; hipChannelFormatDesc channelDesc_2 = hipCreateChannelDesc(); HIP_CHECK(hipMalloc3DArray(&array_2, &channelDesc_2, extent, 0)); hipMemcpy3DPeerParms peerCopyParams{}; peerCopyParams.dstArray = array_2; peerCopyParams.dstPos = make_hipPos(0, 0, 0); peerCopyParams.dstPtr = make_hipPitchedPtr(0, 0, 0, 0); peerCopyParams.srcArray = array_1; peerCopyParams.srcPos = make_hipPos(0, 0, 0); peerCopyParams.srcPtr = make_hipPitchedPtr(0, 0, 0, 0); peerCopyParams.extent = extent; SECTION("Max Destination device id") { peerCopyParams.dstDevice = device_count + 1; HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidDevice); } SECTION("Max Source device id") { peerCopyParams.srcDevice = device_count + 1; HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidDevice); } SECTION("Neg Source device id") { peerCopyParams.srcDevice = -1; HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidDevice); } SECTION("Neg Destination device id") { peerCopyParams.dstDevice = -1; HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidDevice); } SECTION("Source Array as Null") { peerCopyParams.srcArray = nullptr; HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidValue); } SECTION("Destination Array as Null") { peerCopyParams.dstArray = nullptr; HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidValue); } SECTION("Passing Max value to extent") { peerCopyParams.extent = make_hipExtent(std::numeric_limits::max(), std::numeric_limits::max(), std::numeric_limits::max()); HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidValue); } SECTION("Passing width > max width size in extent") { peerCopyParams.extent = make_hipExtent(numW + 1, numH, depth); HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidValue); } SECTION("Passing height > max height size in extent") { peerCopyParams.extent = make_hipExtent(numW, numH + 1, depth); HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidValue); } SECTION("Passing depth > max depth size in extent") { peerCopyParams.extent = make_hipExtent(numW, numH, depth + 1); HIP_CHECK_ERROR(hipMemcpy3DPeer(&peerCopyParams), hipErrorInvalidValue); } SECTION("Memcpy3DPeer Params struct as nullptr") { HIP_CHECK_ERROR(hipMemcpy3D(nullptr), hipErrorInvalidValue); } HIP_CHECK(hipFreeArray(array_1)); HIP_CHECK(hipFreeArray(array_2)); } /** * End doxygen group MemoryTest. * @} */