SWDEV-298757 - Moved long running tests into stress category. (#2461)
このコミットが含まれているのは:
lthakur
GitHub
@@ -17,12 +17,11 @@
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
List of Test cases:
|
||||
1) Unit_hipMallocManaged_Basic
|
||||
2) Unit_hipMallocManaged_MultiSize
|
||||
3) Unit_hipMallocManaged_MultiKrnlHmmAccess
|
||||
4) Unit_hipMallocManaged_KrnlWth2MemTypes
|
||||
/* Test Case Description:
|
||||
1) This testcase verifies the hipMallocManaged basic scenario - supported on
|
||||
all devices
|
||||
2) This testcase verifies the hipMallocManaged basic scenario - supported
|
||||
only on HMM enabled devices
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
@@ -32,12 +31,6 @@
|
||||
|
||||
|
||||
// Kernel functions
|
||||
__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
for (size_t i = index; i < n; i++) {
|
||||
Hmm[i] = Dptr[i] + 10;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) {
|
||||
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
@@ -64,9 +57,6 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr,
|
||||
}
|
||||
}
|
||||
|
||||
// The following variable will be used to get the result of computation
|
||||
// from multiple threads
|
||||
static bool IfTestPassed = true;
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
@@ -93,62 +83,6 @@ static int HmmAttrPrint() {
|
||||
}
|
||||
|
||||
|
||||
static void LaunchKrnl4(size_t NumElms, int InitVal) {
|
||||
int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms)));
|
||||
HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms)));
|
||||
int *Hstptr = reinterpret_cast<int*>(new int[NumElms]);
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
Hstptr[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)),
|
||||
hipMemcpyHostToDevice));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
DataMismatch = 0;
|
||||
KernelMul_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the result
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal * 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
DataMismatch = 0;
|
||||
KernelMulAdd_MngdMem<<<dimGrid, dimBlock, 0, strm>>>(Hmm, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the result
|
||||
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
if (Hmm[i] != (InitVal * 10 * 2 + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n");
|
||||
REQUIRE(false);
|
||||
}
|
||||
delete[] Hstptr;
|
||||
}
|
||||
|
||||
|
||||
|
||||
static size_t N{4 * 1024 * 1024};
|
||||
static unsigned blocksPerCU{6};
|
||||
@@ -241,104 +175,3 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") {
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// The following test case tests the behavior of kernel with a HMM memory and
|
||||
// hipMalloc memory
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
|
||||
IfTestPassed = true;
|
||||
int *Hmm = NULL, *Dptr = NULL, InitVal = 123;
|
||||
size_t NumElms = (1024 * 1024);
|
||||
int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0;
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms));
|
||||
HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms));
|
||||
for (size_t i = 0; i < NumElms; ++i) {
|
||||
Hmm[i] = 0;
|
||||
Hptr[i] = InitVal;
|
||||
}
|
||||
HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms,
|
||||
hipMemcpyHostToDevice));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the results
|
||||
for (size_t k = 0; k < NumElms; ++k) {
|
||||
if (Hmm[k] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch observed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
HIP_CHECK(hipFree(Dptr));
|
||||
delete[] Hptr;
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
// The following test case tests when the same Hmm memory is used for
|
||||
// launching multiple different kernels will results in any issue
|
||||
TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed) {
|
||||
int InitVal = 123, NumElms = (1024 * 1024);
|
||||
LaunchKrnl4(NumElms, InitVal);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
// The following test case allocation, host access, device access of HMM
|
||||
// memory from size 1 to 10KB
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_MultiSize") {
|
||||
IfTestPassed = true;
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
unsigned char *Hmm1 = NULL, *Hmm2 = NULL;
|
||||
int InitVal = 100, blockSize = 64, DataMismatch = 0;
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
dim3 dimBlock(blockSize, 1, 1);
|
||||
for (int i = 1; i < (1024*1024); ++i) {
|
||||
HIP_CHECK(hipMallocManaged(&Hmm1, i));
|
||||
HIP_CHECK(hipMallocManaged(&Hmm2, i));
|
||||
for (int j = 0; j < i; ++j) {
|
||||
Hmm1[j] = InitVal;
|
||||
}
|
||||
dim3 dimGrid((i + blockSize -1)/blockSize, 1, 1);
|
||||
KrnlWth2MemTypesC<<<dimGrid, dimBlock, 0, strm>>>(Hmm2, Hmm1, i);
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
// Verifying the results
|
||||
for (int k = 0; k < i; ++k) {
|
||||
if (Hmm2[k] != (InitVal + 10)) {
|
||||
DataMismatch++;
|
||||
}
|
||||
}
|
||||
if (DataMismatch != 0) {
|
||||
WARN("DataMismatch observed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
DataMismatch = 0;
|
||||
HIP_CHECK(hipFree(Hmm1));
|
||||
HIP_CHECK(hipFree(Hmm2));
|
||||
REQUIRE(IfTestPassed);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -262,79 +262,3 @@ TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") {
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
bool IfTestPassed = true;
|
||||
hipError_t err;
|
||||
void *Hmm = NULL;
|
||||
size_t totalDevMem = 0, freeDevMem = 0;
|
||||
int NumDevs = 0;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
// Testing allocation of extreme and unusual mem values
|
||||
for (int i = 0; i < NumDevs; i++) {
|
||||
HIP_CHECK(hipSetDevice(i));
|
||||
HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem));
|
||||
err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal);
|
||||
if (hipSuccess == err) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
WARN("Observed error while allocating memory on GPU: " << i);
|
||||
WARN(" size 1 with");
|
||||
WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
|
||||
WARN("Error: " << hipGetErrorString(err));
|
||||
IfTestPassed = false;
|
||||
}
|
||||
err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal);
|
||||
if (hipSuccess == err) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
WARN("Observed error while allocating max free memory on GPU: " << i);
|
||||
WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
|
||||
WARN("Error: " << hipGetErrorString(err));
|
||||
IfTestPassed = false;
|
||||
}
|
||||
err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal);
|
||||
if (hipSuccess == err) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
WARN("Observed error while allocating max (free - 1) memory on ");
|
||||
WARN("GPU: " << i);
|
||||
WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n");
|
||||
WARN("Error: " << hipGetErrorString(err));
|
||||
IfTestPassed = false;
|
||||
}
|
||||
err = hipMallocManaged(&Hmm, 1, hipMemAttachHost);
|
||||
if (hipSuccess == err) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
WARN("Observed error while allocating memory size 1 on GPU: " << i);
|
||||
WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
|
||||
WARN("Error: " << hipGetErrorString(err));
|
||||
IfTestPassed = false;
|
||||
}
|
||||
err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost);
|
||||
if (hipSuccess == err) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
WARN("Observed error while allocating max free memory on GPU: " << i);
|
||||
WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
|
||||
WARN("Error: " << hipGetErrorString(err));
|
||||
IfTestPassed = false;
|
||||
}
|
||||
err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost);
|
||||
if (hipSuccess == err) {
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
} else {
|
||||
WARN("Observed error while allocating max (freeDevMem - 1) memory"
|
||||
" on GPU: " << i);
|
||||
WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n");
|
||||
WARN("Error: " << hipGetErrorString(err));
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result");
|
||||
}
|
||||
}
|
||||
|
||||
新しいイシューから参照
ユーザーをブロックする