AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

SWDEV-292637 - [dtest] Catch2 unit and multiprocess tests for Memset3d,HostMalloc and MallocConcurrency tests (#2348)

Kaynağa Gözat 
Change-Id: I9025bc13735c1d9fb0f0811a9c9d6ad304adc134
Bu işleme şunda yer alıyor:
sumanthtg
2021-09-17 12:11:26 +05:30
işlemeyi yapan: GitHub
ebeveyn e3996d3b92
işleme d818f6920d
15 değiştirilmiş dosya ile 2159 ekleme ve 173 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
catch/unit/memory/CMakeLists.txt
+10
Dosyayı Görüntüle
@@ -40,6 +40,11 @@ set(TEST_SRC
hipMemset3D.cc
hipMemset2D.cc
hipMemset2DAsyncMultiThreadAndKernel.cc
hipHostMallocTests.cc
hipMallocConcurrency.cc
hipMemset3DFunctional.cc
hipMemset3DNegative.cc
hipMemset3DRegressMultiThread.cc
)
else()
set(TEST_SRC
@@ -80,6 +85,11 @@ set(TEST_SRC
hipMemset3D.cc
hipMemset2D.cc
hipMemset2DAsyncMultiThreadAndKernel.cc
hipHostMallocTests.cc
hipMallocConcurrency.cc
hipMemset3DFunctional.cc
hipMemset3DNegative.cc
hipMemset3DRegressMultiThread.cc
)
endif()
# Create shared lib of all tests
catch/unit/memory/hipHostMallocTests.cc
+61
Dosyayı Görüntüle
@@ -0,0 +1,61 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Test hipHostMalloc() api with ptr as nullptr and check for return value.
2) Test hipHostMalloc() api with size as max(size_t) and check for OOM error.
3) Test hipHostMalloc() api with flags as max(unsigned int) and validate
return value.
4) Pass size as zero for hipHostMalloc() api and check ptr is reset with
with return value success.
*/
#include <hip_test_common.hh>
/**
* Performs argument validation of hipHostMalloc api.
*/
TEST_CASE("Unit_hipHostMalloc_ArgValidation") {
hipError_t ret;
constexpr size_t allocSize = 1000;
char *ptr;
SECTION("Pass ptr as nullptr") {
ret = hipHostMalloc(static_cast<void **>(nullptr), allocSize);
REQUIRE(ret != hipSuccess);
}
SECTION("Size as max(size_t)") {
ret = hipHostMalloc(&ptr, std::numeric_limits<std::size_t>::max());
REQUIRE(ret != hipSuccess);
}
SECTION("Flags as max(uint)") {
ret = hipHostMalloc(&ptr, allocSize,
std::numeric_limits<unsigned int>::max());
REQUIRE(ret != hipSuccess);
}
SECTION("Pass size as zero and check ptr reset") {
HIP_CHECK(hipHostMalloc(&ptr, 0));
REQUIRE(ptr == nullptr);
}
}
catch/unit/memory/hipMallocConcurrency.cc
+412
Dosyayı Görüntüle
@@ -0,0 +1,412 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Test hipMalloc() api passing zero size and confirming *ptr returning
nullptr. Also pass nullptr to hipFree() api.
2) Pass maximum value of size_t for hipMalloc() api and make sure appropriate
error is returned.
3) Check for hipMalloc() error code, passing invalid/null pointer.
4) Regress hipMalloc()/hipFree() in loop for bigger chunk of allocation
with adequate number of iterations and later test for kernel execution on
default gpu.
5) Regress hipMalloc()/hipFree() in loop while allocating smaller chunks
keeping maximum number of iterations and then run kernel code on default
gpu, perfom data validation.
6) Check hipMalloc() api adaptability when app creates small chunks of memory
continuously, stores it for later use and then frees it at later point
of time.
7) Multithread Scenario : Exercise hipMalloc() api parellely on all gpus from
multiple threads and regress the api.
8) Validate memory usage with hipMemGetInfo() while regressing hipMalloc()
api. Check for any possible memory leaks.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#include <vector>
#include <limits>
#include <atomic>
/* Buffer size for bigger chunks in alloc/free cycles */
static constexpr auto BuffSizeBC = 5*1024*1024;
/* Buffer size for smaller chunks in alloc/free cycles */
static constexpr auto BuffSizeSC = 16;
/* You may change it for individual test.
* But default 100 is for quick return in Jenkin Build */
static constexpr auto NumDiv = 100;
/* Max alloc/free iterations for smaller chunks */
static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv);
/* Max alloc/free iterations for bigger chunks */
static constexpr auto MaxAllocFree_BigChunks = 10000;
/* Max alloc and pool iterations */
static constexpr auto MaxAllocPoolIter = (2000000/NumDiv);
/* Test status shared across threads */
static std::atomic<bool> g_thTestPassed{true};
/**
* Validates data consistency on supplied gpu
*/
static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t prevAvl, prevTot, curAvl, curTot;
bool TestPassed = true;
constexpr auto N = 4 * 1024 * 1024;
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
size_t Nbytes = N * sizeof(int);
HIP_CHECK(hipSetDevice(gpu));
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
UNSCOPED_INFO("Validation PASSED for gpu " << gpu);
} else {
UNSCOPED_INFO("Validation FAILED for gpu " << gpu);
TestPassed = false;
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
// In concurrent calls on one GPU, we cannot verify leaking in this way
UNSCOPED_INFO(
"validateMemoryOnGPU : Memory allocation mismatch observed."
<< "Possible memory leak.");
TestPassed = false;
}
return TestPassed;
}
/**
* Regress memory allocation and free in loop
*/
static bool regressAllocInLoop(int gpu) {
bool TestPassed = true;
size_t tot, avail, ptot, pavail, numBytes;
int i = 0;
int *ptr;
HIP_CHECK(hipSetDevice(gpu));
numBytes = BuffSizeBC;
// Exercise allocation in loop with bigger chunks
for (i = 0; i < MaxAllocFree_BigChunks; i++) {
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
HIP_CHECK(hipMalloc(&ptr, numBytes));
HIP_CHECK(hipMemGetInfo(&avail, &tot));
HIP_CHECK(hipFree(ptr));
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
tot << ", pavail-avail=" << pavail-avail);
TestPassed = false;
break;
}
}
// Exercise allocation in loop with smaller chunks and maximum iters
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
numBytes = BuffSizeSC;
for (i = 0; i < MaxAllocFree_SmallChunks; i++) {
HIP_CHECK(hipMalloc(&ptr, numBytes));
HIP_CHECK(hipFree(ptr));
}
HIP_CHECK(hipMemGetInfo(&avail, &tot));
if ((pavail != avail) || (ptot != tot)) {
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
"Possible memory leak.");
TestPassed &= false;
}
return TestPassed;
}
/**
* Validates data consistency on supplied gpu
* In Multithreaded Environment
*/
static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t prevAvl, prevTot, curAvl, curTot;
bool TestPassed = true;
constexpr auto N = 4 * 1024 * 1024;
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
size_t Nbytes = N * sizeof(int);
HIPCHECK(hipSetDevice(gpu));
HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
UNSCOPED_INFO("Validation PASSED for gpu " << gpu);
} else {
UNSCOPED_INFO("Validation FAILED for gpu " << gpu);
TestPassed = false;
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
// In concurrent calls on one GPU, we cannot verify leaking in this way
UNSCOPED_INFO(
"validateMemoryOnGpuMThread : Memory allocation mismatch observed."
"Possible memory leak.");
TestPassed = false;
}
return TestPassed;
}
/**
* Regress memory allocation and free in loop
* In Multithreaded Environment
*/
static bool regressAllocInLoopMthread(int gpu) {
bool TestPassed = true;
size_t tot, avail, ptot, pavail, numBytes;
int i = 0;
int *ptr;
HIPCHECK(hipSetDevice(gpu));
numBytes = BuffSizeBC;
// Exercise allocation in loop with bigger chunks
for (i = 0; i < MaxAllocFree_BigChunks; i++) {
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
HIPCHECK(hipMalloc(&ptr, numBytes));
HIPCHECK(hipMemGetInfo(&avail, &tot));
HIPCHECK(hipFree(ptr));
if (pavail-avail < numBytes) { // We expect pavail-avail >= numBytes
UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " <<
numBytes << " not matching with hipMemGetInfo - FAIL." << "pavail=" <<
pavail << ", ptot=" << ptot << ", avail=" << avail << ", tot=" <<
tot << ", pavail-avail=" << pavail-avail);
TestPassed = false;
break;
}
}
// Exercise allocation in loop with smaller chunks and maximum iters
HIPCHECK(hipMemGetInfo(&pavail, &ptot));
numBytes = BuffSizeSC;
for (i = 0; i < MaxAllocFree_SmallChunks; i++) {
HIPCHECK(hipMalloc(&ptr, numBytes));
HIPCHECK(hipFree(ptr));
}
HIPCHECK(hipMemGetInfo(&avail, &tot));
if ((pavail != avail) || (ptot != tot)) {
UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." <<
"Possible memory leak.");
TestPassed &= false;
}
return TestPassed;
}
/*
* Thread func to regress alloc and check data consistency
*/
static void threadFunc(int gpu) {
g_thTestPassed = regressAllocInLoopMthread(gpu);
g_thTestPassed = g_thTestPassed & validateMemoryOnGpuMThread(gpu);
UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" <<
g_thTestPassed.load());
}
/* Performs Argument Validation of api */
TEST_CASE("Unit_hipMalloc_ArgumentValidation") {
int *ptr;
hipError_t ret;
SECTION("hipMalloc() when size(0)") {
HIP_CHECK(hipMalloc(&ptr, 0));
// ptr expected to be reset to null ptr
REQUIRE(ptr == nullptr);
}
SECTION("hipFree() when freeing nullptr ") {
ptr = nullptr;
// api should return success and shudnt crash
HIP_CHECK(hipFree(ptr));
}
SECTION("hipMalloc() with invalid argument") {
constexpr auto sizeBytes = 100;
ret = hipMalloc(nullptr, sizeBytes);
REQUIRE(ret != hipSuccess);
}
SECTION("hipMalloc() with max size_t") {
ret = hipMalloc(&ptr, std::numeric_limits<std::size_t>::max());
REQUIRE(ret != hipSuccess);
}
}
/**
* Regress hipMalloc()/hipFree() in loop for bigger chunks and
* smaller chunks of memory allocation
*/
TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") {
int devCnt = 0;
// Get GPU count
HIP_CHECK(hipGetDeviceCount(&devCnt));
REQUIRE(devCnt > 0);
CHECK(regressAllocInLoop(0) == true);
CHECK(validateMemoryOnGPU(0) == true);
}
/**
* Application Behavior Modelling.
* Check hipMalloc() api adaptability when app creates small chunks of memory
* continuously, stores it for later use and then frees it at later point
* of time.
*/
TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") {
size_t avail, tot, pavail, ptot;
bool ret;
hipError_t err;
std::vector<int *> ptrlist;
constexpr auto BuffSize = 10;
int devCnt, *ptr;
// Get GPU count
HIP_CHECK(hipGetDeviceCount(&devCnt));
REQUIRE(devCnt > 0);
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
// Allocate small chunks of memory million times
for (int i = 0; i < MaxAllocPoolIter ; i++) {
if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) {
HIP_CHECK(hipMemGetInfo(&avail, &tot));
INFO("Loop regression pool allocation failure. " <<
"Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " <<
avail/(1024.0*1024.0) << " iter " << i << " error "
<< hipGetErrorString(err));
REQUIRE(false);
}
// Store pointers allocated to emulate memory pool of app
ptrlist.push_back(ptr);
}
// Free ptrs at later point of time
for ( auto &t : ptrlist ) {
HIP_CHECK(hipFree(t));
}
HIP_CHECK(hipMemGetInfo(&avail, &tot));
ret = validateMemoryOnGPU(0);
REQUIRE(ret == true);
REQUIRE(pavail == avail);
REQUIRE(ptot == tot);
}
/**
* Exercise hipMalloc() api parellely on all gpus from
* multiple threads and regress the api.
*/
TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") {
std::vector<std::thread> threadlist;
int devCnt;
// Get GPU count
HIP_CHECK(hipGetDeviceCount(&devCnt));
REQUIRE(devCnt > 0);
for (int i = 0; i < devCnt; i++) {
threadlist.push_back(std::thread(threadFunc, i));
}
for (auto &t : threadlist) {
t.join();
}
REQUIRE(g_thTestPassed == true);
}
catch/unit/memory/hipManagedKeyword.cc
-75
Dosyayı Görüntüle
@@ -1,75 +0,0 @@
/*
Copyright (c) 2021 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.
*/
/*
This testcase verifies the hipManagedKeyword basic scenario
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#define N 1048576
__managed__ float A[N]; // Accessible by ALL CPU and GPU functions !!!
__managed__ float B[N];
__managed__ int x = 0;
__global__ void add(const float *A, float *B) {
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < N; i += stride)
B[i] = A[i] + B[i];
}
__global__ void GPU_func() {
x++;
}
TEST_CASE("Unit_hipManagedKeyword_SingleGpu") {
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
}
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
dim3 dimGrid(numBlocks, 1, 1);
dim3 dimBlock(blockSize, 1, 1);
hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, static_cast<const float*>(A),
static_cast<float*>(B));
hipDeviceSynchronize();
float maxError = 0.0f;
for (int i = 0; i < N; i++)
maxError = fmax(maxError, fabs(B[i]-3.0f));
REQUIRE(maxError == 0.0f);
}
TEST_CASE("Unit_hipManagedKeyword_MultiGpu") {
int numDevices = 0;
hipGetDeviceCount(&numDevices);
for (int i = 0; i < numDevices; i++) {
hipSetDevice(i);
GPU_func<<< 1, 1 >>>();
hipDeviceSynchronize();
}
REQUIRE(x == numDevices);
}
catch/unit/memory/hipMemset3DFunctional.cc
+515
Dosyayı Görüntüle
@@ -0,0 +1,515 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Passing width as 0 in extent, verify hipMemset3D api returns success and
doesn't modify the buffer passed.
2) Passing width as 0 in extent, verify hipMemset3DAsync api returns success
and doesn't modify the buffer passed.
3) Passing height as 0 in extent, verify hipMemset3D api returns success and
doesn't modify the buffer passed.
4) Passing height as 0 in extent, verify hipMemset3DAsync api returns success
and doesn't modify the buffer passed.
5) Passing depth as 0 in extent, verify hipMemset3D api returns success and
doesn't modify the buffer passed.
6) Passing depth as 0 in extent, verify hipMemset3DAsync api returns success
and doesn't modify the buffer passed.
7) When extent passed with width, height and depth all as zeroes, verify
hipMemset3D api returns success and doesn't modify the buffer passed.
8) When extent passed with width, height and depth all as zeroes, verify
hipMemset3DAsync api returns success and doesn't modify the buffer passed.
9) Validate data after performing memory set operation with max memset value
for hipMemset3D api.
10) Validate data after performing memory set operation with max memset value
for hipMemset3DAsync api.
11) Select random slice of 3d array and Memset complete slice with
hipMemset3D api.
12) Select random slice of 3d array and Memset complete slice with
hipMemset3DAsync api.
13) Seek device pitched ptr to desired portion of 3d array and memset the
portion with hipMemset3D api.
14) Seek device pitched ptr to desired portion of 3d array and memset the
portion with hipMemset3DAsync api.
*/
#include <hip_test_common.hh>
/*
* Defines
*/
#define MEMSETVAL 1
#define TESTVAL 2
#define NUMH_EXT 256
#define NUMW_EXT 100
#define DEPTH_EXT 10
#define NUMH_MAX 256
#define NUMW_MAX 256
#define DEPTH_MAX 10
#define ZSIZE_S 32
#define YSIZE_S 32
#define XSIZE_S 32
#define ZSIZE_P 30
#define YSIZE_P 30
#define XSIZE_P 30
#define ZPOS_START 10
#define ZSET_LEN 10
#define ZPOS_END 19
#define YPOS_START 10
#define YSET_LEN 10
#define YPOS_END 19
#define XPOS_START 10
#define XSET_LEN 10
#define XPOS_END 19
/**
* Memset with extent passed and verify data to be intact
*/
static void testMemsetWithExtent(bool bAsync, hipExtent tstExtent) {
hipPitchedPtr devPitchedPtr;
hipError_t ret;
char *A_h;
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
size_t sizeElements = width * numH * depth;
size_t elements = numW* numH* depth;
A_h = reinterpret_cast<char *>(malloc(sizeElements));
REQUIRE(A_h != nullptr);
memset(A_h, 0, sizeElements);
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
if (bAsync) {
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
ret = hipMemset3DAsync(devPitchedPtr, MEMSETVAL, extent, stream);
INFO("testMemsetWithExtent(" << extent.width << "," << extent.height
<< "," << extent.depth << ") memset "
<< MEMSETVAL << ", ret : " << ret);
REQUIRE(ret == hipSuccess);
ret = hipMemset3DAsync(devPitchedPtr, TESTVAL, tstExtent, stream);
INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height
<< "," << tstExtent.depth << ") memset "
<< TESTVAL << "ret : " << ret);
REQUIRE(ret == hipSuccess);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
} else {
ret = hipMemset3D(devPitchedPtr, MEMSETVAL, extent);
INFO("testMemsetWithExtent(" << extent.width << "," << extent.height
<< "," << extent.depth << ") memset "
<< MEMSETVAL << ",ret : " << ret);
REQUIRE(ret == hipSuccess);
ret = hipMemset3D(devPitchedPtr, TESTVAL, tstExtent);
INFO("testMemsetWithExtent(" << tstExtent.width << "," << tstExtent.height
<< "," << tstExtent.depth << ") memset "
<< TESTVAL << ",ret : " << ret);
REQUIRE(ret == hipSuccess);
}
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
myparms.srcPtr = devPitchedPtr;
myparms.extent = extent;
#if HT_NVIDIA
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
#else
myparms.kind = hipMemcpyDeviceToHost;
#endif
HIP_CHECK(hipMemcpy3D(&myparms));
for (size_t i = 0; i < elements; i++) {
if (A_h[i] != MEMSETVAL) {
INFO("testMemsetWithExtent: index:" << i << ",computed:"
<< std::hex << static_cast<int>(A_h[i]) << ",memsetval:"
<< std::hex << MEMSETVAL);
REQUIRE(false);
}
}
HIP_CHECK(hipFree(devPitchedPtr.ptr));
free(A_h);
}
/**
* Validates data after performing memory set operation with max memset value
*/
static void testMemsetMaxValue(bool bAsync) {
hipPitchedPtr devPitchedPtr;
unsigned char *A_h;
int memsetval = std::numeric_limits<unsigned char>::max();
size_t numH = NUMH_MAX, numW = NUMW_MAX, depth = DEPTH_MAX;
size_t width = numW * sizeof(unsigned char);
hipExtent extent = make_hipExtent(width, numH, depth);
size_t sizeElements = width * numH * depth;
size_t elements = numW* numH* depth;
A_h = reinterpret_cast<unsigned char *> (malloc(sizeElements));
REQUIRE(A_h != nullptr);
memset(A_h, 0, sizeElements);
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
if (bAsync) {
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
} else {
HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
}
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
myparms.srcPtr = devPitchedPtr;
myparms.extent = extent;
#if HT_NVIDIA
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
#else
myparms.kind = hipMemcpyDeviceToHost;
#endif
HIP_CHECK(hipMemcpy3D(&myparms));
for (size_t i = 0; i < elements; i++) {
if (A_h[i] != memsetval) {
INFO("testMemsetMaxValue: index:" << i << ",computed:"
<< std::hex << static_cast<int>(A_h[i]) << ",memsetval:"
<< std::hex << memsetval);
REQUIRE(false);
}
}
HIP_CHECK(hipFree(devPitchedPtr.ptr));
free(A_h);
}
/**
* Function seeks device ptr to random slice and performs Memset operation
* on the slice selected.
*/
static void seekAndSet3DArraySlice(bool bAsync) {
char array3D[ZSIZE_S][YSIZE_S][XSIZE_S]{};
dim3 arr_dimensions = dim3(ZSIZE_S, YSIZE_S, XSIZE_S);
hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x,
arr_dimensions.y, arr_dimensions.z);
hipPitchedPtr devicePitchedPointer;
int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL;
HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent));
HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent));
// select random slice for memset
unsigned int seed = time(nullptr);
int slice_index = rand_r(&seed) % ZSIZE_S;
INFO("memset3d for sliceindex " << slice_index);
// Get attributes from device pitched pointer
size_t pitch = devicePitchedPointer.pitch;
size_t slicePitch = pitch * extent.height;
// Point devptr to selected slice
char *devPtrSlice = (reinterpret_cast<char *>(devicePitchedPointer.ptr))
+ slice_index * slicePitch;
hipExtent extentSlice = make_hipExtent(sizeof(char) * arr_dimensions.x,
arr_dimensions.y, 1);
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrSlice, pitch,
arr_dimensions.x, arr_dimensions.y);
if (bAsync) {
// Memset selected slice (Async)
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked,
extentSlice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
} else {
// Memset selected slice
HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, extentSlice));
}
// Copy result back to host buffer
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
arr_dimensions.x, arr_dimensions.y);
myparms.srcPtr = devicePitchedPointer;
myparms.extent = extent;
#if HT_NVIDIA
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
#else
myparms.kind = hipMemcpyDeviceToHost;
#endif
HIP_CHECK(hipMemcpy3D(&myparms));
for (int z = 0; z < ZSIZE_S; z++) {
for (int y = 0; y < YSIZE_S; y++) {
for (int x = 0; x < XSIZE_S; x++) {
if (z == slice_index) {
if (array3D[z][y][x] != memsetval4seeked) {
INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z
<< "," << y << "," << x << ") " << "computed:" << std::hex
<< array3D[z][y][x] << ", memsetval:" << std::hex
<< memsetval4seeked);
REQUIRE(false);
}
} else {
if (array3D[z][y][x] != memsetval) {
INFO("seekAndSet3DArray Slice: mismatch at index: Arr(" << z
<< "," << y << "," << x << ") " << "computed:" << std::hex
<< array3D[z][y][x] << ", memsetval:" << std::hex
<< memsetval);
REQUIRE(false);
}
}
}
}
}
HIP_CHECK(hipFree(devicePitchedPointer.ptr));
}
/**
* Function seeks device ptr to selected portion of 3d array
* and performs Memset operation on the portion.
*/
static void seekAndSet3DArrayPortion(bool bAsync) {
char array3D[ZSIZE_P][YSIZE_P][XSIZE_P]{};
dim3 arr_dimensions = dim3(ZSIZE_P, YSIZE_P, XSIZE_P);
hipExtent extent = make_hipExtent(sizeof(char) * arr_dimensions.x,
arr_dimensions.y, arr_dimensions.z);
hipPitchedPtr devicePitchedPointer;
int memsetval = MEMSETVAL, memsetval4seeked = TESTVAL;
HIP_CHECK(hipMalloc3D(&devicePitchedPointer, extent));
HIP_CHECK(hipMemset3D(devicePitchedPointer, memsetval, extent));
// For memsetting extent/size(10,10,10) in the mid portion of cube(30,30,30),
// seek device ptr to (10,10,10) and then memset 10 bytes across x,y,z axis.
size_t pitch = devicePitchedPointer.pitch;
size_t slicePitch = pitch * extent.height;
int slice_index = ZPOS_START, y = YPOS_START, x = XPOS_START;
// Select 10th slice
char *devPtrSlice = (reinterpret_cast<char *>(devicePitchedPointer.ptr))
+ slice_index * slicePitch;
// Now select row at height as 10
char *current_row = reinterpret_cast<char *>(devPtrSlice + y * pitch);
// Now select index of selected row as 10
char *devPtrIndexed = &current_row[x];
// Make dev Pitchedptr, extent
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrIndexed, pitch,
arr_dimensions.x, arr_dimensions.y);
hipExtent setExtent = make_hipExtent(sizeof(char) * XSET_LEN, YSET_LEN,
ZSET_LEN);
if (bAsync) {
// Memset selected portion (Async)
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemset3DAsync(modDevPitchedPtr, memsetval4seeked,
setExtent, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
} else {
// Memset selected portion
HIP_CHECK(hipMemset3D(modDevPitchedPtr, memsetval4seeked, setExtent));
}
// Copy result back to host buffer
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(array3D, sizeof(char) * arr_dimensions.x,
arr_dimensions.x, arr_dimensions.y);
myparms.srcPtr = devicePitchedPointer;
myparms.extent = extent;
#if HT_NVIDIA
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
#else
myparms.kind = hipMemcpyDeviceToHost;
#endif
HIP_CHECK(hipMemcpy3D(&myparms));
for (int z = 0; z < ZSIZE_P; z++) {
for (int y = 0; y < YSIZE_P; y++) {
for (int x = 0; x < XSIZE_P; x++) {
if ((z >= ZPOS_START && z <= ZPOS_END) &&
(y >= YPOS_START && y <= YPOS_END) &&
(x >= XPOS_START && x <= XPOS_END)) {
if (array3D[z][y][x] != memsetval4seeked) {
INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z
<< "," << y << "," << x << ") " << "computed:" << std::hex
<< array3D[z][y][x] << ", memsetval:" << std::hex
<< memsetval4seeked);
REQUIRE(false);
}
} else {
if (array3D[z][y][x] != memsetval) {
INFO("seekAndSet3DArray Portion: mismatch at index: Arr(" << z
<< "," << y << "," << x << ") " << "computed:" << std::hex
<< array3D[z][y][x] << ", memsetval:" << std::hex
<< memsetval);
REQUIRE(false);
}
}
}
}
}
HIP_CHECK(hipFree(devicePitchedPointer.ptr));
}
/**
* Test Memset3D with different combinations of extent
* taking zero and non-zero fields.
*/
TEST_CASE("Unit_hipMemset3D_MemsetWithExtent") {
hipExtent testExtent;
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
SECTION("Memset with extent width(0)") {
// Memset with extent width(0) and verify data to be intact
testExtent = make_hipExtent(0, numH, depth);
testMemsetWithExtent(0, testExtent);
}
SECTION("Memset with extent height(0)") {
// Memset with extent height(0) and verify data to be intact
testExtent = make_hipExtent(numW, 0, depth);
testMemsetWithExtent(0, testExtent);
}
SECTION("Memset with extent depth(0)") {
// Memset with extent depth(0) and verify data to be intact
testExtent = make_hipExtent(numW, numH, 0);
testMemsetWithExtent(0, testExtent);
}
SECTION("Memset with extent width,height,depth as 0") {
// Memset with extent width,height,depth as 0 and verify data to be intact
testExtent = make_hipExtent(0, 0, 0);
testMemsetWithExtent(0, testExtent);
}
}
/**
* Test Memset3DAsync with different combinations of extent
* taking zero and non-zero fields.
*/
TEST_CASE("Unit_hipMemset3DAsync_MemsetWithExtent") {
hipExtent testExtent;
size_t numH = NUMH_EXT, numW = NUMW_EXT, depth = DEPTH_EXT;
SECTION("Memset with extent width(0)") {
// Memset with extent width(0) and verify data to be intact
testExtent = make_hipExtent(0, numH, depth);
testMemsetWithExtent(1, testExtent);
}
SECTION("Memset with extent height(0)") {
// Memset with extent height(0) and verify data to be intact
testExtent = make_hipExtent(numW, 0, depth);
testMemsetWithExtent(1, testExtent);
}
SECTION("Memset with extent depth(0)") {
// Memset with extent depth(0) and verify data to be intact
testExtent = make_hipExtent(numW, numH, 0);
testMemsetWithExtent(1, testExtent);
}
SECTION("Memset with extent width,height,depth as 0") {
// Memset with extent width,height,depth as 0 and verify data to be intact
testExtent = make_hipExtent(0, 0, 0);
testMemsetWithExtent(1, testExtent);
}
}
/**
* Memset3D with max unsigned char and verify memset operation is success
*/
TEST_CASE("Unit_hipMemset3D_MemsetMaxValue") {
testMemsetMaxValue(0);
}
/**
* Memset3DAsync with max unsigned char and verify memset operation is success
*/
TEST_CASE("Unit_hipMemset3DAsync_MemsetMaxValue") {
testMemsetMaxValue(1);
}
/**
* Seek and set random slice of 3d array, verify memset is success
*/
TEST_CASE("Unit_hipMemset3D_SeekSetSlice") {
seekAndSet3DArraySlice(0);
}
/**
* Seek and set random slice of 3d array with async, verify memset is success
*/
TEST_CASE("Unit_hipMemset3DAsync_SeekSetSlice") {
seekAndSet3DArraySlice(1);
}
/**
* Memset3D selected portion of 3d array
*/
TEST_CASE("Unit_hipMemset3D_SeekSetArrayPortion") {
seekAndSet3DArrayPortion(0);
}
/**
* Memset3DAsync selected portion of 3d array
*/
TEST_CASE("Unit_hipMemset3DAsync_SeekSetArrayPortion") {
seekAndSet3DArrayPortion(1);
}
catch/unit/memory/hipMemset3DNegative.cc
+236
Dosyayı Görüntüle
@@ -0,0 +1,236 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Test hipMemset3D() with uninitialized devPitchedPtr.
2) Test hipMemset3DAsync() with uninitialized devPitchedPtr.
3) Reset devPitchedPtr to zero and check return value for hipMemset3D().
4) Reset devPitchedPtr to zero and check return value for hipMemset3DAsync().
5) Test hipMemset3D() with extent.width as max size_t and keeping height,
depth as valid values.
6) Test hipMemset3DAsync() with extent.width as max size_t and keeping height,
depth as valid values.
7) Test hipMemset3D() with extent.height as max size_t and keeping width,
depth as valid values.
8) Test hipMemset3DAsync() with extent.height as max size_t and keeping width,
depth as valid values.
9) Test hipMemset3D() with extent.depth as max size_t and keeping height,
width as valid values.
10) Test hipMemset3DAsync() with extent.depth as max size_t and keeping height,
width as valid values.
11) Device Ptr out bound and extent(0) passed for hipMemset3D().
12) Device Ptr out bound and extent(0) passed for hipMemset3DAsync().
13) Device Ptr out bound and valid extent passed for hipMemset3D().
14) Device Ptr out bound and valid extent passed for hipMemset3DAsync().
*/
#include <hip_test_common.hh>
TEST_CASE("Unit_hipMemset3D_Negative") {
hipError_t ret;
hipPitchedPtr devPitchedPtr;
constexpr int memsetval = 1;
constexpr size_t numH = 256, numW = 256;
constexpr size_t depth = 10;
constexpr size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
SECTION("Using uninitialized devpitched ptr") {
hipPitchedPtr devPitchedUnPtr;
ret = hipMemset3D(devPitchedUnPtr, memsetval, extent);
REQUIRE(ret != hipSuccess);
}
SECTION("Reset devPitchedPtr to zero") {
hipPitchedPtr rdevPitchedPtr{};
ret = hipMemset3D(rdevPitchedPtr, memsetval, extent);
REQUIRE(ret != hipSuccess);
}
SECTION("Pass extent fields as max size_t") {
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
numH,
depth);
hipExtent extMH = make_hipExtent(width,
std::numeric_limits<std::size_t>::max(),
depth);
hipExtent extMD = make_hipExtent(width,
numH,
std::numeric_limits<std::size_t>::max());
ret = hipMemset3D(devPitchedPtr, memsetval, extMW);
REQUIRE(ret != hipSuccess);
ret = hipMemset3D(devPitchedPtr, memsetval, extMH);
REQUIRE(ret != hipSuccess);
if ((TestContext::get()).isAmd()) {
ret = hipMemset3D(devPitchedPtr, memsetval, extMD);
REQUIRE(ret != hipSuccess);
} else {
WARN("Test is skipped for max depth."
<< "Cuda doesn't check the maximum depth of extent field");
}
}
SECTION("Device Ptr out bound and extent(0) passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
hipExtent extent0{};
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent0);
// api expected to check extent0 and return success before going for ptr.
REQUIRE(ret == hipSuccess);
}
SECTION("Device Ptr out bound and valid extent passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
ret = hipMemset3D(modDevPitchedPtr, memsetval, extent);
REQUIRE(ret != hipSuccess);
}
HIP_CHECK(hipFree(devPitchedPtr.ptr));
}
TEST_CASE("Unit_hipMemset3DAsync_Negative") {
hipError_t ret;
hipPitchedPtr devPitchedPtr;
hipStream_t stream;
constexpr int memsetval = 1;
constexpr size_t numH = 256;
constexpr size_t numW = 256;
constexpr size_t depth = 10;
constexpr size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
SECTION("Using uninitialized devpitched ptr") {
hipPitchedPtr devPitchedUnPtr;
ret = hipMemset3DAsync(devPitchedUnPtr, memsetval, extent, stream);
REQUIRE(ret != hipSuccess);
}
SECTION("Reset devPitchedPtr to zero") {
hipPitchedPtr rdevPitchedPtr{};
ret = hipMemset3DAsync(rdevPitchedPtr, memsetval, extent, stream);
REQUIRE(ret != hipSuccess);
}
SECTION("Pass extent fields as max size_t") {
hipExtent extMW = make_hipExtent(std::numeric_limits<std::size_t>::max(),
numH,
depth);
hipExtent extMH = make_hipExtent(width,
std::numeric_limits<std::size_t>::max(),
depth);
hipExtent extMD = make_hipExtent(width,
numH,
std::numeric_limits<std::size_t>::max());
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMW, stream);
REQUIRE(ret != hipSuccess);
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMH, stream);
REQUIRE(ret != hipSuccess);
if ((TestContext::get()).isAmd()) {
ret = hipMemset3DAsync(devPitchedPtr, memsetval, extMD, stream);
REQUIRE(ret != hipSuccess);
} else {
WARN("Test is skipped for max depth."
<< "Cuda doesn't check the maximum depth of extent field");
}
}
SECTION("Device Ptr out bound and extent(0) passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
hipExtent extent0{};
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent0, stream);
// api expected to check extent0 and return success before going for ptr.
REQUIRE(ret == hipSuccess);
}
SECTION("Device Ptr out bound and valid extent passed for memset") {
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * extent.height;
constexpr auto advanceOffset = 10;
// Point devptr to end of allocated memory
char *devPtrMod = (reinterpret_cast<char *>(devPitchedPtr.ptr))
+ depth * slicePitch;
// Advance devptr further to go out of boundary
devPtrMod = devPtrMod + advanceOffset;
hipPitchedPtr modDevPitchedPtr = make_hipPitchedPtr(devPtrMod, pitch,
numW * sizeof(char), numH);
ret = hipMemset3DAsync(modDevPitchedPtr, memsetval, extent, stream);
REQUIRE(ret != hipSuccess);
}
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipFree(devPitchedPtr.ptr));
}
catch/unit/memory/hipMemset3DRegressMultiThread.cc
+266
Dosyayı Görüntüle
@@ -0,0 +1,266 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Validate Async behavior of hipMemset3DAsync with commands queued
concurrently from multiple threads.
2) Validate hipMemset3DAsync behavior when api is queued along with kernel
function operating on same memory.
3) Perform regression of hipMemset3D api in loop with device memory allocated
on different gpus.
4) Perform regression of hipMemset3DAsync api in loop with device memory
allocated on different gpus.
*/
#include <hip_test_common.hh>
/*
* Defines
*/
#define MAX_REGRESS_ITERS 2
#define MAX_THREADS 10
/**
* kernel function sets device memory with value passed
*/
static __global__ void func_set_value(hipPitchedPtr devicePitchedPointer,
hipExtent extent,
unsigned char val) {
// Index Calculation
size_t x = threadIdx.x + blockDim.x * blockIdx.x;
size_t y = threadIdx.y + blockDim.y * blockIdx.y;
size_t z = threadIdx.z + blockDim.z * blockIdx.z;
// Get attributes from device pitched pointer
char *devicePointer = reinterpret_cast<char *>(devicePitchedPointer.ptr);
size_t pitch = devicePitchedPointer.pitch;
size_t slicePitch = pitch * extent.height;
// Loop over the device buffer
if (z < extent.depth) {
char *current_slice_index = devicePointer + z * slicePitch;
if (y < extent.height) {
// Get data array containing all elements from the current row
char *current_row = reinterpret_cast<char *>(current_slice_index
+ y * pitch);
if (x < extent.width) {
current_row[x] = val;
}
}
}
}
/**
* Thread function queues kernel function and memset cmds
*/
static void threadFunc(hipStream_t stream, hipPitchedPtr devpPtr,
int memsetval, int testval, hipExtent extent, hipMemcpy3DParms myparms) {
// Kernel Launch Configuration
constexpr auto size = 8;
dim3 threadsPerBlock = dim3(size, size, size);
dim3 blocks;
blocks = dim3((extent.width + threadsPerBlock.x - 1) / threadsPerBlock.x,
(extent.height + threadsPerBlock.y - 1) / threadsPerBlock.y,
(extent.depth + threadsPerBlock.z - 1) / threadsPerBlock.z);
hipLaunchKernelGGL(func_set_value, dim3(blocks), dim3(threadsPerBlock), 0,
stream, devpPtr, extent, memsetval);
HIPCHECK(hipMemset3DAsync(devpPtr, testval, extent, stream));
HIPCHECK(hipMemcpy3DAsync(&myparms, stream));
}
/**
* Performs api regression in loop
*/
bool loopRegression(bool bAsync) {
bool testPassed = true;
char *A_h;
constexpr int memsetval = 1;
constexpr size_t numH = 256, numW = 100, depth = 10;
int numGpu = 0, hasPeerAccess = 0;
size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
size_t sizeElements = width * numH * depth;
size_t elements = numW* numH* depth;
std::vector<hipPitchedPtr> devPitchedPtrlist;
hipPitchedPtr pitchedPtr, devpPtr;
A_h = reinterpret_cast<char *>(malloc(sizeElements));
REQUIRE(A_h != nullptr);
memset(A_h, 0, sizeElements);
// Populate hipMemcpy3D parameters
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
myparms.extent = extent;
#if HT_NVIDIA
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
#else
myparms.kind = hipMemcpyDeviceToHost;
#endif
HIP_CHECK(hipGetDeviceCount(&numGpu));
REQUIRE(numGpu > 0);
// Alloc 3D arrays in all GPUs
for (int j = 0; j < numGpu; j++) {
HIP_CHECK(hipSetDevice(j));
HIP_CHECK(hipMalloc3D(&pitchedPtr, extent));
devPitchedPtrlist.push_back(pitchedPtr);
}
for (int itern = 0; itern < MAX_REGRESS_ITERS; itern++) {
// Validate hipMemset3D data consistency in multiple iters
for (int i = 0; i < numGpu; i++) {
for (int j = 0; j < numGpu; j++) {
HIP_CHECK(hipDeviceCanAccessPeer(&hasPeerAccess, i, j));
if (!hasPeerAccess) {
// Skip and continue if no peer access
continue;
}
HIP_CHECK(hipSetDevice(i));
devpPtr = devPitchedPtrlist[j];
HIP_CHECK(hipDeviceEnablePeerAccess(j, 0));
HIP_CHECK(hipMemset3D(devpPtr, 0, extent));
if (bAsync) {
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemset3DAsync(devpPtr, memsetval, extent, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
} else {
HIP_CHECK(hipMemset3D(devpPtr, memsetval, extent));
}
myparms.srcPtr = devpPtr;
memset(A_h, 0, sizeElements);
HIP_CHECK(hipMemcpy3D(&myparms));
for (size_t indx = 0; indx < elements; indx++) {
if (A_h[indx] != memsetval) {
testPassed = false;
printf("RegressIter : mismatch at index:%d computed:%02x, "
"memsetval:%02x\n", static_cast<int>(indx),
static_cast<int>(A_h[indx]), static_cast<int>(memsetval));
break;
}
}
}
}
}
for (int j = 0; j < numGpu; j++) {
HIP_CHECK(hipFree(devPitchedPtrlist[j].ptr));
}
free(A_h);
return testPassed;
}
/**
* Perform regression of hipMemset3D api with device memory allocated
* on different gpus.
*/
TEST_CASE("Unit_hipMemset3D_RegressInLoop") {
bool TestPassed = false;
TestPassed = loopRegression(0);
REQUIRE(TestPassed == true);
}
/**
* Perform regression of hipMemset3DAsync api with device memory allocated
* on different gpus.
*/
TEST_CASE("Unit_hipMemset3DAsync_RegressInLoop") {
bool TestPassed = false;
TestPassed = loopRegression(1);
REQUIRE(TestPassed == true);
}
/**
* Async commands queued concurrently and executed
*/
TEST_CASE("Unit_hipMemset3DAsync_ConcurrencyMthread") {
char *A_h;
constexpr int memsetval = 1, testval = 2;
constexpr size_t numH = 256, numW = 100, depth = 10;
size_t width = numW * sizeof(char);
hipExtent extent = make_hipExtent(width, numH, depth);
size_t sizeElements = width * numH * depth;
size_t elements = numW* numH* depth;
hipPitchedPtr devpPtr;
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMalloc3D(&devpPtr, extent));
A_h = reinterpret_cast<char *>(malloc(sizeElements));
REQUIRE(A_h != nullptr);
memset(A_h, 0, sizeElements);
// Populate hipMemcpy3D parameters
hipMemcpy3DParms myparms{};
myparms.srcPos = make_hipPos(0, 0, 0);
myparms.srcPtr = devpPtr;
myparms.dstPos = make_hipPos(0, 0, 0);
myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
myparms.extent = extent;
#if HT_NVIDIA
myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
#else
myparms.kind = hipMemcpyDeviceToHost;
#endif
std::vector<std::thread> threadlist;
// Queue cmds concurrently from multiple threads on same stream
for (int i = 0; i < MAX_THREADS; i++) {
threadlist.push_back(std::thread(threadFunc, stream, devpPtr, memsetval,
testval, extent, myparms));
}
for (auto &t : threadlist) {
t.join();
}
HIP_CHECK(hipStreamSynchronize(stream));
for (size_t k = 0 ; k < elements ; k++) {
if (A_h[k] != testval) {
CAPTURE(A_h[k], testval, k);
REQUIRE(false);
}
}
HIP_CHECK(hipStreamDestroy(stream));
free(A_h);
HIP_CHECK(hipFree(devpPtr.ptr));
}