From 092279449e2729cfacf7140b3bb9557d43a13c2e Mon Sep 17 00:00:00 2001 From: amd-srinivas1 Date: Tue, 14 Oct 2025 10:47:59 +0530 Subject: [PATCH] SWDEV-546345-[catch2][dtest]-Tests for hipMemSetD2DXX Apis(Memory management) (#896) * SWDEV-546345-Added tests for memsetd2dxx apis * SWDEV-546345-Optimized the code. * SWDEV-546345-Optimized the code. * SWDEV-546345-Addressed review comments * SWDEV-546345-Updated code. --- .../catch/unit/memory/CMakeLists.txt | 6 + .../catch/unit/memory/hipMemsetD2D16.cc | 156 ++++++++++++++++ .../catch/unit/memory/hipMemsetD2D16Async.cc | 171 +++++++++++++++++ .../catch/unit/memory/hipMemsetD2D32.cc | 153 ++++++++++++++++ .../catch/unit/memory/hipMemsetD2D32Async.cc | 167 +++++++++++++++++ .../catch/unit/memory/hipMemsetD2D8.cc | 155 ++++++++++++++++ .../catch/unit/memory/hipMemsetD2D8Async.cc | 173 ++++++++++++++++++ 7 files changed, 981 insertions(+) create mode 100644 projects/hip-tests/catch/unit/memory/hipMemsetD2D16.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemsetD2D16Async.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemsetD2D32.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemsetD2D32Async.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemsetD2D8.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemsetD2D8Async.cc diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 9001ab394b..089973a7c0 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -89,6 +89,12 @@ set(TEST_SRC hipMemcpy2DArrayToArray.cc hipMemcpy3DBatchAsync.cc hipMemcpyBatchAsync.cc + hipMemsetD2D8.cc + hipMemsetD2D8Async.cc + hipMemsetD2D16.cc + hipMemsetD2D16Async.cc + hipMemsetD2D32.cc + hipMemsetD2D32Async.cc ) if(UNIX) diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetD2D16.cc b/projects/hip-tests/catch/unit/memory/hipMemsetD2D16.cc new file mode 100644 index 0000000000..abe0e0a88c --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemsetD2D16.cc @@ -0,0 +1,156 @@ +/* +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, INNCLUDING 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 ANNY 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 +#include +#include +/** + * @addtogroup hipMemsetD2D16 hipMemsetD2D16 + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, + * size_t width, size_t height);` - + * Fills 2D memory range of 'width' 16-bit values synchronously to the specified short value. + * Height specifies numbers of rows to set and dstPitch speicifies the number of bytes between each + * row. + */ +/** + * Test Description + * ------------------------ + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D16.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D16_BasicFunctional") { + constexpr uint16_t memsetval = static_cast(0xDEADBEEF); + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(uint16_t); + size_t sizeElements = numW * numH; + + uint16_t *A_d; + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH, + 2 * sizeof(uint16_t))); + std::vectorA_h(sizeElements, 1); + + HIP_CHECK(hipMemsetD2D16(A_d, pitch_A, memsetval, width, numH)); + HIP_CHECK(hipMemcpy2D(A_h.data(), width, A_d, pitch_A, width, numH, hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < sizeElements; i++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(A_h[i] == memsetval); + } + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Uneven width and Hight 2D Memory. + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D16.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D16_UnEvenRowsCols") { + uint16_t *A_d; + constexpr uint16_t memsetVal = 5; + int rows, cols; + rows = GENERATE(3, 4, 100); + cols = GENERATE(5, 6, 100); + size_t devPitch; + + size_t size = rows * cols; + std::vectorB_h(size, 1); + + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(uint16_t) * cols, + rows, 2 * sizeof(uint16_t))); + + HIP_CHECK(hipMemsetD2D16(A_d, devPitch, memsetVal, sizeof(uint16_t) * cols, rows)); + HIP_CHECK(hipMemcpy2D(B_h.data(), sizeof(uint16_t) * cols, A_d, devPitch, sizeof(uint16_t) * cols, rows, + hipMemcpyDeviceToHost)); + + for (int i = 0; i < rows; i++) { + for (int j = 0; j < cols; j++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << B_h[i * cols + j] + << " memsetval:" << memsetVal); + REQUIRE(B_h[i * cols + j] == memsetVal); + } + } + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Checks function behaviour when provided invalid arguments. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D16.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D16_NegTsts") { + uint16_t* A_d; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t width = numW * sizeof(uint16_t); + size_t devPitch; + constexpr uint16_t memsetval = static_cast(0x26); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, width, numH, + 2 * sizeof(uint16_t))); + SECTION("nullptr destination") { + HIP_CHECK_ERROR(hipMemsetD2D16(nullptr, devPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("OutOfBound destination") { + void* outOfBoundsDst{reinterpret_cast(A_d) + devPitch * numH + 1}; + HIP_CHECK_ERROR(hipMemsetD2D16(outOfBoundsDst, devPitch, memsetval, numW, numH), + hipErrorInvalidValue); + } + SECTION("Dst pointer points to Source Memory") { + uint16_t* B_d; + std::unique_ptr hostPtr; + hostPtr.reset(new uint16_t[numH * width]); + B_d = hostPtr.get(); + HIP_CHECK_ERROR(hipMemsetD2D16(B_d, devPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("Invalid Pitch") { + size_t inValidPitch = 1; + HIP_CHECK_ERROR(hipMemsetD2D16(A_d, inValidPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("Negative Values of Hight, Width") { + HIP_CHECK_ERROR(hipMemsetD2D16(A_d, devPitch, memsetval, numW, -10), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemsetD2D16(A_d, devPitch, memsetval, -10, numH), hipErrorInvalidValue); + } + HIP_CHECK(hipFree(A_d)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetD2D16Async.cc b/projects/hip-tests/catch/unit/memory/hipMemsetD2D16Async.cc new file mode 100644 index 0000000000..4fbd4c79c2 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemsetD2D16Async.cc @@ -0,0 +1,171 @@ +/* +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, INNCLUDING 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 ANNY 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 +#include +#include +/** + * @addtogroup hipMemsetD2D16Async hipMemsetD2D16Async + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, + size_t width, size_t height, hipStream_t stream __dparm(0))` - + * Fills 2D memory range of 'width' 16-bit values asynchronously to the specified short + * value. Height specifies numbers of rows to set and dstPitch speicifies the number of bytes + * between each row. + */ +/** + * Test Description + * ------------------------ + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D16Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D16Async_BasicFunctional") { + constexpr int memsetval = 0x24; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(uint16_t); + size_t sizeElements = numW * numH; + + uint16_t *A_d; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH, + 2 * sizeof(uint16_t))); + std::vectorA_h(sizeElements, 0); + + HIP_CHECK(hipMemsetD2D16Async(A_d, pitch_A, memsetval, width, numH, stream)); + HIP_CHECK(hipMemcpy2DAsync(A_h.data(), width, A_d, pitch_A, width, numH, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t i = 0; i < sizeElements; i++) { + if (A_h[i] != memsetval) { + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(A_h[i] == memsetval); + } + } + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Uneven width and Hight 2D Memory. + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D16Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D16Async_UnEvenRowsCols") { + uint16_t *A_d; + int rows, cols; + constexpr int memsetval = 5; + rows = GENERATE(3, 4, 100); + cols = GENERATE(3, 4, 100); + size_t devPitch; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + size_t size = rows * cols; + + std::vectorB_h(size, 1); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(uint16_t) * cols, + rows, 2 * sizeof(uint16_t))); + + HIP_CHECK(hipMemsetD2D16Async(A_d, devPitch, memsetval, sizeof(uint16_t) * cols, rows, stream)); + HIP_CHECK(hipMemcpy2DAsync(B_h.data(), sizeof(uint16_t) * cols, A_d, devPitch, sizeof(uint16_t) * cols, + rows, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (int i = 0; i < rows; i++) { + for (int j = 0; j < cols; j++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << B_h[i * cols + j] + << " memsetval:" << memsetval); + REQUIRE(B_h[i * cols + j] == memsetval); + } + } + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Checks function behaviour when provided invalid arguments. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D16Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D16Async_NegTsts") { + uint16_t* A_d; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t width = numW * sizeof(uint16_t); + size_t devPitch; + constexpr uint16_t memsetval = static_cast(0x26); + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, width, numH, + 2 * sizeof(uint16_t))); + SECTION("nullptr destination") { + HIP_CHECK_ERROR(hipMemsetD2D16Async(nullptr, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("OutOfBound destination") { + void* outOfBoundsDst{reinterpret_cast(A_d) + devPitch * numH + 1}; + HIP_CHECK_ERROR(hipMemsetD2D16Async(outOfBoundsDst, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Dst pointer points to Source Memory") { + uint16_t* B_d; + std::unique_ptr hostPtr; + hostPtr.reset(new uint16_t[numH * width]); + B_d = hostPtr.get(); + HIP_CHECK_ERROR(hipMemsetD2D16Async(B_d, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Invalid Pitch") { + size_t inValidPitch = 1; + HIP_CHECK_ERROR(hipMemsetD2D16Async(A_d, inValidPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Negative Values of Hight, Width") { + HIP_CHECK_ERROR(hipMemsetD2D16Async(A_d, devPitch, memsetval, numW, -10, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemsetD2D16Async(A_d, devPitch, memsetval, -10, numH, stream), + hipErrorInvalidValue); + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetD2D32.cc b/projects/hip-tests/catch/unit/memory/hipMemsetD2D32.cc new file mode 100644 index 0000000000..3b0dc6abb6 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemsetD2D32.cc @@ -0,0 +1,153 @@ +/* +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, INNCLUDING 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 ANNY 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 +#include +#include +/** + * @addtogroup hipMemsetD2D32 hipMemsetD2D32 + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width, + size_t height)` - + * Fills 2D memory range of 'width' 32-bit values synchronously to thespecified int value. + * Height specifies numbers of rows to set and dstPitch speicifies the number of bytes between each + * row. + */ +/** + * Test Description + * ------------------------ + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D32.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D32_BasicFunctional") { + constexpr int memsetval = 15; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(int); + size_t sizeElements = numW * numH; + int *A_d; + + std::vectorA_h(sizeElements, 1); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH, sizeof(int))); + + HIP_CHECK(hipMemsetD2D32(A_d, pitch_A, memsetval, width, numH)); + HIP_CHECK(hipMemcpy2D(A_h.data(), width, A_d, pitch_A, width, numH, hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < sizeElements; i++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(A_h[i] == memsetval); + } + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Uneven width and Hight 2D Memory. + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D32.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D32_UnEvenRowsCols") { + constexpr int memsetval = 5; + int *A_d; + int rows, cols; + rows = GENERATE(3, 4, 100); + cols = GENERATE(3, 4, 100); + size_t devPitch; + + size_t size = rows * cols; + std::vectorB_h(size, 1); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(int) * cols, rows, + sizeof(int))); + + HIP_CHECK(hipMemsetD2D32(A_d, devPitch, memsetval, sizeof(int) * cols, rows)); + HIP_CHECK(hipMemcpy2D(B_h.data(), sizeof(int) * cols, A_d, devPitch, sizeof(int) * cols, rows, + hipMemcpyDeviceToHost)); + + for (int i = 0; i < rows; i++) { + for (int j = 0; j < cols; j++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << B_h[i * cols + j] + << " memsetval:" << memsetval); + REQUIRE(B_h[i * cols + j] == memsetval); + } + } + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Checks function behaviour when provided invalid arguments. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D32.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D32_NegTsts") { + int* A_d; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t width = numW * sizeof(int); + size_t devPitch; + constexpr int memsetval = static_cast(0x26); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, width, numH, sizeof(int))); + SECTION("nullptr destination") { + HIP_CHECK_ERROR(hipMemsetD2D32(nullptr, devPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("OutOfBound destination") { + void* outOfBoundsDst{reinterpret_cast(A_d) + devPitch * numH + 1}; + HIP_CHECK_ERROR(hipMemsetD2D32(outOfBoundsDst, devPitch, memsetval, numW, numH), + hipErrorInvalidValue); + } + SECTION("Dst pointer points to Source Memory") { + int* B_d; + std::unique_ptr hostPtr; + hostPtr.reset(new int[numH * width]); + B_d = hostPtr.get(); + HIP_CHECK_ERROR(hipMemsetD2D32(B_d, devPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("Invalid Pitch") { + size_t inValidPitch = 1; + HIP_CHECK_ERROR(hipMemsetD2D32(A_d, inValidPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("Negative Values of Hight, Width") { + HIP_CHECK_ERROR(hipMemsetD2D32(A_d, devPitch, memsetval, numW, -10), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemsetD2D32(A_d, devPitch, memsetval, -10, numH), hipErrorInvalidValue); + } + HIP_CHECK(hipFree(A_d)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetD2D32Async.cc b/projects/hip-tests/catch/unit/memory/hipMemsetD2D32Async.cc new file mode 100644 index 0000000000..4f57d74f29 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemsetD2D32Async.cc @@ -0,0 +1,167 @@ +/* +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, INNCLUDING 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 ANNY 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 +#include +#include +/** + * @addtogroup hipMemsetD2D32Async hipMemsetD2D32Async + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height, hipStream_t stream __dparm(0))` - + * Fills 2D memory range of 'width' 32-bit values asynchronously to the specified int + * value. Height specifies numbers of rows to set and dstPitch speicifies the number of bytes + * between each row. + */ +/** + * Test Description + * ------------------------ + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D32Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D32Async_BasicFunctional") { + constexpr int memsetval = 0x24; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(int); + size_t sizeElements = numW * numH; + int *A_d; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH, sizeof(int))); + std::vectorA_h(sizeElements, 1); + + HIP_CHECK(hipMemsetD2D32Async(A_d, pitch_A, memsetval, width, numH, stream)); + HIP_CHECK(hipMemcpy2DAsync(A_h.data(), width, A_d, pitch_A, width, numH, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t i = 0; i < sizeElements; i++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(A_h[i] == memsetval); + } + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Uneven width and Hight 2D Memory. + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D32Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D32Async_UnEvenRowsCols") { + constexpr int memsetval = 5; + int *A_d; + int rows, cols; + rows = GENERATE(3, 4, 100); + cols = GENERATE(3, 4, 100); + size_t devPitch; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t size = rows * cols; + std::vectorB_h(size, 1); + + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(int) * cols, rows, + sizeof(int))); + + HIP_CHECK(hipMemsetD2D32Async(A_d, devPitch, memsetval, sizeof(int) * cols, rows, stream)); + HIP_CHECK(hipMemcpy2DAsync(B_h.data(), sizeof(int) * cols, A_d, devPitch, sizeof(int) * cols, rows, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (int i = 0; i < rows; i++) { + for (int j = 0; j < cols; j++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << B_h[i * cols + j] + << " memsetval:" << memsetval); + REQUIRE(B_h[i * cols + j] == memsetval); + } + } + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Checks function behaviour when provided invalid arguments. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D32Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D32Async_NegTsts") { + int* A_d; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t width = numW * sizeof(int); + size_t devPitch; + constexpr int memsetval = 15; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, width, numH, sizeof(int))); + SECTION("nullptr destination") { + HIP_CHECK_ERROR(hipMemsetD2D32Async(nullptr, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("OutOfBound destination") { + void* outOfBoundsDst{reinterpret_cast(A_d) + devPitch * numH + 1}; + HIP_CHECK_ERROR(hipMemsetD2D32Async(outOfBoundsDst, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Dst pointer points to Source Memory") { + int* B_d; + std::unique_ptr hostPtr; + hostPtr.reset(new int[numH * width]); + B_d = hostPtr.get(); + HIP_CHECK_ERROR(hipMemsetD2D32Async(B_d, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Invalid Pitch") { + size_t inValidPitch = 1; + HIP_CHECK_ERROR(hipMemsetD2D32Async(A_d, inValidPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Negative Values of Hight, Width") { + HIP_CHECK_ERROR(hipMemsetD2D32Async(A_d, devPitch, memsetval, numW, -10, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemsetD2D32Async(A_d, devPitch, memsetval, -10, numH, stream), + hipErrorInvalidValue); + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetD2D8.cc b/projects/hip-tests/catch/unit/memory/hipMemsetD2D8.cc new file mode 100644 index 0000000000..1ee0f1b47e --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemsetD2D8.cc @@ -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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING 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 ANNY 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 +#include +/** + * @addtogroup hipMemsetD2D8 hipMemsetD2D8 + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, + size_t width, size_t height);` - + * Fills 2D memory range of 'width' 8-bit values synchronously to the specified char value. + * Height specifies numbers of rows to set and dstPitch speicifies the number of bytes between each + * row. + */ +/** + * Test Description + * ------------------------ + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D8.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D8_BasicFunctional") { + constexpr char memsetval = 'c'; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = numW * numH; + char *A_d; + HIP_CHECK( + hipMemAllocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH, 4 * sizeof(char))); + std::vectorA_h(sizeElements, 'a'); + HIP_CHECK(hipMemsetD2D8(A_d, pitch_A, memsetval, width, numH)); + HIP_CHECK(hipMemcpy2D(A_h.data(), width, A_d, pitch_A, width, numH, hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < sizeElements; i++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(A_h[i] == memsetval); + } + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Uneven width and Hight 2D Memory. + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D8.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D8_UnEvenRowsCols") { + char *A_d; + int rows, cols; + rows = GENERATE(3, 4, 100); + cols = GENERATE(4, 5, 100); + size_t devPitch; + constexpr char memsetval = 'c'; + size_t size = rows * cols; + std::vectorA_h(size, 'a'); + std::vectorB_h(size, 'a'); + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(char) * cols, rows, + 4 * sizeof(char))); + HIP_CHECK(hipMemcpy2D(A_d, devPitch, A_h.data(), sizeof(char) * cols, sizeof(char) * cols, rows, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipMemsetD2D8(A_d, devPitch, memsetval, sizeof(char) * cols, rows)); + + HIP_CHECK(hipMemcpy2D(B_h.data(), 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++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << B_h[i * cols + j] + << " memsetval:" << memsetval); + REQUIRE(B_h[i * cols + j] == memsetval); + } + } + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Checks function behaviour when provided invalid arguments. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D8.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D8_NegTsts") { + char* A_d; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t width = numW * sizeof(char); + size_t devPitch; + constexpr char memsetval = 'c'; + HIP_CHECK( + hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, width, numH, 4 * sizeof(char))); + SECTION("nullptr destination") { + HIP_CHECK_ERROR(hipMemsetD2D8(nullptr, devPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("OutOfBound destination") { + void* outOfBoundsDst{reinterpret_cast(A_d) + devPitch * numH + 1}; + HIP_CHECK_ERROR(hipMemsetD2D8(outOfBoundsDst, devPitch, memsetval, numW, numH), + hipErrorInvalidValue); + } + SECTION("Dst pointer points to Source Memory") { + char* B_d; + std::unique_ptr hostPtr; + hostPtr.reset(new char[numH * width]); + B_d = hostPtr.get(); + HIP_CHECK_ERROR(hipMemsetD2D8(B_d, devPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("Invalid Pitch") { + size_t inValidPitch = 1; + HIP_CHECK_ERROR(hipMemsetD2D8(A_d, inValidPitch, memsetval, numW, numH), hipErrorInvalidValue); + } + SECTION("Negative Values of Hight, Width") { + HIP_CHECK_ERROR(hipMemsetD2D8(A_d, devPitch, memsetval, numW, -10), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemsetD2D8(A_d, devPitch, memsetval, -10, numH), hipErrorInvalidValue); + } + HIP_CHECK(hipFree(A_d)); +} +/** + * End doxygen group MemoryTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/memory/hipMemsetD2D8Async.cc b/projects/hip-tests/catch/unit/memory/hipMemsetD2D8Async.cc new file mode 100644 index 0000000000..86e1afa8c8 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemsetD2D8Async.cc @@ -0,0 +1,173 @@ +/* +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, INNCLUDING 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 ANNY 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 +#include + +/** + * @addtogroup hipMemsetD2D8Async hipMemsetD2D8Async + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, + unsigned char value, size_t width, + size_t height, hipStream_t stream __dparm(0))` - + * Fills 2D memory range of 'width' 8-bit values asynchronously to the specified char + * value. Height specifies numbers of rows to set and dstPitch speicifies the number + * of bytes between each row. + */ +/** + * Test Description + * ------------------------ + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D8Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D8Async_BasicFunctional") { + constexpr char memsetval = 'c'; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t pitch_A; + size_t width = numW * sizeof(char); + size_t sizeElements = numW * numH; + + char *A_d; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK( + hipMemAllocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH, 4 * sizeof(char))); + std::vectorA_h(sizeElements, 'a'); + + HIP_CHECK(hipMemsetD2D8Async(A_d, pitch_A, memsetval, width, numH, stream)); + HIP_CHECK(hipMemcpy2DAsync(A_h.data(), width, A_d, pitch_A, width, numH, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (size_t i = 0; i < sizeElements; i++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << A_h[i] + << " memsetval:" << memsetval); + REQUIRE(A_h[i] == memsetval); + } + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Uneven width and Hight 2D Memory. + * - Checks that allocated buffers have the expected value + * after setting it to a known constant. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D8Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D8Async_UnEvenRowsCols") { + char *A_d; + int rows, cols; + rows = GENERATE(3, 4, 100); + cols = GENERATE(3, 4, 100); + size_t devPitch; + constexpr char memsetval = 'c'; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t size = rows * cols; + std::vectorA_h(size, 'a'); + std::vectorB_h(size, 'a'); + + HIP_CHECK(hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, sizeof(char) * cols, rows, + 4 * sizeof(char))); + HIP_CHECK(hipMemcpy2DAsync(A_d, devPitch, A_h.data(), sizeof(char) * cols, sizeof(char) * cols, rows, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemsetD2D8Async(A_d, devPitch, memsetval, sizeof(char) * cols, rows, stream)); + HIP_CHECK(hipMemcpy2DAsync(B_h.data(), sizeof(char) * cols, A_d, devPitch, sizeof(char) * cols, rows, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (int i = 0; i < rows; i++) { + for (int j = 0; j < cols; j++) { + INFO("Memset2D mismatch at index:" << i << " computed:" << B_h[i * cols + j] + << " memsetval:" << memsetval); + REQUIRE(B_h[i * cols + j] == memsetval); + } + } + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + * - Checks function behaviour when provided invalid arguments. + * Test source + * ------------------------ + * - catch/unit/memory/hipMemsetD2D8Async.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipMemsetD2D8Async_NegTsts") { + char* A_d; + constexpr size_t numH = 256; + constexpr size_t numW = 256; + size_t width = numW * sizeof(char); + size_t devPitch; + constexpr char memsetval = 'c'; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK( + hipMemAllocPitch(reinterpret_cast(&A_d), &devPitch, width, numH, 4 * sizeof(char))); + SECTION("nullptr destination") { + HIP_CHECK_ERROR(hipMemsetD2D8Async(nullptr, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("OutOfBound destination") { + void* outOfBoundsDst{reinterpret_cast(A_d) + devPitch * numH + 1}; + HIP_CHECK_ERROR(hipMemsetD2D8Async(outOfBoundsDst, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Dst pointer points to Source Memory") { + char* B_d; + std::unique_ptr hostPtr; + hostPtr.reset(new char[numH * width]); + B_d = hostPtr.get(); + HIP_CHECK_ERROR(hipMemsetD2D8Async(B_d, devPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Invalid Pitch") { + size_t inValidPitch = 1; + HIP_CHECK_ERROR(hipMemsetD2D8Async(A_d, inValidPitch, memsetval, numW, numH, stream), + hipErrorInvalidValue); + } + SECTION("Negative Values of Hight, Width") { + HIP_CHECK_ERROR(hipMemsetD2D8Async(A_d, devPitch, memsetval, numW, -10, stream), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemsetD2D8Async(A_d, devPitch, memsetval, -10, numH, stream), + hipErrorInvalidValue); + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipStreamDestroy(stream)); +} +/** + * End doxygen group MemoryTest. + * @} + */