From 460a07f33876fd0151750a2ff65b090ea451bcd9 Mon Sep 17 00:00:00 2001 From: lthakur Date: Thu, 24 Feb 2022 01:28:08 +0530 Subject: [PATCH] SWDEV-298757 - Moved long running tests into stress category. (#2461) --- catch/stress/memory/hipMallocManagedStress.cc | 227 +++++++++++++++++- catch/unit/memory/hipMallocManaged.cc | 177 +------------- catch/unit/memory/hipMallocManagedFlagsTst.cc | 76 ------ 3 files changed, 231 insertions(+), 249 deletions(-) diff --git a/catch/stress/memory/hipMallocManagedStress.cc b/catch/stress/memory/hipMallocManagedStress.cc index 0e7d80edb6..7e992cd152 100644 --- a/catch/stress/memory/hipMallocManagedStress.cc +++ b/catch/stress/memory/hipMallocManagedStress.cc @@ -1,5 +1,5 @@ /* - Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2022 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 @@ -19,6 +19,15 @@ // The following test case allocation, host access, device access of HMM // memory from size 1 to 10KB +/* Test Case Description: + 1) Testing allocation, host access, device access of HMM + memory from size 1 to 10KB + 2) The following test case tests the behavior of kernel with a HMM memory + and hipMalloc memory + 3) The following test case tests when the same Hmm memory is used for + launching multiple different kernels will results in any issue + 4) Testing the allocation of/scenarios around max possible memory + */ #include #include @@ -35,6 +44,86 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr, } static bool IfTestPassed = true; +// 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 KernelMulAdd_MngdMem(int *Hmm, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; + for (size_t i = index; i < n; i += stride) { + Hmm[i] = Hmm[i] * 2 + 10; + } +} + +__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; + for (size_t i = index; i < n; i += stride) { + Hmm[i] = Dptr[i] * 10; + } +} +static bool IfTestPassed = true; + +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(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<<>>(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<<>>(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<<>>(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 int HmmAttrPrint() { int managed = 0; INFO("The following are the attribute values related to HMM for" @@ -104,3 +193,139 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") { } } +// 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<<>>(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"); + } +} + +// Testing the allocation of/scenarios around max possible memory +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"); + } +} diff --git a/catch/unit/memory/hipMallocManaged.cc b/catch/unit/memory/hipMallocManaged.cc index e8bc69997c..d40aa924a1 100644 --- a/catch/unit/memory/hipMallocManaged.cc +++ b/catch/unit/memory/hipMallocManaged.cc @@ -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 @@ -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(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<<>>(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<<>>(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<<>>(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<<>>(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<<>>(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"); - } -} diff --git a/catch/unit/memory/hipMallocManagedFlagsTst.cc b/catch/unit/memory/hipMallocManagedFlagsTst.cc index e0d13b764e..8ac1c37853 100644 --- a/catch/unit/memory/hipMallocManagedFlagsTst.cc +++ b/catch/unit/memory/hipMallocManagedFlagsTst.cc @@ -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"); - } -}