From 9eae1e0f533352af9f2eba009262c2d4177a8402 Mon Sep 17 00:00:00 2001 From: Anton Mitkov Date: Thu, 28 Jul 2022 11:15:42 +0100 Subject: [PATCH 1/9] Added testing for hipMemGetInfo (#2682) --- tests/catch/unit/memory/CMakeLists.txt | 2 + tests/catch/unit/memory/hipMemGetInfo.cc | 429 +++++++++++++++++++++++ 2 files changed, 431 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemGetInfo.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 997572ee6d..853c1453af 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -91,6 +91,7 @@ set(TEST_SRC hipMemPrefetchAsync.cc hipArray.cc hipMemVmm.cc + hipMemGetInfo.cc ) else() set(TEST_SRC @@ -157,6 +158,7 @@ set(TEST_SRC hipPointerGetAttribute.cc hipDrvPtrGetAttributes.cc hipMemPrefetchAsync.cc + hipMemGetInfo.cc ) endif() diff --git a/tests/catch/unit/memory/hipMemGetInfo.cc b/tests/catch/unit/memory/hipMemGetInfo.cc new file mode 100644 index 0000000000..fc91511a3b --- /dev/null +++ b/tests/catch/unit/memory/hipMemGetInfo.cc @@ -0,0 +1,429 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + + +#include +#include +#include + +/* + * This testcase verifies hipMemGetInfo API + * 1. Different memory chunk allocation + * 1.1. hipMalloc + * 1.2. hipMallocArray + * 1.3. hipMalloc3D + * 1.3. hipMalloc3DArray + * 2. Allocation using different threads + * 3. Negative: Invalid args + */ + +struct MinAlloc { + private: + int value; + MinAlloc() { + size_t freeMemInit; + size_t totalMemInit; + + unsigned int* A_mem{nullptr}; + size_t mallocSize{1}; + + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + // allocate 1 byte + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), mallocSize)); + + size_t freeMemRet; + size_t totalMemRet; + // actual allocation should be bigger to reflect the minimum allocation on device + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + REQUIRE(freeMemInit >= freeMemRet); + HIP_CHECK(hipFree(A_mem)); + + // store the size of minimum allocation + value = (freeMemInit - freeMemRet); + } + + public: + static int Get() { + static MinAlloc instance; + return instance.value; + } +}; + +// if the memory being allocated is not divisible by the minimum allocation add an extra minimum +// allocation AddedAllocation = InitialAllocation + (MinAllocation - divisionRemainer) +void fixAllocSize(size_t& allocation) { + REQUIRE(MinAlloc::Get() != 0); + if (allocation % MinAlloc::Get() != 0) { + auto adjustment = allocation % MinAlloc::Get(); + adjustment = MinAlloc::Get() - adjustment; + allocation = allocation + adjustment; + } +} + +// Print information about memory +#define MEMINFO(totalMem, freeMemInit, freeMemRet, usedMem) \ + INFO("Total memory: \t\t\t" << totalMem << "\n" \ + << "Memory used: \t\t\t\t" << freeMemInit - freeMemRet << "\n" \ + << "Free memory after alloc: \t\t" << freeMemRet << "\n" \ + << "Free memory initally: \t\t" << freeMemInit << "\n" \ + << "Memory assumed to be used: \t\t" << usedMem); + + +TEST_CASE("Unit_hipMemGetInfo_DifferentMallocLarge") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + unsigned int* A_mem{nullptr}; + unsigned int* B_mem{nullptr}; + + size_t freeMemRet; + size_t totalMemRet; + int device; + HIP_CHECK(hipGetDevice(&device)); + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + auto totalMemory = prop.totalGlobalMem; + + + // allocate half of free mem + auto Malloc1Size = freeMemInit >> 1; + // if the allocation is not divisible by the MinAllocation + // take into account and add padding + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); + + // allocate an extra quarter of free mem + auto Malloc2Size = Malloc1Size >> 1; + HIP_CHECK(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); + + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size + Malloc2Size); + // check if device property total memory is the same as + // total memory returned from hipMemGetInfo + REQUIRE(totalMemory == totalMemRet); + auto allocSize = Malloc1Size + Malloc2Size; + auto assumedFreeMem = freeMemInit - allocSize; + + REQUIRE(freeMemRet <= assumedFreeMem); + HIP_CHECK(hipFree(A_mem)); + HIP_CHECK(hipFree(B_mem)); +} + +TEST_CASE("Unit_hipMemGetInfo_DifferentMallocSmall") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + unsigned int* A_mem{nullptr}; + size_t freeMemRet; + size_t totalMemRet; + // allocate smaller chunk than minimum + size_t Malloc1Size = 1; + + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); + + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); + + auto assumedFreeMem = freeMemInit - Malloc1Size; + // Free memory should be less than assumed for + // single allocation smaller than min allocation chunk + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFree(A_mem)); +} + +TEST_CASE("Unit_hipMemGetInfo_DifferentMallocMultiSmall") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + unsigned int* A_mem{nullptr}; + unsigned int* B_mem{nullptr}; + size_t freeMemRet; + size_t totalMemRet; + + // Allocate memory that is a quarter of the min allocation + // Expected behaviour is to reuse the min allocation memory + size_t MallocSize = MinAlloc::Get() >> 2; + + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); + HIP_CHECK(hipMalloc(reinterpret_cast(&B_mem), MallocSize)); + + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, MallocSize * 2); + + + auto assumedFreeMem = freeMemInit - (MallocSize * 2); + + // Confirm mem alocation results + REQUIRE(freeMemRet <= assumedFreeMem); + HIP_CHECK(hipFree(A_mem)); + HIP_CHECK(hipFree(B_mem)); +} + +TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_MallocArray", "", int, int4, char) { + // get initial mem data + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + // create and allocate an Array + hipArray_t arrayPtr{}; + + auto bytesPerItem = sizeof(TestType); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + hipExtent extent{}; + extent.width = GENERATE(32, 128, 256, 512, 1024); + + extent.height = GENERATE(0, 32, 128, 256, 512, 1024); + + HIP_CHECK(hipMallocArray(&arrayPtr, &desc, extent.width, extent.height, hipArrayDefault)); + + // check if memory is correct + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + // calculate used memory, take into account 1D array (height = 0) + size_t usedMem = bytesPerItem * extent.width * (extent.height != 0 ? extent.height : 1); + + // ensure we allocate at least the min allocation for the array + MEMINFO(totalMemRet, freeMemInit, freeMemRet, usedMem); + + size_t assumedFreeMem = freeMemInit - usedMem; + + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFreeArray(arrayPtr)); +} + +TEST_CASE("Unit_hipMemGetInfo_Malloc3D") { + // Get initial memory + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + // Allocate 3D object + hipExtent extent{}; + // extent is given in bytes for with + extent.width = GENERATE(32, 128, 256); + extent.height = GENERATE(32, 128, 256); + extent.depth = GENERATE(32, 128, 256); + hipPitchedPtr A_mem{}; + HIP_CHECK(hipMalloc3D(&A_mem, extent)); + + // Get memory after allocation + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + // Verify result + size_t mallocSize = A_mem.pitch * extent.height * extent.depth; + + size_t assumedFreeMem = freeMemInit - mallocSize; + MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); + + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFree(A_mem.ptr)); +} + +TEMPLATE_TEST_CASE("Unit_hipMemGetInfo_Malloc3DArray", "", char, int, int4) { + // Get initial memory + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + // Allocate 3D object + hipArray_t arrayPtr{}; + size_t sizeInBytes = (size_t)sizeof(TestType); + hipChannelFormatDesc desc = hipCreateChannelDesc(); + + int device; + HIP_CHECK(hipGetDevice(&device)); + int allignSize{0}; + hipDeviceGetAttribute(&allignSize, hipDeviceAttributeTextureAlignment, device); + +#if HT_NVIDIA + auto flag = GENERATE(hipArrayDefault, hipArrayLayered, hipArrayCubemap, + hipArrayLayered | hipArrayCubemap); +#else + // hipArrayCubemap not supported on AMD + auto flag = GENERATE(hipArrayDefault, hipArrayLayered); +#endif + + hipExtent extent{}; + extent.width = GENERATE(32, 128, 256, 512); + extent.height = GENERATE(0, 32, 128, 256, 512); + if (flag == hipArrayCubemap) { + // width must be equal to height, and depth must be six. + extent.height = extent.width; + extent.depth = 6; + } else if (flag == hipArrayLayered | hipArrayCubemap) { + // width must be equal to height, and depth must be a multiple six. + extent.height = extent.width; + extent.depth = 6 * GENERATE(4, 8, 16, 32); + } else if (extent.height == 0 && flag != hipArrayLayered) { + // if height = 0 the depth must be 0 unless using hipArrayLayered flag + extent.depth = 0; + } else { + extent.depth = GENERATE(32, 128, 256, 512); + } + + + // Get memory after allocation + auto h = extent.height == 0 ? 1 : extent.height; + auto d = extent.depth == 0 ? 1 : extent.depth; + auto w = extent.width * sizeInBytes; + size_t mallocSize = w * h * d; + + HIP_CHECK(hipMalloc3DArray(&arrayPtr, &desc, extent, flag)); + + // Verify result + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + // Sometimes hipMemGetInfo reports that no new memory has be allocated for testcase + // take this into account + if (freeMemInit == freeMemRet) { + // no new memory allocation has occured verify that memory trying + // to be allocated is less than a min allocation block + MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); + REQUIRE(mallocSize <= static_cast(MinAlloc::Get())); + + } else { + MEMINFO(totalMemRet, freeMemInit, freeMemRet, mallocSize); + size_t assumedFreeMem = freeMemInit - mallocSize; + REQUIRE(freeMemRet <= assumedFreeMem); + } + HIP_CHECK(hipFreeArray(arrayPtr)); +} + + +TEST_CASE("Unit_hipMemGetInfo_ParaLarge") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + unsigned int* A_mem{nullptr}; + unsigned int* B_mem{nullptr}; + + // allocate half of free mem + auto Malloc1Size = freeMemInit >> 1; + // if the allocation is not divisible by the MinAllocation + // take into account and add padding + std::thread t1( + [&] { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)); }); + + // allocate an extra quarter of free mem + auto Malloc2Size = Malloc1Size >> 1; + std::thread t2( + [&] { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&B_mem), Malloc2Size)); }); + + t1.join(); + t2.join(); + HIP_CHECK_THREAD_FINALIZE(); + + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size + Malloc2Size); + auto allocSize = Malloc1Size + Malloc2Size; + REQUIRE(freeMemRet <= freeMemInit - allocSize); + + HIP_CHECK(hipFree(A_mem)); + HIP_CHECK(hipFree(B_mem)); +} + +TEST_CASE("Unit_hipMemGetInfo_ParaSmall") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + unsigned int* A_mem{nullptr}; + // allocate smaller chunk than minimum + size_t Malloc1Size = 1; + + std::thread t1( + [&] { HIP_CHECK_THREAD(hipMalloc(reinterpret_cast(&A_mem), Malloc1Size)) }); + t1.join(); + HIP_CHECK_THREAD_FINALIZE(); + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + MEMINFO(totalMemRet, freeMemInit, freeMemRet, Malloc1Size); + + + auto assumedFreeMem = freeMemInit - Malloc1Size; + // Free memory should be less than assumed for + // single allocation smaller than min allocation chunk + REQUIRE(freeMemRet <= assumedFreeMem); + + HIP_CHECK(hipFree(A_mem)); +} + + +TEST_CASE("Unit_hipMemGetInfo_Negative") { + size_t freeMemInit; + size_t totalMemInit; + HIP_CHECK(hipMemGetInfo(&freeMemInit, &totalMemInit)); + + unsigned int* A_mem{nullptr}; + auto MallocSize = MinAlloc::Get(); + + SECTION("Zero allocation") { + size_t freeMemRet; + size_t totalMemRet; + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), 0)); + HIP_CHECK(hipMemGetInfo(&freeMemRet, &totalMemRet)); + + REQUIRE(freeMemRet == freeMemInit); + } + SECTION("Nullptr as first param passed to hipMemGetInfo") { + size_t* freeMemRet = nullptr; + size_t totalMemRet; + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); + // Segfaults on AMD and returns hipSuccess on Nvidia + HIP_CHECK(hipMemGetInfo(freeMemRet, &totalMemRet)); + } + SECTION("Nullptr as second param passed to hipMemGetInfo") { + size_t freeMemRet; + size_t* totalMemRet = nullptr; + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); + // Segfaults on AMD and returns hipSuccess on Nvidia + HIP_CHECK(hipMemGetInfo(&freeMemRet, totalMemRet)); + } + SECTION("Nullptr as both params passed to hipMemGetInfo") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-135"); + return; +#endif + size_t* freeMemRet = nullptr; + size_t* totalMemRet = nullptr; + HIP_CHECK(hipMalloc(reinterpret_cast(&A_mem), MallocSize)); + // Segfaults on AMD and returns hipSuccess on Nvidia + HIP_CHECK(hipMemGetInfo(freeMemRet, totalMemRet)); + } + + HIP_CHECK(hipFree(A_mem)); +} From 51a945bcfb202809241f938f44a67747eff32db8 Mon Sep 17 00:00:00 2001 From: Finlay Date: Thu, 28 Jul 2022 12:32:59 +0100 Subject: [PATCH 2/9] EXSWCPHIPT-136 Texture Gather tests for hipMalloc3DArray and hipArray3DCreate (#2818) --- tests/catch/unit/memory/hipArray3DCreate.cc | 62 ++++++++++++++++----- tests/catch/unit/memory/hipMalloc3DArray.cc | 43 +++++++++++--- 2 files changed, 82 insertions(+), 23 deletions(-) diff --git a/tests/catch/unit/memory/hipArray3DCreate.cc b/tests/catch/unit/memory/hipArray3DCreate.cc index faa3822b9b..9c11152936 100644 --- a/tests/catch/unit/memory/hipArray3DCreate.cc +++ b/tests/catch/unit/memory/hipArray3DCreate.cc @@ -58,15 +58,17 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, #if HT_AMD desc.Flags = 0; #else - desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); + desc.Flags = GENERATE(0, hipArraySurfaceLoadStore, hipArrayTextureGather); #endif constexpr size_t size = 64; - std::vector extents{ - {size, 0, 0}, // 1D array - {size, size, 0}, // 2D array - {size, size, size} // 3D array + std::vector extents; + extents.reserve(3); + extents.push_back({size, size, 0}); // 2D array + if (desc.Flags != hipArrayTextureGather) { + extents.push_back({size, 0, 0}); // 1D array + extents.push_back({size, size, size}); // 3D array }; for (auto& extent : extents) { @@ -100,8 +102,8 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, us #if HT_AMD desc.Flags = 0; #else - desc.Flags = GENERATE(0, CUDA_ARRAY3D_SURFACE_LDST); - if (desc.Flags == CUDA_ARRAY3D_SURFACE_LDST) { + desc.Flags = GENERATE(0, hipArraySurfaceLoadStore); + if (desc.Flags == hipArraySurfaceLoadStore) { HipTest::HIP_SKIP_TEST("EXSWCPHIPT-58"); return; } @@ -200,7 +202,7 @@ constexpr HIP_ARRAY3D_DESCRIPTOR defaultDescriptor(unsigned int flags, size_t si desc.Depth = size; #if HT_NVIDIA - if (flags == CUDA_ARRAY3D_TEXTURE_GATHER) { + if (flags == hipArrayTextureGather) { desc.Depth = 0; } #endif @@ -247,8 +249,8 @@ TEST_CASE("Unit_hipArray3DCreate_Negative_ZeroHeight") { unsigned int flags = GENERATE(from_range(std::begin(validFlags), std::end(validFlags))); auto desc = defaultDescriptor(flags, 6); #if HT_NVIDIA - std::array exceptions{CUDA_ARRAY3D_LAYERED, - CUDA_ARRAY3D_LAYERED | CUDA_ARRAY3D_SURFACE_LDST}; + std::array exceptions{hipArrayLayered, + hipArrayLayered | hipArraySurfaceLoadStore}; #else std::array exceptions{}; #endif @@ -290,10 +292,9 @@ TEST_CASE("Unit_hipArray3DCreate_Negative_InvalidFlags") { // FIXME: use the same flags for both tests when the values exist for hip #if HT_NVIDIA - unsigned int flags = - GENERATE(0xDEADBEEF, CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST, - CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_CUBEMAP, - CUDA_ARRAY3D_TEXTURE_GATHER | CUDA_ARRAY3D_SURFACE_LDST | CUDA_ARRAY3D_CUBEMAP); + unsigned int flags = GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore, + hipArrayTextureGather | hipArrayCubemap, + hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap); #else unsigned int flags = 0xDEADBEEF; #endif @@ -319,3 +320,36 @@ TEST_CASE("Unit_hipArray3DCreate_Negative_NumericLimit") { testInvalidDescription(desc); } + +// texture gather arrays may only be 2D +TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_Negative_Non2DTextureGather", "", char, uint2, int4, + float2, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); + return; +#endif + using vec_info = vector_info; + DriverContext ctx; + + HIP_ARRAY3D_DESCRIPTOR desc{}; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + desc.Flags = hipArrayTextureGather; + + constexpr size_t size = 64; + + std::array extents{ + make_hipExtent(size, 0, 0), // 1D array + make_hipExtent(size, size, size), // 3D array + }; + + for (auto& extent : extents) { + desc.Width = extent.width; + desc.Height = extent.height; + desc.Depth = extent.depth; + + CAPTURE(desc.Width, desc.Height, desc.Depth); + + testInvalidDescription(desc); + } +} diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index 1a8ec2f0d0..1b7ee27171 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -136,18 +136,26 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, #if HT_AMD const unsigned int flags = hipArrayDefault; #else - const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); + const unsigned int flags = + GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); #endif constexpr size_t size = 64; - hipExtent extent; - SECTION("1D Array") { extent = make_hipExtent(size, 0, 0); } - SECTION("2D Array") { extent = make_hipExtent(size, size, 0); } - SECTION("3D Array") { extent = make_hipExtent(size, size, size); } + std::vector extents; + extents.reserve(3); + extents.push_back({size, size, 0}); // 2D array + if (flags != hipArrayTextureGather) { + extents.push_back({size, 0, 0}); // 1D array + extents.push_back({size, size, size}); // 3D array + }; - HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); - checkArrayIsExpected(array, desc, extent, flags); - HIP_CHECK(hipFreeArray(array)); + for (const auto extent : extents) { + CAPTURE(flags, extent.width, extent.height, extent.depth); + + HIP_CHECK(hipMalloc3DArray(&array, &desc, extent, flags)); + checkArrayIsExpected(array, desc, extent, flags); + HIP_CHECK(hipFreeArray(array)); + } } TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, ushort2, @@ -316,7 +324,7 @@ TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") { HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, makeExtent(flag, s), flag), hipErrorInvalidValue); } -void testInvalidDescription(hipChannelFormatDesc desc){ +void testInvalidDescription(hipChannelFormatDesc desc) { constexpr size_t s = 6; // 6 to keep cubemap happy hipArray_t array; @@ -420,3 +428,20 @@ TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") { HIP_CHECK_ERROR(hipMalloc3DArray(&arrayPtr, &desc, makeExtent(flag, size), flag), hipErrorInvalidValue); } + +// texture gather arrays are only allowed to be 2D +TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_Negative_Non2DTextureGather", "", char, uchar2, short4, + float2, float4) { +#if HT_AMD + HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); + return; +#endif + hipArray_t array; + const auto desc = hipCreateChannelDesc(); + + constexpr unsigned int flags = hipArrayTextureGather; + constexpr size_t size = 64; + const hipExtent extent = GENERATE(make_hipExtent(size, 0, 0), make_hipExtent(size, size, size)); + + HIP_CHECK_ERROR(hipMalloc3DArray(&array, &desc, extent, flags), hipErrorInvalidValue); +} From a922cf68ac628e7a26b4801171aee387eff1354d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Thu, 28 Jul 2022 23:23:45 +0100 Subject: [PATCH 3/9] Fix bug with hipMemsetFunctional tests (#2828) --- .../catch/unit/memory/hipMemsetFunctional.cc | 49 +++++++++---------- 1 file changed, 24 insertions(+), 25 deletions(-) diff --git a/tests/catch/unit/memory/hipMemsetFunctional.cc b/tests/catch/unit/memory/hipMemsetFunctional.cc index 274074266d..d6f152cec9 100644 --- a/tests/catch/unit/memory/hipMemsetFunctional.cc +++ b/tests/catch/unit/memory/hipMemsetFunctional.cc @@ -203,31 +203,30 @@ void partialMemsetTest(T valA, T valB, size_t count, size_t offset, MemsetType m } TEST_CASE("Unit_hipMemsetFunctional_PartialSet_1D") { - for (auto widthOffset = 8; widthOffset <= 8; widthOffset *= 2) { - SECTION("hipMemset - Partial Set") { - partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, false); - } - SECTION("hipMemsetAsync - Partial Set") { - partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, true); - } - SECTION("hipMemsetD8 - Partial Set") { - partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, false); - } - SECTION("hipMemsetD8Async - Partial Set") { - partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, true); - } - SECTION("hipMemsetD16 - Partial Set") { - partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, false); - } - SECTION("hipMemsetD16Async - Partial Set") { - partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, true); - } - SECTION("hipMemsetD32 - Partial Set") { - partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, false); - } - SECTION("hipMemsetD32Async - Partial Set") { - partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, true); - } + auto widthOffset = GENERATE(8, 16, 32, 64, 128, 256, 512, 1024); + SECTION("hipMemset - Partial Set") { + partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, false); + } + SECTION("hipMemsetAsync - Partial Set") { + partialMemsetTest(0x1, 0x42, 1024, widthOffset, hipMemsetTypeDefault, true); + } + SECTION("hipMemsetD8 - Partial Set") { + partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, false); + } + SECTION("hipMemsetD8Async - Partial Set") { + partialMemsetTest(0x1, 0xDE, 1024, widthOffset, hipMemsetTypeD8, true); + } + SECTION("hipMemsetD16 - Partial Set") { + partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, false); + } + SECTION("hipMemsetD16Async - Partial Set") { + partialMemsetTest(0x1, 0xDEAD, 1024, widthOffset, hipMemsetTypeD16, true); + } + SECTION("hipMemsetD32 - Partial Set") { + partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, false); + } + SECTION("hipMemsetD32Async - Partial Set") { + partialMemsetTest(0x1, 0xDEADBEEF, 1024, widthOffset, hipMemsetTypeD32, true); } } From 1873df7bddc51aeeb11a01ac06c76f73efaf889d Mon Sep 17 00:00:00 2001 From: Dylan Angus <61192377+dylan-angus-codeplay@users.noreply.github.com> Date: Fri, 29 Jul 2022 04:05:27 +0100 Subject: [PATCH 4/9] Extending hipMallocManaged tests (#2670) * Extending hipMallocManaged tests * Fixed compilation error * Added tests skips for hipMallocManaged tests on devices that don't support managed memory * Removed unused stream --- tests/catch/include/hip_test_common.hh | 2 +- tests/catch/unit/memory/hipMallocManaged.cc | 194 ++++---- .../unit/memory/hipMallocManagedCommon.hh | 26 ++ .../unit/memory/hipMallocManagedFlagsTst.cc | 417 ++++++++--------- .../memory/hipMallocManaged_MultiScenario.cc | 212 ++++++--- .../unit/memory/hipMallocMngdMultiThread.cc | 419 ++++++++---------- 6 files changed, 636 insertions(+), 634 deletions(-) create mode 100644 tests/catch/unit/memory/hipMallocManagedCommon.hh diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index e97c99bd8f..0fb79a49b3 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -166,7 +166,7 @@ static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlo HIP_CHECK(hipGetDeviceProperties(&props, device)); unsigned blocks = props.multiProcessorCount * blocksPerCU; - if (blocks * threadsPerBlock > N) { + if (blocks * threadsPerBlock < N) { blocks = (N + threadsPerBlock - 1) / threadsPerBlock; } diff --git a/tests/catch/unit/memory/hipMallocManaged.cc b/tests/catch/unit/memory/hipMallocManaged.cc index d40aa924a1..672db36485 100644 --- a/tests/catch/unit/memory/hipMallocManaged.cc +++ b/tests/catch/unit/memory/hipMallocManaged.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 @@ -24,15 +24,14 @@ only on HMM enabled devices */ -#include +#include "hipMallocManagedCommon.hh" #include #include - // Kernel functions -__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { +__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) { @@ -40,7 +39,7 @@ __global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { } } -__global__ void KernelMulAdd_MngdMem(int *Hmm, size_t n) { +__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) { @@ -48,130 +47,93 @@ __global__ void KernelMulAdd_MngdMem(int *Hmm, size_t n) { } } -__global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *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 int HmmAttrPrint() { - int managed = 0; - INFO("The following are the attribute values related to HMM for" - " device 0:\n"); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); - INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeConcurrentManagedAccess, 0)); - INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccess, 0)); - INFO("hipDeviceAttributePageableMemoryAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); - INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" - << managed); - - HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, - 0)); - INFO("hipDeviceAttributeManagedMemory: " << managed); - return managed; -} - - - -static size_t N{4 * 1024 * 1024}; +static size_t numElements{64 * 1024 * 1024}; static unsigned blocksPerCU{6}; static unsigned threadsPerBlock{256}; /* This testcase verifies the hipMallocManaged basic scenario - supported on all devices */ - TEST_CASE("Unit_hipMallocManaged_Basic") { - int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N; - float *A, *B, *C; + auto managed = HmmAttrPrint(); + if (managed != 1) { + WARN( + "GPU doesn't support hipDeviceAttributeManagedMemory attribute so defaulting to system " + "memory."); + } - HIP_CHECK(hipMallocManaged(&A, numElements*sizeof(float))); - HIP_CHECK(hipMallocManaged(&B, numElements*sizeof(float))); - HIP_CHECK(hipMallocManaged(&C, numElements*sizeof(float))); + float *A, *B, *C; + + HIP_CHECK(hipMallocManaged(&A, numElements * sizeof(float))); + HIP_CHECK(hipMallocManaged(&B, numElements * sizeof(float))); + HIP_CHECK(hipMallocManaged(&C, numElements * sizeof(float))); } /* - This testcase verifies the hipMallocManaged basic scenario - supported only on HMM enabled devices + This testcase verifies the hipMallocManaged advanced scenario - supported only on HMM enabled + devices */ - TEST_CASE("Unit_hipMallocManaged_Advanced") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N; - float *A, *B, *C; - - HIP_CHECK(hipMallocManaged(&A, numElements*sizeof(float))); - HIP_CHECK(hipMallocManaged(&B, numElements*sizeof(float))); - HIP_CHECK(hipMallocManaged(&C, numElements*sizeof(float))); - HipTest::setDefaultData(numElements, A, B, C); - - hipDevice_t device = hipCpuDeviceId; - - HIP_CHECK(hipMemAdvise(A, numElements*sizeof(float), - hipMemAdviseSetReadMostly, device)); - HIP_CHECK(hipMemPrefetchAsync(A, numElements*sizeof(float), 0)); - HIP_CHECK(hipMemPrefetchAsync(B, numElements*sizeof(float), 0)); - HIP_CHECK(hipDeviceSynchronize()); - HIP_CHECK(hipMemRangeGetAttribute(&device, sizeof(device), - hipMemRangeAttributeLastPrefetchLocation, - A, numElements*sizeof(float))); - if (device != 0) { - INFO("hipMemRangeGetAttribute error, device = " << device); - } - uint32_t read_only = 0xf; - HIP_CHECK(hipMemRangeGetAttribute(&read_only, sizeof(read_only), - hipMemRangeAttributeReadMostly, - A, numElements*sizeof(float))); - if (read_only != 1) { - SUCCEED("hipMemRangeGetAttribute error, read_only = " << read_only); - } - - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, - numElements); - hipEvent_t event0, event1; - HIP_CHECK(hipEventCreate(&event0)); - HIP_CHECK(hipEventCreate(&event1)); - HIP_CHECK(hipEventRecord(event0, 0)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), - 0, 0, static_cast(A), - static_cast(B), C, numElements); - HIP_CHECK(hipEventRecord(event1, 0)); - HIP_CHECK(hipDeviceSynchronize()); - float time = 0.0f; - HIP_CHECK(hipEventElapsedTime(&time, event0, event1)); - printf("Time %.3f ms\n", time); - float maxError = 0.0f; - HIP_CHECK(hipMemPrefetchAsync(B, numElements*sizeof(float), - hipCpuDeviceId)); - HIP_CHECK(hipDeviceSynchronize()); - device = 0; - HIP_CHECK(hipMemRangeGetAttribute(&device, sizeof(device), - hipMemRangeAttributeLastPrefetchLocation, - A, numElements*sizeof(float))); - if (device != hipCpuDeviceId) { - SUCCEED("hipMemRangeGetAttribute error device = " << device); - } - - for (int i = 0; i < numElements; i++) { - maxError = fmax(maxError, fabs(B[i]-3.0f)); - } - HIP_CHECK(hipFree(A)); - HIP_CHECK(hipFree(B)); - REQUIRE(maxError != 0.0f); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; } -} + float *A, *B, *C; + + HIP_CHECK(hipMallocManaged(&A, numElements * sizeof(float))); + HIP_CHECK(hipMallocManaged(&B, numElements * sizeof(float))); + HIP_CHECK(hipMallocManaged(&C, numElements * sizeof(float))); + HipTest::setDefaultData(numElements, A, B, C); + + hipDevice_t device = hipCpuDeviceId; + + HIP_CHECK(hipMemAdvise(A, numElements * sizeof(float), hipMemAdviseSetReadMostly, device)); + HIP_CHECK(hipMemPrefetchAsync(A, numElements * sizeof(float), 0)); + HIP_CHECK(hipMemPrefetchAsync(B, numElements * sizeof(float), 0)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemRangeGetAttribute(&device, sizeof(device), + hipMemRangeAttributeLastPrefetchLocation, A, + numElements * sizeof(float))); + if (device != 0) { + INFO("hipMemRangeGetAttribute error, device = " << device); + } + uint32_t read_only = 0xf; + HIP_CHECK(hipMemRangeGetAttribute(&read_only, sizeof(read_only), hipMemRangeAttributeReadMostly, + A, numElements * sizeof(float))); + if (read_only != 1) { + SUCCEED("hipMemRangeGetAttribute error, read_only = " << read_only); + } + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); + hipEvent_t event0, event1; + HIP_CHECK(hipEventCreate(&event0)); + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventRecord(event0, 0)); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, + static_cast(A), static_cast(B), C, numElements); + HIP_CHECK(hipEventRecord(event1, 0)); + HIP_CHECK(hipDeviceSynchronize()); + float time = 0.0f; + HIP_CHECK(hipEventElapsedTime(&time, event0, event1)); + printf("Time %.3f ms\n", time); + float maxError = 0.0f; + HIP_CHECK(hipMemPrefetchAsync(B, numElements * sizeof(float), hipCpuDeviceId)); + HIP_CHECK(hipDeviceSynchronize()); + device = 0; + HIP_CHECK(hipMemRangeGetAttribute(&device, sizeof(device), + hipMemRangeAttributeLastPrefetchLocation, A, + numElements * sizeof(float))); + if (device != hipCpuDeviceId) { + SUCCEED("hipMemRangeGetAttribute error device = " << device); + } + + for (size_t i = 0; i < numElements; i++) { + maxError = fmax(maxError, fabs(B[i] - 3.0f)); + } + HIP_CHECK(hipFree(A)); + HIP_CHECK(hipFree(B)); + REQUIRE(maxError != 0.0f); +} diff --git a/tests/catch/unit/memory/hipMallocManagedCommon.hh b/tests/catch/unit/memory/hipMallocManagedCommon.hh new file mode 100644 index 0000000000..cb6c6c7eee --- /dev/null +++ b/tests/catch/unit/memory/hipMallocManagedCommon.hh @@ -0,0 +1,26 @@ +#include + +static int HmmAttrPrint() { + int managed = 0; + INFO( + "The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeConcurrentManagedAccess, 0)); + INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributePageableMemoryAccess, 0)); + INFO("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK( + hipDeviceGetAttribute(&managed, hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" << managed); + + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0)); + INFO("hipDeviceAttributeManagedMemory: " << managed); + if (managed != 1) { + WARN( + "GPU 0 doesn't support hipDeviceAttributeManagedMemory attribute so defaulting to system " + "memory."); + } + return managed; +} \ No newline at end of file diff --git a/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc b/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc index 8ac1c37853..264320c483 100644 --- a/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc +++ b/tests/catch/unit/memory/hipMallocManagedFlagsTst.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 @@ -20,245 +20,210 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +#include "hipMallocManagedCommon.hh" #include // Kernel function -__global__ void MallcMangdFlgTst(int n, float *x, float *y) { +__global__ void MallcMangdFlgTst(int n, float* x, float* y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) - y[i] = x[i] * x[i]; -} - - -// The following function prints info on attributes related to HMM -static int HmmAttrPrint() { - int managed = 0; - INFO("The following are the attribute values related to HMM for" - " device 0:\n"); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); - INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeConcurrentManagedAccess, 0)); - INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccess, 0)); - INFO("hipDeviceAttributePageableMemoryAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); - INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" - << managed); - - HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, - 0)); - INFO("hipDeviceAttributeManagedMemory: " << managed); - return managed; + for (int i = index; i < n; i += stride) y[i] = x[i] * x[i]; } // The following section tests working of hipMallocManaged with flag parameters TEST_CASE("Unit_hipMallocManaged_FlgParam") { - int managed = HmmAttrPrint(); - if (managed == 1) { - std::atomic DataMismatch{0}; - bool IfTestPassed = true; - float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5; - int NumDevs = 0, NUM_ELMS = 4096; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - float *Ad = NULL, *Ah = NULL; - Ah = new float[NUM_ELMS]; - // Testing hipMemAttachGlobal Flag - HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), - hipMemAttachGlobal)); - - // Initializing HmmAG memory - for (int i = 0; i < NUM_ELMS; i++) { - HmmAG[i] = INIT_VAL; - Ah[i] = 0; - } - - int blockSize = 256; - int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; - dim3 dimGrid(numBlocks, 1, 1); - dim3 dimBlock(blockSize, 1, 1); - hipStream_t strm; - for (int i = 0; i < NumDevs; i++) { - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float))); - HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float))); - MallcMangdFlgTst<<>>(NUM_ELMS, HmmAG, Ad); - HIP_CHECK(hipStreamSynchronize(strm)); - HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), - hipMemcpyDeviceToHost)); - for (int j = 0; j < NUM_ELMS; ++j) { - if (Ah[j] != (INIT_VAL * INIT_VAL)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("Data Mismatch observed when kernel launched on"); - WARN(" device: " << i); - IfTestPassed = false; - } - DataMismatch = 0; - - HIP_CHECK(hipFree(Ad)); - HIP_CHECK(hipStreamDestroy(strm)); - } - delete[] Ah; - HIP_CHECK(hipFree(HmmAG)); - - DataMismatch = 0; - HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), - hipMemAttachHost)); - HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), - hipMemAttachHost)); - - // Initializing HmmAH memory - for (int i = 0; i < NUM_ELMS; i++) { - HmmAH1[i] = INIT_VAL; - HmmAH2[i] = 0; - } - for (int i = 0; i < NumDevs; i++) { - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float))); - MallcMangdFlgTst<<>>(NUM_ELMS, - HmmAH1, HmmAH2); - HIP_CHECK(hipStreamSynchronize(strm)); - for (int j = 0; j < NUM_ELMS; ++j) { - if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("Data Mismatch observed when kernel launched on"); - WARN(" device: " << i); - IfTestPassed = false; - } - HIP_CHECK(hipStreamDestroy(strm)); - } - HIP_CHECK(hipFree(HmmAH1)); - HIP_CHECK(hipFree(HmmAH2)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; } + + std::atomic DataMismatch{0}; + bool IfTestPassed = true; + float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5; + int NumDevs = 0, NUM_ELMS = 4096; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + float *Ad = NULL, *Ah = NULL; + Ah = new float[NUM_ELMS]; + // Testing hipMemAttachGlobal Flag + HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), hipMemAttachGlobal)); + + // Initializing HmmAG memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAG[i] = INIT_VAL; + Ah[i] = 0; + } + + int blockSize = 256; + int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + hipStream_t strm; + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float))); + HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, HmmAG, Ad); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), hipMemcpyDeviceToHost)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (Ah[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + DataMismatch = 0; + + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipStreamDestroy(strm)); + } + delete[] Ah; + HIP_CHECK(hipFree(HmmAG)); + + DataMismatch = 0; + HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), hipMemAttachHost)); + HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), hipMemAttachHost)); + + // Initializing HmmAH memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAH1[i] = INIT_VAL; + HmmAH2[i] = 0; + } + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, HmmAH1, HmmAH2); + HIP_CHECK(hipStreamSynchronize(strm)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + HIP_CHECK(hipStreamDestroy(strm)); + } + HIP_CHECK(hipFree(HmmAH1)); + HIP_CHECK(hipFree(HmmAH2)); + REQUIRE(IfTestPassed); } // The following function tests Memory access allocated using hipMallocManaged // in multiple streams TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") { - int managed = HmmAttrPrint(); - if (managed == 1) { - std::atomic DataMismatch{0}; - bool IfTestPassed = true; - float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5; - int NumStrms = 0, MultiDevice = 0, NUM_ELMS = 4096; - HIP_CHECK(hipGetDeviceCount(&MultiDevice)); - if (MultiDevice >= 2) { - HIP_CHECK(hipGetDeviceCount(&NumStrms)); - } else { - NumStrms = 4; - } - hipStream_t **Stream = new hipStream_t*[NumStrms]; - for (int i = 0; i < NumStrms; ++i) { - Stream[i] = reinterpret_cast(malloc(sizeof(hipStream_t))); - } - float *Ad = NULL, *Ah = NULL; - Ah = new float[NUM_ELMS]; - for (int i = 0; i < NumStrms; ++i) { - if (MultiDevice >= 2) { - HIP_CHECK(hipSetDevice(i)); - } - HIP_CHECK(hipStreamCreate(Stream[i])); - } - HIP_CHECK(hipSetDevice(0)); - // Testing hipMemAttachGlobal Flag - HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), - hipMemAttachGlobal)); - - // Initializing HmmAG memory - for (int i = 0; i < NUM_ELMS; i++) { - HmmAG[i] = INIT_VAL; - Ah[i] = 0; - } - - int blockSize = 256; - int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; - dim3 dimGrid(numBlocks, 1, 1); - dim3 dimBlock(blockSize, 1, 1); - for (int i = 0; i < NumStrms; i++) { - if (MultiDevice >= 2) { - HIP_CHECK(hipSetDevice(i)); - } - HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float))); - HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float))); - MallcMangdFlgTst<<>>(NUM_ELMS, - HmmAG, Ad); - HIP_CHECK(hipStreamSynchronize(*(Stream[i]))); - // Validating the results - HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), - hipMemcpyDeviceToHost)); - for (int j = 0; j < NUM_ELMS; ++j) { - if (Ah[j] != (INIT_VAL * INIT_VAL)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("Data Mismatch observed when kernel launched on"); - WARN(" device: " << i); - IfTestPassed = false; - } - DataMismatch = 0; - - HIP_CHECK(hipFree(Ad)); - } - delete[] Ah; - HIP_CHECK(hipFree(HmmAG)); - - DataMismatch = 0; - HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), - hipMemAttachHost)); - HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), - hipMemAttachHost)); - - // Initializing HmmAH memory - for (int i = 0; i < NUM_ELMS; i++) { - HmmAH1[i] = INIT_VAL; - HmmAH2[i] = 0; - } - for (int i = 0; i < NumStrms; i++) { - if (MultiDevice >= 2) { - HIP_CHECK(hipSetDevice(i)); - } - HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float))); - MallcMangdFlgTst<<>>(NUM_ELMS, - HmmAH1, HmmAH2); - HIP_CHECK(hipStreamSynchronize(*(Stream[i]))); - for (int j = 0; j < NUM_ELMS; ++j) { - if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) { - DataMismatch++; - break; - } - } - if (DataMismatch != 0) { - WARN("Data Mismatch observed when kernel launched on"); - WARN(" device: " << i); - IfTestPassed = false; - } - } - - HIP_CHECK(hipFree(HmmAH1)); - HIP_CHECK(hipFree(HmmAH2)); - for (int i = 0; i < NumStrms; ++i) { - HIP_CHECK(hipStreamDestroy(*(Stream[i]))); - } - REQUIRE(IfTestPassed); - } else { - SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; } -} + + std::atomic DataMismatch{0}; + bool IfTestPassed = true; + float *HmmAG = NULL, *HmmAH1 = NULL, *HmmAH2 = NULL, INIT_VAL = 2.5; + int NumStrms = 0, MultiDevice = 0, NUM_ELMS = 4096; + HIP_CHECK(hipGetDeviceCount(&MultiDevice)); + if (MultiDevice >= 2) { + HIP_CHECK(hipGetDeviceCount(&NumStrms)); + } else { + NumStrms = 4; + } + hipStream_t** Stream = new hipStream_t*[NumStrms]; + for (int i = 0; i < NumStrms; ++i) { + Stream[i] = reinterpret_cast(malloc(sizeof(hipStream_t))); + } + float *Ad = NULL, *Ah = NULL; + Ah = new float[NUM_ELMS]; + for (int i = 0; i < NumStrms; ++i) { + if (MultiDevice >= 2) { + HIP_CHECK(hipSetDevice(i)); + } + HIP_CHECK(hipStreamCreate(Stream[i])); + } + HIP_CHECK(hipSetDevice(0)); + // Testing hipMemAttachGlobal Flag + HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), hipMemAttachGlobal)); + // Initializing HmmAG memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAG[i] = INIT_VAL; + Ah[i] = 0; + } + + int blockSize = 256; + int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + for (int i = 0; i < NumStrms; i++) { + if (MultiDevice >= 2) { + HIP_CHECK(hipSetDevice(i)); + } + HIP_CHECK(hipMalloc(&Ad, NUM_ELMS * sizeof(float))); + HIP_CHECK(hipMemset(Ad, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, HmmAG, Ad); + HIP_CHECK(hipStreamSynchronize(*(Stream[i]))); + // Validating the results + HIP_CHECK(hipMemcpy(Ah, Ad, NUM_ELMS * sizeof(float), hipMemcpyDeviceToHost)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (Ah[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + DataMismatch = 0; + + HIP_CHECK(hipFree(Ad)); + } + delete[] Ah; + HIP_CHECK(hipFree(HmmAG)); + + DataMismatch = 0; + HIP_CHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), hipMemAttachHost)); + HIP_CHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), hipMemAttachHost)); + + // Initializing HmmAH memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAH1[i] = INIT_VAL; + HmmAH2[i] = 0; + } + for (int i = 0; i < NumStrms; i++) { + if (MultiDevice >= 2) { + HIP_CHECK(hipSetDevice(i)); + } + HIP_CHECK(hipMemset(HmmAH2, 0, NUM_ELMS * sizeof(float))); + MallcMangdFlgTst<<>>(NUM_ELMS, HmmAH1, HmmAH2); + HIP_CHECK(hipStreamSynchronize(*(Stream[i]))); + for (int j = 0; j < NUM_ELMS; ++j) { + if (HmmAH2[j] != (INIT_VAL * INIT_VAL)) { + DataMismatch++; + break; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on"); + WARN(" device: " << i); + IfTestPassed = false; + } + } + + HIP_CHECK(hipFree(HmmAH1)); + HIP_CHECK(hipFree(HmmAH2)); + for (int i = 0; i < NumStrms; ++i) { + HIP_CHECK(hipStreamDestroy(*(Stream[i]))); + } + REQUIRE(IfTestPassed); +} diff --git a/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc b/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc index bf67d70765..976a4397fc 100644 --- a/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc +++ b/tests/catch/unit/memory/hipMallocManaged_MultiScenario.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 @@ -27,28 +27,81 @@ 6. Multiple Pointers */ -#include +#include "hipMallocManagedCommon.hh" #include #include #include const size_t MAX_GPU{256}; -static size_t N{4*1024*1024}; +static size_t N{4 * 1024 * 1024}; +static unsigned blocksPerCU{6}; +static unsigned threadsPerBlock{256}; #define INIT_VAL 123 /* * Kernel function to perform addition operation. */ -template -__global__ void -vector_sum(T *Ad1, T *Ad2, size_t NUM_ELMTS) { - size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); - size_t stride = blockDim.x * gridDim.x; +template __global__ void vector_sum(T* Ad1, T* Ad2, size_t NUM_ELMTS) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; - for (size_t i = offset; i < NUM_ELMTS; i += stride) { - Ad2[i] = Ad1[i] + Ad1[i]; - } + for (size_t i = offset; i < NUM_ELMTS; i += stride) { + Ad2[i] = Ad1[i] + Ad1[i]; + } +} + +/* + * Kernel function to perform multiplication + */ +__global__ void KernelDouble(float* Hmm, float* dPtr, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < n) { + dPtr[index] = 2 * Hmm[index]; + } +} + +/* + * Host function to perform multiplication + */ +void HostKernelDouble(float* Hmm, float* hPtr, size_t n) { + for (size_t i = 0; i < n; i++) { + hPtr[i] = 2 * Hmm[i]; + } +} + +/* + This testcase verifies the concurrent access of hipMallocManaged Memory on host and device. + */ +TEST_CASE("Unit_hipMallocManaged_HostDeviceConcurrent") { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + + float *Hmm = nullptr, *hPtr = nullptr, *dPtr = nullptr, *resPtr = nullptr; + + hPtr = reinterpret_cast(malloc(N * sizeof(float))); + resPtr = reinterpret_cast(malloc(N * sizeof(float))); + + HIP_CHECK(hipMalloc(&dPtr, N * sizeof(float))); + HIP_CHECK(hipMallocManaged(&Hmm, N * sizeof(float))); + memset(Hmm, 2.0, N * sizeof(float)); + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + std::thread host_thread(HostKernelDouble, Hmm, hPtr, N); + KernelDouble<<>>(Hmm, dPtr, N); + host_thread.join(); + hipMemcpy(resPtr, dPtr, N * sizeof(float), hipMemcpyDeviceToHost); + + for (size_t i = 0; i < N; i++) { + REQUIRE(hPtr[i] == resPtr[i]); + } + + free(hPtr); + HIP_CHECK(hipFree(dPtr)); + HIP_CHECK(hipFree(Hmm)); } // The following Test case tests the following scenario: @@ -57,7 +110,13 @@ vector_sum(T *Ad1, T *Ad2, size_t NUM_ELMTS) { // kernel is launched on acessed chunk of hmm memory // and checks if there are any inconsistencies or access issues TEST_CASE("Unit_hipMallocManaged_MultiChunkSingleDevice") { -std::atomic DataMismatch{0}; + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + + std::atomic DataMismatch{0}; constexpr int Chunks = 4; int Counter = 0; int NUM_ELMS = (1024 * 1024); @@ -74,16 +133,14 @@ std::atomic DataMismatch{0}; Hmm[Counter] = (INIT_VAL + i); } } - const unsigned threadsPerBlock = 256; - const unsigned blocks = (NUM_ELMS + 255)/256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); for (int k = 0; k < Chunks; ++k) { - vector_sum <<>> - (&Hmm[k * NUM_ELMS], Ad[k], NUM_ELMS); + vector_sum + <<>>(&Hmm[k * NUM_ELMS], Ad[k], NUM_ELMS); } HIP_CHECK(hipDeviceSynchronize()); for (int m = 0; m < Chunks; ++m) { - HIP_CHECK(hipMemcpy(Ah, Ad[m], NUM_ELMS * sizeof(float), - hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Ah, Ad[m], NUM_ELMS * sizeof(float), hipMemcpyDeviceToHost)); for (int n = 0; n < NUM_ELMS; ++n) { if (Ah[n] != ((INIT_VAL + m) * 2)) { DataMismatch++; @@ -96,7 +153,7 @@ std::atomic DataMismatch{0}; HIP_CHECK(hipStreamDestroy(stream[i])); } HIP_CHECK(hipFree(Hmm)); - delete [] Ah; + delete[] Ah; } // The following Test case tests the following scenario: @@ -105,10 +162,20 @@ std::atomic DataMismatch{0}; // kernel is launched on acessed chunk of hmm memory // and checks if there are any inconsistencies or access issues TEST_CASE("Unit_hipMallocManaged_MultiChunkMultiDevice") { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + std::atomic DataMismatch{0}; int Counter = 0; int NumDevices = 0; HIP_CHECK(hipGetDeviceCount(&NumDevices)); + if (NumDevices < 2) { + HipTest::HIP_SKIP_TEST("Skipping test because more than one device was not found."); + return; + } unsigned int NUM_ELMS = (1024 * 1024); float *Ad[MAX_GPU], *Hmm = NULL, *Ah = new float[NUM_ELMS]; hipStream_t stream[MAX_GPU]; @@ -124,17 +191,15 @@ TEST_CASE("Unit_hipMallocManaged_MultiChunkMultiDevice") { Hmm[Counter] = INIT_VAL + i; } } - const unsigned threadsPerBlock = 256; - const unsigned blocks = (NUM_ELMS + 255)/256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); for (int Klaunch = 0; Klaunch < NumDevices; ++Klaunch) { HIP_CHECK(hipSetDevice(Klaunch)); - vector_sum <<>> - (&Hmm[Klaunch * NUM_ELMS], Ad[Klaunch], NUM_ELMS); + vector_sum<<>>(&Hmm[Klaunch * NUM_ELMS], + Ad[Klaunch], NUM_ELMS); } HIP_CHECK(hipDeviceSynchronize()); for (int m = 0; m < NumDevices; ++m) { - HIP_CHECK(hipMemcpy(Ah, Ad[m], NUM_ELMS * sizeof(float), - hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Ah, Ad[m], NUM_ELMS * sizeof(float), hipMemcpyDeviceToHost)); for (size_t n = 0; n < NUM_ELMS; ++n) { if (Ah[n] != ((INIT_VAL + m) * 2)) { DataMismatch++; @@ -148,32 +213,38 @@ TEST_CASE("Unit_hipMallocManaged_MultiChunkMultiDevice") { HIP_CHECK(hipStreamDestroy(stream[i])); } HIP_CHECK(hipFree(Hmm)); - delete [] Ah; + delete[] Ah; } // The following tests oversubscription hipMallocManaged() api // Currently disabled. TEST_CASE("Unit_hipMallocManaged_OverSubscription") { - void *A = nullptr; + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + + void* A = nullptr; size_t total = 0, free = 0; HIP_CHECK(hipMemGetInfo(&free, &total)); // ToDo: In case of HMM, memory over-subscription is allowed. Hence, relook // into how out of memory can be tested. // Demanding more mem size than available #if HT_AMD - REQUIRE(hipMallocManaged(&A, (free +1), hipMemAttachGlobal) != hipSuccess); + HIP_CHECK_ERROR(hipMallocManaged(&A, (free + 1), hipMemAttachGlobal), hipErrorOutOfMemory); #endif } // The following test does negative testing of hipMallocManaged() api // by passing invalid values and check if the behavior is as expected TEST_CASE("Unit_hipMallocManaged_Negative") { - void *A; + void* A; size_t total = 0, free = 0; HIP_CHECK(hipMemGetInfo(&free, &total)); SECTION("Nullptr to devPtr") { - REQUIRE(hipMallocManaged(NULL, 1024, hipMemAttachGlobal) != hipSuccess); + HIP_CHECK_ERROR(hipMallocManaged(NULL, 1024, hipMemAttachGlobal), hipErrorInvalidValue); } // cuda api doc says : If size is 0, cudaMallocManaged returns @@ -184,14 +255,14 @@ TEST_CASE("Unit_hipMallocManaged_Negative") { // reset ptr while returning success (to accommodate cuda 11.2 api behavior). SECTION("size 0 with flag hipMemAttachGlobal") { #if HT_AMD - REQUIRE(hipMallocManaged(&A, 0, hipMemAttachGlobal) != hipSuccess); + HIP_CHECK_ERROR(hipMallocManaged(&A, 0, hipMemAttachGlobal), hipErrorInvalidValue); #else - REQUIRE(hipMallocManaged(&A, 0, hipMemAttachHost) == hipSuccess); + HIP_CHECK(hipMallocManaged(&A, 0, hipMemAttachGlobal)); #endif } SECTION("devptr is nullptr with flag hipMemAttachHost") { - REQUIRE(hipMallocManaged(NULL, 1024, hipMemAttachHost) != hipSuccess); + HIP_CHECK_ERROR(hipMallocManaged(NULL, 1024, hipMemAttachHost), hipErrorInvalidValue); } // cuda api doc says : If size is 0, cudaMallocManaged returns @@ -202,32 +273,47 @@ TEST_CASE("Unit_hipMallocManaged_Negative") { // reset ptr while returning success (to accommodate cuda 11.2 api behavior). SECTION("size 0 with flag hipMemAttachHost") { #if HT_AMD - REQUIRE(hipMallocManaged(&A, 0, hipMemAttachHost) != hipSuccess); + HIP_CHECK_ERROR(hipMallocManaged(&A, 0, hipMemAttachHost), hipErrorInvalidValue); #else - REQUIRE(hipMallocManaged(&A, 0, hipMemAttachHost) == hipSuccess); + HIP_CHECK(hipMallocManaged(&A, 0, hipMemAttachHost)); #endif } + SECTION("nullptr to devptr, size 0 and flag 0") { - REQUIRE(hipMallocManaged(NULL, 0, 0) != hipSuccess); + HIP_CHECK_ERROR(hipMallocManaged(NULL, 0, 0), hipErrorInvalidValue); } - SECTION("Numeric value to flag parameter") { - REQUIRE(hipMallocManaged(&A, 1024, 145) != hipSuccess); + SECTION("Invalid flag parameter") { + HIP_CHECK_ERROR(hipMallocManaged(&A, 1024, 145), hipErrorInvalidValue); + } + SECTION("Invalid flag parameter- flag set to 0") { + HIP_CHECK_ERROR(hipMallocManaged(&A, 1024, 0), hipErrorInvalidValue); + } + SECTION("Invalid flag parameter- Both flags set") { + HIP_CHECK_ERROR(hipMallocManaged(&A, 1024, hipMemAttachGlobal | hipMemAttachHost), + hipErrorInvalidValue); } - SECTION("Negative value to size") { - REQUIRE(hipMallocManaged(&A, -10, hipMemAttachGlobal)); + SECTION("Max value to size") { + HIP_CHECK_ERROR(hipMallocManaged(&A, std::numeric_limits::max(), hipMemAttachGlobal), + hipErrorOutOfMemory); } } // Allocate two pointers using hipMallocManaged(), initialize, // then launch kernel using these pointers directly and // later validate the content without using any Memcpy. -TEMPLATE_TEST_CASE("Unit_hipMallocManaged_TwoPointers", "", - int, float, double) { +TEMPLATE_TEST_CASE("Unit_hipMallocManaged_TwoPointers", "", int, float, double) { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + int NumDevices = 0; HIP_CHECK(hipGetDeviceCount(&NumDevices)); TestType *Hmm1 = nullptr, *Hmm2 = nullptr; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); for (int i = 0; i < NumDevices; ++i) { HIP_CHECK(hipSetDevice(i)); @@ -238,10 +324,8 @@ TEMPLATE_TEST_CASE("Unit_hipMallocManaged_TwoPointers", "", Hmm1[m] = m; Hmm2[m] = 0; } - const unsigned threadsPerBlock = 256; - const unsigned blocks = (N + 255)/256; // Kernel launch - vector_sum <<>> (Hmm1, Hmm2, N); + vector_sum<<>>(Hmm1, Hmm2, N); HIP_CHECK(hipDeviceSynchronize()); for (size_t v = 0; v < N; ++v) { if (Hmm2[v] != static_cast(v + v)) { @@ -259,18 +343,28 @@ TEMPLATE_TEST_CASE("Unit_hipMallocManaged_TwoPointers", "", // to all other devices. This include verification and Device two Device // transfers and kernel launch o discover if there any access issues. -TEMPLATE_TEST_CASE("Unit_hipMallocManaged_DeviceContextChange", "", - unsigned char, int, float, double) { +TEMPLATE_TEST_CASE("Unit_hipMallocManaged_DeviceContextChange", "", unsigned char, int, float, + double) { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + std::atomic DataMismatch; - TestType *Ah1 = new TestType[N], *Ah2 = new TestType[N], *Ad = nullptr, - *Hmm = nullptr; + TestType *Ah1 = new TestType[N], *Ah2 = new TestType[N], *Ad = nullptr, *Hmm = nullptr; int NumDevices = 0; HIP_CHECK(hipGetDeviceCount(&NumDevices)); + if (NumDevices < 2) { + HipTest::HIP_SKIP_TEST("Skipping test because more than one device was not found."); + return; + } - for (size_t i =0; i < N; ++i) { + for (size_t i = 0; i < N; ++i) { Ah1[i] = INIT_VAL; Ah2[i] = 0; } + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); for (int Oloop = 0; Oloop < NumDevices; ++Oloop) { DataMismatch = 0; HIP_CHECK(hipSetDevice(Oloop)); @@ -279,8 +373,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocManaged_DeviceContextChange", "", HIP_CHECK(hipSetDevice(Iloop)); HIP_CHECK(hipMalloc(&Ad, N * sizeof(TestType))); // Copy data from host to hipMallocMananged memory and verify - HIP_CHECK(hipMemcpy(Hmm, Ah1, N * sizeof(TestType), - hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Hmm, Ah1, N * sizeof(TestType), hipMemcpyHostToDevice)); for (size_t v = 0; v < N; ++v) { if (Hmm[v] != INIT_VAL) { DataMismatch++; @@ -289,10 +382,8 @@ TEMPLATE_TEST_CASE("Unit_hipMallocManaged_DeviceContextChange", "", REQUIRE(DataMismatch.load() == 0); // Executing D2D transfer with hipMallocManaged memory and verify - HIP_CHECK(hipMemcpy(Ad, Hmm, N * sizeof(TestType), - hipMemcpyDeviceToDevice)); - HIP_CHECK(hipMemcpy(Ah2, Ad, N * sizeof(TestType), - hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Ad, Hmm, N * sizeof(TestType), hipMemcpyDeviceToDevice)); + HIP_CHECK(hipMemcpy(Ah2, Ad, N * sizeof(TestType), hipMemcpyDeviceToHost)); for (size_t k = 0; k < N; ++k) { if (Ah2[k] != INIT_VAL) { DataMismatch++; @@ -300,14 +391,11 @@ TEMPLATE_TEST_CASE("Unit_hipMallocManaged_DeviceContextChange", "", } REQUIRE(DataMismatch.load() == 0); HIP_CHECK(hipMemset(Ad, 0, N * sizeof(TestType))); - const unsigned threadsPerBlock = 256; - const unsigned blocks = (N + 255)/256; // Launching the kernel to check if there is any access issue with // hipMallocManaged memory and local device's memory - vector_sum <<>> (Hmm, Ad, N); - hipDeviceSynchronize(); - HIP_CHECK(hipMemcpy(Ah2, Ad, N * sizeof(TestType), - hipMemcpyDeviceToHost)); + vector_sum<<>>(Hmm, Ad, N); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(Ah2, Ad, N * sizeof(TestType), hipMemcpyDeviceToHost)); for (size_t m = 0; m < N; ++m) { if (Ah2[m] != 246) { DataMismatch++; diff --git a/tests/catch/unit/memory/hipMallocMngdMultiThread.cc b/tests/catch/unit/memory/hipMallocMngdMultiThread.cc index a2f497058c..a163e9a0ab 100644 --- a/tests/catch/unit/memory/hipMallocMngdMultiThread.cc +++ b/tests/catch/unit/memory/hipMallocMngdMultiThread.cc @@ -20,25 +20,25 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #include +#include "hipMallocManagedCommon.hh" #include // Kernel functions -__global__ void HmmMultiThread(int n, float *x, float *y) { +__global__ void HmmMultiThread(int n, float* x, float* y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) - y[i] = x[i] * x[i]; + for (int i = index; i < n; i += stride) y[i] = x[i] * x[i]; } -__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) { +__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_MngdMem123(int *Hmm, int *Dptr, size_t n) { +__global__ void KernelMul_MngdMem123(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) { @@ -47,32 +47,27 @@ __global__ void KernelMul_MngdMem123(int *Hmm, int *Dptr, size_t n) { } - // The following variable is used to determine the failure of test case -static bool IfTestPassed = true; +static bool IfTestPassed = true; -static void LaunchKrnl(int *Hmm1, size_t NumElms, int InitVal, int GpuOrdnl, - int AdviseFlg) { - int *Hmm2 = NULL; +static void LaunchKrnl(int* Hmm1, size_t NumElms, int InitVal, int GpuOrdnl, int AdviseFlg) { + int* Hmm2 = NULL; hipStream_t strm; HIPCHECK(hipSetDevice(GpuOrdnl)); HIPCHECK(hipStreamCreate(&strm)); if (AdviseFlg == 0) { - HIPCHECK(hipMemAdvise(Hmm1 , NumElms * sizeof(int), - hipMemAdviseSetReadMostly, GpuOrdnl)); + HIPCHECK(hipMemAdvise(Hmm1, NumElms * sizeof(int), hipMemAdviseSetReadMostly, GpuOrdnl)); } else if (AdviseFlg == 1) { - HIPCHECK(hipMemAdvise(Hmm1 , NumElms * sizeof(int), - hipMemAdviseSetPreferredLocation, GpuOrdnl)); + HIPCHECK(hipMemAdvise(Hmm1, NumElms * sizeof(int), hipMemAdviseSetPreferredLocation, GpuOrdnl)); } else if (AdviseFlg == 2) { - HIPCHECK(hipMemAdvise(Hmm1 , NumElms * sizeof(int), - hipMemAdviseSetAccessedBy, GpuOrdnl)); + HIPCHECK(hipMemAdvise(Hmm1, NumElms * sizeof(int), hipMemAdviseSetAccessedBy, GpuOrdnl)); } else if (AdviseFlg == 3) { HIPCHECK(hipMemPrefetchAsync(Hmm1, NumElms * sizeof(int), GpuOrdnl, strm)); HIPCHECK(hipStreamSynchronize(strm)); } HIPCHECK(hipMallocManaged(&Hmm2, (sizeof(int) * NumElms))); for (int i = 0; i < 2; ++i) { - KrnlWth2MemTypes<<<((NumElms + 63)/64), 64, 0, strm>>>(Hmm2, Hmm1, NumElms); + KrnlWth2MemTypes<<<((NumElms + 63) / 64), 64, 0, strm>>>(Hmm2, Hmm1, NumElms); HIPCHECK(hipStreamSynchronize(strm)); } // Verifying the result @@ -88,7 +83,7 @@ static void LaunchKrnl(int *Hmm1, size_t NumElms, int InitVal, int GpuOrdnl, } } -static void LaunchKrnl2(int *Hmm, size_t NumElms, int InitVal, int HmmMem) { +static void LaunchKrnl2(int* Hmm, size_t NumElms, int InitVal, int HmmMem) { int *ptr = nullptr, blockSize = 64, *HstPtr = nullptr; hipStream_t strm; HIPCHECK(hipStreamCreate(&strm)); @@ -99,7 +94,7 @@ static void LaunchKrnl2(int *Hmm, size_t NumElms, int InitVal, int HmmMem) { HIPCHECK(hipMallocManaged(&ptr, (sizeof(int) * NumElms))); } dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize - 1) / blockSize, 1, 1); for (int i = 0; i < 2; ++i) { KrnlWth2MemTypes<<>>(ptr, Hmm, NumElms); } @@ -107,8 +102,7 @@ static void LaunchKrnl2(int *Hmm, size_t NumElms, int InitVal, int HmmMem) { // Verifying the result int DataMismatch = 0; if (HmmMem == 0) { - HIPCHECK(hipMemcpy(HstPtr, ptr, (sizeof(int) * NumElms), - hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(HstPtr, ptr, (sizeof(int) * NumElms), hipMemcpyDeviceToHost)); for (size_t i = 0; i < NumElms; ++i) { if (HstPtr[i] != (InitVal + 10)) { DataMismatch++; @@ -127,13 +121,13 @@ static void LaunchKrnl2(int *Hmm, size_t NumElms, int InitVal, int HmmMem) { } } -static void LaunchKrnl3(int *Dptr, size_t NumElms, int InitVal) { +static void LaunchKrnl3(int* Dptr, size_t NumElms, int InitVal) { int *Hmm = NULL, blockSize = 64; hipStream_t strm; HIPCHECK(hipStreamCreate(&strm)); HIPCHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms))); dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize - 1) / blockSize, 1, 1); for (int i = 0; i < 2; ++i) { KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); } @@ -152,14 +146,13 @@ static void LaunchKrnl3(int *Dptr, size_t NumElms, int InitVal) { } -static void LaunchKrnl5(int *Hmm1, size_t NumElms, int InitVal, - int KerneltoLaunch) { +static void LaunchKrnl5(int* Hmm1, size_t NumElms, int InitVal, int KerneltoLaunch) { int *Hmm2 = NULL, blockSize = 64; hipStream_t strm; HIPCHECK(hipStreamCreate(&strm)); HIPCHECK(hipMallocManaged(&Hmm2, (sizeof(int) * NumElms))); dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize - 1) / blockSize, 1, 1); for (int i = 0; i < 2; ++i) { if (KerneltoLaunch == 0) { KrnlWth2MemTypes<<>>(Hmm2, Hmm1, NumElms); @@ -200,8 +193,7 @@ static void TestFlagParamGlobal(int dev) { HIPCHECK(hipSetDevice(dev)); HIPCHECK(hipStreamCreate(&strm)); // Testing hipMemAttachGlobal Flag - HIPCHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), - hipMemAttachGlobal)); + HIPCHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), hipMemAttachGlobal)); // Initializing HmmAG memory for (int i = 0; i < NUM_ELMS; i++) { @@ -246,10 +238,8 @@ static void TestFlagParamHost(int dev) { hipStream_t strm; HIPCHECK(hipSetDevice(dev)); HIPCHECK(hipStreamCreate(&strm)); - HIPCHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), - hipMemAttachHost)); - HIPCHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), - hipMemAttachHost)); + HIPCHECK(hipMallocManaged(&HmmAH1, NUM_ELMS * sizeof(float), hipMemAttachHost)); + HIPCHECK(hipMallocManaged(&HmmAH2, NUM_ELMS * sizeof(float), hipMemAttachHost)); // Initializing HmmAH memory for (int i = 0; i < NUM_ELMS; i++) { HmmAH1[i] = INIT_VAL; @@ -294,77 +284,54 @@ static void AllocateHmmMemory(int flag, int device) { } } - -static int HmmAttrPrint() { - int managed = 0; - INFO("The following are the attribute values related to HMM for" - " device 0:\n"); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); - INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributeConcurrentManagedAccess, 0)); - INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccess, 0)); - INFO("hipDeviceAttributePageableMemoryAccess: " << managed); - HIP_CHECK(hipDeviceGetAttribute(&managed, - hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); - INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" - << managed); - - HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, - 0)); - INFO("hipDeviceAttributeManagedMemory: " << managed); - return managed; -} - TEST_CASE("Unit_hipMallocManaged_MultiThread") { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + IfTestPassed = true; - int NumDevs = 0, managed = 0, ATTACH_GLOBAL = 0, ATTACH_HOST = 1; + int NumDevs = 0, ATTACH_GLOBAL = 0, ATTACH_HOST = 1; int ITERATIONS = 10; - managed = HmmAttrPrint(); - if (managed) { - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - std::vector T1; - std::vector T2; - for (int i = 0; i < NumDevs; ++i) { - for (int j = 0; j < ITERATIONS; ++j) { - T1.push_back(std::thread(TestFlagParamGlobal, i)); - T2.push_back(std::thread(AllocateHmmMemory, ATTACH_GLOBAL, i)); - } - for (auto &t1 : T1) { - if (t1.joinable()) { - t1.join(); - } - } - for (auto &t2 : T2) { - if (t2.joinable()) { - t2.join(); - } + + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + std::vector T1; + std::vector T2; + for (int i = 0; i < NumDevs; ++i) { + for (int j = 0; j < ITERATIONS; ++j) { + T1.push_back(std::thread(TestFlagParamGlobal, i)); + T2.push_back(std::thread(AllocateHmmMemory, ATTACH_GLOBAL, i)); + } + for (auto& t1 : T1) { + if (t1.joinable()) { + t1.join(); } } - T1.clear(); - T2.clear(); - for (int i = 0; i < NumDevs; ++i) { - for (int j = 0; j < ITERATIONS; ++j) { - T1.push_back(std::thread(TestFlagParamHost, i)); - T2.push_back(std::thread(AllocateHmmMemory, ATTACH_HOST, i)); - } - for (auto &t1 : T1) { - if (t1.joinable()) { - t1.join(); - } - } - for (auto &t2 : T2) { - if (t2.joinable()) { - t2.join(); - } + for (auto& t2 : T2) { + if (t2.joinable()) { + t2.join(); + } + } + } + T1.clear(); + T2.clear(); + for (int i = 0; i < NumDevs; ++i) { + for (int j = 0; j < ITERATIONS; ++j) { + T1.push_back(std::thread(TestFlagParamHost, i)); + T2.push_back(std::thread(AllocateHmmMemory, ATTACH_HOST, i)); + } + for (auto& t1 : T1) { + if (t1.joinable()) { + t1.join(); + } + } + for (auto& t2 : T2) { + if (t2.joinable()) { + t2.join(); } } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory" - "attribute. Hence skipping the testing with Pass result.\n"); } REQUIRE(IfTestPassed); } @@ -372,175 +339,169 @@ TEST_CASE("Unit_hipMallocManaged_MultiThread") { // The following test checks what happens when same Hmm memory is used to // launch multiple threads over multiple gpus TEST_CASE("Unit_hipMallocManaged_MGpuMThread") { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } + IfTestPassed = true; int Ngpus = 0; HIP_CHECK(hipGetDeviceCount(&Ngpus)); if (Ngpus < 2) { - WARN("This test needs atleast 2 or more gpus, but the system"); - WARN(" has only " << Ngpus); - WARN(" gpus. Hence skipping the test."); - SUCCEED("\n"); + HipTest::HIP_SKIP_TEST("Skipping test because more than one device was not found."); + return; } - int managed = HmmAttrPrint(); - if (managed == 1) { - int InitVal = 123, *Hmm1 = NULL, NumElms = 4096*4; - HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); - for (int i = 0; i < NumElms; ++i) { - Hmm1[i] = InitVal; - } - std::vector Thrds; - // AdviseFlg=0 for ReadMostly to be applied - // AdviseFlg=1 for PreferredLocation to be applied - // AdviseFlg=2 for AccessedBy to be applied - // AdviseFlg=3 to prefetch the memory to particular gpu - for (int AdviseFlg = 0; AdviseFlg < 4; ++AdviseFlg) { - for (int i = 0; i < Ngpus; ++i) { - Thrds.push_back(std::thread(LaunchKrnl, Hmm1, NumElms, InitVal, i, - AdviseFlg)); - } - for (auto &thr : Thrds) { - if (thr.joinable()) { - thr.join(); - } + int InitVal = 123, *Hmm1 = NULL, NumElms = 4096 * 4; + HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); + for (int i = 0; i < NumElms; ++i) { + Hmm1[i] = InitVal; + } + + std::vector Thrds; + // AdviseFlg=0 for ReadMostly to be applied + // AdviseFlg=1 for PreferredLocation to be applied + // AdviseFlg=2 for AccessedBy to be applied + // AdviseFlg=3 to prefetch the memory to particular gpu + for (int AdviseFlg = 0; AdviseFlg < 4; ++AdviseFlg) { + for (int i = 0; i < Ngpus; ++i) { + Thrds.push_back(std::thread(LaunchKrnl, Hmm1, NumElms, InitVal, i, AdviseFlg)); + } + for (auto& thr : Thrds) { + if (thr.joinable()) { + thr.join(); } } - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); } + REQUIRE(IfTestPassed); } // The following test checks what happens when multiple kernels are launched // with same Hmm memory TEST_CASE("Unit_hipMallocManaged_MultiKrnlComnHmm") { - IfTestPassed = true; - int managed = HmmAttrPrint(); - if (managed == 1) { - int InitVal = 123, *Hmm = NULL, NumElms = 1024*4, TotThrds = 2; - int HmmMem2 = 0, *HstPtr = nullptr; // to indicate the thread that - // hipMalloc() memory has to be used - HstPtr = reinterpret_cast(new int[NumElms]); - HIP_CHECK(hipMalloc(&Hmm, (NumElms * sizeof(int)))); - for (int i = 0; i < NumElms; ++i) { - HstPtr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Hmm, HstPtr, (NumElms * sizeof(int)), - hipMemcpyHostToDevice)); - std::vector Thrds; - for (int i = 0; i < TotThrds; ++i) { - Thrds.push_back(std::thread(LaunchKrnl2, Hmm, NumElms, InitVal, HmmMem2)); - } - - for (auto &thr : Thrds) { - if (thr.joinable()) { - thr.join(); - } - } - delete[] HstPtr; - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; } + + IfTestPassed = true; + + int InitVal = 123, *Hmm = NULL, NumElms = 1024 * 4, TotThrds = 2; + int HmmMem2 = 0, *HstPtr = nullptr; // to indicate the thread that + // hipMalloc() memory has to be used + HstPtr = reinterpret_cast(new int[NumElms]); + HIP_CHECK(hipMalloc(&Hmm, (NumElms * sizeof(int)))); + for (int i = 0; i < NumElms; ++i) { + HstPtr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Hmm, HstPtr, (NumElms * sizeof(int)), hipMemcpyHostToDevice)); + std::vector Thrds; + for (int i = 0; i < TotThrds; ++i) { + Thrds.push_back(std::thread(LaunchKrnl2, Hmm, NumElms, InitVal, HmmMem2)); + } + + for (auto& thr : Thrds) { + if (thr.joinable()) { + thr.join(); + } + } + delete[] HstPtr; } // The following test checks what happens when multiple kernels are launched // with same hipMalloc() memory TEST_CASE("Unit_hipMallocManaged_MultiKrnlComnMalloc") { - IfTestPassed = true; - int managed = HmmAttrPrint(); - if (managed) { - int InitVal = 123, *Dptr = NULL, NumElms = 4096*8, TotThrds = 2; - int *HstPtr = reinterpret_cast(new int[NumElms]); - HIP_CHECK(hipMalloc(&Dptr, (NumElms * sizeof(int)))); - for (int i = 0; i < NumElms; ++i) { - HstPtr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Dptr, HstPtr, (NumElms * sizeof(int)), - hipMemcpyHostToDevice)); - std::vector Thrds; - for (int i = 0; i < TotThrds; ++i) { - Thrds.push_back(std::thread(LaunchKrnl3, Dptr, NumElms, InitVal)); - } - - for (auto &thr : Thrds) { - if (thr.joinable()) { - thr.join(); - } - } - delete[] HstPtr; - HIP_CHECK(hipFree(Dptr)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; } + + IfTestPassed = true; + int InitVal = 123, *Dptr = NULL, NumElms = 4096 * 8, TotThrds = 2; + int* HstPtr = reinterpret_cast(new int[NumElms]); + HIP_CHECK(hipMalloc(&Dptr, (NumElms * sizeof(int)))); + for (int i = 0; i < NumElms; ++i) { + HstPtr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Dptr, HstPtr, (NumElms * sizeof(int)), hipMemcpyHostToDevice)); + std::vector Thrds; + for (int i = 0; i < TotThrds; ++i) { + Thrds.push_back(std::thread(LaunchKrnl3, Dptr, NumElms, InitVal)); + } + + for (auto& thr : Thrds) { + if (thr.joinable()) { + thr.join(); + } + } + delete[] HstPtr; + HIP_CHECK(hipFree(Dptr)); } // The following section tests the scenario wherein multiple threads use their // own stream to launch kernel on common Hmm memory TEST_CASE("Unit_hipMallocManaged_MultiThrdMultiStrm") { - IfTestPassed = true; - int managed = HmmAttrPrint(); - if (managed == 1) { - int NumElms = 4096*4; - int *Hmm1 = NULL, TotlThrds = 4, InitVal = 123; - int HmmMem = 1; // to indicate the thread that Hmm memory need to be - // used inside it - HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); - for (int i = 0; i < NumElms; ++i) { - Hmm1[i] = InitVal; - } - std::vector Thrds; - for (int i = 0; i < TotlThrds; ++i) { - Thrds.push_back(std::thread(LaunchKrnl2, Hmm1, NumElms, InitVal, HmmMem)); - } + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; + } - for (auto &thr : Thrds) { - if (thr.joinable()) { - thr.join(); - } + IfTestPassed = true; + + int NumElms = 4096 * 4; + int *Hmm1 = NULL, TotlThrds = 4, InitVal = 123; + int HmmMem = 1; // to indicate the thread that Hmm memory need to be + // used inside it + HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); + for (int i = 0; i < NumElms; ++i) { + Hmm1[i] = InitVal; + } + std::vector Thrds; + for (int i = 0; i < TotlThrds; ++i) { + Thrds.push_back(std::thread(LaunchKrnl2, Hmm1, NumElms, InitVal, HmmMem)); + } + + for (auto& thr : Thrds) { + if (thr.joinable()) { + thr.join(); } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); } } - // The following section tests the scenario wherein two threads each use // different kernel but common HMM memory TEST_CASE("Unit_hipMallocManaged_TwoKrnlsComnHmmMem") { - IfTestPassed = true; - int managed = HmmAttrPrint(); - if (managed == 1) { - int InitVal = 123, *Dptr = NULL, NumElms = 4096*4, TotThrds = 2; - int *HstPtr = reinterpret_cast(new int[NumElms]); - HIP_CHECK(hipMalloc(&Dptr, (NumElms * sizeof(int)))); - for (int i = 0; i < NumElms; ++i) { - HstPtr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Dptr, HstPtr, (NumElms * sizeof(int)), - hipMemcpyHostToDevice)); - std::vector Thrds; - for (int i = 0; i < TotThrds; ++i) { - Thrds.push_back(std::thread(LaunchKrnl5, Dptr, NumElms, InitVal, i)); - } - - for (auto &thr : Thrds) { - if (thr.joinable()) { - thr.join(); - } - } - delete[] HstPtr; - HIP_CHECK(hipFree(Dptr)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory so skipping test."); + return; } + + IfTestPassed = true; + int InitVal = 123, *Dptr = NULL, NumElms = 4096 * 4, TotThrds = 2; + int* HstPtr = reinterpret_cast(new int[NumElms]); + HIP_CHECK(hipMalloc(&Dptr, (NumElms * sizeof(int)))); + for (int i = 0; i < NumElms; ++i) { + HstPtr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Dptr, HstPtr, (NumElms * sizeof(int)), hipMemcpyHostToDevice)); + std::vector Thrds; + for (int i = 0; i < TotThrds; ++i) { + Thrds.push_back(std::thread(LaunchKrnl5, Dptr, NumElms, InitVal, i)); + } + + for (auto& thr : Thrds) { + if (thr.joinable()) { + thr.join(); + } + } + delete[] HstPtr; + HIP_CHECK(hipFree(Dptr)); } - - From b1619122fd70df2252f589d1a481f4dd5cca3889 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 29 Jul 2022 15:45:57 +0530 Subject: [PATCH 5/9] Fixed typo in Jenkinsfile --- .jenkins/Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.jenkins/Jenkinsfile b/.jenkins/Jenkinsfile index 4f29347af6..3cf3f582a2 100644 --- a/.jenkins/Jenkinsfile +++ b/.jenkins/Jenkinsfile @@ -52,7 +52,7 @@ def hipBuildTest(String backendLabel) { } } timeout(time: 1, unit: 'HOURS') { - stage('HIP Unit Tests - HIT framework ${backendLabel}') { + stage("HIP Unit Tests - HIT framework ${backendLabel}") { dir("${WORKSPACE}/hipamd/build") { sh """#!/usr/bin/env bash set -x @@ -94,7 +94,7 @@ def hipBuildTest(String backendLabel) { } } timeout(time: 1, unit: 'HOURS') { - stage('HIP Unit Tests - Catch2 framework ${backendLabel}') { + stage("HIP Unit Tests - Catch2 framework ${backendLabel}") { dir("${WORKSPACE}/hipamd/build") { sh """#!/usr/bin/env bash set -x From 4d014ce7114596094dc533863fe300b44d88ff41 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 29 Jul 2022 16:31:04 +0530 Subject: [PATCH 6/9] SWDEV-299940 - Needs to Fix Malloc tests (#2832) Change-Id: I21739d88871b5813cf266a6be8eb05ecee487046 --- tests/catch/unit/memory/hipMallocConcurrency.cc | 11 ++++++++++- tests/src/runtimeApi/memory/hipMallocConcurrency.cpp | 11 ++++++++++- 2 files changed, 20 insertions(+), 2 deletions(-) diff --git a/tests/catch/unit/memory/hipMallocConcurrency.cc b/tests/catch/unit/memory/hipMallocConcurrency.cc index 592ce86897..647d38d99d 100644 --- a/tests/catch/unit/memory/hipMallocConcurrency.cc +++ b/tests/catch/unit/memory/hipMallocConcurrency.cc @@ -95,6 +95,14 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { 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); + HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) { + //In concurrent calls on one GPU, we cannot verify leaking in this way + printf("%s : Memory allocation mismatch observed." + "Possible memory leak.\n", __func__); + TestPassed &= false; + } unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -113,10 +121,11 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { TestPassed = false; } + HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); 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)) { + if (!concurOnOneGPU && (curAvl < prevAvl || 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."); diff --git a/tests/src/runtimeApi/memory/hipMallocConcurrency.cpp b/tests/src/runtimeApi/memory/hipMallocConcurrency.cpp index 7115802a8c..4bfd65a7f9 100644 --- a/tests/src/runtimeApi/memory/hipMallocConcurrency.cpp +++ b/tests/src/runtimeApi/memory/hipMallocConcurrency.cpp @@ -109,6 +109,14 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { HIPCHECK(hipSetDevice(gpu)); HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIPCHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl < curAvl|| prevTot != curTot)) { + //In concurrent calls on one GPU, we cannot verify leaking in this way + printf("%s : Memory allocation mismatch observed." + "Possible memory leak.\n", __func__); + TestPassed &= false; + } unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -129,10 +137,11 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { TestPassed &= false; } + HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); 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)) { + if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) { //In concurrent calls on one GPU, we cannot verify leaking in this way printf("%s : Memory allocation mismatch observed." "Possible memory leak.\n", __func__); From 37583e56456764e8a1e09e1cf24ac82d1f4f4100 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Mon, 1 Aug 2022 14:41:20 -0700 Subject: [PATCH 7/9] SWDEV-327563 - Win: Disable failing hipMemGetInfo tests (#2836) --- .../catch/hipTestMain/config/config_amd_windows.json | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 057a8a97f6..36557025fa 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -45,6 +45,15 @@ "Unit_hipGraphNodeGetDependentNodes_Functional", "Unit_hipGraphNodeGetDependentNodes_ParamValidation", "Unit_hipGraphNodeGetDependencies_Functional", - "Unit_hipGraphNodeGetDependencies_ParamValidation" + "Unit_hipGraphNodeGetDependencies_ParamValidation", + "Unit_hipMemGetInfo_DifferentMallocSmall", + "Unit_hipMemGetInfo_MallocArray - int", + "Unit_hipMemGetInfo_MallocArray - int4", + "Unit_hipMemGetInfo_MallocArray - char", + "Unit_hipMemGetInfo_Malloc3D", + "Unit_hipMemGetInfo_Malloc3DArray - char", + "Unit_hipMemGetInfo_Malloc3DArray - int", + "Unit_hipMemGetInfo_Malloc3DArray - int4", + "Unit_hipMemGetInfo_ParaSmall", ] } From 2d90452c9629fd7f3251f7e6099844e9b7ebf572 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 2 Aug 2022 03:53:48 +0530 Subject: [PATCH 8/9] SWDEV-332251 - change catch-amd package name to hip-catch-amd (#2838) 1) catch-amd package name to hip-catch-amd 2) Install to /opt/rocm/test/hip instead of $ROCM_PATH/test/hip Change-Id: Ie466e52a609c516aa14b7fced7552f7fbd66bc3a --- tests/catch/packaging/hip-tests.txt | 37 ++++++++++++++++++++++++++--- 1 file changed, 34 insertions(+), 3 deletions(-) diff --git a/tests/catch/packaging/hip-tests.txt b/tests/catch/packaging/hip-tests.txt index 77b3a0e46c..f934130e34 100644 --- a/tests/catch/packaging/hip-tests.txt +++ b/tests/catch/packaging/hip-tests.txt @@ -20,9 +20,35 @@ # THE SOFTWARE. cmake_minimum_required(VERSION 3.16.8) project(tests) + +MACRO(SUBDIRLIST result curdir) + FILE(GLOB children RELATIVE ${curdir} ${curdir}/*) + SET(dirlist "") + FOREACH(child ${children}) + IF(IS_DIRECTORY ${curdir}/${child}) + LIST(APPEND dirlist ${child}) + ENDIF() + ENDFOREACH() + SET(${result} ${dirlist}) +ENDMACRO() + +SUBDIRLIST(SUBDIRS @PROJECT_BINARY_DIR@/@CATCH_BUILD_DIR@) + +FOREACH(subdir ${SUBDIRS}) + set(CONTENT ${CONTENT} "subdirs(${subdir}) \n") +ENDFOREACH() +# Creating a CTestTestfile so ctest can be executed from @CATCH_BUILD_DIR@ level +# This also helps in executing through jenkins and avoid permission issues +file(WRITE @PROJECT_BINARY_DIR@/@CATCH_BUILD_DIR@/CTestTestfile.cmake + ${CONTENT}) + install(DIRECTORY @PROJECT_BINARY_DIR@/@CATCH_BUILD_DIR@ DESTINATION . - USE_SOURCE_PERMISSIONS) + USE_SOURCE_PERMISSIONS + DIRECTORY_PERMISSIONS OWNER_WRITE OWNER_READ OWNER_EXECUTE + GROUP_WRITE GROUP_READ GROUP_EXECUTE + WORLD_WRITE WORLD_READ WORLD_EXECUTE) + install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake DESTINATION .) @@ -30,8 +56,13 @@ install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake # Packaging steps ############################# set(CPACK_SET_DESTDIR TRUE) -set(CPACK_INSTALL_PREFIX "@ROCM_PATH@/test/hip/") -set(PKG_NAME catch-@HIP_PLATFORM@) +set(CPACK_INSTALL_PREFIX @CPACK_INSTALL_PREFIX@) +if(NOT DEFINED CPACK_INSTALL_PREFIX) + set(CPACK_INSTALL_PREFIX "/opt/rocm/test/hip/") +endif() + + +set(PKG_NAME hip-catch-@HIP_PLATFORM@) set(CPACK_PACKAGE_NAME ${PKG_NAME}) set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [CATCH TESTS]") set(CPACK_PACKAGE_DESCRIPTION "HIP: From eeb10a82389df95a0dd05611eb15a1261839b0c5 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 3 Aug 2022 07:57:33 +0530 Subject: [PATCH 9/9] Fix typo in config_amd_windows.json --- tests/catch/hipTestMain/config/config_amd_windows.json | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 36557025fa..c0ddfbc0a4 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -54,6 +54,6 @@ "Unit_hipMemGetInfo_Malloc3DArray - char", "Unit_hipMemGetInfo_Malloc3DArray - int", "Unit_hipMemGetInfo_Malloc3DArray - int4", - "Unit_hipMemGetInfo_ParaSmall", + "Unit_hipMemGetInfo_ParaSmall" ] }