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.
Этот коммит содержится в:
amd-srinivas1
коммит произвёл
GitHub
GitHub
родитель
cc18890fe8
Коммит
092279449e
@@ -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)
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
/**
|
||||
* @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<uint16_t>(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<void**>(&A_d), &pitch_A, width, numH,
|
||||
2 * sizeof(uint16_t)));
|
||||
std::vector<uint16_t>A_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::vector<uint16_t>B_h(size, 1);
|
||||
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<uint16_t>(0x26);
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<uint16_t*>(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<uint16_t[]> 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.
|
||||
* @}
|
||||
*/
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
/**
|
||||
* @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<void**>(&A_d), &pitch_A, width, numH,
|
||||
2 * sizeof(uint16_t)));
|
||||
std::vector<uint16_t>A_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::vector<uint16_t>B_h(size, 1);
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<uint16_t>(0x26);
|
||||
hipStream_t stream = nullptr;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<uint16_t*>(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<uint16_t[]> 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.
|
||||
* @}
|
||||
*/
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
/**
|
||||
* @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::vector<int>A_h(sizeElements, 1);
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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::vector<int>B_h(size, 1);
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<int>(0x26);
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<int*>(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<int[]> 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.
|
||||
* @}
|
||||
*/
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_kernels.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
/**
|
||||
* @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<void**>(&A_d), &pitch_A, width, numH, sizeof(int)));
|
||||
std::vector<int>A_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::vector<int>B_h(size, 1);
|
||||
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<void**>(&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<int*>(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<int[]> 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.
|
||||
* @}
|
||||
*/
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
/**
|
||||
* @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<void**>(&A_d), &pitch_A, width, numH, 4 * sizeof(char)));
|
||||
std::vector<char>A_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::vector<char>A_h(size, 'a');
|
||||
std::vector<char>B_h(size, 'a');
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<void**>(&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<char*>(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<char[]> 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.
|
||||
* @}
|
||||
*/
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_checkers.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
|
||||
/**
|
||||
* @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<void**>(&A_d), &pitch_A, width, numH, 4 * sizeof(char)));
|
||||
std::vector<char>A_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::vector<char>A_h(size, 'a');
|
||||
std::vector<char>B_h(size, 'a');
|
||||
|
||||
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&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<void**>(&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<char*>(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<char[]> 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.
|
||||
* @}
|
||||
*/
|
||||
Ссылка в новой задаче
Block a user