AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Fix memsetD2XX tests (#1405)

Просмотр исходного кода 
* Fix memsetD2XX tests

* Remove redundant interpret_cast
Этот коммит содержится в:
Danylo Lytovchenko
2025-10-20 11:55:17 +02:00
коммит произвёл GitHub
родитель fd5ad25615
Коммит 61fc256db9
6 изменённых файлов: 92 добавлений и 92 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip-tests/catch/unit/memory/hipMemsetD2D16.cc
+15 -15
Просмотреть файл
@@ -50,20 +50,20 @@ TEST_CASE("Unit_hipMemsetD2D16_BasicFunctional") {
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,
hipDeviceptr_t A_d;
HIP_CHECK(hipMemAllocPitch(&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));
HIP_CHECK(hipMemcpy2D(A_h.data(), width, reinterpret_cast<void *>(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));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* Test Description
@@ -79,7 +79,7 @@ TEST_CASE("Unit_hipMemsetD2D16_BasicFunctional") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D16_UnEvenRowsCols") {
uint16_t *A_d;
hipDeviceptr_t A_d;
constexpr uint16_t memsetVal = 5;
int rows, cols;
rows = GENERATE(3, 4, 100);
@@ -89,11 +89,11 @@ TEST_CASE("Unit_hipMemsetD2D16_UnEvenRowsCols") {
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,
HIP_CHECK(hipMemAllocPitch(&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,
HIP_CHECK(hipMemcpy2D(B_h.data(), sizeof(uint16_t) * cols, reinterpret_cast<void*>(A_d), devPitch, sizeof(uint16_t) * cols, rows,
hipMemcpyDeviceToHost));
for (int i = 0; i < rows; i++) {
@@ -103,7 +103,7 @@ TEST_CASE("Unit_hipMemsetD2D16_UnEvenRowsCols") {
REQUIRE(B_h[i * cols + j] == memsetVal);
}
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* Test Description
@@ -117,27 +117,27 @@ TEST_CASE("Unit_hipMemsetD2D16_UnEvenRowsCols") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D16_NegTsts") {
uint16_t* A_d;
hipDeviceptr_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,
HIP_CHECK(hipMemAllocPitch(&A_d, &devPitch, width, numH,
2 * sizeof(uint16_t)));
SECTION("nullptr destination") {
HIP_CHECK_ERROR(hipMemsetD2D16(nullptr, devPitch, memsetval, numW, numH), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemsetD2D16(NULL, 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),
HIP_CHECK_ERROR(hipMemsetD2D16(reinterpret_cast<hipDeviceptr_t>(outOfBoundsDst), devPitch, memsetval, numW, numH),
hipErrorInvalidValue);
}
SECTION("Dst pointer points to Source Memory") {
uint16_t* B_d;
hipDeviceptr_t B_d;
std::unique_ptr<uint16_t[]> hostPtr;
hostPtr.reset(new uint16_t[numH * width]);
B_d = hostPtr.get();
B_d = reinterpret_cast<hipDeviceptr_t>(hostPtr.get());
HIP_CHECK_ERROR(hipMemsetD2D16(B_d, devPitch, memsetval, numW, numH), hipErrorInvalidValue);
}
SECTION("Invalid Pitch") {
@@ -148,7 +148,7 @@ TEST_CASE("Unit_hipMemsetD2D16_NegTsts") {
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));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* End doxygen group MemoryTest.
projects/hip-tests/catch/unit/memory/hipMemsetD2D16Async.cc
+15 -15
Просмотреть файл
@@ -50,15 +50,15 @@ TEST_CASE("Unit_hipMemsetD2D16Async_BasicFunctional") {
size_t width = numW * sizeof(uint16_t);
size_t sizeElements = numW * numH;
uint16_t *A_d;
hipDeviceptr_t A_d;
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, numH,
HIP_CHECK(hipMemAllocPitch(&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(hipMemcpy2DAsync(A_h.data(), width, reinterpret_cast<void *>(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) {
@@ -68,7 +68,7 @@ TEST_CASE("Unit_hipMemsetD2D16Async_BasicFunctional") {
}
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* Test Description
@@ -84,7 +84,7 @@ TEST_CASE("Unit_hipMemsetD2D16Async_BasicFunctional") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D16Async_UnEvenRowsCols") {
uint16_t *A_d;
hipDeviceptr_t A_d;
int rows, cols;
constexpr int memsetval = 5;
rows = GENERATE(3, 4, 100);
@@ -95,11 +95,11 @@ TEST_CASE("Unit_hipMemsetD2D16Async_UnEvenRowsCols") {
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,
HIP_CHECK(hipMemAllocPitch(&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,
HIP_CHECK(hipMemcpy2DAsync(B_h.data(), sizeof(uint16_t) * cols, reinterpret_cast<void *>(A_d), devPitch, sizeof(uint16_t) * cols,
rows, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (int i = 0; i < rows; i++) {
@@ -110,7 +110,7 @@ TEST_CASE("Unit_hipMemsetD2D16Async_UnEvenRowsCols") {
}
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* Test Description
@@ -124,7 +124,7 @@ TEST_CASE("Unit_hipMemsetD2D16Async_UnEvenRowsCols") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D16Async_NegTsts") {
uint16_t* A_d;
hipDeviceptr_t A_d;
constexpr size_t numH = 256;
constexpr size_t numW = 256;
size_t width = numW * sizeof(uint16_t);
@@ -132,22 +132,22 @@ TEST_CASE("Unit_hipMemsetD2D16Async_NegTsts") {
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,
HIP_CHECK(hipMemAllocPitch(&A_d, &devPitch, width, numH,
2 * sizeof(uint16_t)));
SECTION("nullptr destination") {
HIP_CHECK_ERROR(hipMemsetD2D16Async(nullptr, devPitch, memsetval, numW, numH, stream),
HIP_CHECK_ERROR(hipMemsetD2D16Async(NULL, 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),
HIP_CHECK_ERROR(hipMemsetD2D16Async(reinterpret_cast<hipDeviceptr_t>(outOfBoundsDst), devPitch, memsetval, numW, numH, stream),
hipErrorInvalidValue);
}
SECTION("Dst pointer points to Source Memory") {
uint16_t* B_d;
hipDeviceptr_t B_d;
std::unique_ptr<uint16_t[]> hostPtr;
hostPtr.reset(new uint16_t[numH * width]);
B_d = hostPtr.get();
B_d = reinterpret_cast<hipDeviceptr_t>(hostPtr.get());
HIP_CHECK_ERROR(hipMemsetD2D16Async(B_d, devPitch, memsetval, numW, numH, stream),
hipErrorInvalidValue);
}
@@ -162,7 +162,7 @@ TEST_CASE("Unit_hipMemsetD2D16Async_NegTsts") {
HIP_CHECK_ERROR(hipMemsetD2D16Async(A_d, devPitch, memsetval, -10, numH, stream),
hipErrorInvalidValue);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
HIP_CHECK(hipStreamDestroy(stream));
}
/**
projects/hip-tests/catch/unit/memory/hipMemsetD2D32.cc
+15 -15
Просмотреть файл
@@ -49,20 +49,20 @@ TEST_CASE("Unit_hipMemsetD2D32_BasicFunctional") {
size_t pitch_A;
size_t width = numW * sizeof(int);
size_t sizeElements = numW * numH;
int *A_d;
hipDeviceptr_t 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(hipMemAllocPitch(&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));
HIP_CHECK(hipMemcpy2D(A_h.data(), width, reinterpret_cast<void *>(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));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* Test Description
@@ -79,7 +79,7 @@ TEST_CASE("Unit_hipMemsetD2D32_BasicFunctional") {
*/
TEST_CASE("Unit_hipMemsetD2D32_UnEvenRowsCols") {
constexpr int memsetval = 5;
int *A_d;
hipDeviceptr_t A_d;
int rows, cols;
rows = GENERATE(3, 4, 100);
cols = GENERATE(3, 4, 100);
@@ -87,11 +87,11 @@ TEST_CASE("Unit_hipMemsetD2D32_UnEvenRowsCols") {
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,
HIP_CHECK(hipMemAllocPitch(&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,
HIP_CHECK(hipMemcpy2D(B_h.data(), sizeof(int) * cols, reinterpret_cast<void *>(A_d), devPitch, sizeof(int) * cols, rows,
hipMemcpyDeviceToHost));
for (int i = 0; i < rows; i++) {
@@ -101,7 +101,7 @@ TEST_CASE("Unit_hipMemsetD2D32_UnEvenRowsCols") {
REQUIRE(B_h[i * cols + j] == memsetval);
}
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* Test Description
@@ -115,26 +115,26 @@ TEST_CASE("Unit_hipMemsetD2D32_UnEvenRowsCols") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D32_NegTsts") {
int* A_d;
hipDeviceptr_t 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)));
HIP_CHECK(hipMemAllocPitch(&A_d, &devPitch, width, numH, sizeof(int)));
SECTION("nullptr destination") {
HIP_CHECK_ERROR(hipMemsetD2D32(nullptr, devPitch, memsetval, numW, numH), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemsetD2D32(NULL, 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),
HIP_CHECK_ERROR(hipMemsetD2D32(reinterpret_cast<hipDeviceptr_t>(outOfBoundsDst), devPitch, memsetval, numW, numH),
hipErrorInvalidValue);
}
SECTION("Dst pointer points to Source Memory") {
int* B_d;
hipDeviceptr_t B_d;
std::unique_ptr<int[]> hostPtr;
hostPtr.reset(new int[numH * width]);
B_d = hostPtr.get();
B_d = reinterpret_cast<hipDeviceptr_t>(hostPtr.get());
HIP_CHECK_ERROR(hipMemsetD2D32(B_d, devPitch, memsetval, numW, numH), hipErrorInvalidValue);
}
SECTION("Invalid Pitch") {
@@ -145,7 +145,7 @@ TEST_CASE("Unit_hipMemsetD2D32_NegTsts") {
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));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* End doxygen group MemoryTest.
projects/hip-tests/catch/unit/memory/hipMemsetD2D32Async.cc
+15 -15
Просмотреть файл
@@ -49,14 +49,14 @@ TEST_CASE("Unit_hipMemsetD2D32Async_BasicFunctional") {
size_t pitch_A;
size_t width = numW * sizeof(int);
size_t sizeElements = numW * numH;
int *A_d;
hipDeviceptr_t A_d;
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemAllocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, numH, sizeof(int)));
HIP_CHECK(hipMemAllocPitch(&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(hipMemcpy2DAsync(A_h.data(), width, reinterpret_cast<void*>(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]
@@ -64,7 +64,7 @@ TEST_CASE("Unit_hipMemsetD2D32Async_BasicFunctional") {
REQUIRE(A_h[i] == memsetval);
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* Test Description
@@ -81,7 +81,7 @@ TEST_CASE("Unit_hipMemsetD2D32Async_BasicFunctional") {
*/
TEST_CASE("Unit_hipMemsetD2D32Async_UnEvenRowsCols") {
constexpr int memsetval = 5;
int *A_d;
hipDeviceptr_t A_d;
int rows, cols;
rows = GENERATE(3, 4, 100);
cols = GENERATE(3, 4, 100);
@@ -92,11 +92,11 @@ TEST_CASE("Unit_hipMemsetD2D32Async_UnEvenRowsCols") {
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,
HIP_CHECK(hipMemAllocPitch(&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,
HIP_CHECK(hipMemcpy2DAsync(B_h.data(), sizeof(int) * cols, reinterpret_cast<void*>(A_d), devPitch, sizeof(int) * cols, rows,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (int i = 0; i < rows; i++) {
@@ -107,7 +107,7 @@ TEST_CASE("Unit_hipMemsetD2D32Async_UnEvenRowsCols") {
}
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* Test Description
@@ -121,7 +121,7 @@ TEST_CASE("Unit_hipMemsetD2D32Async_UnEvenRowsCols") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D32Async_NegTsts") {
int* A_d;
hipDeviceptr_t A_d;
constexpr size_t numH = 256;
constexpr size_t numW = 256;
size_t width = numW * sizeof(int);
@@ -129,21 +129,21 @@ TEST_CASE("Unit_hipMemsetD2D32Async_NegTsts") {
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)));
HIP_CHECK(hipMemAllocPitch(&A_d, &devPitch, width, numH, sizeof(int)));
SECTION("nullptr destination") {
HIP_CHECK_ERROR(hipMemsetD2D32Async(nullptr, devPitch, memsetval, numW, numH, stream),
HIP_CHECK_ERROR(hipMemsetD2D32Async(NULL, 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),
HIP_CHECK_ERROR(hipMemsetD2D32Async(reinterpret_cast<hipDeviceptr_t>(outOfBoundsDst), devPitch, memsetval, numW, numH, stream),
hipErrorInvalidValue);
}
SECTION("Dst pointer points to Source Memory") {
int* B_d;
hipDeviceptr_t B_d;
std::unique_ptr<int[]> hostPtr;
hostPtr.reset(new int[numH * width]);
B_d = hostPtr.get();
B_d = reinterpret_cast<hipDeviceptr_t>(hostPtr.get());
HIP_CHECK_ERROR(hipMemsetD2D32Async(B_d, devPitch, memsetval, numW, numH, stream),
hipErrorInvalidValue);
}
@@ -158,7 +158,7 @@ TEST_CASE("Unit_hipMemsetD2D32Async_NegTsts") {
HIP_CHECK_ERROR(hipMemsetD2D32Async(A_d, devPitch, memsetval, -10, numH, stream),
hipErrorInvalidValue);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
HIP_CHECK(hipStreamDestroy(stream));
}
/**
projects/hip-tests/catch/unit/memory/hipMemsetD2D8.cc
+16 -16
Просмотреть файл
@@ -48,19 +48,19 @@ TEST_CASE("Unit_hipMemsetD2D8_BasicFunctional") {
size_t pitch_A;
size_t width = numW * sizeof(char);
size_t sizeElements = numW * numH;
char *A_d;
hipDeviceptr_t A_d;
HIP_CHECK(
hipMemAllocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width, numH, 4 * sizeof(char)));
hipMemAllocPitch(&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));
HIP_CHECK(hipMemcpy2D(A_h.data(), width, reinterpret_cast<void*>(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));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* Test Description
@@ -76,7 +76,7 @@ TEST_CASE("Unit_hipMemsetD2D8_BasicFunctional") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D8_UnEvenRowsCols") {
char *A_d;
hipDeviceptr_t A_d;
int rows, cols;
rows = GENERATE(3, 4, 100);
cols = GENERATE(4, 5, 100);
@@ -85,14 +85,14 @@ TEST_CASE("Unit_hipMemsetD2D8_UnEvenRowsCols") {
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,
HIP_CHECK(hipMemAllocPitch(&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,
HIP_CHECK(hipMemcpy2D(reinterpret_cast<void *>(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,
HIP_CHECK(hipMemcpy2D(B_h.data(), sizeof(char) * cols, reinterpret_cast<void *>(A_d), devPitch, sizeof(char) * cols, rows,
hipMemcpyDeviceToHost));
for (int i = 0; i < rows; i++) {
@@ -102,7 +102,7 @@ TEST_CASE("Unit_hipMemsetD2D8_UnEvenRowsCols") {
REQUIRE(B_h[i * cols + j] == memsetval);
}
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* Test Description
@@ -116,27 +116,27 @@ TEST_CASE("Unit_hipMemsetD2D8_UnEvenRowsCols") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D8_NegTsts") {
char* A_d;
hipDeviceptr_t 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)));
hipMemAllocPitch(&A_d, &devPitch, width, numH, 4 * sizeof(char)));
SECTION("nullptr destination") {
HIP_CHECK_ERROR(hipMemsetD2D8(nullptr, devPitch, memsetval, numW, numH), hipErrorInvalidValue);
HIP_CHECK_ERROR(hipMemsetD2D8(NULL, 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),
HIP_CHECK_ERROR(hipMemsetD2D8(reinterpret_cast<hipDeviceptr_t>( outOfBoundsDst ), devPitch, memsetval, numW, numH),
hipErrorInvalidValue);
}
SECTION("Dst pointer points to Source Memory") {
char* B_d;
hipDeviceptr_t B_d;
std::unique_ptr<char[]> hostPtr;
hostPtr.reset(new char[numH * width]);
B_d = hostPtr.get();
B_d = reinterpret_cast<hipDeviceptr_t>( hostPtr.get() );
HIP_CHECK_ERROR(hipMemsetD2D8(B_d, devPitch, memsetval, numW, numH), hipErrorInvalidValue);
}
SECTION("Invalid Pitch") {
@@ -147,7 +147,7 @@ TEST_CASE("Unit_hipMemsetD2D8_NegTsts") {
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));
HIP_CHECK(hipFree(reinterpret_cast<void*>( A_d )));
}
/**
* End doxygen group MemoryTest.
projects/hip-tests/catch/unit/memory/hipMemsetD2D8Async.cc
+16 -16
Просмотреть файл
@@ -51,15 +51,15 @@ TEST_CASE("Unit_hipMemsetD2D8Async_BasicFunctional") {
size_t width = numW * sizeof(char);
size_t sizeElements = numW * numH;
char *A_d;
hipDeviceptr_t 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)));
hipMemAllocPitch(&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(hipMemcpy2DAsync(A_h.data(), width, reinterpret_cast<void *>( 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]
@@ -67,7 +67,7 @@ TEST_CASE("Unit_hipMemsetD2D8Async_BasicFunctional") {
REQUIRE(A_h[i] == memsetval);
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
}
/**
* Test Description
@@ -83,7 +83,7 @@ TEST_CASE("Unit_hipMemsetD2D8Async_BasicFunctional") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D8Async_UnEvenRowsCols") {
char *A_d;
hipDeviceptr_t A_d;
int rows, cols;
rows = GENERATE(3, 4, 100);
cols = GENERATE(3, 4, 100);
@@ -96,12 +96,12 @@ TEST_CASE("Unit_hipMemsetD2D8Async_UnEvenRowsCols") {
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,
HIP_CHECK(hipMemAllocPitch(&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,
HIP_CHECK(hipMemcpy2DAsync(reinterpret_cast<void*>( 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,
HIP_CHECK(hipMemcpy2DAsync(B_h.data(), sizeof(char) * cols, reinterpret_cast<void*>(A_d), devPitch, sizeof(char) * cols, rows,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (int i = 0; i < rows; i++) {
@@ -112,7 +112,7 @@ TEST_CASE("Unit_hipMemsetD2D8Async_UnEvenRowsCols") {
}
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void *>(A_d)));
}
/**
* Test Description
@@ -126,7 +126,7 @@ TEST_CASE("Unit_hipMemsetD2D8Async_UnEvenRowsCols") {
* - HIP_VERSION >= 7.1
*/
TEST_CASE("Unit_hipMemsetD2D8Async_NegTsts") {
char* A_d;
hipDeviceptr_t A_d;
constexpr size_t numH = 256;
constexpr size_t numW = 256;
size_t width = numW * sizeof(char);
@@ -135,21 +135,21 @@ TEST_CASE("Unit_hipMemsetD2D8Async_NegTsts") {
hipStream_t stream = nullptr;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(
hipMemAllocPitch(reinterpret_cast<void**>(&A_d), &devPitch, width, numH, 4 * sizeof(char)));
hipMemAllocPitch(&A_d, &devPitch, width, numH, 4 * sizeof(char)));
SECTION("nullptr destination") {
HIP_CHECK_ERROR(hipMemsetD2D8Async(nullptr, devPitch, memsetval, numW, numH, stream),
HIP_CHECK_ERROR(hipMemsetD2D8Async(NULL, 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),
HIP_CHECK_ERROR(hipMemsetD2D8Async(reinterpret_cast<hipDeviceptr_t>(outOfBoundsDst), devPitch, memsetval, numW, numH, stream),
hipErrorInvalidValue);
}
SECTION("Dst pointer points to Source Memory") {
char* B_d;
hipDeviceptr_t B_d;
std::unique_ptr<char[]> hostPtr;
hostPtr.reset(new char[numH * width]);
B_d = hostPtr.get();
B_d = reinterpret_cast<hipDeviceptr_t>(hostPtr.get());
HIP_CHECK_ERROR(hipMemsetD2D8Async(B_d, devPitch, memsetval, numW, numH, stream),
hipErrorInvalidValue);
}
@@ -164,7 +164,7 @@ TEST_CASE("Unit_hipMemsetD2D8Async_NegTsts") {
HIP_CHECK_ERROR(hipMemsetD2D8Async(A_d, devPitch, memsetval, -10, numH, stream),
hipErrorInvalidValue);
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
HIP_CHECK(hipStreamDestroy(stream));
}
/**