AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-411880 - Fix test device malloc

Browse Source 
Unit_deviceAllocation_Malloc_ComplexDataType

Change-Id: I8d5c3a9c63e74ffc90ddf6b89729603236a88fcf
This commit is contained in:
Jatin Chaudhary
2024-07-31 16:20:34 +01:00
committed by Rakesh Roy
parent 352932738a
commit f62e3df454
3 changed files with 13 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/hipTestMain/config/config_amd_linux
-4
View File
@@ -1093,10 +1093,6 @@
"Unit_fp8_correctness - double",
"Unit_fp8_vector_basic_conversions",
#endif
#if defined gfx906
"=== SWDEV-419112 Below tests fail in stress test on 29/08/23 ===",
"Unit_deviceAllocation_Malloc_ComplexDataType",
#endif
#if defined gfx908
"=== Below test soft hang in stress test on 29/08/23 ===",
"Unit_hipMultiThreadStreams2",
catch/hipTestMain/config/config_amd_wsl.json
-1
View File
@@ -617,7 +617,6 @@
"Unit_hipDeviceGetPCIBusId_CheckPciBusIDWithLspci",
"Note: TDR",
"Unit_deviceAllocation_InOneThread_AccessInAllThreads",
"Unit_deviceAllocation_Malloc_ComplexDataType",
"Unit_deviceAllocation_Malloc_UnionType",
"Unit_deviceAllocation_New_ComplexDataType",
"Unit_deviceAllocation_New_UnionType",
catch/unit/deviceLib/deviceAllocation.cc
+13 -22
View File
@@ -142,8 +142,7 @@ static __global__ void kerTestAccessInAllThreadsInBlock(T *outputBuf,
* access/modifies it in all threads of block and copies
* data to host and frees the memory in another thread.
*/
static __global__ void kerTestAccessInAllThreads_CmplxStr(int test_type,
int *result) {
static __global__ void kerTestAccessInAllThreads_CmplxStr(int test_type, int* result) {
int myThreadId = threadIdx.x;
int lastThreadId = (blockDim.x - 1);
int myBlockId = blockIdx.x;
@@ -151,14 +150,12 @@ static __global__ void kerTestAccessInAllThreads_CmplxStr(int test_type,
// Allocate memory in thread 0
if (0 == myThreadId) {
if (test_type == TEST_MALLOC_FREE) {
dev_ptr[0][blockIdx.x] =
reinterpret_cast<void*> (malloc(sizeof(struct complexStructure)));
dev_ptr[0][blockIdx.x] = reinterpret_cast<void*>(malloc(sizeof(struct complexStructure)));
} else {
dev_ptr[0][blockIdx.x] =
reinterpret_cast<void*> (new struct complexStructure);
dev_ptr[0][blockIdx.x] = reinterpret_cast<void*>(new struct complexStructure);
}
struct complexStructure *ptr =
reinterpret_cast<struct complexStructure*> (dev_ptr[0][blockIdx.x]);
struct complexStructure* ptr =
reinterpret_cast<struct complexStructure*>(dev_ptr[0][blockIdx.x]);
ptr->alloc_internal_members(test_type, BLOCKSIZE);
}
// All threads wait at this barrier
@@ -168,8 +165,7 @@ static __global__ void kerTestAccessInAllThreads_CmplxStr(int test_type,
printf("Device Allocation Failed in thread = %d \n", myGid);
return;
}
struct complexStructure *ptr =
reinterpret_cast<struct complexStructure*> (dev_ptr[0][blockIdx.x]);
struct complexStructure* ptr = reinterpret_cast<struct complexStructure*>(dev_ptr[0][blockIdx.x]);
if (ptr->sthreadInfo == nullptr) {
printf("Structure Allocation Failed in thread = %d \n", myGid);
return;
@@ -188,12 +184,9 @@ static __global__ void kerTestAccessInAllThreads_CmplxStr(int test_type,
if (lastThreadId == myThreadId) {
int match = 1;
for (int idx = 0; idx < BLOCKSIZE; idx++) {
if ((ptr->sthreadInfo[idx].threadid != idx) ||
(ptr->sthreadInfo[idx].blockid != myBlockId) ||
(ptr->sthreadInfo[idx].ival != INT_MAX) ||
(ptr->sthreadInfo[idx].dval != DBL_MAX) ||
(ptr->sthreadInfo[idx].fval != FLT_MAX) ||
(ptr->sthreadInfo[idx].sval != SHRT_MAX) ||
if ((ptr->sthreadInfo[idx].threadid != idx) || (ptr->sthreadInfo[idx].blockid != myBlockId) ||
(ptr->sthreadInfo[idx].ival != INT_MAX) || (ptr->sthreadInfo[idx].dval != DBL_MAX) ||
(ptr->sthreadInfo[idx].fval != FLT_MAX) || (ptr->sthreadInfo[idx].sval != SHRT_MAX) ||
(ptr->sthreadInfo[idx].cval != SCHAR_MAX)) {
match = 0;
break;
@@ -573,18 +566,16 @@ static bool TestDevMemAllocMulKerMulThrd(int test_type) {
*/
static bool TestMemoryAccessInAllThread_CmplxStr(int test_type) {
int *result_d{nullptr}, *result_h{nullptr};
size_t arraysize = BLOCKSIZE;
result_h = reinterpret_cast<int*> (malloc(sizeof(int) * arraysize));
size_t arraysize = GRIDSIZE;
result_h = reinterpret_cast<int*>(malloc(sizeof(int) * arraysize));
REQUIRE(result_h != nullptr);
HIP_CHECK(hipMalloc(&result_d, (sizeof(int) * arraysize)));
HIP_CHECK(hipMemset(result_d, 0, (sizeof(int) * arraysize)));
// Launch Test Kernel
kerTestAccessInAllThreads_CmplxStr<<<GRIDSIZE, BLOCKSIZE>>>(
test_type, result_d);
kerTestAccessInAllThreads_CmplxStr<<<GRIDSIZE, BLOCKSIZE>>>(test_type, result_d);
HIP_CHECK(hipDeviceSynchronize());
// Copy to host buffer
HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int) * arraysize,
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int) * arraysize, hipMemcpyDefault));
bool bPassed = true;
for (size_t idx = 0; idx < GRIDSIZE; idx++) {
if (result_h[idx] != 1) {