From 52cc6b892fe3669a8752380a72ebc68f41973e59 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 15:12:24 +0530 Subject: [PATCH 1/6] SWDEV-298757 - hipMemAdvise.cc Adding test cases based on hipMemAdvise() api (#2843) Change-Id: I04a580dc54b2f4ac7375b07a68a9343bbf007844 --- tests/catch/unit/memory/CMakeLists.txt | 2 + tests/catch/unit/memory/hipMemAdvise.cc | 962 ++++++++++++++++++++++++ 2 files changed, 964 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemAdvise.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index ce52b2e104..aff4c7ef5e 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -96,6 +96,7 @@ set(TEST_SRC hipMemcpySync.cc hipMemsetSync.cc hipMemsetAsync.cc + hipMemAdvise.cc ) else() set(TEST_SRC @@ -167,6 +168,7 @@ set(TEST_SRC hipMemcpySync.cc hipMemsetSync.cc hipMemsetAsync.cc + hipMemAdvise.cc ) endif() diff --git a/tests/catch/unit/memory/hipMemAdvise.cc b/tests/catch/unit/memory/hipMemAdvise.cc new file mode 100644 index 0000000000..96a0318bb8 --- /dev/null +++ b/tests/catch/unit/memory/hipMemAdvise.cc @@ -0,0 +1,962 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* Test Case Description: + Scenario-1: The following Function Tests the working of flags which can be + assigned to HMM memory using hipMemAdvise() api + Scenario-2: Negative tests on hipMemAdvise() api + Scenario-3: The following function tests various scenarios around the flag + 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api + Scenario-4: The following function tests various scenarios around the flag + 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api + Scenario-5: The following function verifies if assigning of a flag + invalidates the earlier flag which was assigned to the same memory region + using hipMemAdvise() + Scenario-6: The following function tests if peers can set + hipMemAdviseSetAccessedBy flag + on HMM memory prefetched on each of the other gpus + Scenario-7: Set AccessedBy flag and check value returned by + hipMemRangeGetAttribute() It should be -2(same is observed on cuda) + Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the + memory to device 1, then probe for AccessedBy flag using + hipMemRangeGetAttribute() we should still see the said flag is set for + device 0 + Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to + device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should + return 0 + 2) Unset AccessedBy to 0 and set it to device 1 followed by + PreferredLocation to device 1, check for AccessedBy flag using + hipMemRangeGetAttribute() it should return 1 + Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset + AccessedBy, launch kernel. We should not have any access issues + Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation + flag to the allocated memory and launch a kernel. Kernel should get executed + successfully without hang or segfault + Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then + get attribute using the api hipMemRangeGetAttribute() for + hipMemRangeAttributeLastPrefetchLocation the value returned should be -2 + Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch + the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using + hipMemRangeGetAttribute(), we should get 1 + Scenario-14: Allocate HMM memory, set ReadMostly followed by + PreferredLocation, probe for hipMemRangeAttributeReadMostly and + hipMemRangeAttributePreferredLocation + using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. + In other words setting of hipMemRangeAttributePreferredLocation should not + impact hipMemRangeAttributeReadMostly advise to the memory + Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and + launch kernel on all other gpus except 0. This test case may discover any + effect or access denial case arising due to setting ReadMostly only to a + particular gpu +*/ + +#include +#if __linux__ +#include +#include +#include +#endif + +// Kernel function +__global__ void MemAdvseKernel(int n, int *x) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < n) + x[index] = x[index] * x[index]; +} + +// Kernel +__global__ void MemAdvise2(int *Hmm, int n) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + Hmm[i] = Hmm[i] + 10; + } +} + +// Kernel +__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + Hmm1[i] = Hmm[i] + 10; + } +} + + +static bool CheckError(hipError_t err, int LineNo) { + if (err == hipSuccess) { + WARN("Error expected but received hipSuccess at line no.:" << LineNo); + return false; + } else { + return true; + } +} + +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; +} + + +// The following Function Tests the working of flags which can be assigned +// to HMM memory using hipMemAdvise() api +TEST_CASE("Unit_hipMemAdvise_TstFlags") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, *Outpt = nullptr; + int MEM_SIZE = 4*1024, A_CONST = 9999; + float *Hmm = nullptr; + int AttrVal = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + Outpt = new int(NumDevs); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); + // With the following for loop we iterate through each of the Gpus in the + // system set and unset the flags and check the behavior. + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE * 2)); + if (AttrVal != 1) { + WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly, + i)); + + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributeReadMostly, Hmm, + (MEM_SIZE * 2))); + if (AttrVal != 0) { + WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + AttrVal = A_CONST; + // Currently hipMemAdviseSetPreferredLocation and + // hipMemAdviseSetAccessedBy + // flags are resulting in issues: SWDEV-267357 + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, + hipMemAdviseSetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributePreferredLocation, + Hmm, (MEM_SIZE * 2))); + if (AttrVal != i) { + WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n"); + IfTestPassed = false; + } + AttrVal = A_CONST; + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, + hipMemAdviseUnsetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributePreferredLocation, + Hmm, (MEM_SIZE * 2))); + if (AttrVal == i) { + WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation "); + WARN("flag failed!\n"); + IfTestPassed = false; + } + for (int m = 0; m < NumDevs; ++m) { + Outpt[m] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i)); + HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), + hipMemRangeAttributeAccessedBy, Hmm, + (MEM_SIZE * 2))); + if ((Outpt[0]) != i) { + WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + for (int m = 0; m < NumDevs; ++m) { + Outpt[m] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy, + i)); + HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), + hipMemRangeAttributeAccessedBy, Hmm, + (MEM_SIZE * 2))); + if ((Outpt[0]) >= 0) { + WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + } + delete [] Outpt; + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +TEST_CASE("Unit_hipMemAdvise_NegtveTsts") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, MEM_SIZE = 4*1024; + float *Hmm = nullptr; + std::string str; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); +#if HT_AMD + // Passing invalid value(99) device param + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, 99), __LINE__); + + // Passing invalid value(-12) device param + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, -12), __LINE__); +#endif + // Passing NULL as first parameter instead of valid pointer to a memory + IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, 0), __LINE__); + + // Passing 0 for count(2nd param) parameter + IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly, + 0), __LINE__); + + // Passing count much more than actually allocated value + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6, + hipMemAdviseSetReadMostly, 0), __LINE__); + + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function tests various scenarios around the flag +// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api +TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + // Check that when a page fault occurs for the memory region set to devPtr, + // the data is migrated to the destn processor + int MEM_SIZE = 4096, A_CONST = 9999; + int *Hmm = nullptr, NumDevs = 0, dev = A_CONST; + bool IfTestPassed = true; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal)); + for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) { + Hmm[i] = 4; + } + for (int devId = 0; devId < NumDevs; ++devId) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, + hipMemAdviseSetPreferredLocation, devId)); + int NumElms = ((MEM_SIZE * 3)/4); + MemAdvseKernel<<>>(NumElms, Hmm); + int dev = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), + hipMemRangeAttributePreferredLocation, + Hmm, MEM_SIZE * 3)); + if (dev != devId) { + WARN("Memory observed to be not available on expected location\n"); + WARN("line no: " << __LINE__); + WARN("dev: " << dev); + IfTestPassed = false; + } + } + + // Check that when preferred location is set for a memory region, + // data can still be prefetched using hipMemPrefetchAsync + hipStream_t strm; + dev = A_CONST; + for (int devId = 0; devId < NumDevs; ++devId) { + HIP_CHECK(hipSetDevice(devId)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, + hipMemAdviseSetPreferredLocation, devId)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, MEM_SIZE * 3)); + if (dev != devId) { + WARN("Memory reported to be not available at the Prefetched "); + WARN("location with device id: " << devId); + WARN("line no: " << __LINE__); + WARN("dev: " << dev); + IfTestPassed = false; + } + HIP_CHECK(hipStreamDestroy(strm)); + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function tests various scenarios around the flag +// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api + +TEST_CASE("Unit_hipMemAdvise_ReadMostly") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int MEM_SIZE = 4096, A_CONST = 9999; + float *Hmm = nullptr; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE)); + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + Hmm[i] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + // Checking if the data can be read after setting hipMemAdviseSetReadMostly + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + if (Hmm[i] != A_CONST) { + WARN("Didn't find expected value in Hmm memory after setting"); + WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__); + IfTestPassed = false; + } + } + + // Checking if the memory region can be modified + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + Hmm[i] = A_CONST; + } + + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + if (Hmm[i] != A_CONST) { + WARN("Didn't find expected value in Hmm memory after Modification\n"); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + } + } + + int out = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, + Hmm, MEM_SIZE)); + if (out != 1) { + WARN("out value: " << out); + IfTestPassed = false; + } + // Checking the advise attribute after prefetch + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE)); + if (out != 1) { + WARN("Attribute assigned to memory changed after calling "); + WARN("hipMemPrefetchAsync(). line no.: " << __LINE__); + WARN("out value: " << out); + IfTestPassed = false; + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function verifies if assigning of a flag invalidates the +// earlier flag which was assigned to the same memory region using +// hipMemAdvise() +TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int MEM_SIZE = 4*4096, A_CONST = 9999; + float *Hmm = nullptr; + int NumDevs = 0, dev = A_CONST; + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE)); + if (dev != 1) { + WARN("hipMemAdviseSetReadMostly flag did not take affect despite "); + WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); + IfTestPassed = false; + break; + } + dev = A_CONST; + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation, + i)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributePreferredLocation, + Hmm, MEM_SIZE)); + if (dev != i) { + WARN("hipMemAdviseSetPreferredLocation flag did not take affect "); + WARN("despite setting it using hipMemAdvise()\n"); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + break; + } + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i)); + dev = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, + MEM_SIZE)); + if (dev != i) { + WARN("hipMemAdviseSetAccessedBy flag did not take affect despite "); + WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); + IfTestPassed = false; + break; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i)); + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +// The following function tests if peers can set hipMemAdviseSetAccessedBy flag +// on HMM memory prefetched on each of the other gpus +#if HT_AMD +TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;; + int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0; + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + if (NumDevs < 2) { + SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test" + " the scenario. This system has GPUs less than 2"); + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0)); + for (int j = 0; j < NumDevs; ++j) { + if (i == j) + continue; + HIP_CHECK(hipSetDevice(j)); + HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i)); + if (CanAccessPeer) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j)); + for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + Hmm[m] = 4; + } + HIP_CHECK(hipDeviceEnablePeerAccess(i, 0)); + MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>( + (MEM_SIZE/sizeof(int)), Hmm); + HIP_CHECK(hipDeviceSynchronize()); + // Verifying the result + for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + if (Hmm[m] != 16) { + flag = 1; + } + } + if (flag) { + WARN("Didnt get Expected results with device: " << j); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + flag = 0; + } + } + } + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} +#endif + + +/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute() + It should be -2(same is observed on cuda)*/ +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 2*4096)); + if (data != -2) { + WARN("Didnt get expected value!!\n"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to + device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute() + we should still see the said flag is set for device 0*/ +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999, Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus >= 2) { + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); + HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipFree(Hmm)); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + +/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1 + check for AccessedBy flag using hipMemRangeGetAttribute() it should + return 0 + 2) Unset AccessedBy to 0 and set it to device 1 followed by + PreferredLocation to device 1, check for AccessedBy flag using + hipMemRangeGetAttribute() it should return 1*/ + +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999, Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus >= 2) { + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 1) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Set AccessedBy flag to HMM memory launch a kernel and then unset + AccessedBy, launch kernel. We should not have any access issues*/ + +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; + int DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetAccessedBy, 0)); + // Initializing memory + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = InitVal; + } + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + // launching kernel from each one of the gpus + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + 10)) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseUnsetAccessedBy, 0)); + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + (2*10))) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to + the allocated memory and launch a kernel. Kernel should get executed + successfully without hang or segfault*/ +#if __linux__ && HT_AMD +TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { + if ((setenv("HSA_XNACK", "1", 1)) != 0) { + WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); + REQUIRE(false); + } + // The following code block is used to check for gfx906/8 so as to skip if + // any of the gpus available + int fd1[2]; // Used to store two ends of first pipe + pid_t p; + if (pipe(fd1) == -1) { + fprintf(stderr, "Pipe Failed"); + REQUIRE(false); + } + + /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no + GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/ + int GpuId[2] = {0, 0}; + p = fork(); + + if (p < 0) { + fprintf(stderr, "fork Failed"); + REQUIRE(false); + } else if (p > 0) { // parent process + close(fd1[1]); // Close writing end of first pipe + // Wait for child to send a string + wait(NULL); + // Read string from child and close reading end. + read(fd1[0], GpuId, 2 * sizeof(int)); + close(fd1[0]); + if ((GpuId[0] == 1) || (GpuId[0] == 1)) { + WARN("This test is not applicable on MI60 & MI100." + "Skipping the test!!"); + exit(0); + } + } else { // child process + close(fd1[0]); // Close read end of first pipe + hipDeviceProp_t prop; + HIPCHECK(hipGetDeviceProperties(&prop, 0)); + char *p = NULL; + p = strstr(prop.gcnArchName, "gfx906"); + if (p) { + WARN("gfx906 gpu found on this system!!"); + GpuId[0] = 1; + } + p = strstr(prop.gcnArchName, "gfx908"); + if (p) { + WARN("gfx908 gpu found on this system!!"); + GpuId[1] = 1; + } + // Write concatenated string and close writing end + write(fd1[1], GpuId, 2 * sizeof(int)); + close(fd1[1]); + exit(0); + } + int stat = 0; + if (fork() == 0) { + // The below part should be inside fork + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123; + // Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz)); + Mllc = reinterpret_cast(aligned_alloc(4096, 4096*4)); + for (int i = 0; i < NumElms; ++i) { + Mllc[i] = InitVal; + } + hipStream_t strm; + int DataMismatch = 0; + HIP_CHECK(hipStreamCreate(&strm)); + // The following hipMemAdvise() call is made to know if advise on + // aligned_alloc() is causing any issue + HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + for (int i = 0; i < NumElms; ++i) { + if (Mllc[i] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch observed!!"); + exit(9); // 9 for failure + } else { + exit(10); // 10 for Pass result + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } + } else { + wait(&stat); + int Result = WEXITSTATUS(stat); + if (Result != 10) { + REQUIRE(false); + } + } +} +#endif + +/* Allocate Hmm memory, set advise to PreferredLocation and then get + attribute using the api hipMemRangeGetAttribute() for + hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 4096)); + if (data != -2) { + WARN("Didnt receive expected value."); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem + to device1, probe for hipMemRangeAttributeLastPrefetchLocation using + hipMemRangeGetAttribute(), we should get 1*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") { + int NumDevs = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + if (NumDevs >= 2) { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + hipStream_t strm; + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 4096)); + if (data != 1) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } + } else { + SUCCEED("This system has less than 2 gpus hence skipping the test.\n"); + } +} + + +/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe + for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation + using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. + In other words setting of hipMemRangeAttributePreferredLocation should not + impact hipMemRangeAttributeReadMostly advise to the memory*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + 4096)); + if (data != 1) { + WARN("Didnt receive expected value at line: " << data); + REQUIRE(false); + } + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, + 4096)); + if (data != 0) { + WARN("Didnt receive expected value at line: " << data); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + +/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel + on all other gpus except 0. This test case may discover any effect or + access denial case arising due to setting ReadMostly only to a particular + gpu*/ + +TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus < 2) { + SUCCEED("This test needs atleast two gpus to run." + "Hence skipping the test.\n"); + } + int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; + int *Hmm1 = NULL, DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); + // Initializing memory + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = InitVal; + } + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetReadMostly, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); +#if HT_AMD + SECTION("Launch Kernel on all other gpus") { + // launching kernel from each one of the gpus + for (int i = 1; i < Ngpus; ++i) { + DataMismatch = 0; + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); + MemAdvise3<<>>(Hmm, Hmm1, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // verifying the results + for (int j = 0; j < NumElms; ++j) { + if (Hmm1[j] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch is observed with the gpu: " << i); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm1)); + } + } + + SECTION("Launch Kernel on all other gpus and manipulate the content") { + for (int i = 0; i < Ngpus; ++i) { + DataMismatch = 0; + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetReadMostly, i)); + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + } + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + Ngpus * 10)) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + } +#endif + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); + if (data != -2) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + From 981f9076112d7bd78a7263e7ce50ef0e0c190057 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 15:13:19 +0530 Subject: [PATCH 2/6] SWDEV-334508 - CMake create static lib & linker commands for windows. (#2841) Change-Id: I544837a65e47d425a946c666d0b2348aa7a33b15 --- .../2_Cookbook/15_static_library/README.md | 3 +++ .../device_functions/CMakeLists.txt | 22 +++++++++++++++++-- .../host_functions/CMakeLists.txt | 22 ++++++++++++++----- 3 files changed, 40 insertions(+), 7 deletions(-) diff --git a/samples/2_Cookbook/15_static_library/README.md b/samples/2_Cookbook/15_static_library/README.md index bf0123573d..a38dc93451 100644 --- a/samples/2_Cookbook/15_static_library/README.md +++ b/samples/2_Cookbook/15_static_library/README.md @@ -172,5 +172,8 @@ cmake .. make ./test_*.out ``` +It is recommended to use Visual Studio's command prompt for this sample due to requirement of MS Librarian tool - LIB.exe on windows platform. +Override CMAKE_C_COMPILER and CMAKE_CXX_COMPILER to hipcc as Visual Studio's compiler would use cl.exe as default compiler. +i.e. cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=%HIP_PATH% -DCMAKE_C_COMPILER=%HIP_PATH%/bin/hipcc.bat -DCMAKE_CXX_COMPILER=%HIP_PATH%/bin/hipcc.bat .. ## For More Infomation, please refer to the HIP FAQ. diff --git a/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt b/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt index 30a5928389..317b7a819e 100644 --- a/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt +++ b/samples/2_Cookbook/15_static_library/device_functions/CMakeLists.txt @@ -12,6 +12,12 @@ list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) # Find hip find_package(hip REQUIRED) +# For windows, AR is MS Librarian and that is picked by Visual Studio's command prompt. +if (WIN32) + find_program(libpath NAMES lib.exe) + set (CMAKE_AR ${libpath}) +endif() + # Set compiler and linker set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE}) set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) @@ -22,6 +28,11 @@ option(BUILD_SHARED_LIBS "Build as a shared library" OFF) set(CPP_SOURCES hipDevice.cpp) +# For windows, We need to tell cmake how to create static library. +if (WIN32) + set (CMAKE_CXX_CREATE_STATIC_LIBRARY " /out: ") +endif() + # Generate static lib libHipDevice.a add_library(HipDevice STATIC ${CPP_SOURCES}) @@ -35,6 +46,13 @@ set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain2.cpp) add_executable(test_device_static ${TEST_SOURCES}) add_dependencies(test_device_static HipDevice) target_compile_options(test_device_static PRIVATE -fgpu-rdc) -target_link_libraries(test_device_static PRIVATE HipDevice) -target_link_libraries(test_device_static PRIVATE -fgpu-rdc hip::host) + +# For windows, Change in a way to pass lib details +if (WIN32) + target_link_libraries(test_device_static PRIVATE -lHipDevice -L${CMAKE_BINARY_DIR}) +else() + target_link_libraries(test_device_static PRIVATE HipDevice) +endif() + +target_link_libraries(test_device_static PRIVATE -fgpu-rdc amdhip64 amd_comgr) diff --git a/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt b/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt index 347171ed9d..3c7c306866 100644 --- a/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt +++ b/samples/2_Cookbook/15_static_library/host_functions/CMakeLists.txt @@ -12,10 +12,15 @@ list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) # Find hip find_package(hip REQUIRED) +# For windows, AR is MS Librarian and that is pickedby Visual Studio's command prompt. +if (WIN32) + find_program(libpath NAMES lib.exe) + set (CMAKE_AR ${libpath}) +endif() + # Set compiler and linker set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE}) set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) -set(CMAKE_AR ${HIP_HIPCC_EXECUTABLE}) set(CMAKE_BUILD_TYPE Release) # Turn static library generation ON @@ -23,15 +28,17 @@ option(BUILD_SHARED_LIBS "Build as a shared library" OFF) set(CPP_SOURCES hipOptLibrary.cpp) +# For windows, We need to tell cmake how to create static library. +if (WIN32) + set (CMAKE_CXX_CREATE_STATIC_LIBRARY " /out: ") +endif() + # Generate static lib libHipOptLibrary.a. add_library(HipOptLibrary STATIC ${CPP_SOURCES}) # Set-up the correct flags to generate the static library. target_link_libraries(HipOptLibrary PRIVATE --emit-static-lib) target_include_directories(HipOptLibrary PRIVATE /opt/rocm/hsa/include) -get_property(link_libraries TARGET HipOptLibrary PROPERTY LINK_LIBRARIES) -string (REPLACE ";" " " LINK_PROPS "${link_libraries}") -set(CMAKE_CXX_ARCHIVE_CREATE " -o ${LINK_PROPS} ") # Create test executable that uses libHipOptLibrary.a set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain1.cpp) @@ -39,5 +46,10 @@ set(TEST_SOURCES ${CMAKE_SOURCE_DIR}/hipMain1.cpp) add_executable(test_opt_static ${TEST_SOURCES}) add_dependencies(test_opt_static HipOptLibrary) target_link_libraries(test_opt_static PRIVATE -lHipOptLibrary -L${CMAKE_BINARY_DIR}) -target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr hsa-runtime64::hsa-runtime64) + +if (WIN32) + target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr) +else() + target_link_libraries(test_opt_static PRIVATE amdhip64 amd_comgr hsa-runtime64::hsa-runtime64) +endif() From 517c5c7b415d5997b8630b5874425b99798fb59a Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 16:21:15 +0530 Subject: [PATCH 3/6] SWDEV-312831 - Adding device side malloc/new tests (#2635) Following tests are added in this patch 1) Functional Tests 2) Stress Tests 3) Multiprocess Tests Change-Id: Ifd2310fe8f5d555a9af9b87c296fdf6dd4a0a552 --- tests/catch/multiproc/CMakeLists.txt | 1 + .../catch/multiproc/deviceAllocationMproc.cc | 319 ++++ tests/catch/stress/CMakeLists.txt | 1 + .../stress/deviceallocation/CMakeLists.txt | 8 + .../Stress_deviceAllocationStress.cc | 487 ++++++ tests/catch/unit/deviceLib/CMakeLists.txt | 7 + tests/catch/unit/deviceLib/defs.h | 36 + .../catch/unit/deviceLib/deviceAllocCommon.h | 132 ++ .../catch/unit/deviceLib/deviceAllocation.cc | 1494 +++++++++++++++++ .../catch/unit/deviceLib/kerDevAllocMultCO.cc | 39 + .../unit/deviceLib/kerDevAllocSingleKer.cc | 57 + .../catch/unit/deviceLib/kerDevFreeMultCO.cc | 47 + .../catch/unit/deviceLib/kerDevWriteMultCO.cc | 36 + 13 files changed, 2664 insertions(+) create mode 100644 tests/catch/multiproc/deviceAllocationMproc.cc create mode 100644 tests/catch/stress/deviceallocation/CMakeLists.txt create mode 100644 tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc create mode 100644 tests/catch/unit/deviceLib/defs.h create mode 100644 tests/catch/unit/deviceLib/deviceAllocCommon.h create mode 100644 tests/catch/unit/deviceLib/deviceAllocation.cc create mode 100644 tests/catch/unit/deviceLib/kerDevAllocMultCO.cc create mode 100644 tests/catch/unit/deviceLib/kerDevAllocSingleKer.cc create mode 100644 tests/catch/unit/deviceLib/kerDevFreeMultCO.cc create mode 100644 tests/catch/unit/deviceLib/kerDevWriteMultCO.cc diff --git a/tests/catch/multiproc/CMakeLists.txt b/tests/catch/multiproc/CMakeLists.txt index c8da597abd..128384ea3a 100644 --- a/tests/catch/multiproc/CMakeLists.txt +++ b/tests/catch/multiproc/CMakeLists.txt @@ -14,6 +14,7 @@ set(LINUX_TEST_SRC hipMemCoherencyTstMProc.cc hipIpcEventHandle.cc hipIpcMemAccessTest.cc + deviceAllocationMproc.cc ) # the last argument linker libraries is required for this test but optional to the function diff --git a/tests/catch/multiproc/deviceAllocationMproc.cc b/tests/catch/multiproc/deviceAllocationMproc.cc new file mode 100644 index 0000000000..0958cb8781 --- /dev/null +++ b/tests/catch/multiproc/deviceAllocationMproc.cc @@ -0,0 +1,319 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +#ifdef __linux__ +#include +#include +#include +#include +#endif + +#define SIZE 2097152 +// GPU threads +#define BLOCKSIZE 512 +#define GRIDSIZE 256 + +__device__ static char* dev_common_ptr = nullptr; + +/** + * This kernel allocates a memory chunk using malloc(). + */ +static __global__ void kerTestDeviceMalloc(size_t size) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + if (myId == 0) { + dev_common_ptr = reinterpret_cast (malloc(size)); + if (dev_common_ptr == nullptr) { + printf("Device Allocation Failed! \n"); + return; + } + } +} + +/** + * This kernel writes to the memory location allocated in kernel + * kerTestDeviceMalloc or kerTestDeviceNew. + */ +static __global__ void kerTestDeviceWrite() { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + if (dev_common_ptr == nullptr) { + printf("Device Allocation Failed! \n"); + return; + } + *(dev_common_ptr + myId) = SCHAR_MAX; +} + +/** + * This kernel frees the memory chunk allocated in kernel + * kerTestDeviceMalloc using free(). + */ +static __global__ void kerTestDeviceFree(int *result) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + if (myId == 0) { + if (dev_common_ptr != nullptr) { + *result = 1; + for (int idx = 0; idx < (BLOCKSIZE*GRIDSIZE); idx++) { + if (*(dev_common_ptr + myId) != SCHAR_MAX) { + *result = 0; + break; + } + } + free(dev_common_ptr); + } else { + *result = 0; + } + } +} + +/** + * This kernel allocates a memory chunk using new operator. + */ +static __global__ void kerTestDeviceNew(size_t size) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + if (myId == 0) { + dev_common_ptr = new char[size]; + if (dev_common_ptr == nullptr) { + printf("Device Allocation Failed! \n"); + return; + } + } +} + +/** + * This kernel frees the memory chunk allocated in kernel + * kerTestDeviceNew using delete operator. + */ +static __global__ void kerTestDeviceDelete(int *result) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + if (myId == 0) { + if (dev_common_ptr != nullptr) { + *result = 1; + for (int idx = 0; idx < (BLOCKSIZE*GRIDSIZE); idx++) { + if (*(dev_common_ptr + myId) != SCHAR_MAX) { + *result = 0; + break; + } + } + delete[] dev_common_ptr; + } else { + *result = 0; + } + } +} + +/** + * Test device malloc()/new in both Parent and Child Process. + * Allocate SIZE bytes in both parent and child process. Verify + * the allocated size in both parent and child process. + */ +static bool testDeviceAllocMulProc(bool testmalloc) { + int fd[2]; + pid_t childpid; + bool testResult = false; + size_t avail = 0, tot = 0; + // create pipe descriptors + pipe(fd); + // fork process + childpid = fork(); + if (childpid > 0) { // Parent + close(fd[1]); + // Allocate in parent + if (testmalloc) { + kerTestDeviceMalloc<<<1, 1>>>(SIZE); + } else { + kerTestDeviceNew<<<1, 1>>>(SIZE); + } + HIP_CHECK(hipDeviceSynchronize()); + // Check allocated memory size + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if ((tot - avail) < SIZE) { + return false; + } + // parent will wait to read the device cnt + read(fd[0], &testResult, sizeof(testResult)); + // close the read-descriptor + close(fd[0]); + // wait for child exit + wait(NULL); + // At this point the child process exits. + // Ensure that device memory allocated from child is freed. + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if ((tot - avail) < SIZE) { + testResult = false; + } + } else if (!childpid) { // Child + // Wait for hipDeviceSetLimit() completion in parent. + close(fd[0]); + // Allocate in child + if (testmalloc) { + kerTestDeviceMalloc<<<1, 1>>>(SIZE); + } else { + kerTestDeviceNew<<<1, 1>>>(SIZE); + } + HIP_CHECK(hipDeviceSynchronize()); + // Check allocated memory size + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + if ((tot - avail) < SIZE) { + testResult = false; + } else { + testResult = true; + } + // send the value on the write-descriptor: + write(fd[1], &testResult, sizeof(testResult)); + // close the write descriptor: + close(fd[1]); + exit(0); + } + return testResult; +} + +/** + * Test device malloc()/new, write and free()/delete[] + * from both Parent and Child Process. From both Parent and + * Child Process invoke the kernel to allocate memory, the + * kernel to write to the allocated memory and a third kernel + * to verify the memory contents and free it. + */ +static bool testDeviceMemMulProc(bool testmalloc) { + int fd[2]; + bool testResult = false; + pid_t childpid; + int testResultChild = 0; + size_t size = BLOCKSIZE*GRIDSIZE; + // create pipe descriptors + pipe(fd); + // fork process + childpid = fork(); + if (childpid > 0) { // Parent + close(fd[1]); + int *result_d{nullptr}, *result_h{nullptr}; + HIP_CHECK(hipMalloc(&result_d, sizeof(int))); + result_h = reinterpret_cast (malloc(sizeof(int))); + REQUIRE(result_h != nullptr); + // Allocate in parent + if (testmalloc) { + kerTestDeviceMalloc<<<1, 1>>>(size); + } else { + kerTestDeviceNew<<<1, 1>>>(size); + } + // Write + kerTestDeviceWrite<<>>(); + // Free + if (testmalloc) { + kerTestDeviceFree<<<1, 1>>>(result_d); + } else { + kerTestDeviceDelete<<<1, 1>>>(result_d); + } + HIP_CHECK(hipDeviceSynchronize()); + *result_h = 0; + HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int), + hipMemcpyDefault)); + if (*result_h == 0) { + testResult = false; + } else { + testResult = true; + } + // parent will wait to read the device cnt + read(fd[0], &testResultChild, sizeof(int)); + if (testResultChild == 0) { + testResult &= false; + } else { + testResult &= true; + } + // close the read-descriptor + close(fd[0]); + hipFree(result_d); + free(result_h); + // wait for child exit + wait(NULL); + } else if (!childpid) { // Child + // Wait for hipDeviceSetLimit() completion in parent. + close(fd[0]); + int *result_d{nullptr}, *result_h{nullptr}; + HIP_CHECK(hipMalloc(&result_d, sizeof(int))); + result_h = reinterpret_cast (malloc(sizeof(int))); + REQUIRE(result_h != nullptr); + // Allocate in child + if (testmalloc) { + kerTestDeviceMalloc<<<1, 1>>>(size); + } else { + kerTestDeviceNew<<<1, 1>>>(size); + } + // Write + kerTestDeviceWrite<<>>(); + // Free + if (testmalloc) { + kerTestDeviceFree<<<1, 1>>>(result_d); + } else { + kerTestDeviceDelete<<<1, 1>>>(result_d); + } + HIP_CHECK(hipDeviceSynchronize()); + *result_h = 0; + HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int), + hipMemcpyDefault)); + // send the value on the write-descriptor: + write(fd[1], result_h, sizeof(int)); + // close the write descriptor: + close(fd[1]); + hipFree(result_d); + free(result_h); + exit(0); + } + return testResult; +} + +/** + * Multiprocess device side malloc test. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_MultProcess") { + auto res = testDeviceAllocMulProc(true); + REQUIRE(res == true); +} + +/** + * Multiprocess device side new test. + */ +TEST_CASE("Unit_deviceAllocation_New_MultProcess") { + auto res = testDeviceAllocMulProc(false); + REQUIRE(res == true); +} + +/** + * Multiprocess device side malloc, write and free test. + */ +TEST_CASE("Unit_deviceAllocation_MallocFree_MultProcess") { + auto res = testDeviceMemMulProc(true); + REQUIRE(res == true); +} + +/** + * Multiprocess device side new, write and delete test. + */ +TEST_CASE("Unit_deviceAllocation_NewDelete_MultProcess") { + auto res = testDeviceMemMulProc(false); + REQUIRE(res == true); +} diff --git a/tests/catch/stress/CMakeLists.txt b/tests/catch/stress/CMakeLists.txt index 7ecabc21e1..f69313b245 100644 --- a/tests/catch/stress/CMakeLists.txt +++ b/tests/catch/stress/CMakeLists.txt @@ -6,3 +6,4 @@ if(HIP_PLATFORM MATCHES "amd") add_subdirectory(printf) add_subdirectory(stream) endif() +add_subdirectory(deviceallocation) diff --git a/tests/catch/stress/deviceallocation/CMakeLists.txt b/tests/catch/stress/deviceallocation/CMakeLists.txt new file mode 100644 index 0000000000..ee7eb2a580 --- /dev/null +++ b/tests/catch/stress/deviceallocation/CMakeLists.txt @@ -0,0 +1,8 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + Stress_deviceAllocationStress.cc +) + +hip_add_exe_to_target(NAME devalloc + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME stress_test) diff --git a/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc b/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc new file mode 100644 index 0000000000..e60bf66bb4 --- /dev/null +++ b/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc @@ -0,0 +1,487 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include +#include +#include +// Size Macros +#define MEMORY_CHUNK_SIZE (1024*1024) +#define MEMORY_CHUNK_SIZE_ODD (1025*1025) +#define MAXIMUM_CHUNKS (256*1024) +// Subtest Macros +#define NO_ALLOCATION_ONHOST 0 +#define ALLOCATE_ONHOST_HIPMALLOCMANAGED 1 +#define ALLOCATE_ONHOST_HIPMALLOC 2 +// Test Type Macros +#define TEST_MALLOC_FREE 1 +#define TEST_NEW_DELETE 2 +// GPU threads +#define BLOCKSIZE 512 +#define GRIDSIZE 512 +// Test parameters +// Two different loops +#define NUM_OF_LOOP_SINGLE_KER 100000 +#define NUM_OF_LOOP_MULTIPLE_KER 20000 + +// The following flag is defined for platforms (nvidia) +// which honors device memory limit. For AMD this flag +// is disabled and defect is raised. +#if HT_NVIDIA +#define HT_HONORS_DEVICEMEMORY_LIMIT +#endif + +#ifdef HT_HONORS_DEVICEMEMORY_LIMIT +__device__ static char* dev_mem_glob[MAXIMUM_CHUNKS]; +#endif +__device__ static int* dev_mem[GRIDSIZE]; +__device__ static int* dev_common_ptr; + +#ifdef HT_HONORS_DEVICEMEMORY_LIMIT +/** + * This kernel checks kernel allocation of size more than available + * memory. + */ +static __global__ void kerTestDynamicAllocNeg(int test_type, + size_t perThreadSize, + int *ret) { + // Allocate + char* ptr = nullptr; + printf("Memory to allocate in GPU = %zu \n", perThreadSize); + if (test_type == TEST_MALLOC_FREE) { + ptr = reinterpret_cast (malloc(perThreadSize)); + } else { + ptr = new char[perThreadSize]; + } + printf("Allocation Done \n"); + if (ptr == nullptr) { + printf("Allocation Failed. PASSED! \n"); + *ret = 0; + return; + } else { + // Free memory + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } + *ret = -1; + } +} + +/** + * This kernel allocates memory till nullptr is returned. + */ +static __global__ void kerAllocTillExhaust(int test_type, + size_t *total_allocated_mem, + size_t mem_chunk_size) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 of block 0 + if (0 == myId) { + for (int idx = 0; idx < MAXIMUM_CHUNKS; idx++) { + dev_mem_glob[idx] = nullptr; + } + int idx = 0; + if (test_type == TEST_MALLOC_FREE) { + do { + dev_mem_glob[idx] = + reinterpret_cast (malloc(mem_chunk_size)); + if (idx >= MAXIMUM_CHUNKS) { + break; + } + } while (dev_mem_glob[idx++] != nullptr); + } else { + do { + dev_mem_glob[idx] = + reinterpret_cast (new char[mem_chunk_size]); + if (idx >= MAXIMUM_CHUNKS) { + break; + } + } while (dev_mem_glob[idx++] != nullptr); + } + idx = 0; + *total_allocated_mem = 0; + while ((dev_mem_glob[idx] != nullptr) && + (idx < MAXIMUM_CHUNKS)) { + *total_allocated_mem = *total_allocated_mem + mem_chunk_size; + idx++; + } + } +} + +/** + * This kernel deletes the memory. + */ +static __global__ void kerFreeAll(int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + if (0 == myId) { + if (test_type == TEST_MALLOC_FREE) { + int idx = 0; + while (dev_mem_glob[idx] != nullptr) { + free(dev_mem_glob[idx++]); + if (idx >= MAXIMUM_CHUNKS) { + break; + } + } + } else { + int idx = 0; + while (dev_mem_glob[idx] != nullptr) { + delete[] (dev_mem_glob[idx++]); + if (idx >= MAXIMUM_CHUNKS) { + break; + } + } + } + } +} +#endif +/** + * This kernel allocates memory once in thread 0 of each block and + * access this memory in all threads of the block. The memory is + * finally deleted in last thread of each block. + */ +static __global__ void kerBlockLevelMemoryAllocation(int *outputBuf, + int test_type) { + int myThreadId = threadIdx.x, lastThreadId = (blockDim.x - 1); + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 + if (0 == myThreadId) { + if (test_type == TEST_MALLOC_FREE) { + dev_mem[blockIdx.x] = + reinterpret_cast (malloc(blockDim.x*sizeof(int))); + } else { + dev_mem[blockIdx.x] = + reinterpret_cast (new int[blockDim.x]); + } + } + // All threads wait at this barrier + __syncthreads(); + // Check allocated memory in all threads in block before access + if (dev_mem[blockIdx.x] == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + int *ptr = reinterpret_cast (dev_mem[blockIdx.x]); + // Copy to buffer + ptr[myThreadId] = myId; + // All threads wait + __syncthreads(); + // Copy memory to host and free the memory in thread + if (lastThreadId == myThreadId) { + for (size_t idx = 0; idx < blockDim.x; idx++) { + outputBuf[idx + blockDim.x * blockIdx.x] = ptr[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } + } +} + +/** + * This kernel allocates memory in one thread. + */ +static __global__ void kerAlloc(int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 of block 0 + if (0 == myId) { + if (test_type == TEST_MALLOC_FREE) { + dev_common_ptr = + reinterpret_cast (malloc(blockDim.x*gridDim.x*sizeof(int))); + } else { + dev_common_ptr = + reinterpret_cast (new int[blockDim.x*gridDim.x]); + } + } +} + +/** + * This kernel writes to memory allocated in . + */ +static __global__ void kerWrite() { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Check allocated memory in all threads in block before access + if (dev_common_ptr == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + // Copy to buffer + dev_common_ptr[myId] = myId; +} + +/** + * This kernel copies the contents of memory allocated in + * to host and deletes the memory from thread 0. + */ +static __global__ void kerFree(int *outputBuf, int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Check allocated memory in all threads in block before access + if (dev_common_ptr == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + if (0 == myId) { + for (size_t idx = 0; idx < (blockDim.x*gridDim.x); idx++) { + outputBuf[idx] = dev_common_ptr[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(dev_common_ptr); + } else { + delete[] dev_common_ptr; + } + } +} + +#ifdef HT_HONORS_DEVICEMEMORY_LIMIT +/** + * Local function: Launch kerAllocTillExhaust<<<>>> and + * kerFreeAll<<<>>> to test memory allocation till all device + * memory is exhausted. + */ +static bool TestAllocationOfAllAvailableMemory(int test_type, + int category, size_t mem_chunk_size) { + size_t avail1 = 0, avail2 = 0, tot = 0; + constexpr size_t host_alloc = 2147483648; // 2 GB + HIP_CHECK(hipMemGetInfo(&avail1, &tot)); +#if HT_NVIDIA + HIP_CHECK(hipDeviceSetLimit(hipLimitMallocHeapSize, avail1)); +#endif + size_t *tot_alloc_mem_d = nullptr, *tot_alloc_mem_h = nullptr; + tot_alloc_mem_h = + reinterpret_cast (malloc(sizeof(size_t))); + REQUIRE(nullptr != tot_alloc_mem_h); + HIP_CHECK(hipMalloc(&tot_alloc_mem_d, sizeof(size_t))); + REQUIRE(nullptr != tot_alloc_mem_d); + char *devptrHost = nullptr; + if (category == ALLOCATE_ONHOST_HIPMALLOCMANAGED) { + HIP_CHECK(hipMallocManaged(&devptrHost, host_alloc)); + } else if (category == ALLOCATE_ONHOST_HIPMALLOC) { + HIP_CHECK(hipMalloc(&devptrHost, host_alloc)); + } + HIP_CHECK(hipMemGetInfo(&avail1, &tot)); + INFO("Total available memory " << tot); + INFO("Available memory before allocation " << avail1); + // Launch Test Kernel + kerAllocTillExhaust<<<1, 1>>>(test_type, tot_alloc_mem_d, + mem_chunk_size); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(tot_alloc_mem_h, tot_alloc_mem_d, + sizeof(size_t), hipMemcpyDefault)); + HIP_CHECK(hipMemGetInfo(&avail2, &tot)); + kerFreeAll<<<1, 1>>>(test_type); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + bool bPassed = false; + INFO("Available memory after allocation " << avail2); + if (category == NO_ALLOCATION_ONHOST) { + size_t allocated_dev_mem = (tot - avail2); + if (allocated_dev_mem >= *tot_alloc_mem_h) { + bPassed = true; + } + } else if ((category == ALLOCATE_ONHOST_HIPMALLOCMANAGED) || + (category == ALLOCATE_ONHOST_HIPMALLOC)) { + size_t allocated_dev_mem = (tot - avail2 - host_alloc); + if (allocated_dev_mem >= *tot_alloc_mem_h) { + bPassed = true; + } + hipFree(devptrHost); + } + hipFree(tot_alloc_mem_d); + free(tot_alloc_mem_h); + return bPassed; +} +#endif +/** + * Local function: Launch kerBlockLevelMemoryAllocation<<<>>> + * in a loop to stress test allocation and deallocation. + */ +static bool TestMemoryAllocationInLoop(int test_type, + bool isMultikernel = false) { + int *outputVec_d{nullptr}, *outputVec_h{nullptr}; + int arraysize = (BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast (malloc(sizeof(int) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(int) * arraysize))); + bool bPassed = true; + // Launch Test Kernel + int max_index = 0; + if (isMultikernel) { + max_index = NUM_OF_LOOP_MULTIPLE_KER; + } else { + max_index = NUM_OF_LOOP_SINGLE_KER; + } + for (int idx = 0; idx < max_index; idx++) { + if (isMultikernel) { + kerAlloc<<>>(test_type); + kerWrite<<>>(); + kerFree<<>>(outputVec_d, test_type); + } else { + kerBlockLevelMemoryAllocation<<>>(outputVec_d, + test_type); + } + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(int) * arraysize, + hipMemcpyDefault)); + bPassed = true; + for (int idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != idx) { + bPassed = false; + break; + } + } + if (!bPassed) break; + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +#ifdef HT_HONORS_DEVICEMEMORY_LIMIT +/** + * Scenario: Test malloc till nullptr is returned using even chunksize. + */ +TEST_CASE("Stress_deviceAllocation_malloc_Even") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE, + NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE)); +} + +/** + * Scenario: Test malloc till nullptr is returned using odd chunksize. + */ +TEST_CASE("Stress_deviceAllocation_malloc_Odd") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE, + NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE_ODD)); +} + +/** + * Scenario: Test new till nullptr is returned using even chunksize. + */ +TEST_CASE("Stress_deviceAllocation_new_Even") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE, + NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE)); +} + +/** + * Scenario: Test new till nullptr is returned using odd chunksize. + */ +TEST_CASE("Stress_deviceAllocation_new_Odd") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE, + NO_ALLOCATION_ONHOST, MEMORY_CHUNK_SIZE_ODD)); +} + +/** + * Scenario: This test checks device allocation using malloc till nullptr + * is returned. Device memory is also allocated using hipmallocmanaged + * from host. + */ +TEST_CASE("Stress_deviceAllocation_malloc_hipmallocmanaged") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE, + ALLOCATE_ONHOST_HIPMALLOCMANAGED, MEMORY_CHUNK_SIZE)); +} + +/** + * Scenario: This test checks device allocation using new till nullptr + * is returned. Device memory is also allocated using hipmallocmanaged + * from host. + */ +TEST_CASE("Stress_deviceAllocation_new_hipmallocmanaged") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE, + ALLOCATE_ONHOST_HIPMALLOCMANAGED, MEMORY_CHUNK_SIZE)); +} + +/** + * Scenario: This test checks device allocation using malloc till nullptr + * is returned. Device memory is also allocated using hipmalloc from host. + */ +TEST_CASE("Stress_deviceAllocation_malloc_hipmalloc") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_MALLOC_FREE, + ALLOCATE_ONHOST_HIPMALLOC, MEMORY_CHUNK_SIZE)); +} + +/** + * Scenario: This test checks device allocation using new till nullptr + * is returned. Device memory is also allocated using hipmalloc from host. + */ +TEST_CASE("Stress_deviceAllocation_new_hipmalloc") { + REQUIRE(true == TestAllocationOfAllAvailableMemory(TEST_NEW_DELETE, + ALLOCATE_ONHOST_HIPMALLOC, MEMORY_CHUNK_SIZE)); +} + +/** + * Scenario: This test validates device allocation negative scenario + * when size > available memory. + */ +TEST_CASE("Stress_deviceAllocation_Negative") { + int *ret_d{nullptr}, *ret_h{nullptr}; + size_t avail = 0, tot = 0; + HIP_CHECK(hipMemGetInfo(&avail, &tot)); + printf("Available Memory in GPU = %zu \n", avail); + ret_h = reinterpret_cast (malloc(sizeof(int))); + REQUIRE(ret_h != nullptr); + HIP_CHECK(hipMalloc(&ret_d, (sizeof(int)))); + SECTION("Test allocation with malloc") { + kerTestDynamicAllocNeg<<<1, 1>>>(TEST_MALLOC_FREE, (avail + 1), ret_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(ret_h, ret_d, sizeof(int), hipMemcpyDefault)); + REQUIRE(0 == *ret_h); + } + + SECTION("Test allocation with new") { + kerTestDynamicAllocNeg<<<1, 1>>>(TEST_NEW_DELETE, (avail + 1), ret_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(ret_h, ret_d, sizeof(int), hipMemcpyDefault)); + REQUIRE(0 == *ret_h); + } + hipFree(ret_d); + free(ret_h); +} +#endif +/** + * Scenario: This test performs stress test of malloc/free in a loop + * using single kernel. + */ +TEST_CASE("Stress_deviceAllocation_malloc_loop_singlekernel") { + REQUIRE(true == TestMemoryAllocationInLoop(TEST_MALLOC_FREE, false)); +} + +/** + * Scenario: This test performs stress test of new/delete in a loop + * using single kernel. + */ +TEST_CASE("Stress_deviceAllocation_new_loop_singlekernel") { + REQUIRE(true == TestMemoryAllocationInLoop(TEST_NEW_DELETE, false)); +} + +/** + * Scenario: This test performs stress test of malloc/free in a loop + * using multiple kernel. + */ +TEST_CASE("Stress_deviceAllocation_malloc_loop_multkernel") { + REQUIRE(true == TestMemoryAllocationInLoop(TEST_MALLOC_FREE, true)); +} + +/** + * Scenario: This test performs stress test of new/delete in a loop + * using multiple kernel. + */ +TEST_CASE("Stress_deviceAllocation_new_loop_multkernel") { + REQUIRE(true == TestMemoryAllocationInLoop(TEST_NEW_DELETE, true)); +} diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 89027dcbf5..5805e11b5e 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -13,6 +13,7 @@ set(TEST_SRC syncthreadsand.cc syncthreadscount.cc syncthreadsor.cc + deviceAllocation.cc ) # AMD only tests @@ -43,6 +44,11 @@ set(AMD_ARCH_SPEC_TEST_SRC unsafeAtomicAdd_NonCoherent_withunsafeflag.cc ) +add_custom_target(kerDevAllocMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) +add_custom_target(kerDevWriteMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevWriteMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevWriteMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) +add_custom_target(kerDevFreeMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevFreeMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevFreeMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) +add_custom_target(kerDevAllocSingleKer.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocSingleKer.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocSingleKer.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + if(HIP_PLATFORM MATCHES "amd") if (DEFINED OFFLOAD_ARCH_STR) string(FIND ${OFFLOAD_ARCH_STR} "gfx90a" ARCH_CHECK) @@ -78,3 +84,4 @@ elseif(HIP_PLATFORM MATCHES "nvidia") TEST_TARGET_NAME build_tests COMPILE_OPTIONS --Wno-deprecated-declarations) endif() +add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code) diff --git a/tests/catch/unit/deviceLib/defs.h b/tests/catch/unit/deviceLib/defs.h new file mode 100644 index 0000000000..a1b6a5efdf --- /dev/null +++ b/tests/catch/unit/deviceLib/defs.h @@ -0,0 +1,36 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#pragma once + +#define INTERNAL_BUFFER_SIZE 8 +// Test Type +#define TEST_MALLOC_FREE 1 +#define TEST_NEW_DELETE 2 +// Kernel Params +#define BLOCKSIZE 64 +#define GRIDSIZE 32 +// Code Obj +#define DEV_ALLOC_SINGKER_COBJ "kerDevAllocSingleKer.code" +#define DEV_ALLOC_SINGKER_COBJ_FUNC "ker_TestDynamicAllocInAllThreads_CodeObj" +#define DEV_ALLOC_MULCOBJ "kerDevAllocMultCO.code" +#define DEV_WRITE_MULCOBJ "kerDevWriteMultCO.code" +#define DEV_FREE_MULCOBJ "kerDevFreeMultCO.code" +#define DEV_ALLOC_MULCODEOBJ_ALLOC "ker_Alloc_MultCodeObj" +#define DEV_ALLOC_MULCODEOBJ_WRITE "ker_Write_MultCodeObj" +#define DEV_ALLOC_MULCODEOBJ_FREE "ker_Free_MultCodeObj" diff --git a/tests/catch/unit/deviceLib/deviceAllocCommon.h b/tests/catch/unit/deviceLib/deviceAllocCommon.h new file mode 100644 index 0000000000..d5448ddce4 --- /dev/null +++ b/tests/catch/unit/deviceLib/deviceAllocCommon.h @@ -0,0 +1,132 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once +#include +#include +#include +#include +#include +#include "./defs.h" + +static __device__ int* deviceAlloc(int test_type); +static __device__ void deviceWrite(int myId, int *devmem); +static __device__ void deviceFree(int *outputBuf, int *devmem, + int test_type, int myId); + +/** + * Allocation base and derived class to test dynamic allocation. + */ +class baseAlloc{ + public: + virtual __device__ int* alloc(size_t size) = 0; + virtual __device__ void free(int* ptr) = 0; +}; + +class derivedAlloc: public baseAlloc{ + public: + virtual __device__ int* alloc(size_t size) { + return new int[size]; + } + virtual __device__ void free(int* ptr) { + delete ptr; + } +}; + +/** + * Allocation Structure to test dynamic allocation. + */ +struct deviceAllocFunc{ + int* (*alloc)(int); + void (*write)(int, int*); + void (*free)(int*, int*, int, int); +}; + +/** + * Simple Structure to test dynamic allocation. + */ +struct simpleStruct{ + int32_t i; + double d; + float f; + int16_t s; + char c; + int32_t iarr[INTERNAL_BUFFER_SIZE]; + bool operator!=(const struct simpleStruct &inpStr) { + if ((i != inpStr.i) || (d != inpStr.d) || + (f != inpStr.f) || (s != inpStr.s) || (c != inpStr.c)) { + return true; + } + for (int32_t idx = 0; idx < INTERNAL_BUFFER_SIZE; idx++) { + if (iarr[idx] != inpStr.iarr[idx]) { + return true; + } + } + return false; + } +}; + +/** + * Simple Structure containing thread information + */ +struct threadInfo{ + int threadid; + int blockid; + int32_t ival; + double dval; + float fval; + int16_t sval; + char cval; +}; + +/** + * C/C++ Union + */ +union testInfoUnion{ + int32_t ival; + double dval; + float fval; + int16_t sval; + char cval; +}; + +/** + * Complex (nested) Structure to test dynamic allocation using malloc. + */ +struct complexStructure{ + struct threadInfo *sthreadInfo; + __device__ void alloc_internal_members(int test_type, size_t size) { + sthreadInfo = nullptr; + if (test_type == TEST_MALLOC_FREE) { + sthreadInfo = reinterpret_cast( + malloc(size*sizeof(struct threadInfo))); + } else { + sthreadInfo = new struct threadInfo[size]; + } + } + + __device__ void free_internal_members(int test_type) { + if (test_type == TEST_MALLOC_FREE) { + free(sthreadInfo); + } else { + delete[] sthreadInfo; + } + sthreadInfo = nullptr; + } +}; diff --git a/tests/catch/unit/deviceLib/deviceAllocation.cc b/tests/catch/unit/deviceLib/deviceAllocation.cc new file mode 100644 index 0000000000..4982e4f976 --- /dev/null +++ b/tests/catch/unit/deviceLib/deviceAllocation.cc @@ -0,0 +1,1494 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "deviceAllocCommon.h" + +__device__ static void* dev_mem_glob; +__device__ struct deviceAllocFunc allocfunc{&deviceAlloc, &deviceWrite, + &deviceFree}; +__device__ class derivedAlloc classalloc; +constexpr auto num_threads = 5; +static bool thread_results[num_threads]; +__device__ static void* dev_ptr[num_threads][GRIDSIZE]; + +/** + * This kernel allocates and deallocates in every thread + * of every block. + */ +template +static __global__ void kerTestDynamicAllocInAllThread(T *outputBuf, + int test_type, T value, size_t perThreadSize) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + size_t size = 0; + T* ptr = nullptr; + if (test_type == TEST_MALLOC_FREE) { + size = perThreadSize * sizeof(T); + ptr = reinterpret_cast (malloc(size)); + } else { + size = perThreadSize; + ptr = new T[perThreadSize]; + } + if (ptr == nullptr) { + printf("Device Allocation in thread %d Failed! \n", myId); + return; + } + // Set memory + for (size_t idx = 0; idx < perThreadSize; idx++) { + ptr[idx] = value; + } + // Copy to output buffer + for (size_t idx = 0; idx < perThreadSize; idx++) { + outputBuf[myId*perThreadSize + idx] = ptr[idx]; + } + // Free memory + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } +} + +/** + * This kernel allocates and deallocates using virtual functions in every + * thread of every block. + */ +static __global__ void kerTestDynamicAllocVirtualFunc(int *outputBuf, + size_t perThreadSize) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + baseAlloc *palloc = &classalloc; + // Allocate + int* ptr = palloc->alloc(perThreadSize); + + if (ptr == nullptr) { + printf("Device Allocation in thread %d Failed! \n", myId); + return; + } + // Set memory + for (size_t idx = 0; idx < perThreadSize; idx++) { + ptr[idx] = myId; + } + // Copy to output buffer + for (size_t idx = 0; idx < perThreadSize; idx++) { + outputBuf[myId*perThreadSize + idx] = ptr[idx]; + } + // Free memory + palloc->free(ptr); +} + +/** + * This kernel allocates memory in one thread, + * access/modifies it in all threads of block and copies + * data to host and frees the memory in another thread. + */ +template +static __global__ void kerTestAccessInAllThreadsInBlock(T *outputBuf, + int test_type, T value, int host_thr_idx) { + int myThreadId = threadIdx.x, lastThreadId = (blockDim.x - 1); + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 + if (0 == myThreadId) { + if (test_type == TEST_MALLOC_FREE) { + dev_ptr[host_thr_idx][blockIdx.x] = + reinterpret_cast (malloc(blockDim.x*sizeof(T))); + } else { + dev_ptr[host_thr_idx][blockIdx.x] = + reinterpret_cast (new T[blockDim.x]); + } + } + // All threads wait at this barrier + __syncthreads(); + // Check allocated memory in all threads in block before access + if (dev_ptr[host_thr_idx][blockIdx.x] == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + T *ptr = reinterpret_cast (dev_ptr[host_thr_idx][blockIdx.x]); + // Copy to buffer + ptr[myThreadId] = value; + // All threads wait + __syncthreads(); + // Copy memory to host and free the memory in thread + if (lastThreadId == myThreadId) { + for (size_t idx = 0; idx < blockDim.x; idx++) { + outputBuf[idx + blockDim.x * blockIdx.x] = ptr[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } + } +} + +/** + * This kernel allocates a nested structure per block in one thread, + * access/modifies it in all threads of block and copies + * data to host and frees the memory in another thread. + */ +static __global__ void kerTestAccessInAllThreads_CmplxStr(int test_type, + int *result) { + int myThreadId = threadIdx.x; + int lastThreadId = (blockDim.x - 1); + int myBlockId = blockIdx.x; + int myGid = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 + if (0 == myThreadId) { + if (test_type == TEST_MALLOC_FREE) { + dev_ptr[0][blockIdx.x] = + reinterpret_cast (malloc(sizeof(struct complexStructure))); + } else { + dev_ptr[0][blockIdx.x] = + reinterpret_cast (new struct complexStructure); + } + struct complexStructure *ptr = + reinterpret_cast (dev_ptr[0][blockIdx.x]); + ptr->alloc_internal_members(test_type, BLOCKSIZE); + } + // All threads wait at this barrier + __syncthreads(); + // Check allocated memory in all threads in block before access + if (dev_ptr[0][blockIdx.x] == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myGid); + return; + } + struct complexStructure *ptr = + reinterpret_cast (dev_ptr[0][blockIdx.x]); + if (ptr->sthreadInfo == nullptr) { + printf("Structure Allocation Failed in thread = %d \n", myGid); + return; + } + // Copy to buffer + ptr->sthreadInfo[myThreadId].threadid = myThreadId; + ptr->sthreadInfo[myThreadId].blockid = myBlockId; + ptr->sthreadInfo[myThreadId].ival = INT_MAX; + ptr->sthreadInfo[myThreadId].dval = DBL_MAX; + ptr->sthreadInfo[myThreadId].fval = FLT_MAX; + ptr->sthreadInfo[myThreadId].sval = SHRT_MAX; + ptr->sthreadInfo[myThreadId].cval = SCHAR_MAX; + // All threads wait + __syncthreads(); + // Copy memory to host and free the memory in thread + if (lastThreadId == myThreadId) { + int match = 1; + for (int idx = 0; idx < BLOCKSIZE; idx++) { + if ((ptr->sthreadInfo[idx].threadid != idx) || + (ptr->sthreadInfo[idx].blockid != myBlockId) || + (ptr->sthreadInfo[idx].ival != INT_MAX) || + (ptr->sthreadInfo[idx].dval != DBL_MAX) || + (ptr->sthreadInfo[idx].fval != FLT_MAX) || + (ptr->sthreadInfo[idx].sval != SHRT_MAX) || + (ptr->sthreadInfo[idx].cval != SCHAR_MAX)) { + match = 0; + break; + } + } + result[blockIdx.x] = match; + ptr->free_internal_members(test_type); + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete ptr; + } + } +} + +/** + * This kernel allocates a union per block in one thread, + * access/modifies it in all threads of block and copies + * data to host and frees the memory in another thread. + */ +static __global__ void kerTestAccessInAllThreadsForUnion( + testInfoUnion *outputBuf, int test_type) { + int myThreadId = threadIdx.x, lastThreadId = (blockDim.x - 1); + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 + if (0 == myThreadId) { + if (test_type == TEST_MALLOC_FREE) { + dev_ptr[0][blockIdx.x] = + reinterpret_cast (malloc(blockDim.x*sizeof(testInfoUnion))); + } else { + dev_ptr[0][blockIdx.x] = + reinterpret_cast (new testInfoUnion[blockDim.x]); + } + } + // All threads wait at this barrier + __syncthreads(); + // Check allocated memory in all threads in block before access + if (dev_ptr[0][blockIdx.x] == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + testInfoUnion *ptr = + reinterpret_cast (dev_ptr[0][blockIdx.x]); + // Copy to buffer + switch (myId % 5) { + case 0: ptr[myThreadId].ival = INT_MAX; break; + case 1: ptr[myThreadId].dval = DBL_MAX; break; + case 2: ptr[myThreadId].fval = FLT_MAX; break; + case 3: ptr[myThreadId].sval = SHRT_MAX; break; + case 4: ptr[myThreadId].cval = SCHAR_MAX; break; + } + // All threads wait + __syncthreads(); + // Copy memory to host and free the memory in thread + if (lastThreadId == myThreadId) { + for (size_t idx = 0; idx < blockDim.x; idx++) { + outputBuf[idx + blockDim.x * blockIdx.x] = ptr[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } + } +} + +/** + * This kernel allocates memory in one thread. + */ +template +static __global__ void kerAlloc(int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 of block 0 + if (0 == myId) { + if (test_type == TEST_MALLOC_FREE) { + dev_mem_glob = + reinterpret_cast (malloc(blockDim.x*gridDim.x*sizeof(T))); + } else { + dev_mem_glob = + reinterpret_cast (new T[blockDim.x*gridDim.x]); + } + } +} + +/** + * This kernel writes to memory allocated in . + */ +template +static __global__ void kerWrite(T value) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Check allocated memory in all threads in block before access + if (dev_mem_glob == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + T *ptr = reinterpret_cast (dev_mem_glob); + // Copy to buffer + ptr[myId] = value; +} + +/** + * This kernel copies the contents of memory allocated in + * to host and deletes the memory from thread 0. + */ +template +static __global__ void kerFree(T *outputBuf, int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Check allocated memory in all threads in block before access + if (dev_mem_glob == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + + T *ptr = reinterpret_cast (dev_mem_glob); + if (0 == myId) { + for (size_t idx = 0; idx < (blockDim.x*gridDim.x); idx++) { + outputBuf[idx] = ptr[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } + } +} + +/** + * This device function allocates memory in one thread. + */ +static __device__ int* deviceAlloc(int test_type) { + int *ptr = nullptr; + if (test_type == TEST_MALLOC_FREE) { + ptr = + reinterpret_cast (malloc(INTERNAL_BUFFER_SIZE*sizeof(int))); + } else { + ptr = + reinterpret_cast (new int[INTERNAL_BUFFER_SIZE]); + } + return ptr; +} + +/** + * This device function writes to memory allocated in deviceAlloc(). + */ +static __device__ void deviceWrite(int myId, int *devmem) { + // Check allocated memory in all threads in block before access + if (devmem == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + // Copy to buffer + for (size_t idx = 0; idx < INTERNAL_BUFFER_SIZE; idx++) { + devmem[idx] = myId; + } +} + +/** + * This device function copies the contents of memory allocated + * in deviceAlloc() to host and deletes the memory from thread 0. + */ +static __device__ void deviceFree(int *outputBuf, int *devmem, + int test_type, int myId) { + // Check allocated memory in all threads in block before access + if (devmem == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + for (size_t idx = 0; idx < INTERNAL_BUFFER_SIZE; idx++) { + outputBuf[myId*INTERNAL_BUFFER_SIZE + idx] = devmem[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(devmem); + } else { + delete[] devmem; + } +} + +/** + * This kernel invokes __device__ allocation functions via pointers. + */ +static __global__ void kerTestAllocationUsingDevFunc(int *outputBuf, + int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + struct deviceAllocFunc *func = &allocfunc; + int *dev_ptr = nullptr; + dev_ptr = func->alloc(test_type); + func->write(myId, dev_ptr); + func->free(outputBuf, dev_ptr, test_type, myId); +} + +/** + * Local function: Allocate local and device memory from host, + * launches kerTestDynamicAllocInAllThread<<<>>> and copies data back + * to host to validate. + */ +template +static bool TestAllocInAllThread(int test_type, + T value, size_t sizeBufferPerThread) { + T *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (sizeBufferPerThread * BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast (malloc(sizeof(T) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(T) * arraysize))); + // Launch Test Kernel + kerTestDynamicAllocInAllThread<<>>( + outputVec_d, test_type, value, sizeBufferPerThread); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(T) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != value) { + bPassed = false; + break; + } + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Local function: Allocate local and device memory from host, + * launches kerTestAccessInAllThreadsInBlock<<<>>> and copies data back + * to host to validate. + */ +template +static bool TestMemoryAccessInAllThread(int test_type, int thread_idx) { + T *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (BLOCKSIZE * GRIDSIZE); + T data_value = std::numeric_limits::max(); + outputVec_h = reinterpret_cast (malloc(sizeof(T) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(T) * arraysize))); + // Launch Test Kernel + kerTestAccessInAllThreadsInBlock<<>>(outputVec_d, + test_type, data_value, thread_idx); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(T) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != data_value) { + bPassed = false; + break; + } + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Local function: Launch kerAlloc<<<>>> + */ +template +static void runTestMemoryAccessInAllThread(int test_type, int thread_idx) { + thread_results[thread_idx] = TestMemoryAccessInAllThread(test_type, + thread_idx); +} + +/** + * Local function: Launch kerAlloc<<<>>>, kerWrite<<<>>> and kerFree<<<>>> + * to test kernel allocated memory access across multiple kernels and multiple + * streams. + */ +template +static bool TestMemoryAcrossMulKernels(int test_type, + bool multistream = false) { + T *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (BLOCKSIZE * GRIDSIZE); + T data_value = std::numeric_limits::max(); + outputVec_h = reinterpret_cast (malloc(sizeof(T) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(T) * arraysize))); + // Launch Test Kernel + if (multistream) { + hipStream_t stream1, stream2, stream3; + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&stream3)); + kerAlloc<<>>(test_type); + HIP_CHECK(hipStreamSynchronize(stream1)); + kerWrite<<>>(data_value); + HIP_CHECK(hipStreamSynchronize(stream2)); + kerFree<<>>(outputVec_d, test_type); + HIP_CHECK(hipStreamSynchronize(stream3)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipStreamDestroy(stream3)); + } else { + kerAlloc<<>>(test_type); + kerWrite<<>>(data_value); + kerFree<<>>(outputVec_d, test_type); + HIP_CHECK(hipDeviceSynchronize()); + } + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(T) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != data_value) { + bPassed = false; + break; + } + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Local function: Launch kerAlloc<<<>>> + */ +template +static void runKerAlloc(int test_type) { + kerAlloc<<>>(test_type); +} + +/** + * Local function: Launch kerWrite<<<>>> + */ +template +static void runKerWrite(T data_value) { + kerWrite<<>>(data_value); +} + +/** + * Local function: Launch kerFree<<<>>> + */ +template +static void runKerFree(T *outputVec_d, int test_type) { + kerFree<<>>(outputVec_d, test_type); +} + +/** + * Local function: Launch kerAlloc<<<>>>, kerWrite<<<>>> and kerFree<<<>>> + * across multiple threads. + */ +template +static bool TestDevMemAllocMulKerMulThrd(int test_type) { + T *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (BLOCKSIZE * GRIDSIZE); + T data_value = std::numeric_limits::max(); + outputVec_h = reinterpret_cast (malloc(sizeof(T) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(T) * arraysize))); + // Launch all Test Kernel threads + std::thread threadAlloc(runKerAlloc, test_type); + threadAlloc.join(); + std::thread threadWrite(runKerWrite, data_value); + threadWrite.join(); + std::thread threadFree(runKerFree, outputVec_d, test_type); + threadFree.join(); + // Wait for all kernels in device + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(T) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != data_value) { + bPassed = false; + break; + } + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} +/** + * Local function: Allocate local and device memory from host, + * launches kerTestAccessInAllThreads_CmplxStr<<<>>> and copies data back + * to host to validate. + */ +static bool TestMemoryAccessInAllThread_CmplxStr(int test_type) { + int *result_d{nullptr}, *result_h{nullptr}; + size_t arraysize = BLOCKSIZE; + result_h = reinterpret_cast (malloc(sizeof(int) * arraysize)); + REQUIRE(result_h != nullptr); + HIP_CHECK(hipMalloc(&result_d, (sizeof(int) * arraysize))); + HIP_CHECK(hipMemset(result_d, 0, (sizeof(int) * arraysize))); + // Launch Test Kernel + kerTestAccessInAllThreads_CmplxStr<<>>( + test_type, result_d); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(int) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < GRIDSIZE; idx++) { + if (result_h[idx] != 1) { + bPassed = false; + break; + } + } + hipFree(result_d); + free(result_h); + return bPassed; +} + +/** + * Local function: Allocate host and device memory of type union, + * launches kerTestAccessInAllThreadsForUnion<<<>>> and copies data back + * to host to validate. + */ +static bool TestMemoryAccessInAllThread_Union(int test_type) { + testInfoUnion *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast + (malloc(sizeof(testInfoUnion) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, + (sizeof(testInfoUnion) * arraysize))); + // Launch Test Kernel + kerTestAccessInAllThreadsForUnion<<>>(outputVec_d, + test_type); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, + sizeof(testInfoUnion) * arraysize, hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + switch (idx % 5) { + case 0: + if (outputVec_h[idx].ival != INT_MAX) { + bPassed = false; + } + break; + case 1: + if (outputVec_h[idx].dval != DBL_MAX) { + bPassed = false; + } + break; + case 2: + if (outputVec_h[idx].fval != FLT_MAX) { + bPassed = false; + } + break; + case 3: + if (outputVec_h[idx].sval != SHRT_MAX) { + bPassed = false; + } + break; + case 4: + if (outputVec_h[idx].cval != SCHAR_MAX) { + bPassed = false; + } + break; + } + if (bPassed == false) break; + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Local function: Allocate local and device memory from host, + * launches ker_TestDynamicAllocInAllThreads_CodeObj<<<>>> and + * copies data back to host to validate. + */ +static bool TestAlloc_Load_SingleKer_AllocFree(int test_type, + int value, size_t sizeBufferPerThread) { + int *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (sizeBufferPerThread * BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast (malloc(sizeof(int) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(int) * arraysize))); + // Launch Test Kernel + hipModule_t Module; + hipFunction_t Function; + HIP_CHECK(hipModuleLoad(&Module, DEV_ALLOC_SINGKER_COBJ)); + HIP_CHECK(hipModuleGetFunction(&Function, Module, + DEV_ALLOC_SINGKER_COBJ_FUNC)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + struct { + void* _Output_d; + int _test_type; + int _value; + size_t _size; + } args; + args._Output_d = reinterpret_cast(outputVec_d); + args._test_type = test_type; + args._value = value; + args._size = sizeBufferPerThread; + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + HIP_CHECK(hipModuleLaunchKernel(Function, GRIDSIZE, 1, 1, + BLOCKSIZE, 1, 1, 0, + stream, NULL, + reinterpret_cast(&config))); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(int) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != value) { + bPassed = false; + break; + } + } + HIP_CHECK(hipModuleUnload(Module)); + HIP_CHECK(hipStreamDestroy(stream)); + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Local function: Allocate local and device memory from host, + * launches ker_Alloc_MultCodeObj<<<>>>, ker_Write_MultCodeObj<<<>>> and + * ker_Free_MultCodeObj<<<>>> copies data back to host to validate. + */ +static bool TestAlloc_Load_MultKernels(int test_type, + int value) { + int *outputVec_d{nullptr}, *outputVec_h{nullptr}; + int **dev_addr{nullptr}; + size_t arraysize = (BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast (malloc(sizeof(int) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(int) * arraysize))); + HIP_CHECK(hipMalloc(&dev_addr, (sizeof(int*)))); + // Launch Test Kernel + hipModule_t ModuleAlloc, ModuleWrite, ModuleFree; + hipFunction_t FunctionAlloc, FunctionAcess, FunctionFree; + // Load ker_Alloc_MultCodeObj + HIP_CHECK(hipModuleLoad(&ModuleAlloc, DEV_ALLOC_MULCOBJ)); + HIP_CHECK(hipModuleLoad(&ModuleWrite, DEV_WRITE_MULCOBJ)); + HIP_CHECK(hipModuleLoad(&ModuleFree, DEV_FREE_MULCOBJ)); + HIP_CHECK(hipModuleGetFunction(&FunctionAlloc, ModuleAlloc, + DEV_ALLOC_MULCODEOBJ_ALLOC)); + // Load ker_Write_MultCodeObj + HIP_CHECK(hipModuleGetFunction(&FunctionAcess, ModuleWrite, + DEV_ALLOC_MULCODEOBJ_WRITE)); + // Load ker_Free_MultCodeObj + HIP_CHECK(hipModuleGetFunction(&FunctionFree, ModuleFree, + DEV_ALLOC_MULCODEOBJ_FREE)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + struct { + void **__dev_addr; + int _test_type; + } args1; + args1.__dev_addr = reinterpret_cast(dev_addr); + args1._test_type = test_type; + size_t size1 = sizeof(args1); + + void* config1[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args1, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size1, + HIP_LAUNCH_PARAM_END}; + struct { + void **__dev_addr; + int _value; + } args2; + args2.__dev_addr = reinterpret_cast(dev_addr); + args2._value = value; + size_t size2 = sizeof(args2); + + void* config2[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args2, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size2, + HIP_LAUNCH_PARAM_END}; + struct { + void* _output; + void **__dev_addr; + int _test_type; + } args3; + args3._output = reinterpret_cast(outputVec_d); + args3.__dev_addr = reinterpret_cast(dev_addr); + args3._test_type = test_type; + size_t size3 = sizeof(args3); + + void* config3[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args3, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size3, + HIP_LAUNCH_PARAM_END}; + // Launch ker_Alloc_MultCodeObj + HIP_CHECK(hipModuleLaunchKernel(FunctionAlloc, GRIDSIZE, 1, 1, + BLOCKSIZE, 1, 1, 0, + stream, NULL, + reinterpret_cast(&config1))); + // Launch ker_Write_MultCodeObj + HIP_CHECK(hipModuleLaunchKernel(FunctionAcess, GRIDSIZE, 1, 1, + BLOCKSIZE, 1, 1, 0, + stream, NULL, + reinterpret_cast(&config2))); + // Launch ker_Free_MultCodeObj + HIP_CHECK(hipModuleLaunchKernel(FunctionFree, GRIDSIZE, 1, 1, + BLOCKSIZE, 1, 1, 0, + stream, NULL, + reinterpret_cast(&config3))); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(int) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != value) { + bPassed = false; + break; + } + } + HIP_CHECK(hipModuleUnload(ModuleAlloc)); + HIP_CHECK(hipModuleUnload(ModuleWrite)); + HIP_CHECK(hipModuleUnload(ModuleFree)); + HIP_CHECK(hipStreamDestroy(stream)); + hipFree(dev_addr); + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Local function: Launch kerAlloc<<<>>>, kerWrite<<<>>> and kerFree<<<>>> + * to test kernel allocated memory access across multiple kernels using + * hipGraph. + */ +template +static bool TestMemoryAcrossMulKernelsUsingGraph(int test_type) { + T *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (BLOCKSIZE * GRIDSIZE); + T data_value = std::numeric_limits::max(); + outputVec_h = reinterpret_cast (malloc(sizeof(T) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(T) * arraysize))); + // Launch Test Kernels using graph + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create Allocation Kernel Node + hipGraphNode_t kernelnode_1; + hipKernelNodeParams kernelNodeParams1{}; + void* kernelArgs1[] = {reinterpret_cast(&test_type)}; + kernelNodeParams1.func = reinterpret_cast(kerAlloc); + kernelNodeParams1.gridDim = dim3(GRIDSIZE); + kernelNodeParams1.blockDim = dim3(BLOCKSIZE); + kernelNodeParams1.sharedMemBytes = 0; + kernelNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph, nullptr, 0, + &kernelNodeParams1)); + // Create Write Kernel Node + hipGraphNode_t kernelnode_2; + hipKernelNodeParams kernelNodeParams2{}; + void* kernelArgs2[] = {reinterpret_cast(&data_value)}; + kernelNodeParams2.func = reinterpret_cast(kerWrite); + kernelNodeParams2.gridDim = dim3(GRIDSIZE); + kernelNodeParams2.blockDim = dim3(BLOCKSIZE); + kernelNodeParams2.sharedMemBytes = 0; + kernelNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams2.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph, nullptr, 0, + &kernelNodeParams2)); + // Create Free Kernel Node + hipGraphNode_t kernelnode_3; + hipKernelNodeParams kernelNodeParams3{}; + void* kernelArgs3[] = + {&outputVec_d, reinterpret_cast(&test_type)}; + kernelNodeParams3.func = reinterpret_cast(kerFree); + kernelNodeParams3.gridDim = dim3(GRIDSIZE); + kernelNodeParams3.blockDim = dim3(BLOCKSIZE); + kernelNodeParams3.sharedMemBytes = 0; + kernelNodeParams3.kernelParams = reinterpret_cast(kernelArgs3); + kernelNodeParams3.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_3, graph, nullptr, 0, + &kernelNodeParams3)); + // Create Memcpy Node + hipGraphNode_t memcpyD2H; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nullptr, 0, + outputVec_h, outputVec_d, (sizeof(T) * arraysize), + hipMemcpyDeviceToHost)); + // Create dependencies for graph + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode_1, + &kernelnode_2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode_2, + &kernelnode_3, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode_3, + &memcpyD2H, 1)); + // Instantiate and launch the graphs + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != data_value) { + bPassed = false; + break; + } + } + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} +/** + * Local function: Allocate local and device memory from host, + * launches kerTestAllocationUsingDevFunc<<<>>> and copies data back + * to host to validate. + */ +static bool TestAllocInDeviceFunc(int test_type) { + int *outputVec_d{nullptr}, *outputVec_h{nullptr}; + size_t arraysize = (INTERNAL_BUFFER_SIZE * BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast (malloc(sizeof(int) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(int) * arraysize))); + // Launch Test Kernel + kerTestAllocationUsingDevFunc<<>>(outputVec_d, + test_type); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(int) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != (idx / INTERNAL_BUFFER_SIZE)) { + bPassed = false; + break; + } + } + hipFree(outputVec_d); + free(outputVec_h); + return bPassed; +} + +/** + * Scenario: This test validates device allocation and deallocation + * using malloc/free in every gpu thread and block for primitive data + * types like char, short, int etc. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_PrimitiveDataType") { + constexpr size_t sizePerThread = 128; + + // malloc()/free() tests + SECTION("Test char datatype allocation with malloc") { + REQUIRE(true == TestAllocInAllThread(TEST_MALLOC_FREE, + SCHAR_MAX, sizePerThread)); + } + + SECTION("Test short datatype allocation with malloc") { + REQUIRE(true == TestAllocInAllThread(TEST_MALLOC_FREE, + SHRT_MAX, sizePerThread)); + } + + SECTION("Test int datatype allocation with malloc") { + REQUIRE(true == TestAllocInAllThread(TEST_MALLOC_FREE, + INT_MAX, sizePerThread)); + } + + SECTION("Test float datatype allocation with malloc") { + REQUIRE(true == TestAllocInAllThread(TEST_MALLOC_FREE, + FLT_MAX, sizePerThread)); + } + + SECTION("Test double datatype allocation with malloc") { + REQUIRE(true == TestAllocInAllThread(TEST_MALLOC_FREE, + DBL_MAX, sizePerThread)); + } +} + +/** + * Scenario: This test validates device allocation and deallocation + * using new/delete in every gpu thread and block for primitive data + * types like char, short, int etc. + */ +TEST_CASE("Unit_deviceAllocation_New_PerThread_PrimitiveDataType") { + constexpr size_t sizePerThread = 128; + + // new/delete tests + SECTION("Test char datatype allocation with new") { + REQUIRE(true == TestAllocInAllThread(TEST_NEW_DELETE, + SCHAR_MAX, sizePerThread)); + } + + SECTION("Test short datatype allocation with new") { + REQUIRE(true == TestAllocInAllThread(TEST_NEW_DELETE, + SHRT_MAX, sizePerThread)); + } + + SECTION("Test int datatype allocation with new") { + REQUIRE(true == TestAllocInAllThread(TEST_NEW_DELETE, + INT_MAX, sizePerThread)); + } + + SECTION("Test float datatype allocation with new") { + REQUIRE(true == TestAllocInAllThread(TEST_NEW_DELETE, + FLT_MAX, sizePerThread)); + } + + SECTION("Test double datatype allocation with new") { + REQUIRE(true == TestAllocInAllThread(TEST_NEW_DELETE, + DBL_MAX, sizePerThread)); + } +} + +/** + * Scenario: This test validates device allocation and deallocation + * using malloc/free in every gpu thread and block for structure. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_StructDataType") { + constexpr size_t sizePerThread = 64; + struct simpleStruct sampleStr{INT_MAX, DBL_MAX, FLT_MAX, SHRT_MAX, + SCHAR_MAX, {1, 2, 3, 4, 5, 6, 7, 8}}; + REQUIRE(true == TestAllocInAllThread(TEST_MALLOC_FREE, + sampleStr, sizePerThread)); +} + +/** + * Scenario: This test validates device allocation and deallocation + * using new/delete in every gpu thread and block for structure. + */ +TEST_CASE("Unit_deviceAllocation_New_PerThread_StructDataType") { + constexpr size_t sizePerThread = 64; + struct simpleStruct sampleStr{INT_MAX, DBL_MAX, FLT_MAX, SHRT_MAX, + SCHAR_MAX, {1, 2, 3, 4, 5, 6, 7, 8}}; + REQUIRE(true == TestAllocInAllThread(TEST_NEW_DELETE, + sampleStr, sizePerThread)); +} + +/** + * Scenario: This test validates device memory allocation and free + * in 1 thread and access in block for different primitive types like + * char, short, int etc. + */ +TEST_CASE("Unit_deviceAllocation_InOneThread_AccessInAllThreads") { + // malloc()/free() tests + SECTION("Test char datatype allocation with malloc") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_MALLOC_FREE, 0)); + } + + SECTION("Test short datatype allocation with malloc") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_MALLOC_FREE, 0)); + } + + SECTION("Test int datatype allocation with malloc") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_MALLOC_FREE, 0)); + } + + SECTION("Test float datatype allocation with malloc") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_MALLOC_FREE, 0)); + } + + SECTION("Test double datatype allocation with malloc") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_MALLOC_FREE, 0)); + } + + // new/delete tests + SECTION("Test char datatype allocation with new") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_NEW_DELETE, 0)); + } + + SECTION("Test short datatype allocation with new") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_NEW_DELETE, 0)); + } + + SECTION("Test int datatype allocation with new") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_NEW_DELETE, 0)); + } + + SECTION("Test float datatype allocation with new") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_NEW_DELETE, 0)); + } + + SECTION("Test double datatype allocation with new") { + REQUIRE(true == TestMemoryAccessInAllThread(TEST_NEW_DELETE, 0)); + } +} + +/** + * Scenario: This test validates device allocation malloc, access and free + * across multiple kernels for different primitive types like char, short, + * int etc. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_AcrossKernels") { + // malloc()/free() tests + SECTION("Test char datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE)); + } + + SECTION("Test short datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE)); + } + + SECTION("Test int datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE)); + } + + SECTION("Test float datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE)); + } + + SECTION("Test double datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE)); + } +} + +/** + * Scenario: This test validates device new, access and delete + * across multiple kernels for different primitive types like char, short, + * int etc. + */ +TEST_CASE("Unit_deviceAllocation_New_AcrossKernels") { + // new/delete tests + SECTION("Test char datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE)); + } + + SECTION("Test short datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE)); + } + + SECTION("Test int datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE)); + } + + SECTION("Test float datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE)); + } + + SECTION("Test double datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE)); + } +} + +/** + * Scenarios: + * A) This test validates device allocation malloc, access and free + * across multiple kernels for nested structure. + * B) This test also validates memory allocation and deallocation through + * __device__ functions. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_ComplexDataType") { + // malloc()/free() tests + REQUIRE(true == TestMemoryAccessInAllThread_CmplxStr(TEST_MALLOC_FREE)); +} + +/** + * Scenario: + * A) This test validates device allocation malloc, access and free + * across multiple kernels for nested structure. + * B) This test also validates memory allocation and deallocation through + * __device__ functions. + */ +TEST_CASE("Unit_deviceAllocation_New_ComplexDataType") { + // new/delete tests + REQUIRE(true == TestMemoryAccessInAllThread_CmplxStr(TEST_NEW_DELETE)); +} + +/** + * Scenario: This test validates device allocation malloc, access and free + * across multiple kernels for Union data type. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_UnionType") { + // malloc()/free() tests + REQUIRE(true == TestMemoryAccessInAllThread_Union(TEST_MALLOC_FREE)); +} + +/** + * Scenario: This test validates device allocation new, access and delete + * across multiple kernels for Union data type. + */ +TEST_CASE("Unit_deviceAllocation_New_UnionType") { + // new/delete tests + REQUIRE(true == TestMemoryAccessInAllThread_Union(TEST_NEW_DELETE)); +} + +/** + * Scenario: This test validates device allocation and deallocation + * using malloc/free in every gpu thread and block using Single + * Code Object kernel. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_SingleCodeObj") { + constexpr size_t sizePerThread = 128; + + REQUIRE(true == TestAlloc_Load_SingleKer_AllocFree(TEST_MALLOC_FREE, + INT_MAX, sizePerThread)); +} + +/** + * Scenario: This test validates device allocation and deallocation + * using new/delete in every gpu thread and block using Single + * Code Object kernel. + */ +TEST_CASE("Unit_deviceAllocation_New_SingleCodeObj") { + constexpr size_t sizePerThread = 128; + + REQUIRE(true == TestAlloc_Load_SingleKer_AllocFree(TEST_NEW_DELETE, + INT_MAX, sizePerThread)); +} + +#if HT_NVIDIA +/** + * Scenario: This test validates device allocation and deallocation + * using malloc/free in multikernel and multistream environment. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_MultKerMultStrm") { + // malloc()/free() tests + SECTION("Test char datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE, + true)); + } + + SECTION("Test short datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE, + true)); + } + + SECTION("Test int datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE, + true)); + } + + SECTION("Test float datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE, + true)); + } + + SECTION("Test double datatype allocation with malloc") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE, + true)); + } +} + +/** + * Scenario: This test validates device allocation and deallocation + * using new/delete in multikernel and multistream environment. + */ +TEST_CASE("Unit_deviceAllocation_New_PerThread_MultKerMultStrm") { + // new/delete tests + SECTION("Test char datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE, + true)); + } + + SECTION("Test short datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE, + true)); + } + + SECTION("Test int datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE, + true)); + } + + SECTION("Test float datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE, + true)); + } + + SECTION("Test double datatype allocation with new") { + REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE, + true)); + } +} +#endif + +/** + * Scenario: This test validates device allocation and deallocation + * using malloc/free in graph. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_Graph") { + // malloc()/free() tests + SECTION("Test char datatype allocation with malloc") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_MALLOC_FREE)); + } + + SECTION("Test short datatype allocation with malloc") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_MALLOC_FREE)); + } + + SECTION("Test int datatype allocation with malloc") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_MALLOC_FREE)); + } + + SECTION("Test float datatype allocation with malloc") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_MALLOC_FREE)); + } + + SECTION("Test double datatype allocation with malloc") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_MALLOC_FREE)); + } +} + +/** + * Scenario: This test validates device allocation and deallocation + * using new/delete in graph. + */ +TEST_CASE("Unit_deviceAllocation_New_PerThread_Graph") { + // new/delete tests + SECTION("Test char datatype allocation with new") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_NEW_DELETE)); + } + + SECTION("Test short datatype allocation with new") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_NEW_DELETE)); + } + + SECTION("Test int datatype allocation with new") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_NEW_DELETE)); + } + + SECTION("Test float datatype allocation with new") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_NEW_DELETE)); + } + + SECTION("Test double datatype allocation with new") { + REQUIRE(true == + TestMemoryAcrossMulKernelsUsingGraph(TEST_NEW_DELETE)); + } +} + +/** + * Scenario: This test validates device allocation malloc, access and free + * using pointers to device functions. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_DeviceFunc") { + // malloc/free tests + REQUIRE(true == TestAllocInDeviceFunc(TEST_MALLOC_FREE)); +} + +/** + * Scenario: This test validates device allocation new, access and delete + * using pointers to device functions. + */ +TEST_CASE("Unit_deviceAllocation_New_DeviceFunc") { + // new/delete tests + REQUIRE(true == TestAllocInDeviceFunc(TEST_NEW_DELETE)); +} + +/** + * Scenario: This test validates device allocation using vitual functions + */ +TEST_CASE("Unit_deviceAllocation_VirtualFunction") { + int *outputVec_d{nullptr}, *outputVec_h{nullptr}; + constexpr size_t sizeBufferPerThread = 8; + size_t arraysize = (sizeBufferPerThread * BLOCKSIZE * GRIDSIZE); + outputVec_h = reinterpret_cast (malloc(sizeof(int) * arraysize)); + REQUIRE(outputVec_h != nullptr); + HIP_CHECK(hipMalloc(&outputVec_d, (sizeof(int) * arraysize))); + // Launch Test Kernel + kerTestDynamicAllocVirtualFunc<<>>( + outputVec_d, sizeBufferPerThread); + HIP_CHECK(hipDeviceSynchronize()); + // Copy to host buffer + HIP_CHECK(hipMemcpy(outputVec_h, outputVec_d, sizeof(int) * arraysize, + hipMemcpyDefault)); + bool bPassed = true; + for (size_t idx = 0; idx < arraysize; idx++) { + if (outputVec_h[idx] != (idx / sizeBufferPerThread)) { + bPassed = false; + break; + } + } + REQUIRE(true == bPassed); + hipFree(outputVec_d); + free(outputVec_h); +} + +/** + * Scenario: This test validates device allocation malloc, access and free + * across multiple kernels launched using threads. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_MulKernels_MulThreads") { + // malloc()/free() tests + SECTION("Test char datatype allocation with malloc") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_MALLOC_FREE)); + } + + SECTION("Test short datatype allocation with malloc") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_MALLOC_FREE)); + } + + SECTION("Test int datatype allocation with malloc") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_MALLOC_FREE)); + } + + SECTION("Test float datatype allocation with malloc") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_MALLOC_FREE)); + } + + SECTION("Test double datatype allocation with malloc") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_MALLOC_FREE)); + } +} + +/** + * Scenario: This test validates device new, access and delete + * across multiple kernels launched using threads. + */ +TEST_CASE("Unit_deviceAllocation_New_MulKernels_MulThreads") { + // new/delete tests + SECTION("Test char datatype allocation with new") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_NEW_DELETE)); + } + + SECTION("Test short datatype allocation with new") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_NEW_DELETE)); + } + + SECTION("Test int datatype allocation with new") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_NEW_DELETE)); + } + + SECTION("Test float datatype allocation with new") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_NEW_DELETE)); + } + + SECTION("Test double datatype allocation with new") { + REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_NEW_DELETE)); + } +} + +#if HT_AMD +// Scenarios Unit_deviceAllocation_Malloc_SingKernels_MulThreads and +// are failing on NVIDIA platform. +/** + * Scenario: This test validates device allocation malloc, access and free + * in a single kernel launched using threads. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_SingKernels_MulThreads") { + // malloc()/free() tests + std::vector tests; + // Spawn the test threads + for (int idx = 0; idx < num_threads; idx++) { + thread_results[idx] = false; + tests.push_back(std::thread(runTestMemoryAccessInAllThread, + TEST_MALLOC_FREE, idx)); + } + // Wait for all threads to complete + for (std::thread &t : tests) { + t.join(); + } + // Verify All Results + for (int idx = 0; idx < num_threads; idx++) { + REQUIRE(thread_results[idx]); + } +} + +/** + * Scenario: This test validates device new, access and delete + * in a single kernel launched using threads. + */ +TEST_CASE("Unit_deviceAllocation_New_SingKernels_MulThreads") { + // new/delete tests + std::vector tests; + // Spawn the test threads + for (int idx = 0; idx < num_threads; idx++) { + thread_results[idx] = false; + tests.push_back(std::thread(runTestMemoryAccessInAllThread, + TEST_NEW_DELETE, idx)); + } + // Wait for all threads to complete + for (std::thread &t : tests) { + t.join(); + } + // Verify All Results + for (int idx = 0; idx < num_threads; idx++) { + REQUIRE(thread_results[idx]); + } +} +#endif + +/** + * Scenario: This test validates Allocation and Deallocation in multiple + * code object kernels defined in different source files. + */ +TEST_CASE("Unit_deviceAllocation_Malloc_MulCodeObj") { + REQUIRE(true == TestAlloc_Load_MultKernels(TEST_MALLOC_FREE, + INT_MAX)); +} + +/** + * Scenario: This test validates Allocation and Deallocation in multiple + * code object kernels defined in different source files. + */ +TEST_CASE("Unit_deviceAllocation_New_MulCodeObj") { + REQUIRE(true == TestAlloc_Load_MultKernels(TEST_NEW_DELETE, + INT_MAX)); +} diff --git a/tests/catch/unit/deviceLib/kerDevAllocMultCO.cc b/tests/catch/unit/deviceLib/kerDevAllocMultCO.cc new file mode 100644 index 0000000000..f84aabc3c2 --- /dev/null +++ b/tests/catch/unit/deviceLib/kerDevAllocMultCO.cc @@ -0,0 +1,39 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "./defs.h" + +/** + * This kernel allocates memory in thread 0. + */ +extern "C" __global__ void ker_Alloc_MultCodeObj(int **dev_mem, + int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate memory in thread 0 of block 0 + if (0 == myId) { + if (test_type == TEST_MALLOC_FREE) { + *dev_mem = + reinterpret_cast (malloc(blockDim.x*gridDim.x*sizeof(int))); + } else { + *dev_mem = + reinterpret_cast (new int[blockDim.x*gridDim.x]); + } + } +} diff --git a/tests/catch/unit/deviceLib/kerDevAllocSingleKer.cc b/tests/catch/unit/deviceLib/kerDevAllocSingleKer.cc new file mode 100644 index 0000000000..4cda6ce3aa --- /dev/null +++ b/tests/catch/unit/deviceLib/kerDevAllocSingleKer.cc @@ -0,0 +1,57 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "./defs.h" +/** + * This kernel allocates and deallocates memory in every thread. + */ +extern "C" __global__ void ker_TestDynamicAllocInAllThreads_CodeObj( + int *outputBuf, int test_type, int value, + size_t perThreadSize) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Allocate + size_t size = 0; + int* ptr = nullptr; + if (test_type == TEST_MALLOC_FREE) { + size = perThreadSize * sizeof(int); + ptr = reinterpret_cast (malloc(size)); + } else { + size = perThreadSize; + ptr = new int[perThreadSize]; + } + if (ptr == nullptr) { + printf("Device Allocation in thread %d Failed! \n", myId); + return; + } + // Set memory + for (size_t idx = 0; idx < perThreadSize; idx++) { + ptr[idx] = value; + } + // Copy to output buffer + for (size_t idx = 0; idx < perThreadSize; idx++) { + outputBuf[myId*perThreadSize + idx] = ptr[idx]; + } + // Free memory + if (test_type == TEST_MALLOC_FREE) { + free(ptr); + } else { + delete[] ptr; + } +} diff --git a/tests/catch/unit/deviceLib/kerDevFreeMultCO.cc b/tests/catch/unit/deviceLib/kerDevFreeMultCO.cc new file mode 100644 index 0000000000..571ed5d6f6 --- /dev/null +++ b/tests/catch/unit/deviceLib/kerDevFreeMultCO.cc @@ -0,0 +1,47 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "./defs.h" + +/** + * This kernel copies the contents of memory allocated in + * ker_Alloc_MultCodeObj<<<>>> to host and deletes the memory + * from thread 0. + */ +extern "C" __global__ void ker_Free_MultCodeObj(int *outputBuf, + int **dev_mem, int test_type) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Check allocated memory in all threads in block before access + if (*dev_mem == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + + if (0 == myId) { + for (size_t idx = 0; idx < (blockDim.x*gridDim.x); idx++) { + outputBuf[idx] = (*dev_mem)[idx]; + } + if (test_type == TEST_MALLOC_FREE) { + free(*dev_mem); + } else { + delete[] (*dev_mem); + } + } +} diff --git a/tests/catch/unit/deviceLib/kerDevWriteMultCO.cc b/tests/catch/unit/deviceLib/kerDevWriteMultCO.cc new file mode 100644 index 0000000000..c1a1876ce5 --- /dev/null +++ b/tests/catch/unit/deviceLib/kerDevWriteMultCO.cc @@ -0,0 +1,36 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "./defs.h" + +/** + * This kernel writes to memory allocated in ker_Alloc_MultCodeObj<<<>>>. + */ +extern "C" __global__ void ker_Write_MultCodeObj(int **dev_mem, + int value) { + int myId = threadIdx.x + blockDim.x * blockIdx.x; + // Check allocated memory in all threads in block before access + if (*dev_mem == nullptr) { + printf("Device Allocation Failed in thread = %d \n", myId); + return; + } + // Copy to buffer + (*dev_mem)[myId] = value; +} From b7f5db36d20c041a4cebc9849e3e9975e954114a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Mon, 8 Aug 2022 04:59:20 +0100 Subject: [PATCH 4/6] Added testing for hipStreamSynchronize and hipStreamQuery (#2572) --- tests/catch/unit/stream/CMakeLists.txt | 4 + tests/catch/unit/stream/hipStreamQuery.cc | 125 ++++++++++++++ .../catch/unit/stream/hipStreamSynchronize.cc | 156 ++++++++++++++++++ tests/catch/unit/stream/streamCommon.cc | 40 +---- tests/catch/unit/stream/streamCommon.hh | 27 +-- 5 files changed, 290 insertions(+), 62 deletions(-) create mode 100644 tests/catch/unit/stream/hipStreamQuery.cc create mode 100644 tests/catch/unit/stream/hipStreamSynchronize.cc diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 9b829207cf..8518ff07cb 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -14,6 +14,8 @@ set(TEST_SRC hipStreamValue.cc hipStreamWithCUMask.cc hipStreamACb_MultiThread.cc + hipStreamSynchronize.cc + hipStreamQuery.cc hipStreamWaitEvent.cc ) else() @@ -32,6 +34,8 @@ set(TEST_SRC # Fixing would break ABI, to be re-enabled when the fix is made. streamCommon.cc hipStreamValue.cc + hipStreamSynchronize.cc + hipStreamQuery.cc ) # set_source_files_properties(hipStreamAttachMemAsync.cc PROPERTIES COMPILE_FLAGS -std=c++17) diff --git a/tests/catch/unit/stream/hipStreamQuery.cc b/tests/catch/unit/stream/hipStreamQuery.cc new file mode 100644 index 0000000000..ea2e9d0476 --- /dev/null +++ b/tests/catch/unit/stream/hipStreamQuery.cc @@ -0,0 +1,125 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include "streamCommon.hh" + +/** + * @brief Check that querying a stream with no work returns hipSuccess + * + **/ +TEST_CASE("Unit_hipStreamQuery_WithNoWork") { + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamQuery(stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/** + * @brief Check that querying a stream with finished work returns hipSuccess + * + **/ +TEST_CASE("Unit_hipStreamQuery_WithFinishedWork") { + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + + hip::stream::empty_kernel<<>>(); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipStreamQuery(stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +#if !HT_NVIDIA +/** + * @brief Check that submitting work to a destroyed stream sets its status as + * hipErrorContextIsDestroyed + * + * Test removed for Nvidia devices because it returns unexpected error + */ +TEST_CASE("Unit_hipStreamQuery_WithDestroyedStream") { + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorContextIsDestroyed); +} + +/** + * @brief Check that submitting work to an uninitialized stream sets its status as + * hipErrorContextIsDestroyed + * + * Test removed for Nvidia devices because it returns unexpected error + */ +TEST_CASE("Unit_hipStreamQuery_WithUninitializedStream") { + hipStream_t stream{reinterpret_cast(0xFFFF)}; + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorContextIsDestroyed); +} +#endif + +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ + +/** + * @brief Check that submitting work to a stream sets the status of the nullStream to + * hipErrorNotReady + * + */ +TEST_CASE("Unit_hipStreamQuery_SubmitWorkOnStreamAndQueryNullStream") { + { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipStreamQuery(hip::nullStream)); + HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream); + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +/** + * @brief Check that submitting work to the nullStream properly sets its status as + * hipErrorNotReady. + * + */ +TEST_CASE("Unit_hipStreamQuery_NullStreamQuery") { + HIP_CHECK(hipStreamQuery(hip::nullStream)); + HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream); + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + + HIP_CHECK(hipStreamSynchronize(hip::nullStream)); +} + +/** + * @brief Check that querying a stream with pending work returns hipErrorNotReady + * + **/ +TEST_CASE("Unit_hipStreamQuery_WithPendingWork") { + hipStream_t waitingStream{nullptr}; + HIP_CHECK(hipStreamCreate(&waitingStream)); + + HipTest::runKernelForDuration(std::chrono::milliseconds(500), waitingStream); + + HIP_CHECK_ERROR(hipStreamQuery(waitingStream), hipErrorNotReady); + HIP_CHECK(hipStreamSynchronize(waitingStream)); + HIP_CHECK(hipStreamQuery(waitingStream)); + + HIP_CHECK(hipStreamDestroy(waitingStream)); +} +#endif diff --git a/tests/catch/unit/stream/hipStreamSynchronize.cc b/tests/catch/unit/stream/hipStreamSynchronize.cc new file mode 100644 index 0000000000..0555620b1f --- /dev/null +++ b/tests/catch/unit/stream/hipStreamSynchronize.cc @@ -0,0 +1,156 @@ +/* +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 +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include "streamCommon.hh" + +namespace hipStreamSynchronizeTest { + +/** + * @brief Check that hipStreamSynchronize handles empty streams properly. + * + */ +TEST_CASE("Unit_hipStreamSynchronize_EmptyStream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ + +/** + * @brief Check that all work executing in a stream is finished after a call to + * hipStreamSynchronize. + * + */ +TEST_CASE("Unit_hipStreamSynchronize_FinishWork") { + const hipStream_t explicitStream = reinterpret_cast(-1); + hipStream_t stream = GENERATE_COPY(explicitStream, hip::nullStream, hip::streamPerThread); + + const bool isExplicitStream = stream == explicitStream; + if (isExplicitStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + + HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamQuery(stream)); + + if (isExplicitStream) { + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +/** + * @brief Check that synchronizing the nullStream implicitly synchronizes all executing streams. + */ +TEST_CASE("Unit_hipStreamSynchronize_NullStreamSynchronization") { + int totalStreams = 10; + + std::vector streams{}; + + for (int i = 0; i < totalStreams; ++i) { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + streams.push_back(stream); + } + + for (int i = 0; i < totalStreams; ++i) { + HipTest::runKernelForDuration(std::chrono::milliseconds(1000), streams[i]); + } + + for (int i = 0; i < totalStreams; ++i) { + HIP_CHECK_ERROR(hipStreamQuery(streams[i]), hipErrorNotReady); + } + + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + + HIP_CHECK(hipStreamSynchronize(hip::nullStream)); + HIP_CHECK(hipStreamQuery(hip::nullStream)); + + for (int i = 0; i < totalStreams; ++i) { + HIP_CHECK(hipStreamQuery(streams[i])); + } + + for (int i = 0; i < totalStreams; ++i) { + HIP_CHECK(hipStreamDestroy(streams[i])); + } +} + +/** + * @brief Check that synchronizing one stream does implicitly synchronize other streams. + * Check that submiting work to the nullStream does not affect synchronization of other + * streams. Check that querying the nullStream does not affect synchronization of other streams. + */ +TEST_CASE("Unit_hipStreamSynchronize_SynchronizeStreamAndQueryNullStream") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-22"); +#else + + hipStream_t stream1; + hipStream_t stream2; + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + + HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream1); + HipTest::runKernelForDuration(std::chrono::milliseconds(2000), stream2); + + SECTION("Do not use NullStream") {} + SECTION("Submit Kernel to NullStream") { + hip::stream::empty_kernel<<<1, 1, 0, hip::nullStream> > >(); + } + SECTION("Query NullStream") { + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + } + + HIP_CHECK_ERROR(hipStreamQuery(stream1), hipErrorNotReady); + HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady); + + + HIP_CHECK(hipStreamSynchronize(stream1)); + HIP_CHECK(hipStreamQuery(stream1)); + HIP_CHECK_ERROR(hipStreamQuery(stream2), hipErrorNotReady); + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + + HIP_CHECK(hipStreamSynchronize(stream2)); + HIP_CHECK(hipStreamQuery(stream2)); + + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); +#endif +} + +/** + * @brief Check that synchronizing the nullStream also synchronizes the hipStreamPerThread + * special stream. + * + */ +TEST_CASE("Unit_hipStreamSynchronize_NullStreamAndStreamPerThread") { + HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::streamPerThread); + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipErrorNotReady); + HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipErrorNotReady); + HipTest::runKernelForDuration(std::chrono::milliseconds(500), hip::nullStream); + HIP_CHECK(hipStreamSynchronize(hip::nullStream)) + HIP_CHECK_ERROR(hipStreamQuery(hip::streamPerThread), hipSuccess); + HIP_CHECK_ERROR(hipStreamQuery(hip::nullStream), hipSuccess); +} +#endif +} // namespace hipStreamSynchronizeTest \ No newline at end of file diff --git a/tests/catch/unit/stream/streamCommon.cc b/tests/catch/unit/stream/streamCommon.cc index 14ac4000eb..6a2904bbee 100644 --- a/tests/catch/unit/stream/streamCommon.cc +++ b/tests/catch/unit/stream/streamCommon.cc @@ -66,44 +66,8 @@ bool checkStreamFlags_(hipStream_t stream, bool checkFlags = false, unsigned fla inline namespace stream { -__device__ int defaultSemaphore = 0; - -__global__ void signaling_kernel(int* semaphore) { - size_t tid{blockIdx.x * blockDim.x + threadIdx.x}; - if (tid == 0) { - if (semaphore == nullptr) { - atomicAdd(&defaultSemaphore, 1); - } else { - atomicAdd(semaphore, 1); - } - } -} - -__global__ void waiting_kernel(int* semaphore) { - size_t tid{blockIdx.x * blockDim.x + threadIdx.x}; - if (tid == 0) { - if (semaphore == nullptr) { - while (atomicCAS(&defaultSemaphore, 1, 2) == 0) { - } - } else { - while (atomicCAS(semaphore, 1, 2) == 0) { - } - } - } -} - -std::thread startSignalingThread(int* semaphore) { - std::thread signalingThread([semaphore]() { - hipStream_t signalingStream; - HIP_CHECK_THREAD(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking)); - - signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore); - HIP_CHECK_THREAD(hipStreamSynchronize(signalingStream)); - HIP_CHECK_THREAD(hipStreamDestroy(signalingStream)); - }); - - return signalingThread; -} +/* Empty kernel to ensure work finishes on the stream quickly */ +__global__ void empty_kernel() {} bool checkStream(hipStream_t stream) { { // Check default flags diff --git a/tests/catch/unit/stream/streamCommon.hh b/tests/catch/unit/stream/streamCommon.hh index 1d5a1ea958..2017b9298a 100644 --- a/tests/catch/unit/stream/streamCommon.hh +++ b/tests/catch/unit/stream/streamCommon.hh @@ -24,33 +24,12 @@ THE SOFTWARE. namespace hip { inline namespace stream { +/* Empty kernel to ensure work finishes on the stream quickly */ +__global__ void empty_kernel(); + const hipStream_t nullStream = nullptr; const hipStream_t streamPerThread = hipStreamPerThread; -/** - * @brief Kernel that signals a semaphore to change value from 0 to 1. - * - * @param semaphore the semaphore that needs to be signaled. - */ -__global__ void signaling_kernel(int* semaphore = nullptr); - -/** - * @brief Kernel that busy waits until the specified semaphore goes from 0 to 1. - * - * @param semaphore the semaphore to wait for. - */ -__global__ void waiting_kernel(int* semaphore = nullptr); - -/** - * @brief Creates a thread that runs a signaling_kernel on a non-blocking stream. - * hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream. - * You must call HIP_CHECK_THREAD_FINALIZE after joining this thread. - * - * @param semaphore memory location to signal - * @return std::thread thread that has to be joined after the testing is done. - */ -std::thread startSignalingThread(int* semaphore = nullptr); - // Checks stream for valid values of flags and priority bool checkStream(hipStream_t stream); From 485e4677dd1d1cbdb47793f56a3a5a710cbb57df Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 8 Aug 2022 10:37:21 +0530 Subject: [PATCH 5/6] Win: disable compiling unsupported test --- tests/catch/unit/deviceLib/CMakeLists.txt | 22 ++++++++++++++++------ 1 file changed, 16 insertions(+), 6 deletions(-) diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 5805e11b5e..c3f6f43c37 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -13,9 +13,13 @@ set(TEST_SRC syncthreadsand.cc syncthreadscount.cc syncthreadsor.cc - deviceAllocation.cc ) +if(UNIX) + set(TEST_SRC ${TEST_SRC} + deviceAllocation.cc) +endif() + # AMD only tests set(AMD_TEST_SRC unsafeAtomicAdd.cc @@ -44,10 +48,13 @@ set(AMD_ARCH_SPEC_TEST_SRC unsafeAtomicAdd_NonCoherent_withunsafeflag.cc ) -add_custom_target(kerDevAllocMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) -add_custom_target(kerDevWriteMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevWriteMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevWriteMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) -add_custom_target(kerDevFreeMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevFreeMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevFreeMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) -add_custom_target(kerDevAllocSingleKer.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocSingleKer.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocSingleKer.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +if(UNIX) + add_custom_target(kerDevAllocMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevWriteMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevWriteMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevWriteMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevFreeMultCO.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevFreeMultCO.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevFreeMultCO.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(kerDevAllocSingleKer.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/kerDevAllocSingleKer.cc -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/deviceLib/kerDevAllocSingleKer.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) +endif() if(HIP_PLATFORM MATCHES "amd") if (DEFINED OFFLOAD_ARCH_STR) @@ -84,4 +91,7 @@ elseif(HIP_PLATFORM MATCHES "nvidia") TEST_TARGET_NAME build_tests COMPILE_OPTIONS --Wno-deprecated-declarations) endif() -add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code) + +if(UNIX) + add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code) +endif() From 46246c205c0cffb01ce63217fc7d2f047c327f09 Mon Sep 17 00:00:00 2001 From: Finlay Date: Mon, 8 Aug 2022 06:17:42 +0100 Subject: [PATCH 6/6] EXSWCPHIPT-133 - Added support for offsetting to arrays to hipStreamValue tests (#2766) --- tests/catch/unit/stream/CMakeLists.txt | 3 +- .../stream/hipStreamCreateWithPriority.cc | 4 +- tests/catch/unit/stream/hipStreamValue.cc | 617 ++++++++++-------- .../catch/unit/stream/hipStreamWithCUMask.cc | 12 +- 4 files changed, 368 insertions(+), 268 deletions(-) diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 8518ff07cb..7d7c6de46f 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -43,4 +43,5 @@ endif() hip_add_exe_to_target(NAME StreamTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests) + TEST_TARGET_NAME build_tests + COMPILE_OPTIONS -std=c++17) diff --git a/tests/catch/unit/stream/hipStreamCreateWithPriority.cc b/tests/catch/unit/stream/hipStreamCreateWithPriority.cc index 300d78ac3c..9405db8289 100644 --- a/tests/catch/unit/stream/hipStreamCreateWithPriority.cc +++ b/tests/catch/unit/stream/hipStreamCreateWithPriority.cc @@ -509,8 +509,8 @@ bool validateStreamPrioritiesWithEvents() { #define OP(x) \ free(src_h_##x); \ free(dst_h_##x); \ - hipFree(src_d_##x); \ - hipFree(dst_d_##x); + HIP_CHECK(hipFree(src_d_##x)); \ + HIP_CHECK(hipFree(dst_d_##x)); OP(low) OP(normal) OP(high) diff --git a/tests/catch/unit/stream/hipStreamValue.cc b/tests/catch/unit/stream/hipStreamValue.cc index 5eb4321fef..5e70122195 100644 --- a/tests/catch/unit/stream/hipStreamValue.cc +++ b/tests/catch/unit/stream/hipStreamValue.cc @@ -17,276 +17,377 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ #include +#include +#include constexpr unsigned int writeFlag = 0; -#define DEFINE_HIP_STREAM_VALUE(TYPE, BITS, ...) hipStream##TYPE##Value##BITS(__VA_ARGS__) +template auto waitFunc(Args... args) { + if constexpr (std::is_same::value) { + return hipStreamWaitValue32(args...); + } else { + return hipStreamWaitValue64(args...); + } +}; -#define CHECK_HIP_STREAM_VALUE(TYPE, BITS, ...) \ - HIP_CHECK(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__)); +template auto writeFunc(Args... args) { + if constexpr (std::is_same::value) { + return hipStreamWriteValue32(args...); + } else { + return hipStreamWriteValue64(args...); + } +}; -#define NEG_TEST_ERROR_CHECK(TYPE, BITS, errorCode, ...) \ - HIP_CHECK_ERROR(DEFINE_HIP_STREAM_VALUE(TYPE, BITS, __VA_ARGS__), errorCode); +// Random predefined 32 and 64 bit values +using value32_t = std::integral_constant; +using value64_t = std::integral_constant; +template +using testValue = + typename std::conditional::value, value32_t, value64_t>::type; -#if HT_AMD -// Random predefiend 32 and 64 bit values -constexpr uint32_t value32 = 0x70F0F0FF; -constexpr uint64_t value64 = 0x7FFF0000FFFF0000; constexpr uint32_t DATA_INIT = 0x1234; constexpr uint32_t DATA_UPDATE = 0X4321; -template struct TEST_WAIT { - using uintT = typename std::make_unsigned::type; - int compareOp; - uintT mask; - uintT waitValue; - intT signalValueFail; - intT signalValuePass; +template struct TEST_WAIT { + static_assert(std::is_same::value or std::is_same::value, + "only implemented for 32 bit and 64 bit unsigned integers"); + unsigned int compareOp; + UIntT mask = ~static_cast(0); + UIntT waitValue; + UIntT signalValueFail; + UIntT signalValuePass; - TEST_WAIT(int compareOp, uintT waitValue, intT signalValueFail, intT signalValuePass) + TEST_WAIT(unsigned int compareOp, UIntT waitValue, UIntT signalValueFail, UIntT signalValuePass) : compareOp{compareOp}, waitValue{waitValue}, signalValueFail{signalValueFail}, - signalValuePass{signalValuePass} { - mask = static_cast(0xFFFFFFFFFFFFFFFF); - } + signalValuePass{signalValuePass} {} - TEST_WAIT(int compareOp, uintT mask, uintT waitValue, intT signalValueFail, intT signalValuePass) + TEST_WAIT(unsigned int compareOp, UIntT mask, UIntT waitValue, UIntT signalValueFail, + UIntT signalValuePass) : compareOp{compareOp}, mask{mask}, waitValue{waitValue}, signalValueFail{signalValueFail}, signalValuePass{signalValuePass} {} }; -typedef TEST_WAIT TEST_WAIT32; -typedef TEST_WAIT TEST_WAIT64; + +using TEST_WAIT32 = TEST_WAIT; +using TEST_WAIT64 = TEST_WAIT; bool streamWaitValueSupported() { int device_num = 0; HIP_CHECK(hipGetDeviceCount(&device_num)); - int waitValueSupport; for (int device_id = 0; device_id < device_num; ++device_id) { HIP_CHECK(hipSetDevice(device_id)); - waitValueSupport = 0; - HIP_CHECK(hipDeviceGetAttribute(&waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, - device_id)); + int waitValueSupport = 0; + auto getAttributeError = hipDeviceGetAttribute( + &waitValueSupport, hipDeviceAttributeCanUseStreamWaitValue, device_id); + if (getAttributeError != hipSuccess) { + HipTest::HIP_SKIP_TEST("attribute not supported"); + return false; + } if (waitValueSupport == 1) return true; } return false; } -// hipStreamWriteValue Tests -TEST_CASE("Unit_hipStreamValue_Write") { - int64_t* signalPtr; +// The different types of memory that can be used with hipStream[Wait|Write] +enum class PtrType { HostPtr, DevicePtr, DevicePtrToHost, Signal }; - hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); +// Helper class to expose the pointer that is used with hipStream[Write|Wait]Value and also store a +// unique pointer with the deleter to simplify cleanup +// Also includes functions to update and get the value directly +template class TestPtr { + // This stores the memory that must be deleted, as well as the deleter + UniquePtrWithDeleter ptrToDelete; - // Allocate Host Memory - auto hostPtr64 = std::unique_ptr(new uint64_t(1)); - auto hostPtr32 = std::unique_ptr(new uint32_t(1)); + public: + // The pointer that should be used with hipStream[Write|Wait]Value + UIntT* ptr; - // Register Host Memory - HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0)); - HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0)); + TestPtr(UIntT* ptr, UniquePtrWithDeleter ptrToDelete) + : ptrToDelete(std::move(ptrToDelete)), ptr(ptr) {} - // Register Signal Memory - HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory)); - // Initialise Data - *signalPtr = 0x0; - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; - - SECTION("Registered host memory hipStreamWriteValue32") { - INFO("Test writting to registered host pointer using hipStreamWriteValue32"); - HIP_CHECK(hipStreamWriteValue32(stream, hostPtr32.get(), value32, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr32 == value32); + // directly retrieve the value from wherever it was allocated + UIntT getValue(size_t offset = 0) { + if constexpr (type == PtrType::Signal || type == PtrType::HostPtr || + type == PtrType::DevicePtrToHost) { + return ptrToDelete.get()[offset]; + } else { + static_assert(type == PtrType::DevicePtr, "Expected DevicePtr"); + UIntT value; + HIP_CHECK(hipMemcpy(&value, ptr + offset, sizeof(UIntT), hipMemcpyDeviceToHost)); + return value; + } } - SECTION("Registered host memory hipStreamWriteValue64") { - INFO("Test writting to registered host pointer using hipStreamWriteValue32"); - HIP_CHECK(hipStreamWriteValue64(stream, hostPtr64.get(), value64, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr64 == value64); + // directly set the value wherever it was allocated + void setValue(UIntT value, size_t offset = 0) { + if constexpr (type == PtrType::Signal || type == PtrType::DevicePtrToHost || + type == PtrType::HostPtr) { + ptrToDelete.get()[offset] = value; + } else { + // hipMemcpy causes deadlock, so use hipStreamWriteValue + static_assert(type == PtrType::DevicePtr, "Expected DevicePtr"); + hipStream_t stream; + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIP_CHECK(writeFunc(stream, ptr + offset, value, writeFlag)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } } +}; - // Test writting device pointer - void* devicePtr64; - void* devicePtr32; - HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr64, hostPtr64.get(), 0)); - HIP_CHECK(hipHostGetDevicePointer((void**)&devicePtr32, hostPtr32.get(), 0)); - // Reset values - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; +// required for the static assert +template inline constexpr bool AMD_ACTIVE = HT_AMD == 1; - SECTION("Device Memory hipStreamWriteValue32") { - INFO("Test writting to device pointer using hipStreamWriteValue32"); - HIP_CHECK(hipStreamWriteValue32(stream, devicePtr32, value32, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr32 == value32); +template auto allocMem() { + constexpr std::size_t arraySize = 1024; + if constexpr (type == PtrType::Signal) { + static_assert(std::is_same::value, + "signal memory should only be used with 64bit memory"); + + // Allocate Signal Memory + uint64_t* signalPtr{}; + + static_assert(AMD_ACTIVE, + "nvidia backend compiler doesn't like hipExtMallocWithFlags, even in this " + "constexpr branch"); +#if HT_AMD + // 8 is the only acceptable size + HIP_CHECK( + hipExtMallocWithFlags(reinterpret_cast(&signalPtr), 8, hipMallocSignalMemory)); +#endif + + // Init Memory + *signalPtr = 0; + + auto freeStuff = [](uint64_t* sPtr) { HIP_CHECK(hipFree(sPtr)); }; + return TestPtr>{ + signalPtr, std::unique_ptr(signalPtr, freeStuff)}; + } else if constexpr (type == PtrType::DevicePtrToHost) { + auto hostPtr = new UIntT[arraySize]; + + // Register Host Memory + HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0)); + + // Init memory + std::fill(hostPtr, hostPtr + arraySize, 0); + + UIntT* devicePtr; + // Test writing device pointer + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&devicePtr), hostPtr, 0)); + auto freeStuff = [](UIntT* ptr) { + HIP_CHECK(hipHostUnregister(ptr)); + delete[] ptr; + }; + + return TestPtr>{ + devicePtr, std::unique_ptr(hostPtr, freeStuff)}; + } else if constexpr (type == PtrType::HostPtr) { + auto hostPtr = new UIntT[arraySize]; + + // Register Host Memory + HIP_CHECK(hipHostRegister(hostPtr, sizeof(UIntT) * arraySize, 0)); + + // Init memory + std::fill(hostPtr, hostPtr + arraySize, 0); + + auto freeStuff = [](UIntT* ptr) { + HIP_CHECK(hipHostUnregister(ptr)); + delete[] ptr; + }; + + return TestPtr>{ + hostPtr, std::unique_ptr(hostPtr, freeStuff)}; + } else { + static_assert(type == PtrType::DevicePtr, "Expected DevicePtr"); + UIntT* devicePtr; + HIP_CHECK(hipMalloc(&devicePtr, sizeof(UIntT) * arraySize)); + HIP_CHECK(hipMemset(devicePtr, 0, sizeof(UIntT) * arraySize)); + auto freeStuff = [](UIntT* ptr) { HIP_CHECK(hipFree(ptr)); }; + return TestPtr>{ + devicePtr, std::unique_ptr(devicePtr, freeStuff)}; } - - SECTION("Device Memory hipStreamWriteValue64") { - INFO("Test writting to device pointer using hipStreamWriteValue64"); - HIP_CHECK(hipStreamWriteValue64(stream, devicePtr64, value64, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*hostPtr64 == value64); - } - - // Test Writing to Signal Memory - SECTION("Signal Memory hipStreamWriteValue64") { - INFO("Test writting to signal memory using hipStreamWriteValue64"); - HIP_CHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(*signalPtr == value64); - } - - // Cleanup - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipHostUnregister(hostPtr64.get())); - HIP_CHECK(hipHostUnregister(hostPtr32.get())); - HIP_CHECK(hipFree(signalPtr)); } -// hipStreamWaitValue Tests -template -void initData(intT* dataPtr, int64_t* signalPtr, TEST_T tc, std::vector& events) { - // Initialize memory to be waited on - *signalPtr = isBlocking ? tc.signalValueFail : tc.signalValuePass; - - - // Initialize host pointers - dataPtr[0] = DATA_INIT; - dataPtr[1] = DATA_INIT; - - - hipEvent_t firstWriteEvent{nullptr}; - hipEvent_t secondWriteEvent{nullptr}; - HIP_CHECK(hipEventCreate(&firstWriteEvent)); - HIP_CHECK(hipEventCreate(&secondWriteEvent)); - events.push_back(firstWriteEvent); - events.push_back(secondWriteEvent); -} - -template -void syncAndCheckData(hipStream_t stream, intT* dataPtr, int64_t* signalPtr, TEST_T tc, - std::vector& events) { - // Ensure first part of host memory is updated - HIP_CHECK(hipStreamWaitEvent(stream, events[0], 0)); - HIP_ASSERT(dataPtr[0] == DATA_UPDATE); - if (isBlocking) { - // Ensure second part of host memory isn't updated yet - HIP_ASSERT(hipEventQuery(events[1]) == hipErrorNotReady); - HIP_ASSERT(dataPtr[1] == DATA_INIT); - // Update value to release stream - *signalPtr = tc.signalValuePass; +// allows the creation of a list of offsets while avoiding it for signal memory +template constexpr auto get_offsets() { + if constexpr (type == PtrType::Signal) { + return std::array{0}; + } else { + return std::array{0, 1, 2, 3, 31, 1023}; } - - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_ASSERT(hipEventQuery(events[1]) == hipSuccess); - // Finally ensure that second part of host memory is updated - HIP_ASSERT(dataPtr[1] == DATA_UPDATE); } -template void cleanup(hipStream_t& stream, intT* dataPtr, int64_t* signalPtr) { - // Cleanup - HIP_CHECK(hipFree(signalPtr)); - HIP_CHECK(hipHostUnregister(dataPtr)); - HIP_CHECK(hipStreamDestroy(stream)); -} +template struct TestParams { + using UIntType = UIntT; + constexpr static PtrType ptrType = ptrTypeValue; +}; -template void testWait(TEST_T tc) { +#if HT_AMD +TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams)) { +#else +TEMPLATE_TEST_CASE("Unit_hipStreamValue_Write", "", (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams), + (TestParams)) { +#endif +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-126"); + return; +#endif if (!streamWaitValueSupported()) { - UNSCOPED_INFO(" hipStreamWaitValue: not supported on this device , skipping ..."); + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); return; } - // Initialize stream + using UIntT = typename TestType::UIntType; + constexpr auto ptrType = TestType::ptrType; + constexpr auto writeValue = testValue::value; + hipStream_t stream{nullptr}; HIP_CHECK(hipStreamCreate(&stream)); + const auto offsets = get_offsets(); + const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets))); + + CAPTURE(offset); + // Allocate Host Memory - std::unique_ptr dataPtr(new intT(2)); + auto ptr = allocMem(); + UIntT* target = ptr.ptr + offset; + HIP_CHECK(writeFunc(stream, target, writeValue, writeFlag)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(ptr.getValue(offset) == writeValue); - // Register Host Memory - HIP_CHECK(hipHostRegister(&(dataPtr.get()[0]), sizeof(intT), 0)); - HIP_CHECK(hipHostRegister(&(dataPtr.get()[1]), sizeof(intT), 0)); + // Cleanup + HIP_CHECK(hipStreamDestroy(stream)); +} - // Allocate Signal Memory - int64_t* signalPtr; - HIP_CHECK(hipExtMallocWithFlags((void**)&signalPtr, 8, hipMallocSignalMemory)); +template +void syncAndCheckData(hipStream_t stream, UIntT* dataPtr, TestPtr signalPtr, size_t offset, + TEST_WAIT tc, std::array& events) { + // Ensure first part of host memory is updated + HIP_CHECK(hipEventSynchronize(events[0])); + REQUIRE(dataPtr[0] == DATA_UPDATE); - std::vector events; - initData(dataPtr.get(), signalPtr, tc, events); - - if (std::is_same::value) { - CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag) - HIP_CHECK(hipEventRecord(events[0], stream)); - - if (static_cast(tc.mask) != 0xFFFFFFFF) { - CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, static_cast(tc.waitValue), - tc.compareOp, static_cast(tc.mask)); - } else { - CHECK_HIP_STREAM_VALUE(Wait, 32, stream, signalPtr, tc.waitValue, tc.compareOp); - } - - CHECK_HIP_STREAM_VALUE(Write, 32, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag) - } else { - CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag) - HIP_CHECK(hipEventRecord(events[0], stream)); - - if (tc.mask != 0xFFFFFFFFFFFFFFFF) { - CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp, tc.mask); - } else { - CHECK_HIP_STREAM_VALUE(Wait, 64, stream, signalPtr, tc.waitValue, tc.compareOp); - } - - CHECK_HIP_STREAM_VALUE(Write, 64, stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag) + if constexpr (isBlocking) { + // Ensure second part of host memory isn't updated yet + HIP_CHECK_ERROR(hipEventQuery(events[1]), hipErrorNotReady); + REQUIRE(dataPtr[1] == DATA_INIT); + // Update value to release stream + signalPtr.setValue(tc.signalValuePass, offset); } + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipEventQuery(events[1])); + // Finally ensure that second part of host memory is updated + REQUIRE(dataPtr[1] == DATA_UPDATE); +} + + +template +void testWait(TEST_WAIT tc) { + if (!streamWaitValueSupported()) { + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); + return; + } +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-128"); + return; +#endif + + using UIntT = typename TestType::UIntType; + constexpr auto ptrType = TestType::ptrType; + constexpr UIntT defaultMask = ~static_cast(0); + + // Initialize stream + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + + // Allocate Host Memory + auto dataPtr = std::make_unique(2); + // Register Host Memory + HIP_CHECK(hipHostRegister(dataPtr.get(), sizeof(UIntT), 0)); + HIP_CHECK(hipHostRegister(dataPtr.get() + 1, sizeof(UIntT), 0)); + std::fill(dataPtr.get(), dataPtr.get() + 2, DATA_INIT); + + std::array events; + HIP_CHECK(hipEventCreate(&events[0])); + HIP_CHECK(hipEventCreate(&events[1])); + + + const auto offsets = get_offsets(); + const auto offset = GENERATE_COPY(from_range(std::begin(offsets), std::end(offsets))); + + auto waitPtr = allocMem(); + UIntT* const target = waitPtr.ptr + offset; + waitPtr.setValue(isBlocking ? tc.signalValueFail : tc.signalValuePass, offset); + + HIP_CHECK(writeFunc(stream, &(dataPtr.get()[0]), DATA_UPDATE, writeFlag)); + HIP_CHECK(hipEventRecord(events[0], stream)); + + if (tc.mask != defaultMask) { + HIP_CHECK(waitFunc(stream, target, tc.waitValue, tc.compareOp, tc.mask)); + } else { + HIP_CHECK(waitFunc(stream, target, tc.waitValue, tc.compareOp)); + } + + HIP_CHECK(writeFunc(stream, &(dataPtr.get()[1]), DATA_UPDATE, writeFlag)); + HIP_CHECK(hipEventRecord(events[1], stream)); - syncAndCheckData(stream, dataPtr.get(), signalPtr, tc, events); - cleanup(stream, dataPtr.get(), signalPtr); -} -#undef CHECK_HIP_STREAM_VALUE + syncAndCheckData(stream, dataPtr.get(), std::move(waitPtr), offset, tc, events); + // Cleanup + HIP_CHECK(hipEventDestroy(events[0])); + HIP_CHECK(hipEventDestroy(events[1])); + HIP_CHECK(hipHostUnregister(dataPtr.get())); + HIP_CHECK(hipHostUnregister(dataPtr.get() + 1)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead #define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32(suffix, test_t) \ TEST_CASE("Unit_hipStreamValue_Wait32_Blocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, true>(test_t); } \ + SECTION("DevicePtr") { testWait, true>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, true>(test_t); \ + } \ } \ TEST_CASE("Unit_hipStreamValue_Wait32_NonBlocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, false>(test_t); } \ + SECTION("DevicePtr") { testWait, false>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, false>(test_t); \ + } \ } + // Using Mask -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_1", - TEST_WAIT64( // mask will ignore few MSB bits - hipStreamWaitValueGte, 0x0000FFFFFFFFFFFF, - 0x000000007FFF0001, 0x7FFF00007FFF0000, - 0x000000007FFF0001)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte_2", - TEST_WAIT64(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6)) +DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Gte", + TEST_WAIT32(hipStreamWaitValueGte, 0xF, 0x4, 0x3, 0x6)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_1", - TEST_WAIT64( // mask will ignore few MSB bits - hipStreamWaitValueEq, 0x0000FFFFFFFFFFFF, - 0x000000000FFF0001, 0x7FFF00000FFF0000, - 0x7F0000000FFF0001)) + TEST_WAIT32( // mask will ignore few MSB bits + hipStreamWaitValueEq, 0x0000FFFF, 0x00000001, + 0x0FFF0000, 0x0FFF0001)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Eq_2", - TEST_WAIT64(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11)) + TEST_WAIT32(hipStreamWaitValueEq, 0xFF, 0x11, 0x25, 0x11)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_And", - TEST_WAIT64( // mask will discard bits 8 to 11 + TEST_WAIT32( // mask will discard bits 8 to 11 hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_1", - TEST_WAIT64( // mask is set to ignore the sign bit. - hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF, - 0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd, - 0x7FFFFFFFFFFFFdb5)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("Mask_Nor_2", - TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3. - hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76)) // Not Using Mask DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Eq", @@ -299,19 +400,47 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_And", TEST_WAIT32(hipStreamWaitValueAnd, 0x70F0F0F0, 0x0F0F0F0F, 0X1F0F0F0F)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32("NoMask_Nor", - TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA, - static_cast(0x85555555), - static_cast(0x9AAAAAAA))) + TEST_WAIT32(hipStreamWaitValueNor, 0x7AAAAAAA, 0x85555555, + 0x9AAAAAAA)) #undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT32 +#if HT_AMD +// TEMPLATE_TEST_CASE wasn't working within a macro, so sections were used instead #define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \ TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, true>(test_t); } \ + SECTION("DevicePtr") { testWait, true>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, true>(test_t); \ + } \ + SECTION("Signal") { testWait, true>(test_t); } \ } \ TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \ - testWait(test_t); \ + SECTION("HostPtr") { testWait, false>(test_t); } \ + SECTION("DevicePtr") { testWait, false>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, false>(test_t); \ + } \ + SECTION("Signal") { testWait, false>(test_t); } \ } +#else +#define DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64(suffix, test_t) \ + TEST_CASE("Unit_hipStreamValue_Wait64_Blocking_" + std::string(suffix)) { \ + SECTION("HostPtr") { testWait, true>(test_t); } \ + SECTION("DevicePtr") { testWait, true>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, true>(test_t); \ + } \ + } \ + TEST_CASE("Unit_hipStreamValue_Wait64_NonBlocking_" + std::string(suffix)) { \ + SECTION("HostPtr") { testWait, false>(test_t); } \ + SECTION("DevicePtr") { testWait, false>(test_t); } \ + SECTION("DevicePtrToHost") { \ + testWait, false>(test_t); \ + } \ + } +#endif // Using Mask @@ -332,14 +461,6 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Eq_2", DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_And", TEST_WAIT64( // mask will discard bits 8 to 11 hipStreamWaitValueAnd, 0xFF, 0xF4A, 0xF35, 0X02)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_1", - TEST_WAIT64( // mask is set to ignore the sign bit. - hipStreamWaitValueNor, 0x7FFFFFFFFFFFFFFF, - 0x7FFFFFFFFFFFF247, 0x7FFFFFFFFFFFFdbd, - 0x7FFFFFFFFFFFFdb5)) -DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("Mask_Nor_2", - TEST_WAIT64( // mask is set to apply NOR for bits 0 to 3. - hipStreamWaitValueNor, 0xF, 0x7E, 0x7D, 0x76)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Gte", TEST_WAIT64(hipStreamWaitValueGte, 0x7FFFFFFFFFFF0001, @@ -352,94 +473,72 @@ DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_And", 0x0F0F0F0F0F0F0F0F, 0X1F0F0F0F0F0F0F0F)) DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64("NoMask_Nor", TEST_WAIT64(hipStreamWaitValueNor, 0x4724724747247247, - static_cast(0xbddbddbdbddbddbd), - static_cast(0xbddbddbdbddbddb3))) + 0xbddbddbdbddbddbd, 0xbddbddbdbddbddb3)) #undef DEFINE_STREAM_WAIT_VAL_TEST_CASES_INT64 -#endif - // Negative Tests TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") { - #if HT_AMD HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96"); return; #endif + if (!streamWaitValueSupported()) { + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); + return; + } hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); - REQUIRE(stream != nullptr); - // Allocate Host Memory - auto hostPtr32 = std::unique_ptr(new uint32_t(1)); - auto hostPtr64 = std::unique_ptr(new uint64_t(1)); - - // Register Host Memory - HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0)); - HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0)); - - // Set dummy data - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; - - auto compareOp = hipStreamWaitValueGte; + const auto compareOp = hipStreamWaitValueGte; + const auto expectedError = hipErrorInvalidValue; // Memory pointer negative tests - - INFO("Testing Invalid Memory Pointer for hipStreamWriteValue32"); - NEG_TEST_ERROR_CHECK(Write, 32, hipErrorNotSupported, stream, nullptr, 0, writeFlag) - - INFO("Testing Invalid Memory Pointer for hipStreamWriteValue64"); - NEG_TEST_ERROR_CHECK(Write, 64, hipErrorNotSupported, stream, nullptr, 0, writeFlag) - - INFO("Testing Invalid Memory Pointer for hipStreamWaitValue32"); - NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, nullptr, 0, compareOp) - - INFO("Testing Invalid Memory Pointer for hipStreamWaitValue64"); - NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, nullptr, 0, compareOp) + SECTION("Invalid Memory Pointer for hipStreamWriteValue32") { + HIP_CHECK_ERROR(hipStreamWriteValue32(stream, nullptr, 0, writeFlag), expectedError); + } + SECTION("Invalid Memory Pointer for hipStreamWriteValue64") { + HIP_CHECK_ERROR(hipStreamWriteValue64(stream, nullptr, 0, writeFlag), expectedError); + } + SECTION("Invalid Memory Pointer for hipStreamWaitValue32") { + HIP_CHECK_ERROR(hipStreamWaitValue32(stream, nullptr, 0, compareOp), expectedError); + } + SECTION("Invalid Memory Pointer for hipStreamWaitValue32") { + HIP_CHECK_ERROR(hipStreamWaitValue64(stream, nullptr, 0, compareOp), expectedError); + } // Cleanup - HIP_CHECK(hipHostUnregister(hostPtr32.get())); - HIP_CHECK(hipHostUnregister(hostPtr64.get())); HIP_CHECK(hipStreamDestroy(stream)); } -TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag") { +TEMPLATE_TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag", "", uint32_t, uint64_t) { #if HT_AMD HipTest::HIP_SKIP_TEST("EXSWCPHIPT-96"); return; #endif + if (!streamWaitValueSupported()) { + HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device."); + return; + } hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); - REQUIRE(stream != nullptr); // Allocate Host Memory - auto hostPtr32 = std::unique_ptr(new uint32_t(1)); - auto hostPtr64 = std::unique_ptr(new uint64_t(1)); + auto hostPtr = std::make_unique(); // Register Host Memory - HIP_CHECK(hipHostRegister(hostPtr32.get(), sizeof(int32_t), 0)); - HIP_CHECK(hipHostRegister(hostPtr64.get(), sizeof(int64_t), 0)); + HIP_CHECK(hipHostRegister(hostPtr.get(), sizeof(TestType), 0)); // Set dummy data - *hostPtr64 = 0x0; - *hostPtr32 = 0x0; + *hostPtr = 0x0; /* EXSWCPHIPT-96 */ - INFO("Testing Invalid flag for hipStreamWaitValue32"); - NEG_TEST_ERROR_CHECK(Wait, 32, hipErrorNotSupported, stream, hostPtr32.get(), 0, -1) - INFO("Testing Invalid flag for hipStreamWaitValue64"); - NEG_TEST_ERROR_CHECK(Wait, 64, hipErrorNotSupported, stream, hostPtr64.get(), 0, -1) + HIP_CHECK_ERROR(waitFunc(stream, hostPtr.get(), 0, -1), hipErrorInvalidValue); // Cleanup - HIP_CHECK(hipHostUnregister(hostPtr32.get())); - HIP_CHECK(hipHostUnregister(hostPtr64.get())); + HIP_CHECK(hipHostUnregister(hostPtr.get())); HIP_CHECK(hipStreamDestroy(stream)); } - -#undef NEG_TEST_ERROR_CHECK diff --git a/tests/catch/unit/stream/hipStreamWithCUMask.cc b/tests/catch/unit/stream/hipStreamWithCUMask.cc index b65cdc4ca1..ea6b033d9c 100644 --- a/tests/catch/unit/stream/hipStreamWithCUMask.cc +++ b/tests/catch/unit/stream/hipStreamWithCUMask.cc @@ -149,8 +149,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") { HIP_CHECK(hipGetDeviceProperties(&props, 0)); createDefaultCUMask(&defaultCUMask, props.multiProcessorCount); - hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(), - defaultCUMask.data()); + HIP_CHECK(hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(), + defaultCUMask.data())); HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); const unsigned blocks = GRIDSIZE; @@ -244,7 +244,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, streams[0], dA[0], dC[0], N); - hipDeviceSynchronize(); + HIP_CHECK(hipDeviceSynchronize()); auto single_end = std::chrono::steady_clock::now(); std::chrono::duration single_kernel_time = single_end - single_start; @@ -264,7 +264,7 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, streams[np], dA[np], dC[np], N); } - hipDeviceSynchronize(); + HIP_CHECK(hipDeviceSynchronize()); auto all_end = std::chrono::steady_clock::now(); std::chrono::duration all_kernel_time = all_end - all_start; @@ -288,8 +288,8 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { delete [] hA; delete [] hC; for (int np = 0; np < KNumPartition; np++) { - hipFree(dC[np]); - hipFree(dA[np]); + HIP_CHECK(hipFree(dC[np])); + HIP_CHECK(hipFree(dA[np])); HIP_CHECK(hipStreamDestroy(streams[np])); } }