From e108ed9cbf6b7a22927fd9e764a0989ba37c35e7 Mon Sep 17 00:00:00 2001 From: lthakur Date: Fri, 25 Feb 2022 19:39:30 +0530 Subject: [PATCH] SWDEV-298757 - Moving Unit_hipMemPrefetchAsyncOneToAll into stress/memory directory (#2460) --- catch/stress/memory/CMakeLists.txt | 1 + .../memory/hipMemPrftchAsyncStressTst.cc | 139 ++++++++++++++++++ .../unit/memory/hipMemPrefetchAsyncExtTsts.cc | 103 ++----------- 3 files changed, 156 insertions(+), 87 deletions(-) create mode 100644 catch/stress/memory/hipMemPrftchAsyncStressTst.cc diff --git a/catch/stress/memory/CMakeLists.txt b/catch/stress/memory/CMakeLists.txt index 0c2d3fc6aa..da5dfa8182 100644 --- a/catch/stress/memory/CMakeLists.txt +++ b/catch/stress/memory/CMakeLists.txt @@ -3,6 +3,7 @@ set(TEST_SRC memcpy.cc hipMemcpyMThreadMSize.cc hipMallocManagedStress.cc + hipMemPrftchAsyncStressTst.cc ) hip_add_exe_to_target(NAME memory diff --git a/catch/stress/memory/hipMemPrftchAsyncStressTst.cc b/catch/stress/memory/hipMemPrftchAsyncStressTst.cc new file mode 100644 index 0000000000..bc2cb027da --- /dev/null +++ b/catch/stress/memory/hipMemPrftchAsyncStressTst.cc @@ -0,0 +1,139 @@ +/* +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. +*/ + +/* Test Case Description: + The following test allocates a managed memory and prefetch it in + one-to-all and all-to-one fashion followed by kernel launch within available + devices*/ + +#include + +// Kernel function +__global__ void MemPrftchAsyncKernel1(int* Hmm, 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) { + Hmm[i] = Hmm[i] * Hmm[i]; + } +} + +static int HmmAttrPrint() { + int managed = 0; + WARN("The following are the attribute values related to HMM for" + " device 0:\n"); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeDirectManagedMemAccessFromHost, 0)); + WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributeConcurrentManagedAccess, 0)); + WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess: " << managed); + HIP_CHECK(hipDeviceGetAttribute(&managed, + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0)); + WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:" + << managed); + + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, + 0)); + WARN("hipDeviceAttributeManagedMemory: " << managed); + return managed; +} + +static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) { + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipFree(Hmm1)); + HIP_CHECK(hipStreamDestroy(*strm)); +} + + +/* The following test allocates a managed memory and prefetch it in + one-to-all and all-to-one fahsion followed by kernel launch within available + devices*/ +TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + int *Hmm = nullptr, *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4); + int InitVal = 123, NumElms = MemSz/4; + bool IfTestPassed = true; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MemSz)); + HIP_CHECK(hipMallocManaged(&Hmm1, MemSz)); + for (int i = 0; i < NumElms; ++i) { + Hmm1[i] = InitVal; + } + hipStream_t strm; + for (int i = -1; i < NumDevs; ++i) { + HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, 0)); + for (int j = -1; j < NumDevs; ++j) { + if (i == j) { + continue; + } + if (j != -1) { + HIP_CHECK(hipSetDevice(j)); + } + HIP_CHECK(hipStreamCreate(&strm)); + // Prefetching memory from i to j + HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, j, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the result + for (int m = 0; m < NumElms; ++m) { + if (Hmm[m] != (InitVal * InitVal)) { + IfTestPassed = false; + } + } + if (!IfTestPassed) { + ReleaseResource(Hmm, Hmm1, &strm); + INFO("Did not find expected value!"); + REQUIRE(false); + } + // Resetting the values in Hmm + HIP_CHECK(hipMemset(Hmm, 0, MemSz)); + // Prefetching memory from j to i + HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the result + for (int m = 0; m < NumElms; ++m) { + if (Hmm[m] != (InitVal * InitVal)) { + IfTestPassed = false; + } + } + if (!IfTestPassed) { + ReleaseResource(Hmm, Hmm1, &strm); + INFO("Did not find expected value!"); + REQUIRE(false); + } + // Resetting the values in Hmm + HIP_CHECK(hipMemset(Hmm, 0, MemSz)); + HIP_CHECK(hipStreamDestroy(strm)); + } + } + // Releasing the resources in case all the scenarios passed + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipFree(Hmm1)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} diff --git a/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc b/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc index e34e9ec8d0..c5a51d9488 100644 --- a/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc +++ b/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc @@ -17,15 +17,24 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +/* Test Case Description: + 1) Allocate managed memory --> prefetch to gpu 0 + call hipMemAdvise() on the memory and apply the flags ReadMostly, + AccessedBy, and PreferredLocation for gpus other than gpu 0 and verify + the flags using hipMemGetAttribute() + 2) Allocate managed memory --> set AccessedBy using + hipMemAdvise() to gpu1 prefetch the memory to gpu 0 and then query for + AccessedBy using hipMemGetAttribute() and validate if AccessedBy is still + set to gpu1. Similar tests are done with ReadMostly and PreferredLocation + flags + 3) Negative testing with hipMemPrefetchAsync() api + 4) In this test case I am trying to allocate HMM memory + which is not multiple of page Size, but still trying to launch kernel and + see if we are getting values as expected. + */ + #include // Kernel function -__global__ void MemPrftchAsyncKernel(int* Hmm, const int* Hmm1, 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) { - Hmm[i] = Hmm1[i] * Hmm1[i]; - } -} __global__ void MemPrftchAsyncKernel1(int* Hmm, size_t N) { size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); @@ -59,86 +68,6 @@ static int HmmAttrPrint() { return managed; } -static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) { - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipFree(Hmm1)); - HIP_CHECK(hipStreamDestroy(*strm)); -} - -/* The following test allocates a managed memory and prefetch it in - one-to-all and all-to-one fahsion followed by kernel launch within available - devices*/ -TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - int *Hmm = nullptr, *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4); - int InitVal = 123, NumElms = MemSz/4; - bool IfTestPassed = true; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MemSz)); - HIP_CHECK(hipMallocManaged(&Hmm1, MemSz)); - for (int i = 0; i < NumElms; ++i) { - Hmm1[i] = InitVal; - } - hipStream_t strm; - for (int i = -1; i < NumDevs; ++i) { - HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, 0)); - for (int j = -1; j < NumDevs; ++j) { - if (i == j) { - continue; - } - if (j != -1) { - HIP_CHECK(hipSetDevice(j)); - } - HIP_CHECK(hipStreamCreate(&strm)); - // Prefetching memory from i to j - HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, j, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the result - for (int m = 0; m < NumElms; ++m) { - if (Hmm[m] != (InitVal * InitVal)) { - IfTestPassed = false; - } - } - if (!IfTestPassed) { - ReleaseResource(Hmm, Hmm1, &strm); - INFO("Did not find expected value!"); - REQUIRE(false); - } - // Resetting the values in Hmm - HIP_CHECK(hipMemset(Hmm, 0, MemSz)); - // Prefetching memory from j to i - HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the result - for (int m = 0; m < NumElms; ++m) { - if (Hmm[m] != (InitVal * InitVal)) { - IfTestPassed = false; - } - } - if (!IfTestPassed) { - ReleaseResource(Hmm, Hmm1, &strm); - INFO("Did not find expected value!"); - REQUIRE(false); - } - // Resetting the values in Hmm - HIP_CHECK(hipMemset(Hmm, 0, MemSz)); - HIP_CHECK(hipStreamDestroy(strm)); - } - } - // Releasing the resources in case all the scenarios passed - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipFree(Hmm1)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - /* Test Case Description: Allocate managed memory --> prefetch to gpu 0 call hipMemAdvise() on the memory and apply the flags ReadMostly, AccessedBy, and PreferredLocation for gpus other than gpu 0 and verify