From 829cfc1a96d5771744c806375232cbd6199ee862 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 17 May 2022 19:54:28 +0530 Subject: [PATCH] SWDEV-298757 - hipMemPrefetchAsync.cc Adding tests for hipMemPrefetchAsync() api (#2679) Change-Id: Icadd93e62276b660ed0fd9a864e6985a7b03798a --- catch/unit/memory/CMakeLists.txt | 2 + catch/unit/memory/hipMemPrefetchAsync.cc | 124 +++++++++++++++++++++++ 2 files changed, 126 insertions(+) create mode 100644 catch/unit/memory/hipMemPrefetchAsync.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index b4c94f6482..d4e821463c 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -85,6 +85,7 @@ set(TEST_SRC hipPointerGetAttribute.cc hipDrvPtrGetAttributes.cc hipMallocMngdMultiThread.cc + hipMemPrefetchAsync.cc ) else() set(TEST_SRC @@ -148,6 +149,7 @@ set(TEST_SRC hipDrvMemcpy3DAsync.cc hipPointerGetAttribute.cc hipDrvPtrGetAttributes.cc + hipMemPrefetchAsync.cc ) endif() diff --git a/catch/unit/memory/hipMemPrefetchAsync.cc b/catch/unit/memory/hipMemPrefetchAsync.cc new file mode 100644 index 0000000000..6f8fe65c74 --- /dev/null +++ b/catch/unit/memory/hipMemPrefetchAsync.cc @@ -0,0 +1,124 @@ +/* +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 +// Kernel function +__global__ void MemPrftchAsyncKernel(int* C_d, const int* A_d, size_t N) { + size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + size_t stride = hipBlockDim_x * hipGridDim_x; + for (size_t i = offset; i < N; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + + +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 Description: This test prefetches the memory to each of the available + devices and launch kernel followed by result verification + At the end the memory is prefetched to Host and kernel is launched followed + by result verification. +*/ + +TEST_CASE("Unit_hipMemPrefetchAsync") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int A_CONST = 123, MEM_SIZE = (8192 * sizeof(int)); + int *devPtr1 = NULL, *devPtr2 = NULL, NumDevs = 0, flag = 0; + hipStream_t strm; + HIP_CHECK(hipMallocManaged(&devPtr1, MEM_SIZE)); + HIP_CHECK(hipMallocManaged(&devPtr2, MEM_SIZE)); + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + // Initializing the memory + for (uint32_t k = 0; k < (MEM_SIZE/sizeof(int)); ++k) { + devPtr1[k] = A_CONST; + devPtr2[k] = 0; + } + + + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemPrefetchAsync(devPtr1, MEM_SIZE, i, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemPrftchAsyncKernel<<<32, (MEM_SIZE/sizeof(int)/32)>>>(devPtr2, devPtr1, + MEM_SIZE/sizeof(int)); + for (uint32_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + if (devPtr1[m] != (A_CONST * A_CONST)) { + flag = 1; + } + } + HIP_CHECK(hipStreamDestroy(strm)); + if (!flag) { + INFO("Test failed for device: " << i); + IfTestPassed = false; + flag = 0; + } + } + // The memory will be prefetched from last gpu in the system to the host + // memory and kernel is launched followed by result verification. + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemPrefetchAsync(devPtr1, MEM_SIZE, hipCpuDeviceId, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemPrftchAsyncKernel<<<32, (MEM_SIZE/sizeof(int)/32)>>>(devPtr2, devPtr1, + MEM_SIZE/sizeof(int)); + for (uint32_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + if (devPtr1[m] != (A_CONST * A_CONST)) { + flag = 1; + } + } + HIP_CHECK(hipStreamDestroy(strm)); + if (!flag) { + INFO("Failed to prefetch the memory to System space.\n"); + IfTestPassed = false; + flag = 0; + } + + HIP_CHECK(hipFree(devPtr1)); + HIP_CHECK(hipFree(devPtr2)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +}