diff --git a/catch/unit/executionControl/CMakeLists.txt b/catch/unit/executionControl/CMakeLists.txt index 5e30302b8..64216e265 100644 --- a/catch/unit/executionControl/CMakeLists.txt +++ b/catch/unit/executionControl/CMakeLists.txt @@ -12,7 +12,7 @@ if(HIP_PLATFORM MATCHES "amd") hipExtLaunchMultiKernelMultiDevice.cc launch_api.cc hipGetProcAddressLaunchCbExecCtrlApis.cc - ) + hipLaunchKernel_spt.cc) else() # These functions are currently unimplemented on AMD set(TEST_SRC ${TEST_SRC} diff --git a/catch/unit/executionControl/hipLaunchKernel_spt.cc b/catch/unit/executionControl/hipLaunchKernel_spt.cc new file mode 100644 index 000000000..55caaaec5 --- /dev/null +++ b/catch/unit/executionControl/hipLaunchKernel_spt.cc @@ -0,0 +1,152 @@ +/* +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 "execution_control_common.hh" +#include +#include +#include +#include +#include +/** + * @addtogroup hipLaunchKernel_spt hipLaunchKernel_spt + * @{ + * @ingroup ExecutionTest + * `hipError_t hipLaunchKernel_spt(const void* function_address, + dim3 numBlocks, + dim3 dimBlocks, + void** args, + size_t sharedMemBytes __dparm(0), + hipStream_t stream __dparm(0))` - + * C compliant kernel launch API + */ +/** + * Test Description + * ------------------------ + * - Basic test to verify the basic positive behavior of hipLaunchKernel_spt.. + * Test source + * ------------------------ + * - catch\unit\executionControl\hipLaunchKernel_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipLaunchKernel_spt_Positive_Basic") { + SECTION("Kernel with no arguments") { + HIP_CHECK(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, + nullptr, 0, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + SECTION("Kernel with arguments using kernelParams") { + LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); + int* result_ptr = result_dev.ptr(); + void* kernel_args[1] = {&result_ptr}; + HIP_CHECK(hipLaunchKernel_spt(reinterpret_cast(kernel_42), dim3{1, 1, 1}, dim3{1, 1, 1}, + kernel_args, 0, nullptr)); + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } +} +/** + * Test Description + * ------------------------ + * - Basic test to verify the basic functionality with all positive parameters + * of hipLaunchKernel_spt. Test source + * ------------------------ + * - catch\unit\executionControl\hipLaunchKernel_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipLaunchKernel_spt_Positive_Parameters") { + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); + HIP_CHECK(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, + nullptr, 0, nullptr)); + } + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); + HIP_CHECK(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{y, 1, 1}, + nullptr, 0, nullptr)); + } + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); + HIP_CHECK(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{z, 1, 1}, + nullptr, 0, nullptr)); + } +} +/** + * Test Description + * ------------------------ + * - Basic test to verify the negative cases of hipLaunchKernel_spt. + * Test source + * ------------------------ + * - catch\unit\executionControl\hipLaunchKernel_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipLaunchKernel_spt_Negative_Parameters") { + SECTION("f == nullptr") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(nullptr, dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidDeviceFunction); + } + SECTION("gridDim.x == 0") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{0, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + SECTION("gridDim.y == 0") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 0, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + SECTION("gridDim.z == 0") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 0}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + SECTION("blockDim.x == 0") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{0, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + SECTION("blockDim.y == 0") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 0, 1}, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + SECTION("blockDim.z == 0") { + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 0}, nullptr, 0, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid stream") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipLaunchKernel_spt(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, stream), + hipErrorInvalidValue); + } +} +/** + * End doxygen group ExecutionTest. + * @} + */ diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 6beae6009..dd18d53d1 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -106,7 +106,22 @@ if(HIP_PLATFORM MATCHES "amd") hipMemVmm.cc hipArray.cc hipMemcpyDeviceToDeviceNoCU.cc - hipGetProcAddressMemoryApis.cc) + hipGetProcAddressMemoryApis.cc + hipMemcpy2DAsync_spt + hipMemcpy2DFromArrayAsync_spt + hipMemcpy2DFromArray_spt + hipMemcpy2D_spt + hipMemcpy2DToArrayAsync_spt + hipMemcpy2DToArray_spt + hipMemcpy3DAsync_spt + hipMemcpy3D_spt + hipMemcpyAsync_spt + hipMemcpyFromArray_spt + hipMemcpyFromSymbolAsync_spt + hipMemcpyFromSymbol_spt + hipMemcpy_spt + hipMemcpyToSymbolAsync_spt + hipMemcpyToSymbol_spt) if(UNIX) # Should be compiled for NVIDIA as well after EXSWHTEC-346 is addressed # For windows build error occurs undefined symbol: hipPointerSetAttribute diff --git a/catch/unit/memory/hipMemcpy2DAsync_spt.cc b/catch/unit/memory/hipMemcpy2DAsync_spt.cc new file mode 100644 index 000000000..bb724ecc1 --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DAsync_spt.cc @@ -0,0 +1,229 @@ +/* +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 "memcpy2d_tests_common.hh" +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy2DAsync_spt hipMemcpy2DAsync_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy2DAsync_spt(void* dst, size_t dpitch, const void* src, + size_t spitch, size_t width, + size_t height, hipMemcpyKind kind, hipStream_t stream + __dparm(0))` - + * Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Tests to verify the Synchronization Behavior of hipMemcpy2DAsync_spt + * Test source + * ------------------------ + * - unit/memory/hipMemcpy2DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DAsync_spt_Positive_Synchronization_Behavior") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Host to Device") { + Memcpy2DHtoDSyncBehavior(std::bind(hipMemcpy2DAsync_spt, _1, _2, _3, _4, _5, _6, _7, nullptr), + false); + } + SECTION("Device to Pageable Host") { + Memcpy2DDtoHPageableSyncBehavior( + std::bind(hipMemcpy2DAsync_spt, _1, _2, _3, _4, _5, _6, _7, nullptr), true); + } + SECTION("Device to Pinned Host") { + Memcpy2DDtoHPinnedSyncBehavior( + std::bind(hipMemcpy2DAsync_spt, _1, _2, _3, _4, _5, _6, _7, nullptr), false); + } + SECTION("Device to Device") { + Memcpy2DDtoDSyncBehavior(std::bind(hipMemcpy2DAsync_spt, _1, _2, _3, _4, _5, _6, _7, nullptr), + false); + } + SECTION("Host to Host") { + Memcpy2DHtoHSyncBehavior(std::bind(hipMemcpy2DAsync_spt, _1, _2, _3, _4, _5, _6, _7, nullptr), + true); + } +} +/** + * Test Description + * ------------------------ + * - Tests to verify the functionality of hipMemcpy2DAsync_spt + * with positive parameters. + * Test source + * ------------------------ + * - unit/memory/hipMemcpy2DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DAsync_spt_Positive_Parameters") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + constexpr bool async = true; + Memcpy2DZeroWidthHeight( + std::bind(hipMemcpy2DAsync_spt, _1, _2, _3, _4, _5, _6, _7, nullptr)); +} +/** + * Test Description + * ------------------------ + * - Tests to verify the functionality of hipMemcpy2DAsync_spt + * with negative parameters. + * Test source + * ------------------------ + * - unit/memory/hipMemcpy2DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DAsync_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr size_t cols = 128; + constexpr size_t rows = 128; + constexpr auto NegativeTests = [](void* dst, size_t dpitch, const void* src, size_t spitch, + size_t width, size_t height, hipMemcpyKind kind) { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DAsync_spt(nullptr, dpitch, src, spitch, width, height, kind, nullptr), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DAsync_spt(dst, dpitch, nullptr, spitch, width, height, kind, nullptr), + hipErrorInvalidValue); + } + SECTION("dpitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DAsync_spt(dst, width - 1, src, spitch, width, height, kind, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("spitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DAsync_spt(dst, dpitch, src, width - 1, width, height, kind, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("dpitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + HIP_CHECK_ERROR(hipMemcpy2DAsync_spt(dst, static_cast(attr) + 1, src, spitch, width, + height, kind, nullptr), + hipErrorInvalidValue); + } + SECTION("spitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + HIP_CHECK_ERROR(hipMemcpy2DAsync_spt(dst, dpitch, src, static_cast(attr) + 1, width, + height, kind, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpy2DAsync_spt(dst, dpitch, src, spitch, width, height, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + }; + SECTION("Host to device") { + LinearAllocGuard2D device_alloc(cols, rows); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * rows); + NegativeTests(device_alloc.ptr(), device_alloc.pitch(), host_alloc.ptr(), device_alloc.pitch(), + device_alloc.width(), device_alloc.height(), hipMemcpyHostToDevice); + } + SECTION("Device to host") { + LinearAllocGuard2D device_alloc(cols, rows); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * rows); + NegativeTests(host_alloc.ptr(), device_alloc.pitch(), device_alloc.ptr(), device_alloc.pitch(), + device_alloc.width(), device_alloc.height(), hipMemcpyDeviceToHost); + } + SECTION("Host to host") { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, cols * rows * sizeof(int)); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, cols * rows * sizeof(int)); + NegativeTests(dst_alloc.ptr(), cols * sizeof(int), src_alloc.ptr(), cols * sizeof(int), + cols * sizeof(int), rows, hipMemcpyHostToHost); + } + SECTION("Device to device") { + LinearAllocGuard2D src_alloc(cols, rows); + LinearAllocGuard2D dst_alloc(cols, rows); + NegativeTests(dst_alloc.ptr(), dst_alloc.pitch(), src_alloc.ptr(), src_alloc.pitch(), + dst_alloc.width(), dst_alloc.height(), hipMemcpyDeviceToDevice); + } +} +/** + * Test Description + * ------------------------ + * - Basic scenario to trigger capturehipMemcpy2DAsync_spt internal api for + * improved code coverage + * Test source + * ------------------------ + * - unit/memory/hipMemcpy2DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_spt_capturehipMemcpy2DAsync_spt", "", int, float, + double) { + TestType *A_h, *B_h, *A_d; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + int row, col; + row = GENERATE(3, 4, 100); + col = GENERATE(3, 4, 100); + hipStream_t stream; + size_t devPitch; + A_h = reinterpret_cast(malloc(sizeof(TestType) * row * col)); + B_h = reinterpret_cast(malloc(sizeof(TestType) * row * col)); + HIP_CHECK(hipStreamCreate(&stream)); + for (int i = 0; i < row; i++) { + for (int j = 0; j < col; j++) { + B_h[i * col + j] = i * col + j; + } + } + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(TestType) * col, row)); + HIP_CHECK(hipMemcpy2D(A_d, devPitch, B_h, sizeof(TestType) * col, sizeof(TestType) * col, row, + hipMemcpyHostToDevice)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpy2DAsync_spt(A_h, col * sizeof(TestType), A_d, devPitch, col * sizeof(TestType), + row, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (int i = 0; i < row; i++) { + for (int j = 0; j < col; j++) { + REQUIRE(A_h[i * col + j] == B_h[i * col + j]); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); + free(A_h); + free(B_h); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy2DFromArrayAsync_spt.cc b/catch/unit/memory/hipMemcpy2DFromArrayAsync_spt.cc new file mode 100644 index 000000000..601929ea6 --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DFromArrayAsync_spt.cc @@ -0,0 +1,266 @@ +/* +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 "array_memcpy_tests_common.hh" +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy2DFromArrayAsync_spt hipMemcpy2DFromArrayAsync_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, + * hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t + * height, hipMemcpyKind kind, hipStream_t stream __dparm(0));` - Copies data + * between host and device. + */ +/** + * Test Description + * ------------------------ + * - Test basic memcpy between host/device and 2D array with + * hipMemcpy2DFromArrayAsync_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DFromArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_spt_Positive_Default") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); + SECTION("Array to host") { + Memcpy2DHostFromAShell( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToHost, stream), + width, height, stream); + } + SECTION("Array to host with default kind") { + Memcpy2DHostFromAShell( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); + } +} +/** + * Test Description + * ------------------------ + * - Test synchronization behavior for hipMemcpy2DFromArrayAsync_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DFromArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_spt_Positive_Synchronization_Behavior") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Array to host") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + MemcpyAtoHPageableSyncBehavior( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, width * sizeof(int), _2, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost, nullptr), + width, height, false); + MemcpyAtoHPinnedSyncBehavior( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, width * sizeof(int), _2, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost, nullptr), + width, height, false); + } + SECTION("Array to device") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + MemcpyAtoDSyncBehavior(std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + width, height, false); + } +} +/** + * Test Description + * ------------------------ + * - Test that no data is copied when width/height is set to 0 + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DFromArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_spt_Positive_ZeroWidthHeight") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); + const auto width = 16; + const auto height = 16; + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToHost, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight(std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, + 0, 0, height, hipMemcpyDeviceToHost, stream), + width, height, stream); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArrayAsync_spt, _1, _2, _3, 0, 0, 0, height, + hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + } +} +/** + * Test Description + * ------------------------ + * - Test unsuccessful execution of hipMemcpy2DFromArrayAsync_spt api when + * parameters are invalid. Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DFromArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + const unsigned int flag = hipArrayDefault; + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + SECTION("Array to host") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(nullptr, 2 * width * sizeof(int), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(host_alloc.ptr(), 2 * width * sizeof(int), + nullptr, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidHandle); + } + SECTION("dpitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(host_alloc.ptr(), width * sizeof(int) - 10, + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 1, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 0, 1, width * sizeof(int), + height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt( + host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int) + 1, height, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(host_alloc.ptr(), 2 * width * sizeof(int), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height + 1, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt( + host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + } + SECTION("Array to device") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(nullptr, device_alloc.pitch(), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(device_alloc.ptr(), device_alloc.pitch(), + nullptr, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidHandle); + } + SECTION("dpitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(device_alloc.ptr(), width * sizeof(int) - 10, + array_alloc.ptr(), 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 1, 0, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 1, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt( + device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt(device_alloc.ptr(), device_alloc.pitch(), + array_alloc.ptr(), 0, 0, width * sizeof(int), + height + 1, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync_spt( + device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy2DFromArray_spt.cc b/catch/unit/memory/hipMemcpy2DFromArray_spt.cc new file mode 100644 index 000000000..763abffed --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DFromArray_spt.cc @@ -0,0 +1,249 @@ +/* +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 "array_memcpy_tests_common.hh" +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy2DFromArray_spt hipMemcpy2DFromArray_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy2DFromArray( void* dst, size_t dpitch, hipArray_const_t + * src, size_t wOffset, size_t hOffset, size_t width, size_t height, + * hipMemcpyKind kind)` - Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Test basic memcpy between 2D array and host/device with + * hipMemcpy2DFromArray_spt api Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DFromArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_spt_Positive_Default") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); + SECTION("Array to host") { + Memcpy2DHostFromAShell( + std::bind(hipMemcpy2DFromArray_spt, _1, _2, _3, 0, 0, width * sizeof(int), height, + hipMemcpyDeviceToHost), + width, height); + } + SECTION("Array to host with default kind") { + Memcpy2DHostFromAShell(std::bind(hipMemcpy2DFromArray_spt, _1, _2, _3, 0, 0, + width * sizeof(int), height, hipMemcpyDefault), + width, height); + } +} +/** + * Test Description + * ------------------------ + * - Test synchronization behavior for hipMemcpy2DFromArray_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DFromArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_spt_Positive_Synchronization_Behavior") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Array to host") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + MemcpyAtoHPageableSyncBehavior( + std::bind(hipMemcpy2DFromArray_spt, _1, width * sizeof(int), _2, 0, 0, width * sizeof(int), + height, hipMemcpyDeviceToHost), + width, height, true); + MemcpyAtoHPinnedSyncBehavior(std::bind(hipMemcpy2DFromArray_spt, _1, width * sizeof(int), _2, 0, + 0, width * sizeof(int), height, hipMemcpyDeviceToHost), + width, height, true); + } +} +/** + * Test Description + * ------------------------ + * - Test that no data is copied when width/height is set to 0 + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_spt_Positive_ZeroWidthHeight") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 16; + const auto height = 16; + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray_spt, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToHost), + width, height); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray_spt, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToHost), + width, height); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray_spt, _1, _2, _3, 0, 0, width * sizeof(int), 0, + hipMemcpyDeviceToDevice), + width, height); + } + SECTION("Width is 0") { + Memcpy2DFromArrayZeroWidthHeight( + std::bind(hipMemcpy2DFromArray_spt, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToDevice), + width, height); + } + } +} +/** + * Test Description + * ------------------------ + * - Test unsuccessful execution of hipMemcpy2DToArray_spt api when parameters + * are invalid. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DFromArray_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + const unsigned int flag = hipArrayDefault; + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + SECTION("Array to host") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(nullptr, 2 * width * sizeof(int), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), 2 * width * sizeof(int), nullptr, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidHandle); + } + SECTION("dpitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), width * sizeof(int) - 10, array_alloc.ptr(), 0, + 0, width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 1, + 0, width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, + 1, width * sizeof(int), height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, + 0, width * sizeof(int) + 1, height, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, + 0, width * sizeof(int), height + 1, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(host_alloc.ptr(), 2 * width * sizeof(int), array_alloc.ptr(), 0, + 0, width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + } + SECTION("Array to device") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(nullptr, device_alloc.pitch(), array_alloc.ptr(), 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), device_alloc.pitch(), nullptr, 0, 0, + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidHandle); + } + SECTION("dpitch < width") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), width * sizeof(int) - 10, array_alloc.ptr(), + 0, 0, width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 1, + 0, width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, + 1, width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, + 0, width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, + 0, width * sizeof(int), height + 1, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DFromArray_spt(device_alloc.ptr(), device_alloc.pitch(), array_alloc.ptr(), 0, + 0, width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy2DToArrayAsync_spt.cc b/catch/unit/memory/hipMemcpy2DToArrayAsync_spt.cc new file mode 100644 index 000000000..0c1733dff --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DToArrayAsync_spt.cc @@ -0,0 +1,261 @@ +/* +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 "array_memcpy_tests_common.hh" +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy2DToArrayAsync_spt hipMemcpy2DToArrayAsync_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy2DToArrayAsync_spt(hipArray_t dst, size_t wOffset, + size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, + hipMemcpyKind kind, hipStream_t stream __dparm(0))` - + * Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Test basic async memcpy between host/device and 2D array with + * hipMemcpy2DToArrayAsync_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_spt_Positive_Default") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); + SECTION("Host to Array") { + Memcpy2DHosttoAShell( + std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyHostToDevice, stream), + width, height, stream); + } + SECTION("Host to Array with default kind") { + Memcpy2DHosttoAShell( + std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, width * sizeof(int), height, + hipMemcpyDefault, stream), + width, height, stream); + } +} +/** + * Test Description + * ------------------------ + * - Test synchronization behavior for hipMemcpy2DToArrayAsync_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_spt_Positive_Synchronization_Behavior") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Host to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + MemcpyHtoASyncBehavior(std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice, nullptr), + width, height, false); + } + SECTION("Device to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + MemcpyDtoASyncBehavior(std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, + width * sizeof(int), height, hipMemcpyDeviceToDevice, nullptr), + width, height, false); + } +} +/** + * Test Description + * ------------------------ + * - Test that no data is copied when width/height is set to 0 + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_spt_Positive_ZeroWidthHeight") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 16; + const auto height = 16; + const auto stream_type = GENERATE(Streams::nullstream, Streams::perThread, Streams::created); + const StreamGuard stream_guard(stream_type); + const hipStream_t stream = stream_guard.stream(); + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyHostToDevice, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight(std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, + 0, height, hipMemcpyHostToDevice, stream), + width, height, stream); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight(std::bind(hipMemcpy2DToArrayAsync_spt, _1, 0, 0, _2, _3, + 0, height, hipMemcpyDeviceToDevice, stream), + width, height, stream); + } + } +} +/** + * Test Description + * ------------------------ + * - Test unsuccessful execution of hipMemcpy2DToArrayAsync_spt api when + * parameters are invalid. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArrayAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArrayAsync_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + const unsigned int flag = hipArrayDefault; + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + SECTION("Host to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync_spt(nullptr, 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, nullptr, 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + width * sizeof(int) - 10, width * sizeof(int), + height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 1, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), + height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 1, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), + height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int) + 1, + height, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), + height + 1, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + 2 * width * sizeof(int), width * sizeof(int), + height, static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + } + SECTION("Device to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(nullptr, 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, nullptr, + device_alloc.pitch(), width * sizeof(int), height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + width * sizeof(int) - 10, width * sizeof(int), + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 1, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 1, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height, + hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int) + 1, + height, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), + height + 1, hipMemcpyDeviceToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + device_alloc.pitch(), width * sizeof(int), height, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy2DToArray_spt.cc b/catch/unit/memory/hipMemcpy2DToArray_spt.cc new file mode 100644 index 000000000..a8a9cab2b --- /dev/null +++ b/catch/unit/memory/hipMemcpy2DToArray_spt.cc @@ -0,0 +1,244 @@ +/* +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 "array_memcpy_tests_common.hh" +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy2DToArray_spt hipMemcpy2DToArray_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy2DToArray_spt(hipArray_t dst, size_t wOffset, size_t + hOffset, const void* src, size_t spitch, size_t width, size_t height, + hipMemcpyKind kind))` - + * Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Test basic memcpy between host/device and 2D array with + * hipMemcpy2DToArray_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArray_spt_Positive_Default") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(1, 16, 32, 48); + SECTION("Host to Array") { + Memcpy2DHosttoAShell(std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, _3, + width * sizeof(int), height, hipMemcpyHostToDevice), + width, height); + } + SECTION("Host to Array with default kind") { + Memcpy2DHosttoAShell(std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, _3, + width * sizeof(int), height, hipMemcpyDefault), + width, height); + } +} +/** + * Test Description + * ------------------------ + * - Test synchronization behavior for hipMemcpy2DToArray_spt api + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArray_spt_Positive_Synchronization_Behavior") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Host to Array") { + const auto width = GENERATE(16, 32, 48); + const auto height = GENERATE(16, 32, 48); + MemcpyHtoASyncBehavior(std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + width, height, true); + } +} +/** + * Test Description + * ------------------------ + * - Test that no data is copied when width/height is set to 0 + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArray_spt_Positive_ZeroWidthHeight") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 16; + const auto height = 16; + SECTION("Array to host") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyHostToDevice), + width, height); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, _3, 0, height, hipMemcpyHostToDevice), + width, height); + } + } + SECTION("Array to device") { + SECTION("Height is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, _3, width * sizeof(int), 0, + hipMemcpyDeviceToDevice), + width, height); + } + SECTION("Width is 0") { + Memcpy2DToArrayZeroWidthHeight( + std::bind(hipMemcpy2DToArray_spt, _1, 0, 0, _2, _3, 0, height, hipMemcpyDeviceToDevice), + width, height); + } + } +} +/** + * Test Description + * ------------------------ + * - Test unsuccessful execution of hipMemcpy2DToArray_spt api when parameters + * are invalid. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2DToArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2DToArray_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + using namespace std::placeholders; + const auto width = 32; + const auto height = 32; + const auto allocation_size = 2 * width * height * sizeof(int); + const unsigned int flag = hipArrayDefault; + ArrayAllocGuard array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard2D device_alloc(width, height); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + SECTION("Host to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(nullptr, 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, nullptr, 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), + width * sizeof(int) - 10, width * sizeof(int), height, + hipMemcpyHostToDevice), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 1, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 1, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int) + 1, height, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height + 1, hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, host_alloc.ptr(), 2 * width * sizeof(int), + width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + } + SECTION("Device to Array") { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(nullptr, 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidHandle); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, nullptr, device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), + width * sizeof(int) - 10, width * sizeof(int), height, + hipMemcpyDeviceToDevice), + hipErrorInvalidPitchValue); + } + SECTION("Offset + width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 1, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 1, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Width/height overflows") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int) + 1, height, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height + 1, hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + SECTION("Memcpy kind is invalid") { + HIP_CHECK_ERROR( + hipMemcpy2DToArray_spt(array_alloc.ptr(), 0, 0, device_alloc.ptr(), device_alloc.pitch(), + width * sizeof(int), height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + } +} +/** + * End doxygen group MemoryTest. + * @} diff --git a/catch/unit/memory/hipMemcpy2D_spt.cc b/catch/unit/memory/hipMemcpy2D_spt.cc new file mode 100644 index 000000000..d46954e9a --- /dev/null +++ b/catch/unit/memory/hipMemcpy2D_spt.cc @@ -0,0 +1,152 @@ +/* +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 "memcpy2d_tests_common.hh" +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy2D_spt hipMemcpy2D_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy2D_spt(void* dst, size_t dpitch, const void* src, + size_t spitch, size_t width, + size_t height, hipMemcpyKind kind)` - + * Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Test to verify the Synchronization_Behavior of hipMemcpy2D_spt + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2D_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2D_spt_Positive_Synchronization_Behavior") { + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Host to Device") { Memcpy2DHtoDSyncBehavior(hipMemcpy2D_spt, true); } + SECTION("Device to Host") { + Memcpy2DDtoHPageableSyncBehavior(hipMemcpy2D_spt, true); + Memcpy2DDtoHPinnedSyncBehavior(hipMemcpy2D_spt, true); + } + SECTION("Device to Device") { Memcpy2DDtoDSyncBehavior(hipMemcpy2D_spt, true); } + SECTION("Host to Host") { Memcpy2DHtoHSyncBehavior(hipMemcpy2D_spt, true); } +} +/** + * Test Description + * ------------------------ + * - Test to verify the hipMemcpy2D_spt with positive parameters. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2D_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2D_spt_Positive_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr bool async = false; + Memcpy2DZeroWidthHeight(hipMemcpy2D_spt); +} +/** + * Test Description + * ------------------------ + * - Test to verify the hipMemcpy2D_spt with Negative parameters. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy2D_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy2D_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr size_t cols = 128; + constexpr size_t rows = 128; + constexpr auto NegativeTests = [](void* dst, size_t dpitch, const void* src, size_t spitch, + size_t width, size_t height, hipMemcpyKind kind) { + SECTION("dst == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2D_spt(nullptr, dpitch, src, spitch, width, height, kind), + hipErrorInvalidValue); + } + SECTION("src == nullptr") { + HIP_CHECK_ERROR(hipMemcpy2D_spt(dst, dpitch, nullptr, spitch, width, height, kind), + hipErrorInvalidValue); + } + SECTION("dpitch < width") { + HIP_CHECK_ERROR(hipMemcpy2D_spt(dst, width - 1, src, spitch, width, height, kind), + hipErrorInvalidPitchValue); + } + SECTION("spitch < width") { + HIP_CHECK_ERROR(hipMemcpy2D_spt(dst, dpitch, src, width - 1, width, height, kind), + hipErrorInvalidPitchValue); + } + SECTION("dpitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + HIP_CHECK_ERROR( + hipMemcpy2D_spt(dst, static_cast(attr) + 1, src, spitch, width, height, kind), + hipErrorInvalidValue); + } + SECTION("spitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + HIP_CHECK_ERROR( + hipMemcpy2D_spt(dst, dpitch, src, static_cast(attr) + 1, width, height, kind), + hipErrorInvalidValue); + } + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR( + hipMemcpy2D_spt(dst, dpitch, src, spitch, width, height, static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + }; + SECTION("Host to Device") { + LinearAllocGuard2D device_alloc(cols, rows); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * rows); + NegativeTests(device_alloc.ptr(), device_alloc.pitch(), host_alloc.ptr(), device_alloc.pitch(), + device_alloc.width(), device_alloc.height(), hipMemcpyHostToDevice); + } + SECTION("Device to Host") { + LinearAllocGuard2D device_alloc(cols, rows); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * rows); + NegativeTests(host_alloc.ptr(), device_alloc.pitch(), device_alloc.ptr(), device_alloc.pitch(), + device_alloc.width(), device_alloc.height(), hipMemcpyDeviceToHost); + } + SECTION("Host to Host") { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, cols * rows * sizeof(int)); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, cols * rows * sizeof(int)); + NegativeTests(dst_alloc.ptr(), cols * sizeof(int), src_alloc.ptr(), cols * sizeof(int), + cols * sizeof(int), rows, hipMemcpyHostToHost); + } + SECTION("Device to Device") { + LinearAllocGuard2D src_alloc(cols, rows); + LinearAllocGuard2D dst_alloc(cols, rows); + NegativeTests(dst_alloc.ptr(), dst_alloc.pitch(), src_alloc.ptr(), src_alloc.pitch(), + dst_alloc.width(), dst_alloc.height(), hipMemcpyDeviceToDevice); + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy3DAsync_spt.cc b/catch/unit/memory/hipMemcpy3DAsync_spt.cc new file mode 100644 index 000000000..974ce1328 --- /dev/null +++ b/catch/unit/memory/hipMemcpy3DAsync_spt.cc @@ -0,0 +1,225 @@ +/* +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 "memcpy3d_spt_tests_common.hh" +#include +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy3DAsync_spt hipMemcpy3DAsync_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy3DAsync_spt(const struct hipMemcpy3DParms* p, + * hipStream_t stream __dparm(0))` - Copies data between host and device + * asynchronously. + */ +/** + * Test Description + * ------------------------ + * - Basic test to verify the negative parameters of the hipMemcpy3DAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy3DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy3DAsync_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr bool async = true; + constexpr hipExtent extent{128 * sizeof(int), 128, 8}; + constexpr auto NegativeTests = [](hipPitchedPtr dst_ptr, hipPos dst_pos, hipPitchedPtr src_ptr, + hipPos src_pos, hipExtent extent, hipMemcpyKind kind) { + SECTION("dst_ptr.ptr == nullptr") { + hipPitchedPtr invalid_ptr = dst_ptr; + invalid_ptr.ptr = nullptr; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_ptr.ptr == nullptr") { + hipPitchedPtr invalid_ptr = src_ptr; + invalid_ptr.ptr = nullptr; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("dst_ptr.pitch < width") { + hipPitchedPtr invalid_ptr = dst_ptr; + invalid_ptr.pitch = extent.width - 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidPitchValue); + } + SECTION("src_ptr.pitch < width") { + hipPitchedPtr invalid_ptr = src_ptr; + invalid_ptr.pitch = extent.width - 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind), + hipErrorInvalidPitchValue); + } + SECTION("dst_ptr.pitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + hipPitchedPtr invalid_ptr = dst_ptr; + invalid_ptr.pitch = attr; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_ptr.pitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + hipPitchedPtr invalid_ptr = src_ptr; + invalid_ptr.pitch = attr; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("extent.width + dst_pos.x > dst_ptr.pitch") { + hipPos invalid_pos = dst_pos; + invalid_pos.x = dst_ptr.pitch - extent.width + 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("extent.width + src_pos.x > src_ptr.pitch") { + hipPos invalid_pos = src_pos; + invalid_pos.x = src_ptr.pitch - extent.width + 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("dst_pos.y out of bounds") { + hipPos invalid_pos = dst_pos; + invalid_pos.y = 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_pos.y out of bounds") { + hipPos invalid_pos = src_pos; + invalid_pos.y = 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("dst_pos.z out of bounds") { + hipPos invalid_pos = dst_pos; + invalid_pos.z = 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_pos.z out of bounds") { + hipPos invalid_pos = src_pos; + invalid_pos.z = 1; + HIP_CHECK_ERROR( + Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, src_pos, extent, + static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + StreamGuard stream_guard(Streams::created); + HIP_CHECK(hipStreamDestroy(stream_guard.stream())); + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind, + stream_guard.stream()), + hipErrorContextIsDestroyed); + } + }; + SECTION("Host to Device") { + LinearAllocGuard3D device_alloc(extent); + LinearAllocGuard host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.pitch() * device_alloc.height() * device_alloc.depth()); + NegativeTests(device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), + make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(), + device_alloc.height()), + make_hipPos(0, 0, 0), extent, hipMemcpyHostToDevice); + } + SECTION("Device to Host") { + LinearAllocGuard3D device_alloc(extent); + LinearAllocGuard host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.pitch() * device_alloc.height() * device_alloc.depth()); + NegativeTests(make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(), + device_alloc.height()), + make_hipPos(0, 0, 0), device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), extent, + hipMemcpyDeviceToHost); + } + SECTION("Host to Host") { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, + extent.width * extent.height * extent.depth); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, + extent.width * extent.height * extent.depth); + NegativeTests(make_hipPitchedPtr(dst_alloc.ptr(), extent.width, extent.width, extent.height), + make_hipPos(0, 0, 0), + make_hipPitchedPtr(src_alloc.ptr(), extent.width, extent.width, extent.height), + make_hipPos(0, 0, 0), extent, hipMemcpyHostToHost); + } + SECTION("Device to Device") { + LinearAllocGuard3D src_alloc(extent); + LinearAllocGuard3D dst_alloc(extent); + NegativeTests(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), extent, hipMemcpyDeviceToDevice); + } +} +/** + * Test Description + * ------------------------ + * - Basic test to verify the positive parameters of the hipMemcpy3DAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy3DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy3DAsync_spt_Positive_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr bool async = true; + Memcpy3DZeroWidthHeightDepth_spt(Memcpy3DWrapper_spt); +} +/** + * Test Description + * ------------------------ + * - Basic test to verify the data copy between Array with hipMemcpy3DAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy3DAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy3DAsync_spt_Positive_Array") { + CHECK_IMAGE_SUPPORT + constexpr bool async = true; + SECTION("Array from/to Host") { Memcpy3DArrayHostShell(Memcpy3DWrapper_spt); } + SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell(Memcpy3DWrapper_spt); } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy3D_spt.cc b/catch/unit/memory/hipMemcpy3D_spt.cc new file mode 100644 index 000000000..0b56a572c --- /dev/null +++ b/catch/unit/memory/hipMemcpy3D_spt.cc @@ -0,0 +1,204 @@ +/* +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 "memcpy3d_spt_tests_common.hh" +#include +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpy3D_spt hipMemcpy3D_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy3D_spt(const struct hipMemcpy3DParms* p)` - + * Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Basic test to verify the negative parameters of the hipMemcpy3D_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy3D_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy3D_spt_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr hipExtent extent{128 * sizeof(int), 128, 8}; + constexpr auto NegativeTests = [](hipPitchedPtr dst_ptr, hipPos dst_pos, hipPitchedPtr src_ptr, + hipPos src_pos, hipExtent extent, hipMemcpyKind kind) { + SECTION("dst_ptr.ptr == nullptr") { + hipPitchedPtr invalid_ptr = dst_ptr; + invalid_ptr.ptr = nullptr; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_ptr.ptr == nullptr") { + hipPitchedPtr invalid_ptr = src_ptr; + invalid_ptr.ptr = nullptr; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("dst_ptr.pitch < width") { + hipPitchedPtr invalid_ptr = dst_ptr; + invalid_ptr.pitch = extent.width - 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidPitchValue); + } + SECTION("src_ptr.pitch < width") { + hipPitchedPtr invalid_ptr = src_ptr; + invalid_ptr.pitch = extent.width - 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind), + hipErrorInvalidPitchValue); + } + SECTION("dst_ptr.pitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + hipPitchedPtr invalid_ptr = dst_ptr; + invalid_ptr.pitch = attr; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_ptr.pitch > max pitch") { + int attr = 0; + HIP_CHECK(hipDeviceGetAttribute(&attr, hipDeviceAttributeMaxPitch, 0)); + hipPitchedPtr invalid_ptr = src_ptr; + invalid_ptr.pitch = attr; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("extent.width + dst_pos.x > dst_ptr.pitch") { + hipPos invalid_pos = dst_pos; + invalid_pos.x = dst_ptr.pitch - extent.width + 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("extent.width + src_pos.x > src_ptr.pitch") { + hipPos invalid_pos = src_pos; + invalid_pos.x = src_ptr.pitch - extent.width + 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("dst_pos.y out of bounds") { + hipPos invalid_pos = dst_pos; + invalid_pos.y = 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_pos.y out of bounds") { + hipPos invalid_pos = src_pos; + invalid_pos.y = 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("dst_pos.z out of bounds") { + hipPos invalid_pos = dst_pos; + invalid_pos.z = 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, invalid_pos, src_ptr, src_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("src_pos.z out of bounds") { + hipPos invalid_pos = src_pos; + invalid_pos.z = 1; + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, invalid_pos, extent, kind), + hipErrorInvalidValue); + } + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(Memcpy3DWrapper_spt(dst_ptr, dst_pos, src_ptr, src_pos, extent, + static_cast(-1)), + hipErrorInvalidMemcpyDirection); + } + }; + SECTION("Host to Device") { + LinearAllocGuard3D device_alloc(extent); + LinearAllocGuard host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.pitch() * device_alloc.height() * device_alloc.depth()); + NegativeTests(device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), + make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(), + device_alloc.height()), + make_hipPos(0, 0, 0), extent, hipMemcpyHostToDevice); + } + SECTION("Device to Host") { + LinearAllocGuard3D device_alloc(extent); + LinearAllocGuard host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.pitch() * device_alloc.height() * device_alloc.depth()); + NegativeTests(make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(), device_alloc.width(), + device_alloc.height()), + make_hipPos(0, 0, 0), device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), extent, + hipMemcpyDeviceToHost); + } + SECTION("Host to Host") { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, + extent.width * extent.height * extent.depth); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, + extent.width * extent.height * extent.depth); + NegativeTests(make_hipPitchedPtr(dst_alloc.ptr(), extent.width, extent.width, extent.height), + make_hipPos(0, 0, 0), + make_hipPitchedPtr(src_alloc.ptr(), extent.width, extent.width, extent.height), + make_hipPos(0, 0, 0), extent, hipMemcpyHostToHost); + } + SECTION("Device to Device") { + LinearAllocGuard3D src_alloc(extent); + LinearAllocGuard3D dst_alloc(extent); + NegativeTests(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), extent, hipMemcpyDeviceToDevice); + } +} +/** + * Test Description + * ------------------------ + * - Basic test to verify the positive parameters of the hipMemcpy3D_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy3D_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy3D_spt_Positive_Parameters") { + CHECK_IMAGE_SUPPORT + constexpr bool async = false; + Memcpy3DZeroWidthHeightDepth_spt(Memcpy3DWrapper_spt); +} +/** + * Test Description + * ------------------------ + * - Basic test to verify the data copy between Array with hipMemcpy3D_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpy3D_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy3D_spt_Positive_Array") { + CHECK_IMAGE_SUPPORT + constexpr bool async = false; + SECTION("Array from/to Host") { Memcpy3DArrayHostShell(Memcpy3DWrapper_spt); } + SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell(Memcpy3DWrapper_spt); } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpyAsync_spt.cc b/catch/unit/memory/hipMemcpyAsync_spt.cc new file mode 100644 index 000000000..2540ff801 --- /dev/null +++ b/catch/unit/memory/hipMemcpyAsync_spt.cc @@ -0,0 +1,163 @@ +/* +Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include +#include +#include +#include +/** + * @addtogroup hipMemcpyAsync_spt hipMemcpyAsync_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpyAsync_spt(void* dst, const void* src, size_t sizeBytes, + hipMemcpyKind kind, hipStream_t stream __dparm(0))` - + * Copy data from src to dst asynchronously. + */ +/** + * Test Description + * ------------------------ + * - Basic test to verify the Synchronization Behavior of hipMemcpyAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyAsync_spt_Positive_Synchronization_Behavior") { + using namespace std::placeholders; + HIP_CHECK(hipDeviceSynchronize()); + SECTION("Host pinned memory to device memory") { + MemcpyHPinnedtoDSyncBehavior( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyHostToDevice, nullptr), false); + } + SECTION("Device memory to pageable host memory") { + MemcpyDtoHPageableSyncBehavior( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyDeviceToHost, nullptr), true); + } + SECTION("Device memory to pinned host memory") { + MemcpyDtoHPinnedSyncBehavior( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyDeviceToHost, nullptr), false); + } + SECTION("Device memory to device memory") { + MemcpyDtoDSyncBehavior( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyDeviceToDevice, nullptr), false); + } + SECTION("Device memory to device Memory No CU") { + MemcpyDtoDSyncBehavior( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyDeviceToDeviceNoCU, nullptr), false); + } + SECTION("Host memory to host memory") { + MemcpyHtoHSyncBehavior(std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyHostToHost, nullptr), + true); + MemcpyDtoHPinnedSyncBehavior( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyHostToHost, nullptr), true); + } +} +/** + * Test Description + * ------------------------ + * - Basic test to verify negative test cases of hipMemcpyAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyAsync_spt_Negative_Parameters") { + using namespace std::placeholders; + constexpr auto InvalidStream = [] { + StreamGuard sg(Streams::created); + return sg.stream(); + }; + SECTION("Host to device") { + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); + MemcpyCommonNegativeTests( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyHostToDevice, nullptr), + device_alloc.ptr(), host_alloc.ptr(), kPageSize); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(device_alloc.ptr(), host_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(device_alloc.ptr(), host_alloc.ptr(), kPageSize, + hipMemcpyHostToDevice, InvalidStream()), + hipErrorContextIsDestroyed); + } + } + SECTION("Device to host") { + LinearAllocGuard device_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, kPageSize); + MemcpyCommonNegativeTests( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyDeviceToHost, nullptr), host_alloc.ptr(), + device_alloc.ptr(), kPageSize); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(host_alloc.ptr(), device_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(host_alloc.ptr(), device_alloc.ptr(), kPageSize, + hipMemcpyDeviceToHost, InvalidStream()), + hipErrorContextIsDestroyed); + } + } + SECTION("Host to host") { + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, kPageSize); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, kPageSize); + MemcpyCommonNegativeTests( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyHostToHost, nullptr), dst_alloc.ptr(), + src_alloc.ptr(), kPageSize); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, + hipMemcpyHostToHost, InvalidStream()), + hipErrorContextIsDestroyed); + } + } + SECTION("Device to device") { + LinearAllocGuard src_alloc(LinearAllocs::hipMalloc, kPageSize); + LinearAllocGuard dst_alloc(LinearAllocs::hipMalloc, kPageSize); + MemcpyCommonNegativeTests( + std::bind(hipMemcpyAsync_spt, _1, _2, _3, hipMemcpyDeviceToDevice, nullptr), + dst_alloc.ptr(), src_alloc.ptr(), kPageSize); + SECTION("Invalid MemcpyKind") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(src_alloc.ptr(), dst_alloc.ptr(), kPageSize, + static_cast(-1), nullptr), + hipErrorInvalidMemcpyDirection); + } + SECTION("Invalid stream") { + HIP_CHECK_ERROR(hipMemcpyAsync_spt(dst_alloc.ptr(), src_alloc.ptr(), kPageSize, + hipMemcpyDeviceToDevice, InvalidStream()), + hipErrorContextIsDestroyed); + } + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpyFromArray_spt.cc b/catch/unit/memory/hipMemcpyFromArray_spt.cc new file mode 100644 index 000000000..4ef0347ef --- /dev/null +++ b/catch/unit/memory/hipMemcpyFromArray_spt.cc @@ -0,0 +1,117 @@ +/* +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 hipMemcpyFromArray_spt hipMemcpyFromArray_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t srcArray, + size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ +/** + * Test Description + * ------------------------ + * - Basic test to verify the functionality of hipMemcpyFromArray_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyFromArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyFromArray_spt_Basic_Postive") { + size_t width = 64; + size_t height = 1; + const int N = width * height; + int value = 10; + int* hostMem = reinterpret_cast(malloc(N * sizeof(int))); + REQUIRE(hostMem != nullptr); + for (int i = 0; i < N; i++) { + hostMem[i] = value; + } + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpyToArray(array, 0, 0, hostMem, N * sizeof(int), hipMemcpyHostToDevice)); + int* hostMemory = reinterpret_cast(malloc(N * sizeof(int))); + REQUIRE(hostMemory != nullptr); + HIP_CHECK(hipMemcpyFromArray_spt(hostMemory, array, 0, 0, N * sizeof(int), hipMemcpyDefault)); + for (int i = 0; i < N; i++) { + if (hostMemory[i] != value) { + REQUIRE(false); + } + } + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); +} +/** + * Test Description + * ------------------------ + * - Negative tests of hipMemcpyFromArray_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyFromArray_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyFromArray_spt_NegativeTests") { + size_t width = 64; + size_t height = 1; + const int N = width * height; + int value = 10; + int* hostMem = reinterpret_cast(malloc(N * sizeof(int))); + REQUIRE(hostMem != nullptr); + for (int i = 0; i < N; i++) { + hostMem[i] = value; + } + hipArray_t array = nullptr; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + unsigned int flags = hipArrayDefault; + HIP_CHECK(hipMallocArray(&array, &desc, width, height, flags)); + REQUIRE(array != nullptr); + HIP_CHECK(hipMemcpyToArray(array, 0, 0, hostMem, N * sizeof(int), hipMemcpyHostToDevice)); + int* hostMemory = reinterpret_cast(malloc(N * sizeof(int))); + REQUIRE(hostMemory != nullptr); + SECTION("Destination Array as nullptr") { + HIP_CHECK_ERROR(hipMemcpyFromArray_spt(nullptr, array, 0, 0, N * sizeof(int), hipMemcpyDefault), + hipErrorInvalidValue); + } + SECTION("Source Array as nullptr") { + HIP_CHECK_ERROR( + hipMemcpyFromArray_spt(hostMemory, nullptr, 0, 0, N * sizeof(int), hipMemcpyDefault), + hipErrorInvalidValue); + } + SECTION("Invalid Size") { + HIP_CHECK_ERROR(hipMemcpyFromArray_spt(hostMemory, array, 0, 0, -3, hipMemcpyDefault), + hipErrorInvalidValue); + } + free(hostMem); + free(hostMemory); + HIP_CHECK(hipFreeArray(array)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpyFromSymbolAsync_spt.cc b/catch/unit/memory/hipMemcpyFromSymbolAsync_spt.cc new file mode 100644 index 000000000..53df5bed0 --- /dev/null +++ b/catch/unit/memory/hipMemcpyFromSymbolAsync_spt.cc @@ -0,0 +1,158 @@ +/* +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 +__device__ int devSymbol[10]; +__constant__ int constSymbol[10]; +TEST_CASE("Unit_hipMemcpyFromSymbolAsync_spt_Negative") { + SECTION("Invalid Src Ptr") { + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync_spt(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid Dst Ptr") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync_spt(&result, nullptr, sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidSymbol); + } + SECTION("Invalid Size") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, + 0, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid Offset") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid Direction") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidMemcpyDirection); + } +} +TEST_CASE("Unit_hipMemcpyFromSymbolAsync_spt_PositiveTests") { + enum StreamTestType { NullStream = 0, StreamPerThread, CreatedStream, NoStream }; + /* Test type NoStream - Use Sync variants, else use async variants */ + auto streamType = GENERATE(StreamTestType::NoStream, StreamTestType::NullStream, + StreamTestType::StreamPerThread, StreamTestType::CreatedStream); + hipStream_t stream{nullptr}; + if (streamType == StreamTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + INFO("Stream :: " << streamType); + SECTION("Singular Value") { + int set{42}; + int result{0}; + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &set, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(result == set); + } + SECTION("Array Values") { + constexpr size_t size{10}; + int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2}; + int result[size] = {0}; + + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, sizeof(int) * size, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } + SECTION("Offset'ed Values") { + constexpr size_t size{10}; + constexpr size_t offset = 5 * sizeof(int); + int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2}; + int result[size] = {0}; + + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, offset, 0, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set + 5, offset, offset, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync_spt(result, HIP_SYMBOL(devSymbol), offset, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync_spt(result + 5, HIP_SYMBOL(devSymbol), offset, offset, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } +} +/** + * Test Description + * ------------------------ + * - Basic functional testcase to trigger capturehipMemcpyToSymbolAsync + * and capturehipMemcpyFromSymbolAsync internal apis to improve + * code coverage. + * Test source + * ------------------------ + * - unit/memory/hipMemcpyFromSymbol.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMemcpyFromSymbolAsync_spt_capturehipMemcpyToFromSymbolAsync") { + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + int A_h = 0, B_h = 42; + // Start Capturing + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + SECTION("__constant__ symbol") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(constSymbol), &B_h, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync_spt(&A_h, HIP_SYMBOL(constSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + } + SECTION("__device__ symbol") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &B_h, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync_spt(&A_h, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + } + // End Capture + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + // Create and Launch Executable Graphs + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(A_h == B_h); + HIP_CHECK(hipGraphExecDestroy(graphExec)) + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpyFromSymbol_spt.cc b/catch/unit/memory/hipMemcpyFromSymbol_spt.cc new file mode 100644 index 000000000..c9b335ac5 --- /dev/null +++ b/catch/unit/memory/hipMemcpyFromSymbol_spt.cc @@ -0,0 +1,122 @@ +/*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 hipMemcpyFromSymbol_spt hipMemcpyFromSymbol_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpyFromSymbol_spt(void* dst, const void* symbol, + size_t sizeBytes, size_t offset __dparm(0), + hipMemcpyKind kind + __dparm(hipMemcpyDeviceToHost))` - + * Copies data from the given symbol on the device. + */ +__device__ int devSymbol[10]; +__constant__ int constSymbol[10]; +/** + * Test Description + * ------------------------ + * - Basic test to check the negative cases of hipMemcpyFromSymbol_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyFromSymbol_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyFromSymbol_spt_Negative") { + SECTION("Invalid Src Ptr") { + HIP_CHECK_ERROR(hipMemcpyFromSymbol_spt(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Invalid Dst Ptr") { + int result{0}; + HIP_CHECK_ERROR( + hipMemcpyFromSymbol_spt(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidSymbol); + } + SECTION("Invalid Size") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Invalid Offset") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + } + SECTION("Invalid Direction") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyHostToDevice), + hipErrorInvalidMemcpyDirection); + } +} +/** + * Test Description + * ------------------------ + * - Test Verifies hipMemcpyFromSymbol_spt for simple use case + * For single value From Symbol + * For Array Values From Symbol + * For Array Values with offset From Symbol + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyFromSymbol_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyFromSymbol_spt_Sync") { + SECTION("Singular Value") { + int set{42}; + int result{0}; + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &set, sizeof(int))); + HIP_CHECK(hipMemcpyFromSymbol_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int))); + REQUIRE(result == set); + } + SECTION("Array Values") { + constexpr size_t size{10}; + int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2}; + int result[size] = {0}; + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, sizeof(int) * size)); + HIP_CHECK(hipMemcpyFromSymbol_spt(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } + SECTION("Offset'ed Values") { + constexpr size_t size{10}; + constexpr size_t offset = 5 * sizeof(int); + int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2}; + int result[size] = {0}; + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, offset)); + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set + 5, offset, offset)); + HIP_CHECK(hipMemcpyFromSymbol_spt(result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpyToSymbolAsync_spt.cc b/catch/unit/memory/hipMemcpyToSymbolAsync_spt.cc new file mode 100644 index 000000000..496ea84bc --- /dev/null +++ b/catch/unit/memory/hipMemcpyToSymbolAsync_spt.cc @@ -0,0 +1,185 @@ +/* +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 +__device__ int devSymbol[10]; +__constant__ int constSymbol[10]; +/** + * @addtogroup hipMemcpyToSymbolAsync_spt hipMemcpyToSymbolAsync_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpyToSymbolAsync_spt(const void* symbol, const void* src, + size_t sizeBytes, size_t offset, + hipMemcpyKind kind, hipStream_t stream + __dparm(0))` - + * Copies data to the given symbol on the device asynchronously. + */ +/** + * Test Description + * ------------------------ + * - Basic test to check the negative cases of hipMemcpyToSymbolAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyToSymbolAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyToSymbolAsync_spt_Negative") { + SECTION("Invalid Src Ptr") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync_spt(nullptr, &result, sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidSymbol); + } + SECTION("Invalid Dst Ptr") { + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid Size") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid Offset") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + } + SECTION("Invalid Direction") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidMemcpyDirection); + } +} +/** + * Test Description + * ------------------------ + * - Basic test to check the positive cases of hipMemcpyToSymbolAsync_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyToSymbolAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyToSymbolAsync_spt_PositiveTest") { + enum StreamTestType { NullStream = 0, StreamPerThread, CreatedStream, NoStream }; + auto streamType = GENERATE(StreamTestType::NoStream, StreamTestType::NullStream, + StreamTestType::StreamPerThread, StreamTestType::CreatedStream); + hipStream_t stream{nullptr}; + if (streamType == StreamTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + INFO("Stream :: " << streamType); + SECTION("Singular Value") { + int set{42}; + int result{0}; + HIP_CHECK(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), &set, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(result == set); + } + SECTION("Array Values") { + constexpr size_t size{10}; + int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2}; + int result[size] = {0}; + HIP_CHECK(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), set, sizeof(int) * size, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } + SECTION("Offset'ed Values") { + constexpr size_t size{10}; + constexpr size_t offset = 5 * sizeof(int); + int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2}; + int result[size] = {0}; + HIP_CHECK(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), set, offset, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), set + 5, offset, offset, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(result, HIP_SYMBOL(devSymbol), offset, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(result + 5, HIP_SYMBOL(devSymbol), offset, offset, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } +} +/** + * Test Description + * ------------------------ + * - Basic functional testcase to trigger capturehipMemcpyToSymbolAsync + * and capturehipMemcpyToSymbolAsync internal apis to improve + * code coverage. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyToSymbolAsync_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyToSymbolAsync_spt_capturehipMemcpyToSymbolAsync_spt") { + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + int A_h = 0, B_h = 42; + // Start Capturing + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + SECTION("__constant__ symbol") { + HIP_CHECK(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(constSymbol), &B_h, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(&A_h, HIP_SYMBOL(constSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + } + SECTION("__device__ symbol") { + HIP_CHECK(hipMemcpyToSymbolAsync_spt(HIP_SYMBOL(devSymbol), &B_h, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(&A_h, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + } + // End Capture + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + // Create and Launch Executable Graphs + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(A_h == B_h); + HIP_CHECK(hipGraphExecDestroy(graphExec)) + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpyToSymbol_spt.cc b/catch/unit/memory/hipMemcpyToSymbol_spt.cc new file mode 100644 index 000000000..dfc5d1e9c --- /dev/null +++ b/catch/unit/memory/hipMemcpyToSymbol_spt.cc @@ -0,0 +1,118 @@ +/*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 hipMemcpyToSymbol_spt hipMemcpyToSymbol_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpyToSymbol_spt(const void* symbol, const void* src, + size_t sizeBytes, size_t offset __dparm(0), + hipMemcpyKind kind __dparm(hipMemcpyHostToDevice))` + - + * Copies data from the given symbol on the device. + */ +__device__ int devSymbol[10]; +__constant__ int constSymbol[10]; +/** + * Test Description + * ------------------------ + * - Basic test to check the negative cases of hipMemcpyToSymbol_spt. + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyToSymbol_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyToSymbol_spt_Negative") { + int result{0}; + SECTION("Invalid Src Ptr") { + HIP_CHECK_ERROR(hipMemcpyToSymbol_spt(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidSymbol); + } + SECTION("Invalid Dst Ptr") { + HIP_CHECK_ERROR(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Invalid Size") { + HIP_CHECK_ERROR(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Invalid Offset") { + HIP_CHECK_ERROR(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + SECTION("Invalid Direction") { + HIP_CHECK_ERROR(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, + hipMemcpyDeviceToHost), + hipErrorInvalidMemcpyDirection); + } +} +/** + * Test Description + * ------------------------ + * - Test Verifies hipMemcpyToSymbol_spt for simple use case + * For single value To Symbol + * For Array Values To Symbol + * For Array Values with offset To Symbol + * Test source + * ------------------------ + * - catch\unit\memory\hipMemcpyToSymbol_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpyToSymbol_spt_Sync") { + SECTION("Singular Value") { + int set{42}; + int result{0}; + HIP_CHECK(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), &set, sizeof(int))); + HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int))); + REQUIRE(result == set); + } + SECTION("Array Values") { + constexpr size_t size{10}; + int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2}; + int result[size] = {0}; + HIP_CHECK(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), set, sizeof(int) * size)); + HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } + SECTION("Offset'ed Values") { + constexpr size_t size{10}; + constexpr size_t offset = 5 * sizeof(int); + int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2}; + int result[size] = {0}; + HIP_CHECK(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), set, offset)); + HIP_CHECK(hipMemcpyToSymbol_spt(HIP_SYMBOL(devSymbol), set + 5, offset, offset)); + HIP_CHECK(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemcpy_spt.cc b/catch/unit/memory/hipMemcpy_spt.cc new file mode 100644 index 000000000..4e83a09ea --- /dev/null +++ b/catch/unit/memory/hipMemcpy_spt.cc @@ -0,0 +1,228 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#define LEN8 8 * 4 +#define LEN9 9 * 4 +#define LEN10 10 * 4 +#define LEN11 11 * 4 +#define LEN12 12 * 4 +__global__ void MemCpy8(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 8, In + tid * 8, 8); +} +__global__ void MemCpy9(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 9, In + tid * 9, 9); +} +__global__ void MemCpy10(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 10, In + tid * 10, 10); +} +__global__ void MemCpy11(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 11, In + tid * 11, 11); +} +__global__ void MemCpy12(uint8_t* In, uint8_t* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memcpy(Out + tid * 12, In + tid * 12, 12); +} +__global__ void MemSet8(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 8, 1, 8); +} +__global__ void MemSet9(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 9, 1, 9); +} +__global__ void MemSet10(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 10, 1, 10); +} +__global__ void MemSet11(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 11, 1, 11); +} +__global__ void MemSet12(uint8_t* In) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + memset(In + tid * 12, 1, 12); +} +/** + * @addtogroup hipMemcpy_spt + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemcpy_spt(void* dst, const void* src, size_t sizeBytes, + * hipMemcpyKind kind)` - Copy data from src to dst. + */ +/** + * Test Description + * ------------------------ + * - Test case to check memcpy and memset via kernel call. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemcpy_spt.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemcpy_spt_MemoryOperationsViaKernels") { + uint8_t *A, *Ad, *B, *Bd, *C, *Cd; + A = new uint8_t[LEN8]; + B = new uint8_t[LEN8]; + C = new uint8_t[LEN8]; + for (uint32_t i = 0; i < LEN8; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN8)); + HIP_CHECK(hipMalloc(&Bd, LEN8)); + HIP_CHECK(hipMalloc(&Cd, LEN8)); + HIP_CHECK(hipMemcpy_spt(Ad, A, LEN8, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy8, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet8, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy_spt(B, Bd, LEN8, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy_spt(C, Cd, LEN8, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN8; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + SECTION("MemCpySet1") { + A = new uint8_t[LEN9]; + B = new uint8_t[LEN9]; + C = new uint8_t[LEN9]; + for (uint32_t i = 0; i < LEN9; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN9)); + HIP_CHECK(hipMalloc(&Bd, LEN9)); + HIP_CHECK(hipMalloc(&Cd, LEN9)); + HIP_CHECK(hipMemcpy_spt(Ad, A, LEN9, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy9, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet9, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy_spt(B, Bd, LEN9, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy_spt(C, Cd, LEN9, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN9; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } + SECTION("MemCpySet2") { + A = new uint8_t[LEN10]; + B = new uint8_t[LEN10]; + C = new uint8_t[LEN10]; + for (uint32_t i = 0; i < LEN10; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN10)); + HIP_CHECK(hipMalloc(&Bd, LEN10)); + HIP_CHECK(hipMalloc(&Cd, LEN10)); + HIP_CHECK(hipMemcpy_spt(Ad, A, LEN10, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy10, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet10, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy_spt(B, Bd, LEN10, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy_spt(C, Cd, LEN10, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN10; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } + SECTION("MemCpySet3") { + A = new uint8_t[LEN11]; + B = new uint8_t[LEN11]; + C = new uint8_t[LEN11]; + for (uint32_t i = 0; i < LEN11; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN11)); + HIP_CHECK(hipMalloc(&Bd, LEN11)); + HIP_CHECK(hipMalloc(&Cd, LEN11)); + HIP_CHECK(hipMemcpy_spt(Ad, A, LEN11, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy11, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet11, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy_spt(B, Bd, LEN11, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy_spt(C, Cd, LEN11, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN11; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } + SECTION("MemCpySet4") { + A = new uint8_t[LEN12]; + B = new uint8_t[LEN12]; + C = new uint8_t[LEN12]; + for (uint32_t i = 0; i < LEN12; i++) { + A[i] = i; + B[i] = 0; + C[i] = 0; + } + HIP_CHECK(hipMalloc(&Ad, LEN12)); + HIP_CHECK(hipMalloc(&Bd, LEN12)); + HIP_CHECK(hipMalloc(&Cd, LEN12)); + HIP_CHECK(hipMemcpy_spt(Ad, A, LEN12, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(MemCpy12, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Ad, Bd); + hipLaunchKernelGGL(MemSet12, dim3(2, 1, 1), dim3(2, 1, 1), 0, 0, Cd); + HIP_CHECK(hipMemcpy_spt(B, Bd, LEN12, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy_spt(C, Cd, LEN12, hipMemcpyDeviceToHost)); + for (uint32_t i = 0; i < LEN12; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(C[i] == 1); + } + delete[] A; + delete[] B; + delete[] C; + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); + HIP_CHECK(hipFree(Cd)); + } +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/catch/unit/memory/hipMemset.cc b/catch/unit/memory/hipMemset.cc index 20bde2904..1b9681764 100644 --- a/catch/unit/memory/hipMemset.cc +++ b/catch/unit/memory/hipMemset.cc @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2021-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 @@ -30,50 +30,52 @@ Testcase Scenarios : #include - +#define NUM_OF_THREADS 10 // Table with unique number of elements and memset values. // (N, memsetval, memsetD32val, memsetD16val, memsetD8val) typedef std::tuple tupletype; -static constexpr std::initializer_list tableItems { - std::make_tuple((4*1024*1024), 0x42, 0xDEADBEEF, 0xDEAD, 0xDE), - std::make_tuple((10) , 0x42, 0x101 , 0x10, 0x1), - std::make_tuple((10013) , 0x5a, 0xDEADBEEF, 0xDEAD, 0xDE), - std::make_tuple((256*1024*1024), 0xa6, 0xCAFEBABE, 0xCAFE, 0xCA) - }; +static constexpr std::initializer_list tableItems{ + std::make_tuple((4 * 1024 * 1024), 0x42, 0xDEADBEEF, 0xDEAD, 0xDE), + std::make_tuple((10), 0x42, 0x101, 0x10, 0x1), + std::make_tuple((10013), 0x5a, 0xDEADBEEF, 0xDEAD, 0xDE), + std::make_tuple((256 * 1024 * 1024), 0xa6, 0xCAFEBABE, 0xCAFE, 0xCA)}; enum MemsetType { hipMemsetTypeDefault, + hipMemsetTypeDefaultSpt, hipMemsetTypeD8, hipMemsetTypeD16, hipMemsetTypeD32 }; -template -static bool testhipMemset(T *A_h, T *A_d, T memsetval, enum MemsetType type, - size_t numElements) { +template +static bool testhipMemset(T* A_h, T* A_d, T memsetval, enum MemsetType type, size_t numElements) { size_t Nbytes = numElements * sizeof(T); bool testResult = true; constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr. HIP_CHECK(hipMalloc(&A_d, Nbytes)); - A_h = reinterpret_cast (malloc(Nbytes)); + A_h = reinterpret_cast(malloc(Nbytes)); REQUIRE(A_h != nullptr); - for (int offset = MAX_OFFSET; offset >= 0; offset --) { + for (int offset = MAX_OFFSET; offset >= 0; offset--) { if (type == hipMemsetTypeDefault) { HIP_CHECK(hipMemset(A_d + offset, memsetval, numElements - offset)); + } else if (type == hipMemsetTypeDefaultSpt) { +#if HT_AMD + HIP_CHECK(hipMemset_spt(A_d + offset, memsetval, numElements - offset)); +#else + HIP_CHECK(hipMemset(A_d + offset, memsetval, numElements - offset)); +#endif } else if (type == hipMemsetTypeD8) { - HIP_CHECK(hipMemsetD8((hipDeviceptr_t)(A_d + offset), memsetval, - numElements - offset)); + HIP_CHECK(hipMemsetD8((hipDeviceptr_t)(A_d + offset), memsetval, numElements - offset)); } else if (type == hipMemsetTypeD16) { - HIP_CHECK(hipMemsetD16((hipDeviceptr_t)(A_d + offset), memsetval, - numElements - offset)); + HIP_CHECK(hipMemsetD16((hipDeviceptr_t)(A_d + offset), memsetval, numElements - offset)); } else if (type == hipMemsetTypeD32) { - HIP_CHECK(hipMemsetD32((hipDeviceptr_t)(A_d + offset), memsetval, - numElements - offset)); + HIP_CHECK(hipMemsetD32((hipDeviceptr_t)(A_d + offset), memsetval, numElements - offset)); } HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); @@ -92,9 +94,9 @@ static bool testhipMemset(T *A_h, T *A_d, T memsetval, enum MemsetType type, } -template -static bool testhipMemsetAsync(T *A_h, T *A_d, T memsetval, - enum MemsetType type, size_t numElements) { +template +static bool testhipMemsetAsync(T* A_h, T* A_d, T memsetval, enum MemsetType type, + size_t numElements) { size_t Nbytes = numElements * sizeof(T); bool testResult = true; constexpr auto MAX_OFFSET = 3; // To memset on unaligned ptr. @@ -102,25 +104,29 @@ static bool testhipMemsetAsync(T *A_h, T *A_d, T memsetval, HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipMalloc(&A_d, Nbytes)); - A_h = reinterpret_cast (malloc(Nbytes)); + A_h = reinterpret_cast(malloc(Nbytes)); REQUIRE(A_h != nullptr); - for (int offset = MAX_OFFSET; offset >= 0; offset --) { + for (int offset = MAX_OFFSET; offset >= 0; offset--) { if (type == hipMemsetTypeDefault) { - HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset, - stream)); - + HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset, stream)); + } else if (type == hipMemsetTypeDefaultSpt) { +#if HT_AMD + HIP_CHECK(hipMemsetAsync_spt(A_d + offset, memsetval, numElements - offset, stream)); +#else + HIP_CHECK(hipMemsetAsync(A_d + offset, memsetval, numElements - offset, stream)); +#endif } else if (type == hipMemsetTypeD8) { - HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)(A_d + offset), memsetval, - numElements - offset, stream)); + HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)(A_d + offset), memsetval, numElements - offset, + stream)); } else if (type == hipMemsetTypeD16) { - HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)(A_d + offset), memsetval, - numElements - offset, stream)); + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)(A_d + offset), memsetval, numElements - offset, + stream)); } else if (type == hipMemsetTypeD32) { - HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)(A_d + offset), memsetval, - numElements - offset, stream)); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)(A_d + offset), memsetval, numElements - offset, + stream)); } HIP_CHECK(hipStreamSynchronize(stream)); @@ -154,7 +160,7 @@ TEST_CASE("Unit_hipMemset_SetMemoryWithOffset") { bool ret; std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) = - GENERATE(table(tableItems)); + GENERATE(table(tableItems)); SECTION("Memset with hipMemsetTypeDefault") { @@ -162,7 +168,13 @@ TEST_CASE("Unit_hipMemset_SetMemoryWithOffset") { ret = testhipMemset(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N); REQUIRE(ret == true); } - +#if HT_AMD + SECTION("Memset with hipMemsetTypeDefaultSpt") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemset(cA_h, cA_d, memsetval, hipMemsetTypeDefaultSpt, N); + REQUIRE(ret == true); + } +#endif SECTION("Memset with hipMemsetTypeD32") { int32_t *iA_d{nullptr}, *iA_h{nullptr}; ret = testhipMemset(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N); @@ -196,7 +208,7 @@ TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") { bool ret; std::tie(N, memsetval, memsetD32val, memsetD16val, memsetD8val) = - GENERATE(table(tableItems)); + GENERATE(table(tableItems)); SECTION("Memset with hipMemsetTypeDefault") { @@ -204,7 +216,13 @@ TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") { ret = testhipMemsetAsync(cA_h, cA_d, memsetval, hipMemsetTypeDefault, N); REQUIRE(ret == true); } - +#if HT_AMD + SECTION("Memset with hipMemsetTypeDefaultSpt") { + char *cA_d{nullptr}, *cA_h{nullptr}; + ret = testhipMemsetAsync(cA_h, cA_d, memsetval, hipMemsetTypeDefaultSpt, N); + REQUIRE(ret == true); + } +#endif SECTION("Memset with hipMemsetTypeD32") { int32_t *iA_d{nullptr}, *iA_h{nullptr}; ret = testhipMemsetAsync(iA_h, iA_d, memsetD32val, hipMemsetTypeD32, N); @@ -230,20 +248,27 @@ TEST_CASE("Unit_hipMemsetAsync_SetMemoryWithOffset") { TEST_CASE("Unit_hipMemset_SmallBufferSizes") { char *A_d, *A_h; constexpr int memsetval = 0x24; - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); auto numElements = GENERATE(range(1, 4)); int numBytes = numElements * sizeof(char); HIP_CHECK(hipMalloc(&A_d, numBytes)); - A_h = reinterpret_cast (malloc(numBytes)); - + A_h = reinterpret_cast(malloc(numBytes)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset(A_d, memsetval, numBytes)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset_spt(A_d, memsetval, numBytes)); +#else + HIP_CHECK(hipMemset(A_d, memsetval, numBytes)); +#endif + } HIP_CHECK(hipMemset(A_d, memsetval, numBytes)); HIP_CHECK(hipMemcpy(A_h, A_d, numBytes, hipMemcpyDeviceToHost)); for (int i = 0; i < numBytes; i++) { if (A_h[i] != memsetval) { - INFO("Mismatch at index:" << i << " computed:" << A_h[i] - << " memsetval:" << memsetval); + INFO("Mismatch at index:" << i << " computed:" << A_h[i] << " memsetval:" << memsetval); REQUIRE(false); } } @@ -259,17 +284,17 @@ TEST_CASE("Unit_hipMemset_SmallBufferSizes") { TEST_CASE("Unit_hipMemset_2AsyncOperations") { std::vector v; v.resize(2048); - float* p2, *p3; - HIP_CHECK(hipMalloc(reinterpret_cast(&p2), 4096 + 4096*2)); - p3 = p2+2048; + float *p2, *p3; + HIP_CHECK(hipMalloc(reinterpret_cast(&p2), 4096 + 4096 * 2)); + p3 = p2 + 2048; hipStream_t s; HIP_CHECK(hipStreamCreate(&s)); - HIP_CHECK(hipMemsetAsync(p2, 0, 32*32*4, s)); - HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s)); + HIP_CHECK(hipMemsetAsync(p2, 0, 32 * 32 * 4, s)); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32 * 32, s)); HIP_CHECK(hipStreamSynchronize(s)); for (int i = 0; i < 256; ++i) { - HIP_CHECK(hipMemsetAsync(p2, 0, 32*32*4, s)); - HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s)); + HIP_CHECK(hipMemsetAsync(p2, 0, 32 * 32 * 4, s)); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32 * 32, s)); } HIP_CHECK(hipStreamSynchronize(s)); HIP_CHECK(hipDeviceSynchronize()); @@ -282,3 +307,71 @@ TEST_CASE("Unit_hipMemset_2AsyncOperations") { HIP_CHECK(hipFree(p2)); HIP_CHECK(hipStreamDestroy(s)); } + +/** + * Thread functions. + */ +#if HT_AMD +static void memsetSptTest(char* devBuf, char* hostBuf, size_t N, char val) { + HIP_CHECK(hipMemset_spt(devBuf, val, N)); + HIP_CHECK(hipMemcpy(hostBuf, devBuf, N, hipMemcpyDeviceToHost)); +} + +static void memsetSptAsyncTest(char* devBuf, char* hostBuf, size_t N, char val) { + HIP_CHECK(hipMemsetAsync_spt(devBuf, val, N, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemcpy(hostBuf, devBuf, N, hipMemcpyDeviceToHost)); +} + +/** + * Test multiple memset spt operations in parallel. + */ +TEST_CASE("Unit_hipMemset_spt_paralleloperations") { + size_t N = NUM_OF_THREADS * 1024 * 1024; + std::vector v(N); + std::vector testThread; + char* dev_buf[NUM_OF_THREADS]; + char val = 'a'; + for (int idx = 0; idx < NUM_OF_THREADS; idx++) { + HIP_CHECK(hipMalloc(&dev_buf[idx], N * sizeof(char))); + testThread.emplace_back( + std::thread(memsetSptTest, dev_buf[idx], &v[idx * 1024 * 1024], (1024 * 1024), val)); + } + // Wait for completion of all threads + for (int idx = 0; idx < NUM_OF_THREADS; idx++) { + testThread[idx].join(); + } + // Validate the output values + REQUIRE(true == std::all_of(v.begin(), v.end(), [](char val) { return val == 'a'; })); + // free all resources + for (int idx = 0; idx < NUM_OF_THREADS; idx++) { + HIP_CHECK(hipFree(dev_buf[idx])); + } +} + +/** + * Test multiple hipMemsetAsync_spt operations in parallel. + */ +TEST_CASE("Unit_hipMemsetAsync_spt_paralleloperations") { + size_t N = NUM_OF_THREADS * 1024 * 1024; + std::vector v(N); + std::vector testThread; + char* dev_buf[NUM_OF_THREADS]; + char val = 'a'; + for (int idx = 0; idx < NUM_OF_THREADS; idx++) { + HIP_CHECK(hipMalloc(&dev_buf[idx], N * sizeof(char))); + testThread.emplace_back( + std::thread(memsetSptAsyncTest, dev_buf[idx], &v[idx * 1024 * 1024], (1024 * 1024), val)); + } + // Wait for completion of all threads + for (int idx = 0; idx < NUM_OF_THREADS; idx++) { + testThread[idx].join(); + } + // Validate the output values + REQUIRE(true == std::all_of(v.begin(), v.end(), [](char val) { return val == 'a'; })); + // free all resources + for (int idx = 0; idx < NUM_OF_THREADS; idx++) { + HIP_CHECK(hipFree(dev_buf[idx])); + } +} +#endif diff --git a/catch/unit/memory/hipMemset2D.cc b/catch/unit/memory/hipMemset2D.cc index 975929850..cd6f3acae 100644 --- a/catch/unit/memory/hipMemset2D.cc +++ b/catch/unit/memory/hipMemset2D.cc @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2021-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 @@ -32,24 +32,20 @@ // (width2D, height2D, memsetWidth, memsetHeight) typedef std::tuple tupletype; -static constexpr std::initializer_list tableItems { - std::make_tuple(20, 20, 20, 20), - std::make_tuple(10, 10, 4, 4), - std::make_tuple(100, 100, 20, 40), - std::make_tuple(256, 256, 39, 19), - std::make_tuple(100, 100, 20, 0), - std::make_tuple(100, 100, 0, 20), - std::make_tuple(100, 100, 0, 0), - }; - - +static constexpr std::initializer_list tableItems{ + std::make_tuple(20, 20, 20, 20), std::make_tuple(10, 10, 4, 4), + std::make_tuple(100, 100, 20, 40), std::make_tuple(256, 256, 39, 19), + std::make_tuple(100, 100, 20, 0), std::make_tuple(100, 100, 0, 20), + std::make_tuple(100, 100, 0, 0), +}; +enum MemsetType { hipMemsetTypeDefault, hipMemsetTypeDefaultSpt }; /** * Basic Functionality of hipMemset2D */ TEST_CASE("Unit_hipMemset2D_BasicFunctional") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); constexpr int memsetval = 0x24; constexpr size_t numH = 256; constexpr size_t numW = 256; @@ -59,23 +55,28 @@ TEST_CASE("Unit_hipMemset2D_BasicFunctional") { size_t elements = numW * numH; char *A_d, *A_h; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, - numH)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH)); A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); for (size_t i = 0; i < elements; i++) { A_h[i] = 1; } - - HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); - HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, - hipMemcpyDeviceToHost)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset2D_spt(A_d, pitch_A, memsetval, numW, numH)); +#else + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); +#endif + } + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost)); for (size_t i = 0; i < elements; i++) { if (A_h[i] != memsetval) { - INFO("Memset2D mismatch at index:" << i << " computed:" - << A_h[i] << " memsetval:" << memsetval); + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); REQUIRE(false); } } @@ -90,7 +91,7 @@ TEST_CASE("Unit_hipMemset2D_BasicFunctional") { */ TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); constexpr int memsetval = 0x26; constexpr size_t numH = 256; constexpr size_t numW = 256; @@ -100,26 +101,31 @@ TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") { size_t elements = numW * numH; char *A_d, *A_h; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, - width, numH)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH)); A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); for (size_t i = 0; i < elements; i++) { - A_h[i] = 1; + A_h[i] = 1; } - hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset2DAsync_spt(A_d, pitch_A, memsetval, numW, numH, stream)); +#else + HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); +#endif + } HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, - hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost)); - for (size_t i=0; i < elements; i++) { + for (size_t i = 0; i < elements; i++) { if (A_h[i] != memsetval) { - INFO("Memset2DAsync mismatch at index:" << i << " computed:" - << A_h[i] << " memsetval:" << memsetval); + INFO("Memset2DAsync mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); REQUIRE(false); } } @@ -135,7 +141,7 @@ TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") { */ TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); int width2D, height2D; int memsetWidth, memsetHeight; char *A_d, *A_h; @@ -143,13 +149,12 @@ TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") { constexpr int memsetval = 0x26; std::tie(width2D, height2D, memsetWidth, memsetHeight) = - GENERATE(table(tableItems)); + GENERATE(table(tableItems)); size_t width = width2D * sizeof(char); size_t sizeElements = width * height2D; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, - width, height2D)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, height2D)); A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); @@ -158,18 +163,23 @@ TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") { A_h[index] = 'c'; } - INFO("2D Dimension: Width:" << width2D << " Height:" << height2D << - " MemsetWidth:" << memsetWidth << " MemsetHeight:" << memsetHeight); - - HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight)); - HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, width2D, height2D, - hipMemcpyDeviceToHost)); + INFO("2D Dimension: Width:" << width2D << " Height:" << height2D << " MemsetWidth:" << memsetWidth + << " MemsetHeight:" << memsetHeight); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset2D_spt(A_d, pitch_A, memsetval, memsetWidth, memsetHeight)); +#else + HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, memsetWidth, memsetHeight)); +#endif + } + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, width2D, height2D, hipMemcpyDeviceToHost)); for (int row = 0; row < memsetHeight; row++) { for (int column = 0; column < memsetWidth; column++) { if (A_h[(row * width) + column] != memsetval) { - INFO("A_h[" << row << "][" << column << "]" << - " didnot match " << memsetval); + INFO("A_h[" << row << "][" << column << "]" << " didnot match " << memsetval); REQUIRE(false); } } @@ -180,22 +190,23 @@ TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") { } /** -* Test Description -* ------------------------ -* - Basic functional testcase for triggering capturehipMemset2DAsync internal -* API to improve code coverage -* Test source -* ------------------------ -* - unit/memory/hipMemset2D.cc -* Test requirements -* ------------------------ -* - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + * - Basic functional testcase for triggering capturehipMemset2DAsync internal + * API to improve code coverage + * Test source + * ------------------------ + * - unit/memory/hipMemset2D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipMemset2DAsync_capturehipMemset2DAsync") { char *A_h, *B_h, *A_d; hipGraph_t graph{nullptr}; hipGraphExec_t graphExec{nullptr}; int rows, cols; + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); rows = GENERATE(3, 4, 100); cols = GENERATE(3, 4, 100); hipStream_t stream; @@ -209,15 +220,21 @@ TEST_CASE("Unit_hipMemset2DAsync_capturehipMemset2DAsync") { A_h[i * cols + j] = 'a'; } } - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &devPitch, - sizeof(char) * cols, rows)); - HIP_CHECK(hipMemcpy2D(A_d, devPitch, A_h, sizeof(char) * cols, - sizeof(char) * cols, rows, hipMemcpyHostToDevice)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(char) * cols, rows)); + HIP_CHECK(hipMemcpy2D(A_d, devPitch, A_h, sizeof(char) * cols, sizeof(char) * cols, rows, + hipMemcpyHostToDevice)); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); - HIP_CHECK(hipMemset2DAsync(A_d, devPitch, 'b', sizeof(char) * cols, rows, - stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset2DAsync(A_d, devPitch, 'b', sizeof(char) * cols, rows, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset2DAsync_spt(A_d, devPitch, 'b', sizeof(char) * cols, rows, stream)); +#else + HIP_CHECK(hipMemset2DAsync(A_d, devPitch, 'b', sizeof(char) * cols, rows, stream)); +#endif + } HIP_CHECK(hipStreamEndCapture(stream, &graph)); HIP_CHECK(hipDeviceSynchronize()); @@ -225,8 +242,8 @@ TEST_CASE("Unit_hipMemset2DAsync_capturehipMemset2DAsync") { HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2D(B_h, sizeof(char) * cols, A_d, devPitch, - sizeof(char) * cols, rows, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy2D(B_h, sizeof(char) * cols, A_d, devPitch, sizeof(char) * cols, rows, + hipMemcpyDeviceToHost)); for (int i = 0; i < rows; i++) { for (int j = 0; j < cols; j++) { diff --git a/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc b/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc index f70736da8..b56d49242 100644 --- a/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc +++ b/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2021-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 @@ -15,7 +15,7 @@ * 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. -*/ + */ /** Testcase Scenarios : @@ -35,14 +35,21 @@ #define NUM_H 256 #define NUM_W 256 +enum MemsetType { hipMemsetTypeDefault, hipMemsetTypeDefaultSpt }; - -void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch, - size_t width, hipStream_t stream) { +void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch, size_t width, + hipStream_t stream, enum MemsetType type) { constexpr int memsetval = 0x22; - HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream)); - HIPCHECK(hipMemcpy2DAsync(A_h, width, A_d, pitch, NUM_W, NUM_H, - hipMemcpyDeviceToHost, stream)); + if (type == hipMemsetTypeDefault) { + HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream)); + } else { +#if HT_AMD + HIPCHECK(hipMemset2DAsync_spt(A_d, pitch, memsetval, NUM_W, NUM_H, stream)); +#else + HIPCHECK(hipMemset2DAsync(A_d, pitch, memsetval, NUM_W, NUM_H, stream)); +#endif + } + HIPCHECK(hipMemcpy2DAsync(A_h, width, A_d, pitch, NUM_W, NUM_H, hipMemcpyDeviceToHost, stream)); } @@ -51,7 +58,7 @@ void queueJobsForhipMemset2DAsync(char* A_d, char* A_h, size_t pitch, */ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); constexpr auto N = 4 * 1024 * 1024; constexpr auto blocksPerCU = 6; // to hide latency constexpr auto threadsPerBlock = 256; @@ -65,39 +72,41 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { int validateCount{}; blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, - width, NUM_H)); - HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, - width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, width, NUM_H)); A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); B_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(B_h != nullptr); - HIP_CHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, - width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, width, NUM_H)); for (size_t i = 0; i < elements; i++) { B_h[i] = i; } - HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, - hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, hipMemcpyHostToDevice)); SECTION("Using User created stream") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); for (size_t k = 0; k < ITER; k++) { - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream, B_d, C_d, elements); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream, + B_d, C_d, elements); HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, - hipMemcpyDeviceToHost)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset2DAsync_spt(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); +#else + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); +#endif + } + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, hipMemcpyDeviceToHost)); - for (size_t p = 0 ; p < elements ; p++) { + for (size_t p = 0; p < elements; p++) { if (A_h[p] == memsetval) { - validateCount+= 1; + validateCount += 1; } } } @@ -105,19 +114,25 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { } SECTION("Using hipStreamPerThread stream") { for (size_t k = 0; k < ITER; k++) { - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, hipStreamPerThread, B_d, C_d, elements); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + hipStreamPerThread, B_d, C_d, elements); HIP_CHECK(hipGetLastError()); HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, - hipStreamPerThread)); - HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, - hipMemcpyDeviceToHost)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, hipStreamPerThread)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset2DAsync_spt(C_d, pitch_C, memsetval, NUM_W, NUM_H, hipStreamPerThread)); +#else + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, hipStreamPerThread)); +#endif + } + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, hipMemcpyDeviceToHost)); - for (size_t p = 0 ; p < elements ; p++) { + for (size_t p = 0; p < elements; p++) { if (A_h[p] == memsetval) { - validateCount+= 1; + validateCount += 1; } } } @@ -125,8 +140,11 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { REQUIRE(static_cast(validateCount) == (ITER * elements)); - HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); - free(A_h); free(B_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); + free(A_h); + free(B_h); } @@ -135,7 +153,7 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { */ TEST_CASE("Unit_hipMemset2DAsync_MultiThread") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); constexpr auto memPerThread = 200; constexpr int memsetval = 0x22; char *A_d, *A_h, *B_d, *B_h, *C_d; @@ -153,52 +171,49 @@ TEST_CASE("Unit_hipMemset2DAsync_MultiThread") { return; } - std::thread *t = new std::thread[thread_count]; + std::thread* t = new std::thread[thread_count]; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, - width, NUM_H)); - HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, - width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&B_d), &pitch_B, width, NUM_H)); A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); B_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(B_h != nullptr); - HIP_CHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, - width, NUM_H)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&C_d), &pitch_C, width, NUM_H)); - for (size_t i = 0 ; i < elements ; i++) { + for (size_t i = 0; i < elements; i++) { B_h[i] = i; } - HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, - hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, hipMemcpyHostToDevice)); HIP_CHECK(hipStreamCreate(&stream)); - for (int i = 0 ; i < ITER ; i++) { - for (size_t k = 0 ; k < thread_count; k++) { - if (k%2) { - t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, A_h, pitch_A, - width, stream); + for (int i = 0; i < ITER; i++) { + for (size_t k = 0; k < thread_count; k++) { + if (k % 2) { + t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, A_h, pitch_A, width, stream, type); } else { - t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, B_h, pitch_A, - width, stream); + t[k] = std::thread(queueJobsForhipMemset2DAsync, A_d, B_h, pitch_A, width, stream, type); } } - for (size_t j = 0 ; j < thread_count; j++) { + for (size_t j = 0; j < thread_count; j++) { t[j].join(); } HIP_CHECK(hipStreamSynchronize(stream)); - for (size_t k = 0 ; k < elements ; k++) { + for (size_t k = 0; k < elements; k++) { if ((A_h[k] == memsetval) && (B_h[k] == memsetval)) { - validateCount+= 1; + validateCount += 1; } } } REQUIRE(static_cast(validateCount) == (ITER * elements)); - HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); - free(A_h); free(B_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); + free(A_h); + free(B_h); HIP_CHECK(hipStreamDestroy(stream)); delete[] t; diff --git a/catch/unit/memory/hipMemset3D.cc b/catch/unit/memory/hipMemset3D.cc index 105fe3e68..33460960b 100644 --- a/catch/unit/memory/hipMemset3D.cc +++ b/catch/unit/memory/hipMemset3D.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021i-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 @@ -26,14 +26,13 @@ THE SOFTWARE. #include - - +enum MemsetType { hipMemsetTypeDefault, hipMemsetTypeDefaultSpt }; /** * Basic Functional test of hipMemset3D */ TEST_CASE("Unit_hipMemset3D_BasicFunctional") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); constexpr int memsetval = 0x22; constexpr size_t numH = 256; constexpr size_t numW = 256; @@ -41,23 +40,31 @@ TEST_CASE("Unit_hipMemset3D_BasicFunctional") { size_t width = numW * sizeof(char); size_t sizeElements = width * numH * depth; size_t elements = numW * numH * depth; - char *A_h; + char* A_h; hipExtent extent = make_hipExtent(width, numH, depth); hipPitchedPtr devPitchedPtr; HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); - A_h = reinterpret_cast(malloc(sizeElements)); + A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); for (size_t i = 0; i < elements; i++) { - A_h[i] = 1; + A_h[i] = 1; + } + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3D_spt(devPitchedPtr, memsetval, extent)); +#else + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); +#endif } - HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); hipMemcpy3DParms myparms{}; myparms.srcPos = make_hipPos(0, 0, 0); myparms.dstPos = make_hipPos(0, 0, 0); - myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); myparms.srcPtr = devPitchedPtr; myparms.extent = extent; #if HT_NVIDIA @@ -68,11 +75,11 @@ TEST_CASE("Unit_hipMemset3D_BasicFunctional") { HIP_CHECK(hipMemcpy3D(&myparms)); for (size_t i = 0; i < elements; i++) { - if (A_h[i] != memsetval) { - INFO("Memset3D mismatch at index:" << i << " computed:" - << A_h[i] << " memsetval:" << memsetval); - REQUIRE(false); - } + if (A_h[i] != memsetval) { + INFO("Memset3D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(false); + } } HIP_CHECK(hipFree(devPitchedPtr.ptr)); free(A_h); @@ -83,7 +90,7 @@ TEST_CASE("Unit_hipMemset3D_BasicFunctional") { */ TEST_CASE("Unit_hipMemset3DAsync_BasicFunctional") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); constexpr int memsetval = 0x22; constexpr size_t numH = 256; constexpr size_t numW = 256; @@ -93,24 +100,32 @@ TEST_CASE("Unit_hipMemset3DAsync_BasicFunctional") { size_t elements = numW * numH * depth; hipExtent extent = make_hipExtent(width, numH, depth); hipPitchedPtr devPitchedPtr; - char *A_h; + char* A_h; HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); - A_h = reinterpret_cast(malloc(sizeElements)); + A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); for (size_t i = 0; i < elements; i++) { - A_h[i] = 1; + A_h[i] = 1; } hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3DAsync_spt(devPitchedPtr, memsetval, extent, stream)); +#else + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); +#endif + } HIP_CHECK(hipStreamSynchronize(stream)); hipMemcpy3DParms myparms{}; myparms.srcPos = make_hipPos(0, 0, 0); myparms.dstPos = make_hipPos(0, 0, 0); - myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); myparms.srcPtr = devPitchedPtr; myparms.extent = extent; #if HT_NVIDIA @@ -121,30 +136,31 @@ TEST_CASE("Unit_hipMemset3DAsync_BasicFunctional") { HIP_CHECK(hipMemcpy3D(&myparms)); for (size_t i = 0; i < elements; i++) { - if (A_h[i] != memsetval) { - INFO("Memset3DAsync mismatch at index:" << i << " computed:" - << A_h[i] << " memsetval:" << memsetval); - REQUIRE(false); - } + if (A_h[i] != memsetval) { + INFO("Memset3DAsync mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(false); + } } HIP_CHECK(hipFree(devPitchedPtr.ptr)); free(A_h); } /** -* Test Description -* ------------------------ -* - Basic scenario to trigger capturehipMemset3DAsync internal -* api for improved code coverage -* Test source -* ------------------------ -* - unit/memory/hipMemset3D.cc -* Test requirements -* ------------------------ -* - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + * - Basic scenario to trigger capturehipMemset3DAsync internal + * api for improved code coverage + * Test source + * ------------------------ + * - unit/memory/hipMemset3D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipMemset3DAsync_capturehipMemset3DAsync") { - char *A_h; + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + char* A_h; hipPitchedPtr A_d; hipGraph_t graph{nullptr}; hipGraphExec_t graphExec{nullptr}; @@ -160,7 +176,15 @@ TEST_CASE("Unit_hipMemset3DAsync_capturehipMemset3DAsync") { HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); - HIP_CHECK(hipMemset3DAsync(A_d, 'a', extent, stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3DAsync(A_d, 'a', extent, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3DAsync_spt(A_d, 'a', extent, stream)); +#else + HIP_CHECK(hipMemset3DAsync(A_d, 'a', extent, stream)); +#endif + } HIP_CHECK(hipStreamEndCapture(stream, &graph)); HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); @@ -176,7 +200,7 @@ TEST_CASE("Unit_hipMemset3DAsync_capturehipMemset3DAsync") { params.kind = hipMemcpyDeviceToHost; HIP_CHECK(hipMemcpy3D(¶ms)); for (int i = 0; i < (row * col * dep); i++) { - REQUIRE(A_h[i]=='a'); + REQUIRE(A_h[i] == 'a'); } HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); diff --git a/catch/unit/memory/hipMemset3DFunctional.cc b/catch/unit/memory/hipMemset3DFunctional.cc index aae255e04..9de0a435b 100644 --- a/catch/unit/memory/hipMemset3DFunctional.cc +++ b/catch/unit/memory/hipMemset3DFunctional.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021-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 @@ -85,58 +85,95 @@ Testcase Scenarios : #define XSET_LEN 10 #define XPOS_END 19 - +enum MemsetType { hipMemsetTypeDefault, hipMemsetTypeDefaultSpt }; /** * Memset with extent passed and verify data to be intact */ -static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) { +static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent, enum MemsetType type) { hipPitchedPtr devPitchedPtr; hipError_t ret; - char *A_h; + char* A_h; size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; size_t width = numW * sizeof(char); hipExtent extent = make_hipExtent(width, numH, depth); size_t sizeElements = width * numH * depth; - size_t elements = numW* numH* depth; + size_t elements = numW * numH * depth; - A_h = reinterpret_cast(malloc(sizeElements)); + A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); memset(A_h, 0, sizeElements); HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); if (bAsync) { - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream); - INFO("testMemsetWithExtent(" << extent.width << "," << extent.height - << "," << extent.depth << ") memset " - << MEMSETVAL << ", ret : " << ret); - REQUIRE(ret == hipSuccess); - - ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream); - INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height - << "," << tstExtent.depth << ") memset " - << TESTVAL << "ret : " << ret); - REQUIRE(ret == hipSuccess); - + /*hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream));*/ + if (type == hipMemsetTypeDefault) { + ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height << "," << extent.depth + << ") memset " << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height << "," + << tstExtent.depth << ") memset " << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); + } else { +#if HT_AMD + ret = hipMemset3DAsync_spt(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height << "," << extent.depth + << ") memset " << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + ret = hipMemset3DAsync_spt(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height << "," + << tstExtent.depth << ") memset " << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); +#else + ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height << "," << extent.depth + << ") memset " << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height << "," + << tstExtent.depth << ") memset " << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); +#endif + } HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); + // HIP_CHECK(hipStreamDestroy(stream)); } else { - ret = hipMemset3D(devPitchedPtr, MEMSETVAL, extent); - INFO("testMemsetWithExtent(" << extent.width << "," << extent.height - << "," << extent.depth << ") memset " - << MEMSETVAL << ",ret : " << ret); - REQUIRE(ret == hipSuccess); - - ret = hipMemset3D(devPitchedPtr, TESTVAL, tstExtent); - INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height - << "," << tstExtent.depth << ") memset " - << TESTVAL << ",ret : " << ret); - REQUIRE(ret == hipSuccess); + if (type == hipMemsetTypeDefault) { + ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height << "," << extent.depth + << ") memset " << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height << "," + << tstExtent.depth << ") memset " << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); + } else { +#if HT_AMD + ret = hipMemset3DAsync_spt(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height << "," << extent.depth + << ") memset " << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + ret = hipMemset3DAsync_spt(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height << "," + << tstExtent.depth << ") memset " << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); +#else + ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream); + INFO("testMemsetWithExtent(" << extent.width << "," << extent.height << "," << extent.depth + << ") memset " << MEMSETVAL << ", ret : " << ret); + REQUIRE(ret == hipSuccess); + ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream); + INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height << "," + << tstExtent.depth << ") memset " << TESTVAL << "ret : " << ret); + REQUIRE(ret == hipSuccess); +#endif + } } - - + HIP_CHECK(hipStreamDestroy(stream)); hipMemcpy3DParms myparms{}; myparms.srcPos = make_hipPos(0, 0, 0); myparms.dstPos = make_hipPos(0, 0, 0); @@ -153,9 +190,9 @@ static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) { for (size_t i = 0; i < elements; i++) { if (A_h[i] != MEMSETVAL) { - INFO("testMemsetWithExtent: index:" << i << ",computed:" - << std::hex << static_cast(A_h[i]) << ",memsetval:" - << std::hex << MEMSETVAL); + INFO("testMemsetWithExtent: index:" << i << ",computed:" << std::hex + << static_cast(A_h[i]) << ",memsetval:" << std::hex + << MEMSETVAL); REQUIRE(false); } } @@ -168,17 +205,17 @@ static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) { /** * Validates data after performing memory set operation with max memset value */ -static void testMemsetMaxValue(bool bAsync) { +static void testMemsetMaxValue(bool bAsync, enum MemsetType type) { hipPitchedPtr devPitchedPtr; - unsigned char *A_h; + unsigned char* A_h; int memsetval = std::numeric_limits::max(); size_t numH = NUMH_MAX, numW = NUMW_MAX, depth = DEPTH_MAX; size_t width = numW * sizeof(unsigned char); hipExtent extent = make_hipExtent(width, numH, depth); size_t sizeElements = width * numH * depth; - size_t elements = numW* numH* depth; + size_t elements = numW * numH * depth; - A_h = reinterpret_cast (malloc(sizeElements)); + A_h = reinterpret_cast(malloc(sizeElements)); REQUIRE(A_h != nullptr); memset(A_h, 0, sizeElements); @@ -187,16 +224,40 @@ static void testMemsetMaxValue(bool bAsync) { SECTION("Using user created stream") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3DAsync_spt(devPitchedPtr, memsetval, extent, stream)); +#else + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); +#endif + } HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipStreamDestroy(stream)); } SECTION("Using hipStreamPerThread") { - HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, hipStreamPerThread)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, hipStreamPerThread)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3DAsync_spt(devPitchedPtr, memsetval, extent, hipStreamPerThread)); +#else + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, hipStreamPerThread)); +#endif + } HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); } } else { - HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3D_spt(devPitchedPtr, memsetval, extent)); +#else + HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); +#endif + } } hipMemcpy3DParms myparms{}; @@ -215,9 +276,8 @@ static void testMemsetMaxValue(bool bAsync) { for (size_t i = 0; i < elements; i++) { if (A_h[i] != memsetval) { - INFO("testMemsetMaxValue: index:" << i << ",computed:" - << std::hex << static_cast(A_h[i]) << ",memsetval:" - << std::hex << memsetval); + INFO("testMemsetMaxValue: index:" << i << ",computed:" << std::hex << static_cast(A_h[i]) + << ",memsetval:" << std::hex << memsetval); REQUIRE(false); } } @@ -229,17 +289,24 @@ static void testMemsetMaxValue(bool bAsync) { * Function seeks device ptr to random slice and performs Memset operation * on the slice selected. */ -static void seekAndSet3DArraySlice(bool bAsync) { +static void seekAndSet3DArraySlice(bool bAsync, enum MemsetType type) { char array3D[ZSIZE_S][YSIZE_S][XSIZE_S]{}; dim3 arr_dimensions = dim3(ZSIZE_S, YSIZE_S, XSIZE_S); - hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x, - arr_dimensions.y, arr_dimensions.z); + hipExtent extent = + make_hipExtent(sizeof(char) * arr_dimensions.x, arr_dimensions.y, arr_dimensions.z); hipPitchedPtr devicePitchedPointer; int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL; HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent)); - HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); - + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3D_spt(devicePitchedPointer, memsetval, extent)); +#else + HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); +#endif + } // select random slice for memset unsigned int seed = time(nullptr); int slice_index = HipTest::RAND_R(&seed) % ZSIZE_S; @@ -251,32 +318,46 @@ static void seekAndSet3DArraySlice(bool bAsync) { size_t slicePitch = pitch * extent.height; // Point devptr to selected slice - char *devPtrSlice = (reinterpret_cast(devicePitchedPointer.ptr)) - + slice_index * slicePitch; - hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x, - arr_dimensions.y, 1); - hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrSlice, pitch, - arr_dimensions.x, arr_dimensions.y); + char* devPtrSlice = + (reinterpret_cast(devicePitchedPointer.ptr)) + slice_index * slicePitch; + hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x, arr_dimensions.y, 1); + hipPitchedPtr modDevPitchedPtr = + make_hipPitchedPtr(devPtrSlice, pitch, arr_dimensions.x, arr_dimensions.y); if (bAsync) { // Memset selected slice (Async) hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, - extentSlice, stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, extentSlice, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3DAsync_spt(modDevPitchedPtr, memsetval4seeked, extentSlice, stream)); +#else + HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, extentSlice, stream)); +#endif + } HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipStreamDestroy(stream)); } else { // Memset selected slice - HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3D_spt(modDevPitchedPtr, memsetval4seeked, extentSlice)); +#else + HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice)); +#endif + } } // Copy result back to host buffer hipMemcpy3DParms myparms{}; myparms.srcPos = make_hipPos(0, 0, 0); myparms.dstPos = make_hipPos(0, 0, 0); - myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, - arr_dimensions.x, arr_dimensions.y); + myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, arr_dimensions.x, + arr_dimensions.y); myparms.srcPtr = devicePitchedPointer; myparms.extent = extent; #if HT_NVIDIA @@ -292,18 +373,18 @@ static void seekAndSet3DArraySlice(bool bAsync) { for (int x = 0; x < XSIZE_S; x++) { if (z == slice_index) { if (array3D[z][y][x] != memsetval4seeked) { - INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z - << "," << y << "," << x << ") " << "computed:" << std::hex - << array3D[z][y][x] << ", memsetval:" << std::hex - << memsetval4seeked); + INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" + << z << "," << y << "," << x << ") " + << "computed:" << std::hex << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval4seeked); REQUIRE(false); } } else { if (array3D[z][y][x] != memsetval) { - INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z - << "," << y << "," << x << ") " << "computed:" << std::hex - << array3D[z][y][x] << ", memsetval:" << std::hex - << memsetval); + INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" + << z << "," << y << "," << x << ") " + << "computed:" << std::hex << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval); REQUIRE(false); } } @@ -318,17 +399,24 @@ static void seekAndSet3DArraySlice(bool bAsync) { * Function seeks device ptr to selected portion of 3d array * and performs Memset operation on the portion. */ -static void seekAndSet3DArrayPortion(bool bAsync) { +static void seekAndSet3DArrayPortion(bool bAsync, enum MemsetType type) { char array3D[ZSIZE_P][YSIZE_P][XSIZE_P]{}; dim3 arr_dimensions = dim3(ZSIZE_P, YSIZE_P, XSIZE_P); - hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x, - arr_dimensions.y, arr_dimensions.z); + hipExtent extent = + make_hipExtent(sizeof(char) * arr_dimensions.x, arr_dimensions.y, arr_dimensions.z); hipPitchedPtr devicePitchedPointer; int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL; HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent)); - HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); - + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3D_spt(devicePitchedPointer, memsetval, extent)); +#else + HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent)); +#endif + } // For memsetting extent/size(10,10,10) in the mid portion of cube(30,30,30), // seek device ptr to (10,10,10) and then memset 10 bytes across x,y,z axis. size_t pitch = devicePitchedPointer.pitch; @@ -336,40 +424,54 @@ static void seekAndSet3DArrayPortion(bool bAsync) { int slice_index = ZPOS_START, y = YPOS_START, x = XPOS_START; // Select 10th slice - char *devPtrSlice = (reinterpret_cast(devicePitchedPointer.ptr)) - + slice_index * slicePitch; + char* devPtrSlice = + (reinterpret_cast(devicePitchedPointer.ptr)) + slice_index * slicePitch; // Now select row at height as 10 - char *current_row = reinterpret_cast(devPtrSlice + y * pitch); + char* current_row = reinterpret_cast(devPtrSlice + y * pitch); // Now select index of selected row as 10 - char *devPtrIndexed = ¤t_row[x]; + char* devPtrIndexed = ¤t_row[x]; // Make dev Pitchedptr, extent - hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrIndexed, pitch, - arr_dimensions.x, arr_dimensions.y); - hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN, - ZSET_LEN); + hipPitchedPtr modDevPitchedPtr = + make_hipPitchedPtr(devPtrIndexed, pitch, arr_dimensions.x, arr_dimensions.y); + hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN, ZSET_LEN); if (bAsync) { // Memset selected portion (Async) hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, - setExtent, stream)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, setExtent, stream)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3DAsync_spt(modDevPitchedPtr, memsetval4seeked, setExtent, stream)); +#else + HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked, setExtent, stream)); +#endif + } HIP_CHECK(hipStreamSynchronize(stream)); HIP_CHECK(hipStreamDestroy(stream)); } else { // Memset selected portion - HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent)); + if (type == hipMemsetTypeDefault) { + HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent)); + } else { +#if HT_AMD + HIP_CHECK(hipMemset3D_spt(modDevPitchedPtr, memsetval4seeked, setExtent)); +#else + HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent)); +#endif + } } // Copy result back to host buffer hipMemcpy3DParms myparms{}; myparms.srcPos = make_hipPos(0, 0, 0); myparms.dstPos = make_hipPos(0, 0, 0); - myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, - arr_dimensions.y, arr_dimensions.z); + myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x, arr_dimensions.y, + arr_dimensions.z); myparms.srcPtr = devicePitchedPointer; myparms.extent = extent; #if HT_NVIDIA @@ -383,24 +485,23 @@ static void seekAndSet3DArrayPortion(bool bAsync) { for (int z = 0; z < ZSIZE_P; z++) { for (int y = 0; y < YSIZE_P; y++) { for (int x = 0; x < XSIZE_P; x++) { - if ((z >= ZPOS_START && z <= ZPOS_END) && - (y >= YPOS_START && y <= YPOS_END) && + if ((z >= ZPOS_START && z <= ZPOS_END) && (y >= YPOS_START && y <= YPOS_END) && (x >= XPOS_START && x <= XPOS_END)) { if (array3D[z][y][x] != memsetval4seeked) { - INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z - << "," << y << "," << x << ") " << "computed:" << std::hex - << array3D[z][y][x] << ", memsetval:" << std::hex - << memsetval4seeked); + INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" + << z << "," << y << "," << x << ") " + << "computed:" << std::hex << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval4seeked); REQUIRE(false); } } else { - if (array3D[z][y][x] != memsetval) { - INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z - << "," << y << "," << x << ") " << "computed:" << std::hex - << array3D[z][y][x] << ", memsetval:" << std::hex - << memsetval); + if (array3D[z][y][x] != memsetval) { + INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" + << z << "," << y << "," << x << ") " + << "computed:" << std::hex << array3D[z][y][x] << ", memsetval:" << std::hex + << memsetval); REQUIRE(false); - } + } } } } @@ -410,39 +511,38 @@ static void seekAndSet3DArrayPortion(bool bAsync) { } - /** * Test Memset3D with different combinations of extent * taking zero and non-zero fields. */ TEST_CASE("Unit_hipMemset3D_MemsetWithExtent") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); hipExtent testExtent; size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; SECTION("Memset with extent width(0)") { // Memset with extent width(0) and verify data to be intact testExtent = make_hipExtent(0, numH, depth); - testMemsetWithExtent(0, testExtent); + testMemsetWithExtent(0, testExtent, type); } SECTION("Memset with extent height(0)") { // Memset with extent height(0) and verify data to be intact testExtent = make_hipExtent(numW, 0, depth); - testMemsetWithExtent(0, testExtent); + testMemsetWithExtent(0, testExtent, type); } SECTION("Memset with extent depth(0)") { // Memset with extent depth(0) and verify data to be intact testExtent = make_hipExtent(numW, numH, 0); - testMemsetWithExtent(0, testExtent); + testMemsetWithExtent(0, testExtent, type); } SECTION("Memset with extent width,height,depth as 0") { // Memset with extent width,height,depth as 0 and verify data to be intact testExtent = make_hipExtent(0, 0, 0); - testMemsetWithExtent(0, testExtent); + testMemsetWithExtent(0, testExtent, type); } } @@ -453,32 +553,32 @@ TEST_CASE("Unit_hipMemset3D_MemsetWithExtent") { */ TEST_CASE("Unit_hipMemset3DAsync_MemsetWithExtent") { CHECK_IMAGE_SUPPORT - + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); hipExtent testExtent; size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT; SECTION("Memset with extent width(0)") { // Memset with extent width(0) and verify data to be intact testExtent = make_hipExtent(0, numH, depth); - testMemsetWithExtent(1, testExtent); + testMemsetWithExtent(1, testExtent, type); } SECTION("Memset with extent height(0)") { // Memset with extent height(0) and verify data to be intact testExtent = make_hipExtent(numW, 0, depth); - testMemsetWithExtent(1, testExtent); + testMemsetWithExtent(1, testExtent, type); } SECTION("Memset with extent depth(0)") { // Memset with extent depth(0) and verify data to be intact testExtent = make_hipExtent(numW, numH, 0); - testMemsetWithExtent(1, testExtent); + testMemsetWithExtent(1, testExtent, type); } SECTION("Memset with extent width,height,depth as 0") { // Memset with extent width,height,depth as 0 and verify data to be intact testExtent = make_hipExtent(0, 0, 0); - testMemsetWithExtent(1, testExtent); + testMemsetWithExtent(1, testExtent, type); } } @@ -487,8 +587,8 @@ TEST_CASE("Unit_hipMemset3DAsync_MemsetWithExtent") { */ TEST_CASE("Unit_hipMemset3D_MemsetMaxValue") { CHECK_IMAGE_SUPPORT - - testMemsetMaxValue(0); + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + testMemsetMaxValue(0, type); } /** @@ -496,8 +596,8 @@ TEST_CASE("Unit_hipMemset3D_MemsetMaxValue") { */ TEST_CASE("Unit_hipMemset3DAsync_MemsetMaxValue") { CHECK_IMAGE_SUPPORT - - testMemsetMaxValue(1); + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + testMemsetMaxValue(1, type); } /** @@ -505,8 +605,8 @@ TEST_CASE("Unit_hipMemset3DAsync_MemsetMaxValue") { */ TEST_CASE("Unit_hipMemset3D_SeekSetSlice") { CHECK_IMAGE_SUPPORT - - seekAndSet3DArraySlice(0); + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + seekAndSet3DArraySlice(0, type); } /** @@ -514,8 +614,8 @@ TEST_CASE("Unit_hipMemset3D_SeekSetSlice") { */ TEST_CASE("Unit_hipMemset3DAsync_SeekSetSlice") { CHECK_IMAGE_SUPPORT - - seekAndSet3DArraySlice(1); + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + seekAndSet3DArraySlice(1, type); } /** @@ -523,8 +623,8 @@ TEST_CASE("Unit_hipMemset3DAsync_SeekSetSlice") { */ TEST_CASE("Unit_hipMemset3D_SeekSetArrayPortion") { CHECK_IMAGE_SUPPORT - - seekAndSet3DArrayPortion(0); + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + seekAndSet3DArrayPortion(0, type); } /** @@ -532,6 +632,6 @@ TEST_CASE("Unit_hipMemset3D_SeekSetArrayPortion") { */ TEST_CASE("Unit_hipMemset3DAsync_SeekSetArrayPortion") { CHECK_IMAGE_SUPPORT - - seekAndSet3DArrayPortion(1); + enum MemsetType type = GENERATE(hipMemsetTypeDefault, hipMemsetTypeDefaultSpt); + seekAndSet3DArrayPortion(1, type); } diff --git a/catch/unit/memory/memcpy3d_spt_tests_common.hh b/catch/unit/memory/memcpy3d_spt_tests_common.hh new file mode 100644 index 000000000..a7a9aea03 --- /dev/null +++ b/catch/unit/memory/memcpy3d_spt_tests_common.hh @@ -0,0 +1,155 @@ +/* +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. +*/ +#pragma clang diagnostic ignored "-Wmissing-field-initializers" +#include +#include +#include +#include +#include +using PtrVariant = std::variant; +static inline hipMemcpy3DParms GetMemcpy3DParms_spt(PtrVariant dst_ptr, hipPos dst_pos, + PtrVariant src_ptr, hipPos src_pos, + hipExtent extent, hipMemcpyKind kind) { + hipMemcpy3DParms parms = {0}; + if (std::holds_alternative(dst_ptr)) { + parms.dstArray = std::get(dst_ptr); + } else { + parms.dstPtr = std::get(dst_ptr); + } + parms.dstPos = dst_pos; + if (std::holds_alternative(src_ptr)) { + parms.srcArray = std::get(src_ptr); + } else { + parms.srcPtr = std::get(src_ptr); + } + parms.srcPos = src_pos; + parms.extent = extent; + parms.kind = kind; + return parms; +} +template +hipError_t Memcpy3DWrapper_spt(PtrVariant dst_ptr, hipPos dst_pos, PtrVariant src_ptr, + hipPos src_pos, hipExtent extent, hipMemcpyKind kind, + hipStream_t stream = nullptr) { + auto parms = GetMemcpy3DParms_spt(dst_ptr, dst_pos, src_ptr, src_pos, extent, kind); + if constexpr (async) { + return hipMemcpy3DAsync_spt(&parms, stream); + } else { + return hipMemcpy3D_spt(&parms); + } +} +template +void Memcpy3DZeroWidthHeightDepth_spt(F memcpy_func, const hipStream_t stream = nullptr) { + constexpr hipExtent extent{127 * sizeof(int), 128, 8}; + const auto [width_mult, height_mult, depth_mult] = + GENERATE(std::make_tuple(0, 1, 1), std::make_tuple(1, 0, 1), std::make_tuple(1, 1, 0)); + SECTION("Device to Host") { + LinearAllocGuard3D device_alloc(extent); + LinearAllocGuard host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.width() * device_alloc.height() * device_alloc.depth()); + std::fill_n(host_alloc.ptr(), + device_alloc.width_logical() * device_alloc.height() * device_alloc.depth(), 42); + HIP_CHECK(hipMemset3D(device_alloc.pitched_ptr(), 1, device_alloc.extent())); + HIP_CHECK(memcpy_func( + make_hipPitchedPtr(host_alloc.ptr(), device_alloc.width(), device_alloc.width(), + device_alloc.height()), + make_hipPos(0, 0, 0), device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), + make_hipExtent(device_alloc.width() * width_mult, device_alloc.height() * height_mult, + device_alloc.depth() * depth_mult), + hipMemcpyDeviceToHost, stream)); + if constexpr (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + ArrayFindIfNot(host_alloc.ptr(), static_cast(42), + device_alloc.width_logical() * device_alloc.height() * device_alloc.depth()); + } + SECTION("Device to Device") { + LinearAllocGuard3D src_alloc(extent); + LinearAllocGuard3D dst_alloc(extent); + LinearAllocGuard host_alloc( + LinearAllocs::hipHostMalloc, dst_alloc.width() * dst_alloc.height() * dst_alloc.depth()); + HIP_CHECK(hipMemset3D(src_alloc.pitched_ptr(), 1, src_alloc.extent())); + HIP_CHECK(hipMemset3D(dst_alloc.pitched_ptr(), 42, dst_alloc.extent())); + HIP_CHECK( + memcpy_func(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), + make_hipExtent(dst_alloc.width() * width_mult, dst_alloc.height() * height_mult, + dst_alloc.depth() * depth_mult), + hipMemcpyDeviceToDevice, stream)); + if constexpr (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + HIP_CHECK(Memcpy3DWrapper_spt(make_hipPitchedPtr(host_alloc.ptr(), dst_alloc.width(), + dst_alloc.width(), dst_alloc.height()), + make_hipPos(0, 0, 0), dst_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToHost)); + ArrayFindIfNot(host_alloc.ptr(), static_cast(42), + dst_alloc.width_logical() * dst_alloc.height()); + } + SECTION("Host to Device") { + LinearAllocGuard3D device_alloc(extent); + LinearAllocGuard src_host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.width() * device_alloc.height() * device_alloc.depth()); + LinearAllocGuard dst_host_alloc( + LinearAllocs::hipHostMalloc, + device_alloc.width() * device_alloc.height() * device_alloc.depth()); + std::fill_n(src_host_alloc.ptr(), + device_alloc.width_logical() * device_alloc.height() * device_alloc.depth(), 1); + HIP_CHECK(hipMemset3D(device_alloc.pitched_ptr(), 42, device_alloc.extent())); + HIP_CHECK(memcpy_func( + device_alloc.pitched_ptr(), make_hipPos(0, 0, 0), + make_hipPitchedPtr(src_host_alloc.ptr(), device_alloc.width(), device_alloc.width(), + device_alloc.height()), + make_hipPos(0, 0, 0), + make_hipExtent(device_alloc.width() * width_mult, device_alloc.height() * height_mult, + device_alloc.depth() * depth_mult), + hipMemcpyHostToDevice, stream)); + if constexpr (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + HIP_CHECK(Memcpy3DWrapper_spt(make_hipPitchedPtr(dst_host_alloc.ptr(), device_alloc.width(), + device_alloc.width(), device_alloc.height()), + make_hipPos(0, 0, 0), device_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), device_alloc.extent(), + hipMemcpyDeviceToHost)); + ArrayFindIfNot(dst_host_alloc.ptr(), static_cast(42), + device_alloc.width_logical() * device_alloc.height()); + } + SECTION("Host to Host") { + const auto alloc_size = extent.width * extent.height * extent.depth; + LinearAllocGuard src_alloc(LinearAllocs::hipHostMalloc, alloc_size); + LinearAllocGuard dst_alloc(LinearAllocs::hipHostMalloc, alloc_size); + std::fill_n(src_alloc.ptr(), alloc_size, 1); + std::fill_n(dst_alloc.ptr(), alloc_size, 42); + HIP_CHECK( + memcpy_func(make_hipPitchedPtr(dst_alloc.ptr(), extent.width, extent.width, extent.height), + make_hipPos(0, 0, 0), + make_hipPitchedPtr(src_alloc.ptr(), extent.width, extent.width, extent.height), + make_hipPos(0, 0, 0), + make_hipExtent(extent.width * width_mult, extent.height * height_mult, + extent.depth * depth_mult), + hipMemcpyHostToHost, stream)); + if constexpr (should_synchronize) { + HIP_CHECK(hipStreamSynchronize(stream)); + } + ArrayFindIfNot(dst_alloc.ptr(), static_cast(42), alloc_size); + } +}