From 7f520ae104d1803490bbec8b2cf38c7e53ef5935 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 23 Feb 2022 03:49:05 +0530 Subject: [PATCH 01/30] SWDEV-301947 - Add new enum for boost CU count (#2490) Change-Id: I39b96218edbdf8a77473be8fa524999bb7df036a Co-authored-by: Saleel Kudchadker --- include/hip/hip_runtime_api.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 7d4d58ec24..eca5b9ff63 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -445,6 +445,8 @@ typedef enum hipDeviceAttribute_t { ///< hipStreamWaitValue64(), '0' otherwise. hipDeviceAttributeImageSupport, ///< '1' if Device supports image, '0' otherwise. + hipDeviceAttributeMultiprocessorBoostCount, ///< All available boost compute units for the device + hipDeviceAttributeAmdSpecificEnd = 19999, hipDeviceAttributeVendorSpecificBegin = 20000, // Extended attributes for vendors From 990f546b387b19d5514ce50d15d04edfef8dd5cf Mon Sep 17 00:00:00 2001 From: lthakur Date: Wed, 23 Feb 2022 09:00:16 +0530 Subject: [PATCH 02/30] SWDEV-298757 - hipMemRangeGetAttribute.cc Adding test cases for hipMemRangeGetAttribute() api (#2459) --- tests/catch/unit/memory/CMakeLists.txt | 4 + .../unit/memory/hipMemRangeGetAttribute.cc | 408 ++++++++++++++++++ 2 files changed, 412 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemRangeGetAttribute.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index b450921321..95cf97d736 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -35,6 +35,8 @@ set(TEST_SRC hipMemPrefetchAsyncExtTsts.cc hipMemAdviseMmap.cc hipMemCoherencyTst.cc + hipMallocManaged.cc + hipMemRangeGetAttribute.cc ) else() set(TEST_SRC @@ -69,6 +71,8 @@ set(TEST_SRC hipMallocManagedFlagsTst.cc hipMemPrefetchAsyncExtTsts.cc hipMemAdviseMmap.cc + hipMallocManaged.cc + hipMemRangeGetAttribute.cc ) endif() diff --git a/tests/catch/unit/memory/hipMemRangeGetAttribute.cc b/tests/catch/unit/memory/hipMemRangeGetAttribute.cc new file mode 100644 index 0000000000..64fc8e6be4 --- /dev/null +++ b/tests/catch/unit/memory/hipMemRangeGetAttribute.cc @@ -0,0 +1,408 @@ +/* +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, 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. +*/ + +/* Test Case Description: + Scenario-1: The following function tests the count parameter(last param) to + hipMemRangeGetAttribute api by passing possible extreme values. + Curently the only way to test if count param working properly is to verify + the first parameter of hipMemRangeGetAttribute() api has value 1 stored + + Scenario-2: This test case checks the behavior of hipMemRangeGetAttribute() with + AccessedBy flag is consistent with cuda's counter part + + Scenario-3: Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise + AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the + memory and probe the for the flags which are set earlier using + hipMemRangeGetAttribute() api for the full size(4*PageSz). + + + Scenario-4: The following scenarios tests that probing the attributes which are not set + by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should + not result in a crash + + Scenario-5: The following scenario is a simple test which does the following: + Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then + probe LastPrefetchLocation attribute using hipMemRangeGetAttribute + + Scenario-6: The following Test Case does negative tests on hipMemRangeGetAttribute()*/ + +#include +#include +#ifdef __linux__ + #include + #include +#endif + +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 count parameter(last param) to +// hipMemRangeGetAttribute api by passing possible extreme values. +// Curently the only way to test if count param working properly is to verify +// the first parameter of hipMemRangeGetAttribute() api has value 1 stored +TEST_CASE("Unit_hipMemRangeGetAttribute_TstCountParam") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + int MEM_SIZE = 4096, RND_NUM = 9999, FLG_READMOSTLY_ENBLD = 1; + bool IfTestPassed = true; + int data = RND_NUM, *devPtr = nullptr; + size_t TotGpuMem, TotGpuFreeMem; + HIP_CHECK(hipMemGetInfo(&TotGpuFreeMem, &TotGpuMem)); + + HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(reinterpret_cast(&data), + sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE)); + if (data != FLG_READMOSTLY_ENBLD) { + WARN("hipMemRangeGetAttribute() api didnt return expected value!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipFree(devPtr)); + HIP_CHECK(hipMallocManaged(&devPtr, TotGpuFreeMem, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, TotGpuFreeMem, hipMemAdviseSetReadMostly, + 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, TotGpuFreeMem)); + + if (data != FLG_READMOSTLY_ENBLD) { + WARN("hipMemRangeGetAttribute() api didnt return expected value!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipFree(devPtr)); + HIP_CHECK(hipMallocManaged(&devPtr, (TotGpuFreeMem - 1), + hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, (TotGpuFreeMem - 1), + hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, (TotGpuFreeMem - 1))); + + if (data != FLG_READMOSTLY_ENBLD) { + WARN("hipMemRangeGetAttribute() api didnt return expected value!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipFree(devPtr)); + + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* The following Test Case does negative tests on hipMemRangeGetAttribute()*/ + +TEST_CASE("Unit_hipMemRangeGetAttribute_NegativeTests") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + int MEM_SIZE = 4096, RND_NUM = 9999; + float *devPtr = nullptr; + int NumDevs; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + int data = RND_NUM; + int *OutData = new int[NumDevs]; + for (int m = 0; m < NumDevs; ++m) { + OutData[m] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal)); + HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + + // checking the behavior with dataSize 0 + SECTION("checking the behavior with dataSize 0") { + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, 0, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behavior with dataSize > 4 and even + SECTION("checking the behavior with dataSize > 4 and even") { + REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 6, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behavior with dataSize > 4 and odd + SECTION("checking the behavior with dataSize > 4 and odd") { + REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 7, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behavior with dataSize which is not multiple of 4 + SECTION("checking the behavior with dataSize which is not multiple of 4") { + REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 27, + hipMemRangeAttributeReadMostly, + devPtr, MEM_SIZE), __LINE__)); + } + // checking the behaviour with devPtr(4th param) as NULL + SECTION("checking the behaviour with devPtr(4th param) as NULL") { + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + NULL, MEM_SIZE), __LINE__)); + } + // checking the behaviour with count(5th param) as 0 + SECTION("checking the behaviour with count(5th param) as 0") { + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + devPtr, 0), __LINE__)); + } + // checking the behavior with invalid attribute (3rd param) as 0 + // as it is attribute hence avoiding the negative tests with 3rd param + + // checking the behaviour of the api with ptr allocated using + // hipHostMalloc + void *ptr = nullptr; + SECTION("Checking behavior with hipHostMalloc ptr") { + HIP_CHECK(hipHostMalloc(&ptr, MEM_SIZE, 0)); + REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, + ptr, MEM_SIZE), __LINE__)); + HIP_CHECK(hipHostFree(ptr)); + } + HIP_CHECK(hipFree(devPtr)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* This test case checks the behavior of hipMemRangeGetAttribute() with + AccessedBy flag is consistent with cuda's counter part*/ +TEST_CASE("Unit_hipMemRangeGetAttribute_AccessedBy1") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096, RND_NUM = 999; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[Ngpus] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + Hmm, MEM_SZ)); + if (OutData[0] != 0) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + for (int i = 1; i < Ngpus; ++i) { + if (OutData[i] != -2) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + } + if (Ngpus >= 2) { + for (int i = 0; i < Ngpus; ++i) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, i)); + } + // checking the behavior with dataSize less than the number of gpus + // This should not result in segfault. + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*(Ngpus-1), + hipMemRangeAttributeAccessedBy, + Hmm, MEM_SZ)); + // OutData should have stored the gpu ordinals for which AccessedBy is + // assigned except for the last element which should have -2 stored + // so as to be consistent with cuda's behavior + for (int i = 0; i < (Ngpus - 1); ++i) { + if (OutData[i] != i) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + } + if (OutData[Ngpus - 1] != -2) { + WARN("Didn't receive expected value at line: " << __LINE__); + REQUIRE(false); + } + } + HIP_CHECK(hipFree(Hmm)); + delete[] OutData; + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + + +/* Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise + AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the + memory and probe the for the flags which are set earlier using + hipMemRangeGetAttribute() api for the full size(4*PageSz).*/ +/* Need to discuss the difference in behavior w.r.t cuda*/ + +TEST_CASE("Unit_hipMemRangeGetAttribte_3") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096*4, RND_NUM = 999; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[Ngpus] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + (Hmm), MEM_SZ)); + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetReadMostly, 0)); + // The Api called below should not fail + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeReadMostly, + (Hmm), MEM_SZ)); + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetPreferredLocation, 0)); + // The api called below should not fail + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributePreferredLocation, + (Hmm), MEM_SZ)); + HIP_CHECK(hipFree(Hmm)); + delete[] OutData; + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* The following scenarios tests that probing the attributes which are not set + by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should + not result in a crash*/ + +TEST_CASE("Unit_hipMemRangeGetAttribute_4") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, PageSz = 4096, Ngpus, RND_NUM = 999; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[i] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, 4*PageSz)); + SECTION("Set ReadMostly & probe other flags") { + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributePreferredLocation, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetReadMostly, 0)); + } + SECTION("Set AccessedBy & probe other flags") { + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeReadMostly, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributePreferredLocation, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetAccessedBy, 0)); + } + SECTION("Set AccessedBy & probe other flags") { + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetPreferredLocation, + 0)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeReadMostly, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus, + hipMemRangeAttributeAccessedBy, + Hmm, 4*PageSz)); + HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetPreferredLocation, + 0)); + } + HIP_CHECK(hipFree(Hmm)); + delete[] OutData; + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* The following scenario is a simple test which does the following: + Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then + probe LastPrefetchLocation attribute using hipMemRangeGetAttribute*/ + +TEST_CASE("Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0, *Hmm = NULL, RND_NUM = 999; + size_t PageSz = 4096; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + + int *OutData = new int[Ngpus]; + for (int i = 0; i < Ngpus; ++i) { + OutData[Ngpus] = RND_NUM; + } + HIP_CHECK(hipMallocManaged(&Hmm, PageSz*4)); + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, PageSz*4, 0, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(OutData, 4, + hipMemRangeAttributeLastPrefetchLocation, + Hmm, PageSz*4)); + HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipFree(Hmm)); + if (OutData[0] != 0) { + WARN("Didnt receive expected value at line: " << __LINE__); + delete[] OutData; + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + From 6e696a028bcb11cbc41668711704e335dc893dae Mon Sep 17 00:00:00 2001 From: lthakur Date: Thu, 24 Feb 2022 01:28:08 +0530 Subject: [PATCH 03/30] SWDEV-298757 - Moved long running tests into stress category. (#2461) --- .../stress/memory/hipMallocManagedStress.cc | 227 +++++++++++++++++- tests/catch/unit/memory/hipMallocManaged.cc | 177 +------------- .../unit/memory/hipMallocManagedFlagsTst.cc | 76 ------ 3 files changed, 231 insertions(+), 249 deletions(-) diff --git a/tests/catch/stress/memory/hipMallocManagedStress.cc b/tests/catch/stress/memory/hipMallocManagedStress.cc index 0e7d80edb6..7e992cd152 100644 --- a/tests/catch/stress/memory/hipMallocManagedStress.cc +++ b/tests/catch/stress/memory/hipMallocManagedStress.cc @@ -1,5 +1,5 @@ /* - Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -19,6 +19,15 @@ // The following test case allocation, host access, device access of HMM // memory from size 1 to 10KB +/* Test Case Description: + 1) Testing allocation, host access, device access of HMM + memory from size 1 to 10KB + 2) The following test case tests the behavior of kernel with a HMM memory + and hipMalloc memory + 3) The following test case tests when the same Hmm memory is used for + launching multiple different kernels will results in any issue + 4) Testing the allocation of/scenarios around max possible memory + */ #include #include @@ -35,6 +44,86 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr, } static bool IfTestPassed = true; +// Kernel functions +__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + for (size_t i = index; i < n; i++) { + Hmm[i] = Dptr[i] + 10; + } +} + +__global__ void KernelMulAdd_MngdMem(int *Hmm, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; + for (size_t i = index; i < n; i += stride) { + Hmm[i] = Hmm[i] * 2 + 10; + } +} + +__global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { + size_t index = blockIdx.x * blockDim.x + threadIdx.x; + size_t stride = blockDim.x * gridDim.x; + for (size_t i = index; i < n; i += stride) { + Hmm[i] = Dptr[i] * 10; + } +} +static bool IfTestPassed = true; + +static void LaunchKrnl4(size_t NumElms, int InitVal) { + int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms))); + HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms))); + int *Hstptr = reinterpret_cast(new int[NumElms]); + for (size_t i = 0; i < NumElms; ++i) { + Hstptr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)), + hipMemcpyHostToDevice)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + for (size_t i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n"); + REQUIRE(false); + } + DataMismatch = 0; + KernelMul_MngdMem<<>>(Hmm, Dptr, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the result + for (size_t i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal * 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); + REQUIRE(false); + } + DataMismatch = 0; + KernelMulAdd_MngdMem<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the result + + for (size_t i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal * 10 * 2 + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); + REQUIRE(false); + } + delete[] Hstptr; +} + static int HmmAttrPrint() { int managed = 0; INFO("The following are the attribute values related to HMM for" @@ -104,3 +193,139 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") { } } +// The following test case tests the behavior of kernel with a HMM memory and +// hipMalloc memory + +TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") { + IfTestPassed = true; + int *Hmm = NULL, *Dptr = NULL, InitVal = 123; + size_t NumElms = (1024 * 1024); + int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0; + int managed = HmmAttrPrint(); + if (managed == 1) { + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms)); + HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms)); + for (size_t i = 0; i < NumElms; ++i) { + Hmm[i] = 0; + Hptr[i] = InitVal; + } + HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms, + hipMemcpyHostToDevice)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // Verifying the results + for (size_t k = 0; k < NumElms; ++k) { + if (Hmm[k] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch observed!\n"); + IfTestPassed = false; + } + + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipFree(Dptr)); + delete[] Hptr; + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +// The following test case tests when the same Hmm memory is used for +// launching multiple different kernels will results in any issue +TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") { + int managed = HmmAttrPrint(); + if (managed) { + int InitVal = 123, NumElms = (1024 * 1024); + LaunchKrnl4(NumElms, InitVal); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// Testing the allocation of/scenarios around max possible memory +TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") { + int managed = HmmAttrPrint(); + if (managed == 1) { + bool IfTestPassed = true; + hipError_t err; + void *Hmm = NULL; + size_t totalDevMem = 0, freeDevMem = 0; + int NumDevs = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + // Testing allocation of extreme and unusual mem values + for (int i = 0; i < NumDevs; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem)); + err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating memory on GPU: " << i); + WARN(" size 1 with"); + WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max free memory on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max (free - 1) memory on "); + WARN("GPU: " << i); + WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, 1, hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating memory size 1 on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max free memory on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost); + if (hipSuccess == err) { + HIP_CHECK(hipFree(Hmm)); + } else { + WARN("Observed error while allocating max (freeDevMem - 1) memory" + " on GPU: " << i); + WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); + WARN("Error: " << hipGetErrorString(err)); + IfTestPassed = false; + } + } + REQUIRE(IfTestPassed); + } else { + SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); + } +} diff --git a/tests/catch/unit/memory/hipMallocManaged.cc b/tests/catch/unit/memory/hipMallocManaged.cc index e8bc69997c..d40aa924a1 100644 --- a/tests/catch/unit/memory/hipMallocManaged.cc +++ b/tests/catch/unit/memory/hipMallocManaged.cc @@ -17,12 +17,11 @@ THE SOFTWARE. */ -/* - List of Test cases: - 1) Unit_hipMallocManaged_Basic - 2) Unit_hipMallocManaged_MultiSize - 3) Unit_hipMallocManaged_MultiKrnlHmmAccess - 4) Unit_hipMallocManaged_KrnlWth2MemTypes +/* Test Case Description: + 1) This testcase verifies the hipMallocManaged basic scenario - supported on + all devices + 2) This testcase verifies the hipMallocManaged basic scenario - supported + only on HMM enabled devices */ #include @@ -32,12 +31,6 @@ // Kernel functions -__global__ void KrnlWth2MemTypes(int *Hmm, int *Dptr, size_t n) { - size_t index = blockIdx.x * blockDim.x + threadIdx.x; - for (size_t i = index; i < n; i++) { - Hmm[i] = Dptr[i] + 10; - } -} __global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { size_t index = blockIdx.x * blockDim.x + threadIdx.x; @@ -64,9 +57,6 @@ __global__ void KrnlWth2MemTypesC(unsigned char *Hmm, unsigned char *Dptr, } } -// The following variable will be used to get the result of computation -// from multiple threads -static bool IfTestPassed = true; static int HmmAttrPrint() { int managed = 0; @@ -93,62 +83,6 @@ static int HmmAttrPrint() { } -static void LaunchKrnl4(size_t NumElms, int InitVal) { - int *Hmm = NULL, *Dptr = NULL, blockSize = 64, DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, (sizeof(int) * NumElms))); - HIP_CHECK(hipMalloc(&Dptr, (sizeof(int) * NumElms))); - int *Hstptr = reinterpret_cast(new int[NumElms]); - for (size_t i = 0; i < NumElms; ++i) { - Hstptr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Dptr, Hstptr, (NumElms * sizeof(int)), - hipMemcpyHostToDevice)); - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); - KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - for (size_t i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - INFO("Data Mismatch observed after the Kernel: KrnlWth2MemTypes!!\n"); - REQUIRE(false); - } - DataMismatch = 0; - KernelMul_MngdMem<<>>(Hmm, Dptr, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the result - for (size_t i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal * 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); - REQUIRE(false); - } - DataMismatch = 0; - KernelMulAdd_MngdMem<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the result - - for (size_t i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal * 10 * 2 + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - INFO("Data Mismatch observedafter the Kernel: KernelMul_MngdMem!!\n"); - REQUIRE(false); - } - delete[] Hstptr; -} - - static size_t N{4 * 1024 * 1024}; static unsigned blocksPerCU{6}; @@ -241,104 +175,3 @@ TEST_CASE("Unit_hipMallocManaged_Advanced") { } } - -// The following test case tests the behavior of kernel with a HMM memory and -// hipMalloc memory - -TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") { - IfTestPassed = true; - int *Hmm = NULL, *Dptr = NULL, InitVal = 123; - size_t NumElms = (1024 * 1024); - int *Hptr = new int[NumElms], blockSize = 64, DataMismatch = 0; - int managed = HmmAttrPrint(); - if (managed == 1) { - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms)); - HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms)); - for (size_t i = 0; i < NumElms; ++i) { - Hmm[i] = 0; - Hptr[i] = InitVal; - } - HIP_CHECK(hipMemcpy(Dptr, Hptr, sizeof(int) * NumElms, - hipMemcpyHostToDevice)); - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); - KrnlWth2MemTypes<<>>(Hmm, Dptr, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the results - for (size_t k = 0; k < NumElms; ++k) { - if (Hmm[k] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!\n"); - IfTestPassed = false; - } - - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipFree(Dptr)); - delete[] Hptr; - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -// The following test case tests when the same Hmm memory is used for -// launching multiple different kernels will results in any issue -TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") { - int managed = HmmAttrPrint(); - if (managed) { - int InitVal = 123, NumElms = (1024 * 1024); - LaunchKrnl4(NumElms, InitVal); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -// The following test case allocation, host access, device access of HMM -// memory from size 1 to 10KB - -TEST_CASE("Unit_hipMallocManaged_MultiSize") { - IfTestPassed = true; - int managed = HmmAttrPrint(); - if (managed == 1) { - unsigned char *Hmm1 = NULL, *Hmm2 = NULL; - int InitVal = 100, blockSize = 64, DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - dim3 dimBlock(blockSize, 1, 1); - for (int i = 1; i < (1024*1024); ++i) { - HIP_CHECK(hipMallocManaged(&Hmm1, i)); - HIP_CHECK(hipMallocManaged(&Hmm2, i)); - for (int j = 0; j < i; ++j) { - Hmm1[j] = InitVal; - } - dim3 dimGrid((i + blockSize -1)/blockSize, 1, 1); - KrnlWth2MemTypesC<<>>(Hmm2, Hmm1, i); - HIP_CHECK(hipStreamSynchronize(strm)); - // Verifying the results - for (int k = 0; k < i; ++k) { - if (Hmm2[k] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!\n"); - IfTestPassed = false; - } - DataMismatch = 0; - HIP_CHECK(hipFree(Hmm1)); - HIP_CHECK(hipFree(Hmm2)); - REQUIRE(IfTestPassed); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} diff --git a/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc b/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc index e0d13b764e..8ac1c37853 100644 --- a/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc +++ b/tests/catch/unit/memory/hipMallocManagedFlagsTst.cc @@ -262,79 +262,3 @@ TEST_CASE("Unit_hipMallocManaged_AccessMultiStream") { } } -TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") { - int managed = HmmAttrPrint(); - if (managed == 1) { - bool IfTestPassed = true; - hipError_t err; - void *Hmm = NULL; - size_t totalDevMem = 0, freeDevMem = 0; - int NumDevs = 0; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - // Testing allocation of extreme and unusual mem values - for (int i = 0; i < NumDevs; i++) { - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMemGetInfo(&freeDevMem, &totalDevMem)); - err = hipMallocManaged(&Hmm, 1, hipMemAttachGlobal); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating memory on GPU: " << i); - WARN(" size 1 with"); - WARN(" hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachGlobal); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max free memory on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachGlobal); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max (free - 1) memory on "); - WARN("GPU: " << i); - WARN(" using hipMallocManaged() api with flag 'hipMemAttachGlobal'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, 1, hipMemAttachHost); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating memory size 1 on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, freeDevMem, hipMemAttachHost); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max free memory on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - err = hipMallocManaged(&Hmm, (freeDevMem - 1), hipMemAttachHost); - if (hipSuccess == err) { - HIP_CHECK(hipFree(Hmm)); - } else { - WARN("Observed error while allocating max (freeDevMem - 1) memory" - " on GPU: " << i); - WARN(" with hipMallocManaged() api with flag 'hipMemAttachHost'\n"); - WARN("Error: " << hipGetErrorString(err)); - IfTestPassed = false; - } - } - REQUIRE(IfTestPassed); - } else { - SUCCEED("Gpu doesnt support HMM! Hence skipping the test with PASS result"); - } -} From 389db38b03dd1bf0fd458eb59ae92986ec04d804 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 10:22:58 +0530 Subject: [PATCH 04/30] SWDEV-323472 - Add hipThreadExchangeStreamCaptureMode (#2519) Change-Id: I1c610833bf3e21e4ba311a89cec0e8866b9164bd Co-authored-by: Christophe Paquot --- include/hip/hip_runtime_api.h | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index eca5b9ff63..0d376d0673 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -4401,6 +4401,18 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t size_t numDependencies, unsigned int flags __dparm(0)); +/** + * @brief Swaps the stream capture mode of a thread. + * + * @param [in] mode - Pointer to mode value to swap with the current mode + * @returns #hipSuccess, #hipErrorInvalidValue + * + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + * + */ +hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode); + /** * @brief Creates a graph * From 73b520a2e0eebdfce08ba495fc017f346fa673a0 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 16:39:07 +0530 Subject: [PATCH 05/30] SWDEV-273235 - build hip samples in Windows (#2513) Change-Id: I90f06ef058a3a8b54726b86d72f36b2ba7f94d4c --- cmake/FindHIP.cmake | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index 75806e1b5f..d56b20f26e 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -50,8 +50,8 @@ mark_as_advanced(HIP_HOST_COMPILATION_CPP) get_filename_component(_IMPORT_PREFIX "${CMAKE_CURRENT_LIST_DIR}/../" REALPATH) -# HIP is supported on Linux only -if(UNIX AND NOT APPLE AND NOT CYGWIN) +# HIP is currently not supported for apple +if(NOT APPLE) # Search for HIP installation if(NOT HIP_ROOT_DIR) # Search in user specified path first @@ -94,7 +94,6 @@ if(UNIX AND NOT APPLE AND NOT CYGWIN) # Now search in default paths find_program(HIP_HIPCC_EXECUTABLE hipcc) endif() - mark_as_advanced(HIP_HIPCC_EXECUTABLE) # Find HIPCONFIG executable find_program( @@ -113,7 +112,12 @@ if(UNIX AND NOT APPLE AND NOT CYGWIN) # Now search in default paths find_program(HIP_HIPCONFIG_EXECUTABLE hipconfig) endif() + if(NOT UNIX) + set(HIP_HIPCONFIG_EXECUTABLE "${HIP_HIPCONFIG_EXECUTABLE}.bat") + set(HIP_HIPCC_EXECUTABLE "${HIP_HIPCC_EXECUTABLE}.bat") + endif() mark_as_advanced(HIP_HIPCONFIG_EXECUTABLE) + mark_as_advanced(HIP_HIPCC_EXECUTABLE) # Find HIPCC_CMAKE_LINKER_HELPER executable find_program( From 302b0e8e125e0952236df21c0b01269920228b53 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 16:40:14 +0530 Subject: [PATCH 06/30] SWDEV-306122 - [catch2][dtest] hipGraphClone API (#2510) Added Negative, functional and multithreaded tests for hipGraphClone API Change-Id: I0733b469c1f1f0e897ed70c49b6ad3756a35f1b5 --- tests/catch/unit/graph/CMakeLists.txt | 3 +- tests/catch/unit/graph/hipGraphClone.cc | 319 ++++++++++++++++++++++++ 2 files changed, 321 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphClone.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 557ddeca06..d6d6f3bafe 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2021 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -25,6 +25,7 @@ set(TEST_SRC hipGraph.cc hipSimpleGraphWithKernel.cc hipGraphAddMemcpyNode.cc + hipGraphClone.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphClone.cc b/tests/catch/unit/graph/hipGraphClone.cc new file mode 100644 index 0000000000..4a6cbc6ec9 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphClone.cc @@ -0,0 +1,319 @@ +/* +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, 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. +*/ + +/* +Testcase Scenarios of hipGraphClone API: + +Negative: + +1. Pass nullptr to cloned graph +2. pass nullptr to original graph + +Functional: + +1. Clone the graph,Instantiate and execute the cloned graph +2. Clone the graph and modify the original graph and ensure that the + cloned graph is not modified +3. Create graph on one GPU device and clone it from peer GPU device +4. Create graph in one thread and clone it from multiple threads. +*/ + +#include +#include +#include + +#define NUM_THREADS 10 + +/* This test covers the negative scenarios of + hipGraphClone API */ + +TEST_CASE("Unit_hipGraphClone_Negative") { + SECTION("Passing nullptr to Cloned graph") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + REQUIRE(hipGraphClone(nullptr, graph) == hipErrorInvalidValue); + HIP_CHECK(hipGraphDestroy(graph)); + } + + SECTION("Passing nullptr to original graph") { + hipGraph_t clonedGraph; + REQUIRE(hipGraphClone(&clonedGraph, nullptr) == hipErrorInvalidValue); + } +} +/* +This function creates the graph with dependencies +then performs device context change and clones the cloned graph +Executes the cloned graph and validates the result +*/ +void hipGraphClone_DeviceContextChange() { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph, clonedgraph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memcpyH2D_A, memcpyD2H_A; + int *A_d{nullptr}, *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, N, false); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A, graph, nullptr, 0, B_h, A_d, + Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &memcpyD2H_A, 1)); + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Instantiate and launch the original graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + for (size_t i = 0; i < N; i++) { + if (A_h[i] != B_h[i]) { + INFO("Validation failed A_h[i] " << A_h[i] << " B_h[i] " << B_h[i]); + REQUIRE(false); + } + } + HipTest::freeArrays(A_d, nullptr, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} +/* +This function does the following +1. Creates the graph with multiple dependencies + clones the graph and validates the result. +2. Creates the graph, clones the graph and modifies + the existing graph and execute the cloned graph + to ensure that cloned graph is not modified +*/ +void hipGraphClone_Func(bool ModifyOrigGraph = false) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, clonedgraph; + hipGraphNode_t memset_A, memset_B, memsetKer_C; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C, memcpyD2D_C, + memcpyD2H_C_new; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + hipMemsetParams memsetParams{}; + int memsetVal{}; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, + &memsetParams)); + + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = + reinterpret_cast(HipTest::memsetReverse); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, + &kernelNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + + if (ModifyOrigGraph) { + // Modify Original graph by adding new dependency + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0, + C_d, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C_new, graph, nullptr, 0, + C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2D_C, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2D_C, + &memcpyD2H_C_new, 1)); + + // Instantiate and launch the original graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + for (size_t i= 0; i < NElem; i++) { + if (C_h[i] != B_h[i]) { + INFO("Validation failed C_h is " << C_h[i] << + "B_h is " << B_h[i]); + REQUIRE(false); + } + } + } + + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies following scenarios +1. Clones the graph and verify the result +2. Clones the graph, Modify the original graph and + validate the result of the cloned graph +3. Device context change for cloned graph +*/ +TEST_CASE("Unit_hipGraphClone_Functional") { + SECTION("hipGraphClone Basic Functionality") { + hipGraphClone_Func(); + } + SECTION("hipGraphClone Modify Original graph") { + hipGraphClone_Func(true); + } + + SECTION("hipGraphClone Device context change") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + hipGraphClone_DeviceContextChange(); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } + } +} + +/* +This testcase creates the graph with dependencies +then creates multiple threads and clones the graph +in each thread and executes the cloned graph +hipGraphClone is failing in CUDA in multi threaded +scenario so excluded for nvidia +*/ +#if HT_AMD +TEST_CASE("Unit_hipGraphClone_MultiThreaded") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphNode_t memcpyH2D_A, memcpyD2H_A; + int *A_d{nullptr}, *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, N, false); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_A, graph, nullptr, 0, B_h, A_d, + Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &memcpyD2H_A, 1)); + std::vector threads; + auto lambdaFunc = [&](){ + hipGraph_t clonedgraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, + nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + for (size_t i = 0; i < N; i++) { + if (A_h[i] != B_h[i]) { + INFO("Validation failed A_h[i] " << A_h[i] << " B_h[i] " << B_h[i]); + REQUIRE(false); + } + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + }; + for (int i = 0; i < NUM_THREADS; i++) { + std::thread t(lambdaFunc); + threads.push_back(std::move(t)); + } + for (auto &t : threads) { + t.join(); + } + HipTest::freeArrays(A_d, nullptr, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} +#endif From 11271fe362bb14c99ac26d1b2f9045cb5031c88a Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 16:41:35 +0530 Subject: [PATCH 07/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphInstantiateWithFlags API (#2508) Added Negative and functional tests for hipGraphInstantiateWithFlags API Change-Id: Ife82615f4aec3c5df90e4640a7cc4381709306d9 Co-authored-by: DURGESH KROTTAPALLI --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphInstantiateWithFlags.cc | 304 ++++++++++++++++++ 2 files changed, 305 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index d6d6f3bafe..ae8445a104 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -26,6 +26,7 @@ set(TEST_SRC hipSimpleGraphWithKernel.cc hipGraphAddMemcpyNode.cc hipGraphClone.cc + hipGraphInstantiateWithFlags.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc new file mode 100644 index 0000000000..b8fabb69ee --- /dev/null +++ b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc @@ -0,0 +1,304 @@ +/* +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. +*/ + +/* +hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph, unsigned long long flags); +Testcase Scenarios of hipGraphInstantiateWithFlags API: + +Negative: +1) Pass nullptr to pGraphExec +2) Pass nullptr to graph +4) Pass invalid flag + +Functional: + +1) Create dependencies graph and instantiate the graph +2) Create graph in one GPU device and instantiate, launch in peer GPU device +3) Create stream capture graph and instantite the graph +4) Create stream capture graph in one GPU device and instantite the graph launch + in peer GPU device + +Mapping is missing for NVIDIA platform hence skipping the testcases +*/ + + +#include +#include +#include + +constexpr size_t N = 1000000; +#if HT_AMD +/* This test covers the negative scenarios of + hipGraphInstantiateWithFlags API */ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") { +#if HT_NVIDIA + SECTION("Passing nullptr pGraphExec") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + REQUIRE(hipGraphInstantiateWithFlags(nullptr, + graph, 0) == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph") { + hipGraphExec_t graphExec; + REQUIRE(hipGraphInstantiateWithFlags(&graphExec, + nullptr, 0) == hipErrorInvalidValue); + } + + SECTION("Passing Invalid flag") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphExec_t graphExec; + REQUIRE(hipGraphInstantiateWithFlags(&graphExec, graph, 10) != hipSuccess); + } +#endif +} +/* +This function verifies the following scenarios +1. Creates dependency graph, Instantiates the graph with flags and verifies it +2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device +*/ +void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memset_A, memset_B, memsetKer_C; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + hipMemsetParams memsetParams{}; + int memsetVal{}; + size_t NElem{N}; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, + &memsetParams)); + + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = + reinterpret_cast(HipTest::memsetReverse); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, + &kernelNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + if (ctxt_change) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* +This function verifies the following scenarios +1. Creates stream capture graph, Instantiates the graph with flags and verifies it +2. Creates graph on one GPU-1 device and instantiates the graph on peer GPU device +*/ +void GraphInstantiateWithFlags_StreamCapture(bool deviceContextChg = false) { + float *A_d, *C_d; + float *A_h, *C_h; + size_t Nbytes = N * sizeof(float); + hipStream_t stream; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + + // Fill with Phi + i + for (size_t i = 0; i < N; i++) { + A_h[i] = 1.618f + i; + } + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + + HIP_CHECK(hipStreamCreate(&stream)); + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + + if (deviceContextChg) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + + // Validate end capture is successful + REQUIRE(graph != nullptr); + HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0)); + REQUIRE(graphExec != nullptr); + + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + UNSCOPED_INFO("A and C not matching at " << i); + REQUIRE(false); + } + } + + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating dependency graph and instantiate, launching and verifying +the result +*/ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") { + GraphInstantiateWithFlags_DependencyGraph(); +} +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating dependency graph on GPU-0 and instantiate, launching and verifying +the result on GPU-1 +*/ +#if HT_NVIDIA +TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + GraphInstantiateWithFlags_DependencyGraph(true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating capture graph and instantiate, launching and verifying +the result +*/ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCapture") { + GraphInstantiateWithFlags_StreamCapture(); +} + +/* +This testcase verifies hipGraphInstantiateWithFlags API +by creating capture graph on GPU-0 and instantiate, launching and verifying +the result on GPU-1 +*/ +TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + GraphInstantiateWithFlags_StreamCapture(true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif From e21231649d43ece78c89a0d0a9a56467b07834b0 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 16:54:07 +0530 Subject: [PATCH 08/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphAddHostNode API (#2507) Added Negative and functional tests for hipGraphAddHostNode API Change-Id: I7237cdf62daa492d94b1f67fcb9df8b2260d8052 Co-authored-by: DURGESH KROTTAPALLI --- tests/catch/unit/graph/CMakeLists.txt | 1 + tests/catch/unit/graph/hipGraphAddHostNode.cc | 312 ++++++++++++++++++ 2 files changed, 313 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphAddHostNode.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index ae8445a104..1143ddda4e 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -27,6 +27,7 @@ set(TEST_SRC hipGraphAddMemcpyNode.cc hipGraphClone.cc hipGraphInstantiateWithFlags.cc + hipGraphAddHostNode.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddHostNode.cc b/tests/catch/unit/graph/hipGraphAddHostNode.cc new file mode 100644 index 0000000000..596a57d755 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddHostNode.cc @@ -0,0 +1,312 @@ +/* +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. +*/ + +/** +Testcase Scenarios of hipGraphAddHostNode API: + +Functional: +1. Creates graph, Adds HostNode which updates the variable and validates the result +2. Create graph, Add Graphnodes and clones the graph. Add Hostnode to the cloned graph + and validate the result +3. Creates graph which performs the square of number in the kernel function and the result + is validated in the callback function of hipGraphAddHostNode API + +Negative: + +1) Pass pGraphNode as nullptr and verify api doesn’t crash, returns error code. +2) Pass graph as nullptr and verify api doesn’t crash, returns error code. +3) Pass pNodeParams as nullptr and verify api doesn’t crash, returns error code. +4) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. +*/ + +#include +#include + +#define SIZE 1024 + +static int *B_h; +static int *D_h; + +static void callbackfunc(void *A_h) { + int *A = reinterpret_cast(A_h); + for (int i = 0; i < SIZE; i++) { + A[i] = i; + } +} + +static void __global__ vector_square(int *B_d, int *D_d) { + for (int i = 0; i < SIZE; i++) { + D_d[i] = B_d[i] * B_d[i]; + } +} +static void vectorsquare_callback(void* ptr) { + // The callback func is not working with zero parameters + // Temporary fix for adding the below 2 lines and ticket + // has been raised for the same. + int *A = reinterpret_cast(ptr); + A++; + for (int i = 0; i < SIZE; i++) { + if (D_h[i] != B_h[i] * B_h[i]) { + INFO("Validation failed " << D_h[i] << B_h[i]); + REQUIRE(false); + } + } +} +/* +This testcase verifies the negative scenarios of +hipGraphAddHostNode API +*/ +TEST_CASE("Unit_hipGraphAddHostNode_Negative") { + constexpr size_t N = 1024; + hipGraph_t graph; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphAddHostNode(nullptr, graph, + nullptr, + 0, &hostParams) == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphAddHostNode(&hostNode, nullptr, + nullptr, + 0, &hostParams) == hipErrorInvalidValue); + } + +#if HT_NVIDIA + SECTION("Passing nullptr to host params") { + REQUIRE(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, nullptr) == hipErrorInvalidValue); + } +#endif + + SECTION("Passing nullptr to host func") { + hostParams.fn = nullptr; + REQUIRE(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams) == hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies hipGraphAddHostNode API in cloned graph +Creates graph, Add graph nodes and clone the graph +Add HostNode to the cloned graph and validate the result +*/ +TEST_CASE("Unit_hipGraphAddHostNode_ClonedGraphwithHostNode") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyH2D_C, + memcpyD2H_AC; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + hipGraph_t clonedgraph; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, clonedgraph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (A_h[i] != static_cast(i)) { + INFO("Validation failed i " << i << "C_h[i] "<< C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies the square of number by +creating graph, Add kernel node which does the square +of number and the result is validated byhipGrahAddHostNode API +*/ +TEST_CASE("Unit_hipGraphAddHostNode_VectorSquare") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *A_h{nullptr}, *B_d{nullptr}, *D_d{nullptr}; + int *param = reinterpret_cast(sizeof(int));; + HipTest::initArrays(&A_d, &B_d, &D_d, + &A_h, &B_h, &D_h, N, false); + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_B, memcpyH2D_D, memcpyD2H_D, kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = vectorsquare_callback; + hostParams.userData = param; + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, + 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_D, graph, nullptr, + 0, D_d, D_h, + Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgs2[] = {&B_d, &D_d}; + kernelNodeParams.func = reinterpret_cast(vector_square); + kernelNodeParams.gridDim = dim3(1); + kernelNodeParams.blockDim = dim3(1); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_D, graph, nullptr, + 0, D_h, D_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_D, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, + &memcpyD2H_D, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_D, + &hostNode, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HipTest::freeArrays(A_d, B_d, D_d, A_h, B_h, D_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} +/* +This testcase verifies the following scenario +Create graph, calls the host function and updates +the parameters in the callback function and +validates it. +*/ +TEST_CASE("Unit_hipGraphAddHostNode_BasicFunc") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC, + &hostNode, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (A_h[i] != static_cast(i)) { + INFO("Validation failed i " << i << "A_h[i] "<< A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} From d305e341730580499536545d7d2a5c8a9d099684 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:56:56 +0530 Subject: [PATCH 09/30] SWDEV-240806 - Added API hipLaunchHostFunc and updated graph test for host graph node (#2512) Change-Id: Idf759064946e503a2b3cc2a38f1a7cb049012688 --- include/hip/hip_runtime_api.h | 13 +++++++++ tests/src/runtimeApi/graph/hipGraph.cpp | 35 ++++++++++++++++++++++--- 2 files changed, 45 insertions(+), 3 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 0d376d0673..b8b511d816 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -4401,6 +4401,19 @@ hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t size_t numDependencies, unsigned int flags __dparm(0)); +/** + * @brief Enqueues a host function call in a stream. + * + * @param [in] stream - stream to enqueue work to. + * @param [in] fn - function to call once operations enqueued preceeding are complete. + * @param [in] userData - User-specified data to be passed to the function. + * @returns #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorInvalidValue, + * #hipErrorNotSupported + * @warning : This API is marked as beta, meaning, while this is feature complete, + * it is still open to changes and may have outstanding issues. + */ +hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData); + /** * @brief Swaps the stream capture mode of a thread. * diff --git a/tests/src/runtimeApi/graph/hipGraph.cpp b/tests/src/runtimeApi/graph/hipGraph.cpp index 427e583181..fc32512208 100644 --- a/tests/src/runtimeApi/graph/hipGraph.cpp +++ b/tests/src/runtimeApi/graph/hipGraph.cpp @@ -105,6 +105,25 @@ bool hipWithoutGraphs(float* inputVec_h, float* inputVec_d, double* outputVec_d, return true; } +typedef struct callBackData { + const char* fn_name; + double* data; +} callBackData_t; +double result_gpu = 0.0; +void myHostNodeCallback(void* data) { + static int iter = 0; + iter++; + // Check status of GPU after stream operations are done + callBackData_t* tmp = (callBackData_t*)(data); + // checkCudaErrors(tmp->status); + double* result = (double*)(tmp->data); + char* function = (char*)(tmp->fn_name); + if (iter == GRAPH_LAUNCH_ITERATIONS) + printf("[%s] Host callback final reduced sum = %lf\n", function, *result); + result_gpu = *result; + *result = 0.0; // reset the result +} + bool hipGraphsUsingStreamCapture(float* inputVec_h, float* inputVec_d, double* outputVec_d, double* result_d, size_t inputSize, size_t numOfBlocks) { hipStream_t stream1, stream2, stream3, streamForGraph; @@ -237,6 +256,16 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, nodeDependencies.clear(); nodeDependencies.push_back(memcpyNode); hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0}; + hostParams.fn = myHostNodeCallback; + callBackData_t hostFnData; + hostFnData.data = &result_h; + hostFnData.fn_name = "hipGraphsManual"; + hostParams.userData = &hostFnData; + + HIPCHECK(hipGraphAddHostNode(&hostNode, graph, nodeDependencies.data(), nodeDependencies.size(), + &hostParams)); + hipGraphExec_t graphExec; hipGraphNode_t* nodes = NULL; size_t numNodes = 0; @@ -266,8 +295,8 @@ bool hipGraphsManual(float* inputVec_h, float* inputVec_d, double* outputVec_d, for (int i = 0; i < inputSize; i++) { result_h_cpu += inputVec_h[i]; } - if (result_h_cpu != result_h) { - printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_h); + if (result_h_cpu != result_gpu) { + printf("Final reduced sum = %lf %lf\n", result_h_cpu, result_gpu); return false; } return true; @@ -304,4 +333,4 @@ int main(int argc, char** argv) { failed("Failed during hipGraph with capture\n"); } passed(); -} \ No newline at end of file +} From 7d8d41e45bfdd49aae3d2b9a2cec09e993c7ec83 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:57:27 +0530 Subject: [PATCH 10/30] SWDEV-306122 - [catch2][dtest] hipGraph tests for hipGraphAddMemcpyNode api (#2498) Change-Id: I65088f4107158e320d54b5e54a47d4ad8ce5d85f --- .../catch/unit/graph/hipGraphAddMemcpyNode.cc | 270 +++++++++++++++++- 1 file changed, 262 insertions(+), 8 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc index 5fcf1ccc40..6fcb3fb90e 100644 --- a/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc @@ -19,18 +19,22 @@ THE SOFTWARE. /** Testcase Scenarios : - 1) Add multiple Memcpy nodes to graph and verify node execution is - working as expected. +1) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D). +Memcpy nodes are added and assigned to default device. +2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and verify the results. +3) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D). +Memcpy nodes are added and assigned to Peer device. +4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and verify the results. */ #include #include -/** - * Functional Test adds memcpy nodes of types H2D, D2D and D2H to graph - * and verifies execution sequence by launching graph. - */ -TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { +#define ZSIZE 32 +#define YSIZE 32 +#define XSIZE 32 + +void validateMemcpyNode3DArray(bool peerAccess = false) { constexpr int width{10}, height{10}, depth{10}; hipArray *devArray1, *devArray2; hipChannelFormatKind formatKind = hipChannelFormatKindSigned; @@ -42,6 +46,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { hipStream_t streamForGraph; hipGraphExec_t graphExec; + HIP_CHECK(hipSetDevice(0)); int *hData = reinterpret_cast(malloc(size)); int *hOutputData = reinterpret_cast(malloc(size)); @@ -69,6 +74,12 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { make_hipExtent(width, height, depth), hipArrayDefault)); HIP_CHECK(hipGraphCreate(&graph, 0)); + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + // Host to Device memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); myparams.srcPos = make_hipPos(0, 0, 0); @@ -79,7 +90,6 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { myparams.dstArray = devArray1; myparams.kind = hipMemcpyHostToDevice; - HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); dependencies.push_back(memcpyNode); @@ -126,3 +136,247 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { free(hData); free(hOutputData); } + +void validateMemcpyNode2DArray(bool peerAccess = false) { + int harray2D[YSIZE][XSIZE]{}; + int harray2Dres[YSIZE][XSIZE]{}; + constexpr int width{XSIZE}, height{YSIZE}; + hipArray *devArray1, *devArray2; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparams; + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize 2D object + for (int i = 0; i < YSIZE; i++) { + for (int j = 0; j < XSIZE; j++) { + harray2D[i][j] = i + j + 1; + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + // Allocate 2D device array by passing depth(0) + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc, + make_hipExtent(width, height, 0), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, + make_hipExtent(width, height, 0), hipArrayDefault)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + + // Host to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, 1); + myparams.srcPtr = make_hipPitchedPtr(harray2D, width * sizeof(int), + width, height); + myparams.dstArray = devArray1; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); + dependencies.push_back(memcpyNode); + + // Device to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.srcArray = devArray1; + myparams.dstArray = devArray2; + myparams.extent = make_hipExtent(width, height, 1); + myparams.kind = hipMemcpyDeviceToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + dependencies.clear(); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, 1); + myparams.dstPtr = make_hipPitchedPtr(harray2Dres, width * sizeof(int), + width, height); + myparams.srcArray = devArray2; + myparams.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate result + for (int i = 0; i < YSIZE; i++) { + for (int j = 0; j < XSIZE; j++) { + if (harray2D[i][j] != harray2Dres[i][j]) { + INFO("harray2D: " << harray2D[i][j] << "harray2Dres: " + << harray2Dres[i][j] << " mismatch at (i,j) : " << i << j); + REQUIRE(false); + } + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + hipFreeArray(devArray1); + hipFreeArray(devArray2); +} + +void validateMemcpyNode1DArray(bool peerAccess = false) { + int harray1D[XSIZE]{}; + int harray1Dres[XSIZE]{}; + constexpr int width{XSIZE}; + hipArray *devArray1, *devArray2; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparams; + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize 1D object + for (int i = 0; i < XSIZE; i++) { + harray1D[i] = i + 1; + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + // Allocate 1D device array by passing depth(0), height(0) + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc, + make_hipExtent(width, 0, 0), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, + make_hipExtent(width, 0, 0), hipArrayDefault)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + + // Host to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, 1, 1); + myparams.srcPtr = make_hipPitchedPtr(harray1D, width * sizeof(int), + width, 1); + myparams.dstArray = devArray1; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); + dependencies.push_back(memcpyNode); + + // Device to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.srcArray = devArray1; + myparams.dstArray = devArray2; + myparams.extent = make_hipExtent(width, 1, 1); + myparams.kind = hipMemcpyDeviceToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + dependencies.clear(); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, 1, 1); + myparams.dstPtr = make_hipPitchedPtr(harray1Dres, width * sizeof(int), + width, 1); + myparams.srcArray = devArray2; + myparams.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate result + for (int i = 0; i < XSIZE; i++) { + if (harray1D[i] != harray1Dres[i]) { + INFO("harray1D: " << harray1D[i] << " harray1Dres: " << harray1Dres[i] + << " mismatch at : " << i); + REQUIRE(false); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + hipFreeArray(devArray1); + hipFreeArray(devArray2); +} + +/** + * Basic Functional Tests adds memcpy nodes of types H2D, D2D and D2H to graph + * and verifies execution sequence by launching graph on default device. + * Tests also verify memcpy node addition with 1D, 2D and 3D objects. + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode_BasicFunctional") { + SECTION("Memcpy with 3D array on default device") { + validateMemcpyNode3DArray(); + } + + SECTION("Memcpy with 2D array on default device") { + validateMemcpyNode2DArray(); + } + + SECTION("Memcpy with 1D array on default device") { + validateMemcpyNode1DArray(); + } +} + +/** + * Peer access tests adds and assigns memcpy nodes of types H2D, D2D and D2H + * to peer device. Memory allocations happen on device(0) and memcpy operations + * are performed from device(1). + * Tests also verify memcpy node addition with 1D, 2D and 3D objects. + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") { + int numDevices{}, peerAccess{}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + } + + if (!peerAccess) { + WARN("Skipping test as peer device access is not found!"); + return; + } + + SECTION("Memcpy with 3D array on peer device") { + validateMemcpyNode3DArray(true); + } + + SECTION("Memcpy with 2D array on peer device") { + validateMemcpyNode2DArray(true); + } + + SECTION("Memcpy with 1D array on peer device") { + validateMemcpyNode1DArray(true); + } +} From d18b4b8848937ee7750187f03df1295e29c34b7e Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:58:04 +0530 Subject: [PATCH 11/30] SWDEV-314080 - Tested All hipMemcpy**() apis with hipStreamPerThread stream obj (#2496) Change-Id: I8f429eb0cc3be2e4d62c76ccb8c1510c56a1e143 --- tests/catch/unit/deviceLib/CMakeLists.txt | 1 + .../unit/deviceLib/hipTestDeviceSymbol.cc | 141 ++++++++++++++++++ tests/catch/unit/memory/hipMemcpy2DAsync.cc | 44 ++++-- .../unit/memory/hipMemcpy2DFromArrayAsync.cc | 16 +- .../unit/memory/hipMemcpy2DToArrayAsync.cc | 17 ++- tests/catch/unit/memory/hipMemcpy3DAsync.cc | 10 +- tests/catch/unit/memory/hipMemcpyPeerAsync.cc | 19 ++- 7 files changed, 219 insertions(+), 29 deletions(-) create mode 100644 tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 347e6d73c9..2f6819f43f 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -10,6 +10,7 @@ set(TEST_SRC popc.cc ldg.cc threadfence_system.cc + hipTestDeviceSymbol.cc ) # skipped for windows compiler issue - Illegal instruction detected diff --git a/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc b/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc new file mode 100644 index 0000000000..501a0805cb --- /dev/null +++ b/tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc @@ -0,0 +1,141 @@ +/* +Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + + +/* Test Case Description: Calling hipMemcpyTo/FromSymbolAsync() using user + declared stream obj and hipStreamPerThread*/ + +#include +#define NUM 1024 +#define SIZE 1024 * 4 + +__device__ int globalIn[NUM]; +__device__ int globalOut[NUM]; + +__global__ void Assign(int* Out) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Out[tid] = globalIn[tid]; + globalOut[tid] = globalIn[tid]; +} + +__device__ __constant__ int globalConst[NUM]; + +__global__ void checkAddress(int* addr, bool* out) { + *out = (globalConst == addr); +} + +TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { + int *A, *Am, *B, *Ad, *C, *Cm; + A = new int[NUM]; + B = new int[NUM]; + C = new int[NUM]; + for (int i = 0; i < NUM; i++) { + A[i] = -1 * i; + B[i] = 0; + C[i] = 0; + } + + HIP_CHECK(hipMalloc((void**)&Ad, SIZE)); + HIP_CHECK(hipHostMalloc((void**)&Am, SIZE)); + HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE)); + for (int i = 0; i < NUM; i++) { + Am[i] = -1 * i; + Cm[i] = 0; + } + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + for (int i = 0; i < NUM; i++) { + assert(Am[i] == B[i]); + assert(Am[i] == Cm[i]); + } + + for (int i = 0; i < NUM; i++) { + A[i] = -2 * i; + B[i] = 0; + } + + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, + hipMemcpyHostToDevice)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost)); + for (int i = 0; i < NUM; i++) { + assert(A[i] == B[i]); + assert(A[i] == C[i]); + } + + for (int i = 0; i < NUM; i++) { + A[i] = -3 * i; + B[i] = 0; + } + SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } + for (int i = 0; i < NUM; i++) { + assert(A[i] == B[i]); + assert(A[i] == C[i]); + } + + bool *checkOkD; + bool checkOk = false; + size_t symbolSize = 0; + int *symbolAddress; + HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst))); + HIP_CHECK(hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst))); + HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool))); + hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD); + HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(checkOkD)); + HIP_ASSERT(checkOk); + HIP_ASSERT((symbolSize == SIZE)); + + HIP_CHECK(hipHostFree(Am)); + HIP_CHECK(hipHostFree(Cm)); + HIP_CHECK(hipFree(Ad)); + delete[] A; + delete[] B; + delete[] C; +} diff --git a/tests/catch/unit/memory/hipMemcpy2DAsync.cc b/tests/catch/unit/memory/hipMemcpy2DAsync.cc index c5733a095e..10924687df 100644 --- a/tests/catch/unit/memory/hipMemcpy2DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy2DAsync.cc @@ -83,22 +83,40 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_Host&PinnedMem", "" // Initialize the data HipTest::setDefaultData(NUM_W*NUM_H, A_h, B_h, C_h); + SECTION("Calling Async apis with stream object created by user") { + // Host to Device + HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyHostToDevice, stream)); - // Host to Device - HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), - COLUMNS*sizeof(TestType), ROWS, - hipMemcpyHostToDevice, stream)); + // Performs D2D on same GPU device + HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, + pitch_A, COLUMNS*sizeof(TestType), + ROWS, hipMemcpyDeviceToDevice, stream)); - // Performs D2D on same GPU device - HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, - pitch_A, COLUMNS*sizeof(TestType), - ROWS, hipMemcpyDeviceToDevice, stream)); + // hipMemcpy2DAsync Device to Host + HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling Async apis with hipStreamPerThread") { + // Host to Device + HIP_CHECK(hipMemcpy2DAsync(A_d, pitch_A, A_h, COLUMNS*sizeof(TestType), + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyHostToDevice, hipStreamPerThread)); - // hipMemcpy2DAsync Device to Host - HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, - COLUMNS*sizeof(TestType), ROWS, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + // Performs D2D on same GPU device + HIP_CHECK(hipMemcpy2DAsync(B_d, pitch_B, A_d, + pitch_A, COLUMNS*sizeof(TestType), + ROWS, hipMemcpyDeviceToDevice, hipStreamPerThread)); + + // hipMemcpy2DAsync Device to Host + HIP_CHECK(hipMemcpy2DAsync(B_h, COLUMNS*sizeof(TestType), B_d, pitch_B, + COLUMNS*sizeof(TestType), ROWS, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } // Validating the result REQUIRE(HipTest::checkArray(A_h, B_h, COLUMNS, ROWS) == true); diff --git a/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc index a8c9ac2944..6603243585 100644 --- a/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc @@ -62,10 +62,18 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_Basic") { HIP_CHECK(hipMemcpy2DToArray(A_d, 0, 0, hData, width, width, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, - 0, 0, width, NUM_H, - hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpy2DFromArrayAsync() with user declared stream obj") { + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy2DFromArrayAsync() with hipStreamPerThread") { + HIP_CHECK(hipMemcpy2DFromArrayAsync(A_h, width, A_d, + 0, 0, width, NUM_H, + hipMemcpyDeviceToHost, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } REQUIRE(HipTest::checkArray(A_h, hData, NUM_W, NUM_H) == true); // Cleaning the memory diff --git a/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc b/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc index fecc031f52..5145f28aa1 100644 --- a/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy2DToArrayAsync.cc @@ -58,11 +58,18 @@ TEST_CASE("Unit_hipMemcpy2DToArrayAsync_Basic") { HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault)); HipTest::setDefaultData(width*NUM_H, A_h, hData, nullptr); HIP_CHECK(hipStreamCreate(&stream)); - - HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, - width, NUM_H, - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpy2DToArrayAsync() with user declared stream obj") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy2DToArrayAsync() with hipStreamPerThread") { + HIP_CHECK(hipMemcpy2DToArrayAsync(A_d, 0, 0, hData, width, + width, NUM_H, + hipMemcpyHostToDevice, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } HIP_CHECK(hipMemcpy2DFromArray(A_h, width, A_d, 0, 0, width, NUM_H, hipMemcpyDeviceToHost)); diff --git a/tests/catch/unit/memory/hipMemcpy3DAsync.cc b/tests/catch/unit/memory/hipMemcpy3DAsync.cc index fb2fa8e1dc..6ebc480844 100644 --- a/tests/catch/unit/memory/hipMemcpy3DAsync.cc +++ b/tests/catch/unit/memory/hipMemcpy3DAsync.cc @@ -602,8 +602,14 @@ void Memcpy3DAsync::simple_Memcpy3DAsync() { #else myparms.kind = hipMemcpyHostToDevice; #endif - REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpy3DAsync() using user declared stream obj") { + REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpy3DAsync() using hipStreamPerThread") { + REQUIRE(hipMemcpy3DAsync(&myparms, hipStreamPerThread) == hipSuccess); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } // Array to Array memset(&myparms, 0x0, sizeof(hipMemcpy3DParms)); diff --git a/tests/catch/unit/memory/hipMemcpyPeerAsync.cc b/tests/catch/unit/memory/hipMemcpyPeerAsync.cc index c2e67bb802..c3a0b31501 100644 --- a/tests/catch/unit/memory/hipMemcpyPeerAsync.cc +++ b/tests/catch/unit/memory/hipMemcpyPeerAsync.cc @@ -149,11 +149,20 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") { // Copying data from GPU-0 to GPU-1 and performing vector addition HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, - stream)); - HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, - stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + SECTION("Calling hipMemcpyPerAsync() using user defined stream obj") { + HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, + stream)); + HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") { + HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, + hipStreamPerThread)); + HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, + hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, static_cast(X_d), static_cast(Y_d), Z_d, numElements*sizeof(int)); From fe2bb1bf3eb5dfd1b698014b190fd1b26e22336f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:58:17 +0530 Subject: [PATCH 12/30] SWDEV-298757 - hipMemRangeGetAttribute.cc Adding test cases for hipMemRangeGetAttribute() api (#2495) Change-Id: I6dc6476383799963dd2a079b4a074b2da6fe8cca From bdc8f76306dffd36acaa4b36d628f79159b907a8 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:58:38 +0530 Subject: [PATCH 13/30] SWDEV-314080 - Testing All hipMemsetxx() apis with hipStreamPerThread (#2494) Change-Id: I452edb7d0047146b22500750bf529a59ab7f99d5 --- .../hipMemset2DAsyncMultiThreadAndKernel.cc | 49 +++++++++++++------ .../unit/memory/hipMemset3DFunctional.cc | 18 ++++--- .../unit/memory/hipMemsetAsyncAndKernel.cc | 35 +++++++++---- 3 files changed, 71 insertions(+), 31 deletions(-) diff --git a/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc b/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc index 479e163c26..89390bb908 100644 --- a/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc +++ b/tests/catch/unit/memory/hipMemset2DAsyncMultiThreadAndKernel.cc @@ -61,7 +61,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { size_t elements = NUM_W * NUM_H; unsigned blocks{}; int validateCount{}; - hipStream_t stream; blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, @@ -81,21 +80,42 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { } HIP_CHECK(hipMemcpy2D(B_d, width, B_h, pitch_B, NUM_W, NUM_H, hipMemcpyHostToDevice)); - HIP_CHECK(hipStreamCreate(&stream)); + SECTION("Using User created stream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + for (size_t k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, B_d, C_d, elements); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, + stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); + for (size_t p = 0 ; p < elements ; p++) { + if (A_h[p] == memsetval) { + validateCount+= 1; + } + } + } + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Using hipStreamPerThread stream") { + for (size_t k = 0; k < ITER; k++) { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, hipStreamPerThread, B_d, C_d, elements); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, + hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, + hipMemcpyDeviceToHost)); - for (size_t k = 0; k < ITER; k++) { - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream, B_d, C_d, elements); - - HIP_CHECK(hipMemset2DAsync(C_d, pitch_C, memsetval, NUM_W, NUM_H, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipMemcpy2D(A_h, width, C_d, pitch_C, NUM_W, NUM_H, - hipMemcpyDeviceToHost)); - - for (size_t p = 0 ; p < elements ; p++) { - if (A_h[p] == memsetval) { - validateCount+= 1; + for (size_t p = 0 ; p < elements ; p++) { + if (A_h[p] == memsetval) { + validateCount+= 1; + } } } } @@ -104,7 +124,6 @@ TEST_CASE("Unit_hipMemset2DAsync_WithKernel") { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(B_d)); HIP_CHECK(hipFree(C_d)); free(A_h); free(B_h); - HIP_CHECK(hipStreamDestroy(stream)); } diff --git a/tests/catch/unit/memory/hipMemset3DFunctional.cc b/tests/catch/unit/memory/hipMemset3DFunctional.cc index b47798487b..35d5696bc4 100644 --- a/tests/catch/unit/memory/hipMemset3DFunctional.cc +++ b/tests/catch/unit/memory/hipMemset3DFunctional.cc @@ -184,11 +184,17 @@ static void testMemsetMaxValue(bool bAsync) { HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); if (bAsync) { - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); + SECTION("Using user created stream") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Using hipStreamPerThread") { + HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } } else { HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent)); } @@ -236,7 +242,7 @@ static void seekAndSet3DArraySlice(bool bAsync) { // select random slice for memset unsigned int seed = time(nullptr); - int slice_index = HipTest::RAND_R(&seed) % ZSIZE_S; + int slice_index = rand_r(&seed) % ZSIZE_S; INFO("memset3d for sliceindex " << slice_index); diff --git a/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc b/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc index adbd4a3964..270561f811 100644 --- a/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc +++ b/tests/catch/unit/memory/hipMemsetAsyncAndKernel.cc @@ -21,7 +21,6 @@ * Test for checking order of execution of device kernel and * hipMemsetAsync apis on all gpus */ - #include #include #include @@ -83,16 +82,26 @@ class MemSetKernelTest { } }; -static bool testhipMemsetAsyncWithKernel() { +static bool testhipMemsetAsyncWithKernel(bool UseStrmPerThrd) { MemSetKernelTest obj; constexpr char memsetval = 0x42; obj.memAllocate(memsetval); - for (int k = 0 ; k < ITER ; k++) { - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); - HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); - HIP_CHECK(hipStreamSynchronize(obj.stream)); + for (int k = 0 ; k < ITER ; ++k) { + if (UseStrmPerThrd) { // will use hipStreamPerThread stream object + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, hipStreamPerThread, obj.B_d, + obj.C_d, N); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal, N, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + } else { + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, + N); + HIP_CHECK(hipMemsetAsync(obj.C_d , obj.memSetVal , N , obj.stream)); + HIP_CHECK(hipStreamSynchronize(obj.stream)); + } HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); obj.validateExecutionOrder(); @@ -109,7 +118,7 @@ static bool testhipMemsetD32AsyncWithKernel() { hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, obj.stream, obj.B_d, obj.C_d, N); HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)obj.C_d , obj.memSetVal, - N, obj.stream)); + N, obj.stream)); HIP_CHECK(hipStreamSynchronize(obj.stream)); HIP_CHECK(hipMemcpy(obj.A_h, obj.C_d, obj.Nbytes, hipMemcpyDeviceToHost)); @@ -161,7 +170,7 @@ static bool testhipMemsetD8AsyncWithKernel() { */ TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") { int numDevices = 0; - bool ret; + bool ret, UseStrmPerThrd = false; blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -172,7 +181,13 @@ TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") { HIP_CHECK(hipSetDevice(devNum)); SECTION("hipMemsetAsync With Kernel") { - ret = testhipMemsetAsyncWithKernel(); + UseStrmPerThrd = false; + ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd); + REQUIRE(ret == true); + } + SECTION("hipMemsetAsync With Kernel using hipStreamPerThread stream obj") { + UseStrmPerThrd = true; + ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd); REQUIRE(ret == true); } From c10b08bc3430881fe8d9d053b3a9225edb5e0536 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:59:00 +0530 Subject: [PATCH 14/30] SWDEV-298757 - hipHmmOvrSubscriptionTst.cc adding HMM over subscription test (#2493) Change-Id: I4728a329e673cbb94656a28b266cdcd046a07144 --- tests/catch/unit/memory/CMakeLists.txt | 2 + .../unit/memory/hipHmmOvrSubscriptionTst.cc | 213 ++++++++++++++++++ 2 files changed, 215 insertions(+) create mode 100644 tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 95cf97d736..23bbd07aae 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -37,6 +37,7 @@ set(TEST_SRC hipMemCoherencyTst.cc hipMallocManaged.cc hipMemRangeGetAttribute.cc + hipHmmOvrSubscriptionTst.cc ) else() set(TEST_SRC @@ -73,6 +74,7 @@ set(TEST_SRC hipMemAdviseMmap.cc hipMallocManaged.cc hipMemRangeGetAttribute.cc + hipHmmOvrSubscriptionTst.cc ) endif() diff --git a/tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc b/tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc new file mode 100644 index 0000000000..de5d48ee4b --- /dev/null +++ b/tests/catch/unit/memory/hipHmmOvrSubscriptionTst.cc @@ -0,0 +1,213 @@ +/* +Copyright (c) 2021-Present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* Test Case Description: This test case tests the working of OverSubscription + feature which is part of HMM.*/ + +#include +#ifdef __linux__ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#endif +#include + +#define INIT_VAL 2.5 +#define NUM_ELMS 268435456 // 268435456 * 4 = 1GB +#define ITERATIONS 10 +#define ONE_GB 1024 * 1024 * 1024 + +static void GetTotGpuMem(int *TotMem); +static void DisplayHmmFlgs(int *Signal); +// Kernel function +__global__ void Square(int n, float *x) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + x[i] = x[i] + 10; + } +} + +static void OneGBMemTest(int dev) { + int DataMismatch = 0; + float *HmmAG = nullptr; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + // Testing hipMemAttachGlobal Flag + HIP_CHECK(hipMallocManaged(&HmmAG, NUM_ELMS * sizeof(float), + hipMemAttachGlobal)); + + // Initializing HmmAG memory + for (int i = 0; i < NUM_ELMS; i++) { + HmmAG[i] = INIT_VAL; + } + + int blockSize = 256; + int numBlocks = (NUM_ELMS + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + HIP_CHECK(hipSetDevice(dev)); + for (int i = 0; i < ITERATIONS; ++i) { + Square<<>>(NUM_ELMS, HmmAG); + } + HIP_CHECK(hipStreamSynchronize(strm)); + for (int j = 0; j < NUM_ELMS; ++j) { + if (HmmAG[j] != (INIT_VAL + ITERATIONS * 10)) { + DataMismatch++; + break; + } + } + if (DataMismatch != 0) { + WARN("Data Mismatch observed when kernel launched on device: " << dev); + REQUIRE(false); + } + HIP_CHECK(hipFree(HmmAG)); + HIP_CHECK(hipStreamDestroy(strm)); +} + +static void GetTotGpuMem(int *TotMem) { + size_t FreeMem, TotGpuMem; + HIP_CHECK(hipMemGetInfo(&FreeMem, &TotGpuMem)); + TotMem[0] = (TotGpuMem/(ONE_GB)); + TotMem[1] = 1; +} + +static void DisplayHmmFlgs(int *Signal) { + 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); + + // Checking for Vega20 or MI100 + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + char *p = NULL; + p = strstr(prop.gcnArchName, "gfx906"); + if (p) { + WARN("This system has MI60 gpu hence OverSubscription test will be"); + WARN(" skipped"); + Signal[2] = 1; + } + p = strstr(prop.gcnArchName, "gfx908"); + if (p) { + WARN("This system has MI100 gpu hence OverSubscription test will be"); + WARN(" skipped"); + Signal[2] = 1; + } + Signal[1] = managed; + Signal[0] = 1; +} + +TEST_CASE("Unit_HMM_OverSubscriptionTst") { + int HmmEnabled = 0; + // The following Shared Mem is to get Max GPU Mem + // The size requested is for three ints + // 1) To get Max GPU Mem in GB + // 2) To Signal parent that req. info is available to consume + // 3) To know if MI60 or MI100 gpu are there in the system + key_t key = ftok("shmTotMem", 66); + int shmid = shmget(key, (3 * sizeof(int)), 0666|IPC_CREAT); + int *TotGpuMem = reinterpret_cast(shmat(shmid, NULL, 0)); + TotGpuMem[0] = 0; TotGpuMem[1] = 0; + // The following function DisplayHmmFlgs() displays the flag values related + // to HMM and also sends us ManagedMemory attribute value + if (fork() == 0) { + DisplayHmmFlgs(TotGpuMem); + exit(1); + } + while (TotGpuMem[0] == 0) { + sleep(2); + } + // The following if block will skip test if either of MI60 or MI100 is found + if (TotGpuMem[2] == 1) { + SUCCEED("Test is skipped!!"); + REQUIRE(true); + } else { + HmmEnabled = TotGpuMem[1]; + + // Re-setting the shared memory values for further usage + TotGpuMem[0] = 0; + TotGpuMem[1] = 0; + + std::list PidLst; + // The following function gets the MaxGpu memory in GBs and also launches + // OverSubscription test + if (HmmEnabled) { + if ((setenv("HSA_XNACK", "1", 1)) != 0) { + WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); + REQUIRE(false); + } + if (fork() == 0) { + GetTotGpuMem(TotGpuMem); + } + while (TotGpuMem[1] == 0) { + sleep(2); + } + int NumGB = TotGpuMem[0], TotalThreads = (NumGB + 10); + WARN("Launching " << TotalThreads); + WARN(" processes to test OverSubscription."); + pid_t pid; + for (int k = 0; k < TotalThreads; ++k) { + pid = fork(); + PidLst.push_back(pid); + if (pid == 0) { + OneGBMemTest(0); + exit(10); + } + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } + int status; + for (pid_t pd : PidLst) { + waitpid(pd, &status, 0); + if (!(WIFEXITED(status))) { + REQUIRE(false); + } + } + } + shmdt(TotGpuMem); + shmctl(shmid, IPC_RMID, NULL); +} From 46d9f76702b440991377dacaa084aa9a700d046c Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Thu, 24 Feb 2022 05:29:18 -0800 Subject: [PATCH 15/30] SWDEV-1 - hipcc/hipconfig path detection in bat file (#2488) Change-Id: I075148087bc07b5044a3360f1636d07372b16a1d --- bin/hipcc.bat | 2 +- bin/hipconfig.bat | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/bin/hipcc.bat b/bin/hipcc.bat index a19ce9d80f..7a3c25ba8d 100644 --- a/bin/hipcc.bat +++ b/bin/hipcc.bat @@ -1,2 +1,2 @@ -@IF DEFINED HIP_PATH (set HIPCC="%HIP_PATH%/bin/hipcc") ELSE (set HIPCC="%~dp0/hipcc") +set HIPCC="%~dp0/hipcc" @perl %HIPCC% %* diff --git a/bin/hipconfig.bat b/bin/hipconfig.bat index 9ba2d4205c..47d9500c71 100644 --- a/bin/hipconfig.bat +++ b/bin/hipconfig.bat @@ -1,2 +1,2 @@ -@IF DEFINED HIP_PATH (set HIPCONFIG="%HIP_PATH%/bin/hipconfig") ELSE (set HIPCONFIG="%~dp0/hipconfig") +set HIPCONFIG="%~dp0/hipconfig" @perl %HIPCONFIG% %* From 1dca3b328fae080ebdd18d432a199358c3a7ee9c Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:43:14 +0530 Subject: [PATCH 16/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphAddMemcpyNodeFromSymbol API (#2506) Added Negative and functional tests for hipGraphAddMemcpyNodeFromSymbol API Change-Id: Ic65045ca56ae8949efd7690e62f788ff888cd893 --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphAddMemcpyNodeFromSymbol.cc | 435 ++++++++++++++++++ 2 files changed, 436 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 1143ddda4e..1e57939a81 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -28,6 +28,7 @@ set(TEST_SRC hipGraphClone.cc hipGraphInstantiateWithFlags.cc hipGraphAddHostNode.cc + hipGraphAddMemcpyNodeFromSymbol.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc new file mode 100644 index 0000000000..7d9cbaffed --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc @@ -0,0 +1,435 @@ +/* +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. +*/ + +/** +Testcase Scenarios of hipGraphAddMemcpyNodeFromSymbol API: + +Functional : + +1. Allocate global symbol memory, add the MemcpyNodeFromSymbol + node to the graph and verify for different memory kinds +2. Allocate const memory add the MemcpyNodeFromSymbol node to + the graph and verify for different memory kinds +3. Allocate global symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +4. Allocate const symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +5. Allocate global memory, Add MemcpyFromSymbolNode,KernelNode and memcpynode and validating + the behaviour + +Negative : + +1) Pass nullptr to graph node +2) Pass nullptr to graph +3) Pass nullptr to dependencies +4) Pass invalid numDependencies +5) Pass nullptr to dst +6) Pass nullptr to symbol +7) Pass invalid count +8) Pass offset+count greater than allocated size +9) Pass unintialized graph +*/ + +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE]; +__device__ int globalOut[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + +__global__ void MemcpyFromSymbolKernel(int* B_d) { + for (int i = 0 ; i < SIZE; i++) { + globalIn[i] = B_d[i]; + } +} + +/* This testcase verifies negative scenarios of + hipGraphAddMemcpyNodeFromSymbol API */ +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + +#if HT_NVIDIA + hipGraphNode_t memcpyFromSymbolNode; + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, nullptr, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(nullptr, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing size > 1 and dependencies as nullptr") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + nullptr, + 1, + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing invalid dependencies size") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + 10, + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to dst") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + nullptr, + HIP_SYMBOL(globalIn), Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to source") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidSymbol); + } + + SECTION("Passing offset+size > max size") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 10, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing Max count") { + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + std::numeric_limits::max(), 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Pass Unintialized graph") { + hipGraph_t unint_graph; + REQUIRE(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, unint_graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } +#endif + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This function is used to verify the following scenarios +1. Create global variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeFromSymbol API in GPU-1 and validate the result +2. Allocate global memory, Create dependency graph and validate the result on GPU-0 +3. Allocate global const memory, Create dependency graph and validate the result on GPU-0 +4. Create global const variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeFromSymbol API in GPU-1 and validate the result +*/ + +void hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(bool device_ctxchg = false, + bool const_device_var = + false) { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + if (device_ctxchg) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalConst), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + + } else { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + + // Adding MemcpyNodeFromSymbol + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalConst), + Nbytes, 0, + hipMemcpyDeviceToHost)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToHost)); + } + + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies allocating global symbol memory, +add the MemcpyNodeFromSymbol node to the graph and +erifying the result +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory") { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(false, false); +} + +/* +This testcase verifies allocating global const symbol memory, +add the MemcpyNodeFromSymbol node to the graph and +verifying the result +*/ + +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory") { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(false, true); +} + +/* +This testcase verifies allocating global symbol memory and device variables +in GPU-0 and add the MemcpyNodeFromSymbol node to the graph and +verifying the result in GPU-1 +*/ +#if HT_NVIDIA +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(true, false); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} + +/* +This testcase verifies allocating global const symbol memory and device variables +in GPU-0 and add the MemcpyNodeFromSymbol node to the graph and +verifying the result in GPU-1 +*/ + +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); + if (canAccessPeer) { + hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(true, true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif +/* +This testcaser verifies allocating global memory, +Add MemcpyFromSymbolNode,KernelNode and memcpynode and validating +the behaviour +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel") { + constexpr size_t Nbytes = SIZE * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE); + hipGraphNode_t memcpyfromsymbolkernel, memcpyD2H_B; + hipKernelNodeParams kernelNodeParams{}; + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_d, + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyFromSymbolNode); + + // Adding Kernel node + void* kernelArgs1[] = {&B_d}; + kernelNodeParams.func = + reinterpret_cast(MemcpyFromSymbolKernel); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memcpyfromsymbolkernel, graph, + dependencies.data(), dependencies.size(), + &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(memcpyfromsymbolkernel); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(), + dependencies.size(), B_h, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} From 97273320b166568d2292f3c7afc23a8d839c26b2 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:43:51 +0530 Subject: [PATCH 17/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphChildGraphNodeGetGraph API (#2505) Added Negative and functional tests for hipGraphChildGraphNodeGetGraph API Change-Id: I93c905687b684448297bde203414f9e6cc1f454a --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphChildGraphNodeGetGraph.cc | 163 ++++++++++++++++++ 2 files changed, 164 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 1e57939a81..b2aed70084 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -29,6 +29,7 @@ set(TEST_SRC hipGraphInstantiateWithFlags.cc hipGraphAddHostNode.cc hipGraphAddMemcpyNodeFromSymbol.cc + hipGraphChildGraphNodeGetGraph.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc b/tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc new file mode 100644 index 0000000000..4a975dd3c0 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphChildGraphNodeGetGraph.cc @@ -0,0 +1,163 @@ +/*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. +*/ + +/** +Testcase Scenarios of hipGraphChildGraphNodeGetGraph API: + +Functional Scenarios: +1. Get the child graph node from the original graph and execute it + +Negative Scenarios: +1. Pass nullptr to graph +2. Pass nullptr to graphnode +3. Pass uninitialized graph node +4. Pass orginial graph node instead of child graph node +**/ + +#include +#include +#include + +/* +This testcase verifies the following scenario +Create graph, add multiple child nodes and gets the +graph of one of the child nodes using hipGraphChildGraphNodeGetGraph API +executes it and validates the results +*/ +TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + size_t NElem{N}; + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, childgraph1, childgraph2; + hipGraphExec_t graphExec; + hipKernelNodeParams kernelNodeParams{}; + hipGraphNode_t kernel_vecAdd; + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1, + childGraphNode2, memcpyD2H_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphCreate(&childgraph1, 0)); + HIP_CHECK(hipGraphCreate(&childgraph2, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph1, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, childgraph2, nullptr, + 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, + nullptr, 0, childgraph1)); + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode2, graph, + nullptr, 0, childgraph2)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, + 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode1, + &childGraphNode2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &childGraphNode2, + &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + hipGraph_t Getgraph; + HIP_CHECK(hipGraphChildGraphNodeGetGraph(childGraphNode1, &Getgraph)); + // Instantiate and launch the child graph + HIP_CHECK(hipGraphInstantiate(&graphExec, Getgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + + // Verify child graph execution result + HIP_CHECK(hipMemcpy(C_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < N; i++) { + if (A_h[i] != C_h[i]) { + INFO("Validation failed " << A_h[i] << C_h[i]); + REQUIRE(false); + } + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(childgraph2)); + HIP_CHECK(hipGraphDestroy(childgraph1)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies the negative scenarios +of hipGraphChildGraphNodeGetGraph API +*/ +TEST_CASE("Unit_hipGraphChildGraphNodeGetGraph_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph, childgraph1; + int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, childGraphNode1; + HIP_CHECK(hipGraphCreate(&childgraph1, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, childgraph1, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, + nullptr, 0, childgraph1)); + + hipGraph_t Getgraph; + SECTION("nullptr to child node") { + REQUIRE((hipGraphChildGraphNodeGetGraph(nullptr, &Getgraph)) + == hipErrorInvalidValue); + } +#if HT_NVIDIA + SECTION("nullptr to graph") { + REQUIRE((hipGraphChildGraphNodeGetGraph(childGraphNode1, nullptr)) + == hipErrorInvalidValue); + } + + SECTION("Passing parent instead of child graph node") { + REQUIRE((hipGraphChildGraphNodeGetGraph(memcpyH2D_A, &Getgraph)) + == hipErrorInvalidValue); + } + + SECTION("Passing unintialized node") { + hipGraphNode_t unint_node; + REQUIRE((hipGraphChildGraphNodeGetGraph(unint_node, &Getgraph)) + == hipErrorInvalidValue); + } +#endif + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphDestroy(childgraph1)); +} From 11b71138cd3b5ceb75bbacbf538ea789ea2326af Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:44:23 +0530 Subject: [PATCH 18/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphNodeFindInClone API (#2504) Added Negative and functional tests for hipGraphNodeFindInClone API Change-Id: I62382f6e5fcb5c1ca62fca7326efb8e6fc0ab624 --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../unit/graph/hipGraphNodeFindInClone.cc | 241 ++++++++++++++++++ 2 files changed, 242 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphNodeFindInClone.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index b2aed70084..14a2dbf40f 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -30,6 +30,7 @@ set(TEST_SRC hipGraphAddHostNode.cc hipGraphAddMemcpyNodeFromSymbol.cc hipGraphChildGraphNodeGetGraph.cc + hipGraphNodeFindInClone.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphNodeFindInClone.cc b/tests/catch/unit/graph/hipGraphNodeFindInClone.cc new file mode 100644 index 0000000000..174c0ebcb9 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphNodeFindInClone.cc @@ -0,0 +1,241 @@ +/* +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. +*/ + +/* +Testcase Scenarios of hipGraphNodeFindInClone API: + +Negative: + +1) Pass nullptr to graph node +2) pass nullptr to original graph node +3) pass nullptr to clonedGraph +4) Pass original graph in place of the cloned graph +5) Pass invalid originalNode +6) Destroy the graph node in the original graph + and try to get the deleted graph node + from the cloned graph +7) Clone the graph,Add node to Original graph + and try to find the original node in the cloned graph + + +Functional: + +1) Get the graph node from the cloned graph corresponding to the original node +2) Create and clone the graph, modify the original graph and clone the graph again, + then try to find the newly added graph node from the cloned graph + +*/ + +#include +#include +#include +#include + + +/* This test covers the negative scenarios of + hipGraphNodeFindInClone API */ + +TEST_CASE("Unit_hipGraphNodeFindInClone_Negative") { + hipGraph_t graph; + hipGraph_t clonedgraph; + hipGraphNode_t graphnode, newnode; + hipGraphNode_t clonedgraphnode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + + int *A_d, *A_h, *B_d, *B_h; + HipTest::initArrays(&A_d, &B_d, nullptr, &A_h, + &B_h, nullptr, 1024, false); + HIP_CHECK(hipGraphAddMemcpyNode1D(&graphnode, graph, nullptr, 0, A_d, A_h, + 1024, hipMemcpyHostToDevice)); + // Cloned the graph + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&newnode, graph, nullptr, 0, B_d, B_h, + 1024, hipMemcpyHostToDevice)); + + SECTION("Passing nullptr to Cloned graph") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, nullptr) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to original graph") { + REQUIRE(hipGraphNodeFindInClone(nullptr, graphnode, clonedgraph) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, nullptr, clonedgraph) + == hipErrorInvalidValue); + } +#if HT_NVIDIA + SECTION("Pass uncloned graph") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, graph) + == hipErrorInvalidValue); + } + + SECTION("Destroy the graph node and find in cloned graph") { + HIP_CHECK(hipGraphDestroyNode(graphnode)); + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, graphnode, + clonedgraph) + == hipErrorInvalidValue); + } +#endif + + SECTION("Pass invalid original graphnode") { + hipGraphNode_t unintialized_graphnode{nullptr}; + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, unintialized_graphnode, + graph) + == hipErrorInvalidValue); + } + + SECTION("Find node in cloned graph which is only present in original graph") { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, newnode, + clonedgraph) == hipErrorInvalidValue); + } + + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); +} + + +void hipGraphNodeFindInClone_Func(bool ModifyOrigGraph = false) { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph, clonedgraph; + hipGraphNode_t memset_A, memset_B, memsetKer_C; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C, memcpyD2D_C, + memcpyD2H_C_new; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipMemsetParams memsetParams{}; + int memsetVal{}; + size_t NElem{N}; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(B_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0, + &memsetParams)); + + void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast(&NElem)}; + kernelNodeParams.func = + reinterpret_cast(HipTest::memsetReverse); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0, + &kernelNodeParams)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + + if (ModifyOrigGraph) { + // Cloned the graph + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + // Modify Original graph by adding new dependency + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D_C, graph, nullptr, 0, + C_d, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C_new, graph, nullptr, 0, + C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2D_C, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2D_C, + &memcpyD2H_C_new, 1)); + } + // Cloned the graph + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + hipGraphNode_t clonedgraphnode; + if (ModifyOrigGraph) { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, + memcpyD2H_C_new, clonedgraph) + == hipSuccess); + } else { + REQUIRE(hipGraphNodeFindInClone(&clonedgraphnode, + memcpyH2D_A, clonedgraph) + == hipSuccess); + } + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(clonedgraph)); +} + +TEST_CASE("Unit_hipGraphNodeFindInClone_Functional") { + SECTION("hipGraphNodeFindInClone Basic Functionality") { + hipGraphNodeFindInClone_Func(); + } + SECTION("hipGraphNodeFindInClone Modify Original graph") { + hipGraphNodeFindInClone_Func(true); + } +} From 16edb6c7d2ae4f24eb91718dd3924edeac2a1936 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:44:54 +0530 Subject: [PATCH 19/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphExecHostNodeSetParams API (#2502) Added Negative and functional tests for hipGraphExecHostNodeSetParams API Change-Id: I56ae94cff564fc0888fa6f0398ecb913cf4b9eed --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphExecHostNodeSetParams.cc | 276 ++++++++++++++++++ 2 files changed, 277 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 14a2dbf40f..48340d488a 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -31,6 +31,7 @@ set(TEST_SRC hipGraphAddMemcpyNodeFromSymbol.cc hipGraphChildGraphNodeGetGraph.cc hipGraphNodeFindInClone.cc + hipGraphExecHostNodeSetParams.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc b/tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc new file mode 100644 index 0000000000..6f07b23269 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecHostNodeSetParams.cc @@ -0,0 +1,276 @@ +/* +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. +*/ + +/** +Testcase Scenarios of hipGraphExecHostNodeSetParams API: + +Functional: +1. Creates graph, Adds HostNode, update hostNode params using hipGraphExecHostNodeSetParams API + and validates the result +2. Create graph, Add Graphnodes and clones the graph. Add Hostnode to the cloned graph, update + hostNode params using hipGraphExecHostNodeSetParams API and validate the result + +Negative: + +1) Pass hGraphExec as nullptr and verify api doen't crash, returns error code. +2) Pass node as nullptr and verify api doen't crash, returns error code. +3) Pass pNodeParams as nullptr and verify api doesn't crash, returns error code. +3) Pass hipHostNodeParams::hipHostFn_t as nullptr and verify api doesn't crash, returns error code. +4) Pass unintialized host params and verify api doesn't crash, returns error code. +5) Pass unintialized graph and verify api doesn't crash, returns error code. +6) Pass nullptr to hostfunc and verify api doesn't crash, returns error code. +*/ + +#include +#include + +#define SIZE 1024 + +void callbackfunc(void *A_h) { + int *A = reinterpret_cast(A_h); + for (int i = 0; i < SIZE; i++) { + A[i] = i; + } +} + +void callbackfunc_setparams(void *B_h) { + int *B = reinterpret_cast(B_h); + for (int i = 0; i < SIZE; i++) { + B[i] = i * i; + } +} + +/* +This testcase verifies the negative scenarios of +hipGraphExecHostNodeSetParams API +*/ +TEST_CASE("Unit_hipGraphExecHostNodeSetParams_Negative") { +constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC, + &hostNode, 1)); + + hipHostNodeParams sethostParams = {0, 0}; + sethostParams.fn = callbackfunc_setparams; + sethostParams.userData = C_h; + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); +#if HT_NVIDIA + SECTION("Passing nullptr to graphExec") { + REQUIRE(hipGraphExecHostNodeSetParams(nullptr, hostNode, &sethostParams) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to hostParams") { + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, nullptr) + == hipErrorInvalidValue); + } +#endif + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, nullptr, &sethostParams) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to host func") { + sethostParams.fn = nullptr; + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams) + == hipErrorInvalidValue); + } + + + SECTION("Passing unintialized hostParams") { + hipHostNodeParams unintParams = {0, 0}; + REQUIRE(hipGraphExecHostNodeSetParams(graphExec, hostNode, &unintParams) + == hipErrorInvalidValue); + } + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies hipGraphExecHostNodeSetParams API in cloned graph +Creates graph, Add graph nodes and clone the graph +Add HostNode to the cloned graph,update the host params using +hipGraphExecHostNodeSetParams API and validates the result +*/ +TEST_CASE("Unit_hipGraphExecHostNodeSetParams_ClonedGraphwithHostNode") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyH2D_C, + memcpyD2H_AC; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + hipGraph_t clonedgraph; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, clonedgraph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + + hipHostNodeParams sethostParams = {0, 0}; + sethostParams.fn = callbackfunc_setparams; + sethostParams.userData = C_h; + + + // Instantiate and launch the cloned graph + HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (C_h[i] != static_cast(i * i)) { + INFO("Validation failed i " << i << "C_h[i] "<< C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* +This testcase verifies the following scenario +Create graph, Adds host node to the graph, +updates the host params using hipGraphExecHostNodeSetParams API +and validates the result +*/ +TEST_CASE("Unit_hipGraphExecHostNodeSetParams_BasicFunc") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph; + hipGraphExec_t graphExec; + int *A_d{nullptr}, *C_d{nullptr}; + int *A_h{nullptr}, *C_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, &C_d, + &A_h, nullptr, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D_A, memcpyD2H_AC, memcpyH2D_C; + hipStream_t streamForGraph; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, + 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_C, graph, nullptr, + 0, C_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_AC, graph, nullptr, + 0, A_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + hipGraphNode_t hostNode; + hipHostNodeParams hostParams = {0, 0}; + hostParams.fn = callbackfunc; + hostParams.userData = A_h; + HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, + nullptr, + 0, &hostParams)); + + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_C, + &memcpyD2H_AC, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H_AC, + &hostNode, 1)); + + hipHostNodeParams sethostParams = {0, 0}; + sethostParams.fn = callbackfunc_setparams; + sethostParams.userData = C_h; + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphExecHostNodeSetParams(graphExec, hostNode, &sethostParams)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify execution result + for (size_t i = 0; i < N; i++) { + if (C_h[i] != static_cast(i * i)) { + INFO("Validation failed i " << i << "C_h[i] "<< C_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, C_d, A_h, nullptr, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} From 6c34fc554de30d7c48e99ddd5a259f59cf655641 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:45:30 +0530 Subject: [PATCH 20/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphAddMemcpyNodeToSymbol API (#2501) Added Negative and functional tests for hipGraphAddMemcpyNodeToSymbol API Change-Id: I2cae4de9977c8d2ecbdd6ec6868514cd6e1c6efd --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphAddMemcpyNodeToSymbol.cc | 400 ++++++++++++++++++ 2 files changed, 401 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 48340d488a..a87c229549 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -32,6 +32,7 @@ set(TEST_SRC hipGraphChildGraphNodeGetGraph.cc hipGraphNodeFindInClone.cc hipGraphExecHostNodeSetParams.cc + hipGraphAddMemcpyNodeToSymbol.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc new file mode 100644 index 0000000000..5a5b08f77e --- /dev/null +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc @@ -0,0 +1,400 @@ +/* +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, 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. +*/ + +/** +Testcase Scenarios of hipGraphAddMemcpyNodeToSymbol API: + +Functional : + +1. Allocate global symbol memory, add the MemcpyNodeToSymbol + node to the graph and verify for different memory kinds +2. Allocate const memory add the MemcpyNodeToSymbol node to + the graph and verify for different memory kinds +3. Allocate global symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +4. Allocate const symbol memory and device memory in GPU-0 + and perform MemcpyToSymbol from peer GPU by adding it to the graph node. +5. Allocate global memory, Add MemcpyToSymbolNode,KernelNode and memcpynode and validating + the behaviour + +Negative : + +1) Pass nullptr to graph node +2) Pass nullptr to graph +3) Pass nullptr to dependencies +4) Pass invalid numDependencies +5) Pass nullptr to dst +6) Pass nullptr to symbol +7) Pass invalid count +8) Pass offset+count greater than allocated size +9) Pass unintialized graph +*/ + +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + +__global__ void MemcpyToSymbolKernel(int* B_d) { + for (int i = 0 ; i < SIZE; i++) { + B_d[i] = globalIn[i]; + } +} + +/* This testcase verifies negative scenarios of + hipGraphAddMemcpyNodeToSymbol API */ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphNode_t memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); +#if HT_NVIDIA + hipGraphNode_t memcpyToSymbolNode; + SECTION("Passing nullptr to graph") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, nullptr, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_h, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to graph node") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(nullptr, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing size > 1 and dependencies as nullptr") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + nullptr, + 1, + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing invalid dependencies size") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + 10, + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing nullptr to dst") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + nullptr, + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidSymbol); + } + + SECTION("Passing nullptr to source") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing offset+size > max size") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 10, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Passing Max count") { + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, + std::numeric_limits::max(), 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } + + SECTION("Pass Unintialized graph") { + hipGraph_t unint_graph; + REQUIRE(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, unint_graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, + Nbytes, 0, + hipMemcpyDeviceToDevice) + == hipErrorInvalidValue); + } +#endif + + HipTest::freeArrays(A_d, nullptr, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This function is used to verify the following scenarios +1. Create global variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeToSymbol API in GPU-1 and validate the result +2. Allocate global memory, Create dependency graph and validate the result on GPU-0 +3. Allocate global const memory, Create dependency graph and validate the result on GPU-0 +4. Create global const variable, allocate Memory in GPU-0 and create dependency graph of + hipGraphAddMemcpyNodeToSymbol API in GPU-1 and validate the result +*/ +void hipGraphAddMemcpyNodeToSymbol_GlobalMemory(bool device_ctxchg = false, + bool const_device_var = false) { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, nullptr, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + if (device_ctxchg) { + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipDeviceEnablePeerAccess(0, 0)); + } + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + // Adding MemcpyNodeToSymbol + + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalConst), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + } + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + // Adding MemcpyNodeFromSymbol + if (const_device_var) { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalConst), + Nbytes, 0, hipMemcpyDeviceToHost)); + } else { + HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph, + dependencies.data(), + dependencies.size(), + B_h, + HIP_SYMBOL(globalIn), + Nbytes, 0, hipMemcpyDeviceToHost)); + } + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, nullptr, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} +/* +This testcase verifies allocating global symbol memory, +add the MemcpyNodeToSymbol node to the graph and +erifying the result +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemory") { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(false, false); +} + +/* +This testcase verifies allocating global const symbol memory, +add the MemcpyNodeToSymbol node to the graph and +verifying the result +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemory") { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(false, true); +} + +#if HT_NVIDIA +/* +This testcase verifies allocating global symbol memory and device variables +in GPU-0 and add the MemcpyNodeToSymbol node to the graph and +verifying the result in GPU-1 +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + if (canAccessPeer) { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(true, false); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +/* +This testcase verifies allocating global const symbol memory and device variables +in GPU-0 and add the MemcpyNodeToSymbol node to the graph and +verifying the result in GPU-1 +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemoryPeerDevice") { + int numDevices = 0; + int canAccessPeer = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + if (canAccessPeer) { + hipGraphAddMemcpyNodeToSymbol_GlobalMemory(true, true); + } else { + SUCCEED("Machine does not seem to have P2P"); + } + } else { + SUCCEED("skipped the testcase as no of devices is less than 2"); + } +} +#endif +/* +This testcaser verifies allocating global memory, +Add MemcpyToSymbolNode,KernelNode and memcpynode and validating +the behaviour +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel") { + constexpr size_t Nbytes = SIZE * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, SIZE); + hipGraphNode_t memcpytosymbolkernel, memcpyD2H_B; + hipKernelNodeParams kernelNodeParams{}; + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + dependencies.clear(); + dependencies.push_back(memcpyToSymbolNode); + + // Adding Kernel node + void* kernelArgs1[] = {&B_d}; + kernelNodeParams.func = + reinterpret_cast(MemcpyToSymbolKernel); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&memcpytosymbolkernel, graph, + dependencies.data(), dependencies.size(), + &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(memcpytosymbolkernel); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(), + dependencies.size(), B_h, B_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, 0)); + + // Validating the result + for (int i = 0; i < SIZE; i++) { + if (B_h[i] != A_h[i]) { + WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + + HipTest::freeArrays(A_d, B_d, nullptr, + A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} From 411110ebc9e87ea3c736085b092700f3590092f5 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:46:30 +0530 Subject: [PATCH 21/30] SWDEV-306122 - [catch2][dtest] hipGraphExecMemsetNodeSetParams added Functional and Negative test scenario (#2500) Change-Id: Ia6cc644e59985bdee8a7b5ebccf6ced711dcfc7b --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphExecMemsetNodeSetParams.cc | 201 ++++++++++++++++++ 2 files changed, 202 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index a87c229549..34fbb28daf 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -33,6 +33,7 @@ set(TEST_SRC hipGraphNodeFindInClone.cc hipGraphExecHostNodeSetParams.cc hipGraphAddMemcpyNodeToSymbol.cc + hipGraphExecMemsetNodeSetParams.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc b/tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc new file mode 100644 index 0000000000..7e51eb296b --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecMemsetNodeSetParams.cc @@ -0,0 +1,201 @@ +/* +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. +*/ + +/** +Testcase Scenarios : +Functional- +1) Instantiate a graph with memset node, obtain executable graph and update the + hipMemsetParams node params with set. Make sure they are taking effect. +Negative- +1) Pass hGraphExec as nullptr and verify api returns error code. +2) Pass graph node as nullptr and verify api returns error code. +3) Pass different hipGraphNode_t which was not used in graphExec and verify api returns error code. +4) Pass Pass different Graph which was not used in graphExec and verify api returns error code. +5) Pass pNodeParams as nullptr and verify api returns error code. +6) Pass pNodeParams as empty structure object and verify api returns error code. +7) Pass hipMemsetParams::dst as nullptr, api should return error code. +8) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code. +9) Pass hipMemsetParams::height as zero and check api should return error code. +*/ + +#include +#include + +/* Test verifies hipGraphExecMemsetNodeSetParams API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + char *devData, *hOutputData; + + HIP_CHECK(hipMalloc(&devData, Nbytes)); + hOutputData = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, 0, Nbytes); + + hipGraph_t graph; + hipError_t ret; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams mParams{}; + memset(&mParams, 0, sizeof(mParams)); + mParams.dst = reinterpret_cast(devData); + mParams.value = val; + mParams.pitch = 0; + mParams.elementSize = sizeof(char); + mParams.width = Nbytes; + mParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &mParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + SECTION("Pass hGraphExec as nullptr") { + ret = hipGraphExecMemsetNodeSetParams(nullptr, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hGraphNode as nullptr") { + ret = hipGraphExecMemsetNodeSetParams(graphExec, nullptr, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass different hGraphNode which was not used in graphExec") { + hipGraphNode_t memsetNode1{}; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode1, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass different Graph which was not used in graphExec") { + hipGraph_t graph1; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph1, nullptr, 0, &mParams)); + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK(hipGraphDestroy(graph1)); + } + SECTION("Pass pNodeParams as nullptr") { + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_NVIDIA + SECTION("Pass pNodeParams as empty structure object") { + hipMemsetParams mParmTemp{}; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParmTemp); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + SECTION("Pass hipMemsetParams::dst as nullptr") { + mParams.dst = nullptr; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_NVIDIA + SECTION("Pass hipMemsetParams::element size other than 1, 2, or 4") { + mParams.dst = reinterpret_cast(devData); + mParams.elementSize = 9; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass hipMemsetParams::height as zero") { + mParams.elementSize = sizeof(char); + mParams.height = 0; + ret = hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, &mParams); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + + free(hOutputData); + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} + +/* Test verifies hipGraphExecMemsetNodeSetParams API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphExecMemsetNodeSetParams_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(char); + constexpr size_t val = 0; + constexpr size_t updateVal = 2; + char *devData, *devData1, *hOutputData, *hOutputData1; + + HIP_CHECK(hipMalloc(&devData, Nbytes)); + HIP_CHECK(hipMalloc(&devData1, Nbytes)); + hOutputData = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hOutputData != nullptr); + memset(hOutputData, updateVal, Nbytes); + hOutputData1 = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hOutputData1 != nullptr); + memset(hOutputData1, 0, Nbytes); + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + hipGraphNode_t memsetNode; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemsetParams memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData); + memsetParams.value = val; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + + std::vector dependencies; + dependencies.push_back(memsetNode); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(devData1); + memsetParams.value = updateVal; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + + REQUIRE(hipSuccess == hipGraphExecMemsetNodeSetParams(graphExec, memsetNode, + &memsetParams)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipMemcpy(hOutputData1, devData1, Nbytes, hipMemcpyDeviceToHost)); + HipTest::checkArray(hOutputData, hOutputData1, Nbytes, 1); + + free(hOutputData); + free(hOutputData1); + HIP_CHECK(hipFree(devData)); + HIP_CHECK(hipFree(devData1)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); +} From e30b950817f79dd7d9f6aeb707f4a74df70b773e Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:47:07 +0530 Subject: [PATCH 22/30] SWDEV-316467 - tests to verify hipGraphMemcpyNodeSetParamsToSymbol negative scenarios (#2499) Change-Id: I38c7fe6ea060ce8f4c2ffe3ddd6ef927da2e864a --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../hipGraphMemcpyNodeSetParamsToSymbol.cc | 107 ++++++++++++++++++ 2 files changed, 108 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 34fbb28daf..aa23776294 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -34,6 +34,7 @@ set(TEST_SRC hipGraphExecHostNodeSetParams.cc hipGraphAddMemcpyNodeToSymbol.cc hipGraphExecMemsetNodeSetParams.cc + hipGraphMemcpyNodeSetParamsToSymbol.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc new file mode 100644 index 0000000000..804c502a0c --- /dev/null +++ b/tests/catch/unit/graph/hipGraphMemcpyNodeSetParamsToSymbol.cc @@ -0,0 +1,107 @@ +#include +#include +#include +#define SIZE 256 + +__device__ int globalIn[SIZE], globalOut[SIZE]; +__device__ __constant__ int globalConst[SIZE]; + + +/* This testcase verifies negative scenarios of + hipGraphMemcpyNodeSetParamsToSymbol API */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParamsToSymbol_Negative") { + constexpr size_t Nbytes = SIZE * sizeof(int); + int *A_d{nullptr}, *B_d{nullptr}; + int *A_h{nullptr}, *B_h{nullptr}; + HipTest::initArrays(&A_d, &B_d, nullptr, + &A_h, &B_h, nullptr, SIZE, false); + + hipGraph_t graph; + hipError_t ret; + hipGraphNode_t memcpyToSymbolNode, memcpyH2D_A; + std::vector dependencies; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Adding MemcpyNode + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyH2D_A); + + HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph, + dependencies.data(), + dependencies.size(), + HIP_SYMBOL(globalIn), + A_d, Nbytes, 0, + hipMemcpyDeviceToDevice)); + + SECTION("Pass GraphNode as nullptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(nullptr, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass symbol ptr as nullptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + nullptr, + B_d, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidSymbol == ret); + } + SECTION("Pass src ptr as nullptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + nullptr, Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count as zero") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, 0, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + + SECTION("Pass count more than allocated size for source and dstn ptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, Nbytes+8, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass offset+count greater than allocated size") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + B_d, Nbytes, 10, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass same symbol pointer as source ptr and destination ptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + HIP_SYMBOL(globalIn), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass 2 different symbol pointer as source ptr and dstn ptr") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + HIP_SYMBOL(globalOut), + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Copy from host ptr to device ptr but pass kind as different") { + ret = hipGraphMemcpyNodeSetParamsToSymbol(memcpyToSymbolNode, + HIP_SYMBOL(globalIn), + A_h, + Nbytes, 0, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + + HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); + HIP_CHECK(hipGraphDestroy(graph)); +} \ No newline at end of file From c093f29d32d0a015def353453f8d00556ea30bf8 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:55:21 +0530 Subject: [PATCH 23/30] SWDEV-312141 - Renaming of enum to avoid shadow variable compilation warning (#2487) Change-Id: I0579d8b3e4175f14ad28f9f4b263fa0a792a9120 --- include/hip/hip_bfloat16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hip_bfloat16.h b/include/hip/hip_bfloat16.h index 642090b4f5..eb99206632 100644 --- a/include/hip/hip_bfloat16.h +++ b/include/hip/hip_bfloat16.h @@ -56,7 +56,7 @@ struct hip_bfloat16 enum truncate_t { - truncate + truncate_0 }; __host__ __device__ hip_bfloat16() = default; From 610b859b348dfc19942f0b832ec2ab87e35962b1 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Thu, 24 Feb 2022 20:57:04 +0530 Subject: [PATCH 24/30] SWDEV-306940 - Fix hip errors/warnings by CppCheck (#2486) Change-Id: I91f5e97b0c7ca80b6d21fa66d24e3cf9fde716a9 --- samples/0_Intro/module_api/launchKernelHcc.cpp | 3 --- samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp | 2 +- samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 4 ++-- samples/1_Utils/hipCommander/ResultDatabase.cpp | 2 +- samples/1_Utils/hipCommander/hipCommander.cpp | 12 ++++++------ samples/2_Cookbook/14_gpu_arch/gpuarch.cpp | 11 ++++------- .../15_static_library/device_functions/hipMain2.cpp | 4 ++-- 7 files changed, 16 insertions(+), 22 deletions(-) diff --git a/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp index 39d3bd31d2..464f6d8851 100644 --- a/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -69,9 +69,6 @@ int main() { HIP_CHECK(hipModuleLoad(&Module, fileName)); HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); - uint32_t len = LEN; - uint32_t one = 1; - struct { void* _Ad; void* _Bd; diff --git a/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp b/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp index 581f742834..e094f70d07 100644 --- a/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp +++ b/samples/1_Utils/hipBusBandwidth/ResultDatabase.cpp @@ -401,7 +401,6 @@ void ResultDatabase::DumpCsv(string fileName) { // **************************************************************************** bool ResultDatabase::IsFileEmpty(string fileName) { - bool fileEmpty; ifstream file(fileName.c_str()); @@ -409,6 +408,7 @@ bool ResultDatabase::IsFileEmpty(string fileName) { if (!file.good()) { return true; } else { + bool fileEmpty; fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof()); file.close(); diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 8032bd0a20..79d80ce317 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -252,9 +252,9 @@ void RunBenchmark_H2D(ResultDatabase& resultDB) { case MallocUnpinned: if (p_alignedhost) { - delete[] hostMem; - } else { free(hostMem); + } else { + delete[] hostMem; } break; diff --git a/samples/1_Utils/hipCommander/ResultDatabase.cpp b/samples/1_Utils/hipCommander/ResultDatabase.cpp index 1b1ee3a70d..51ced81fae 100644 --- a/samples/1_Utils/hipCommander/ResultDatabase.cpp +++ b/samples/1_Utils/hipCommander/ResultDatabase.cpp @@ -393,7 +393,6 @@ void ResultDatabase::DumpCsv(string fileName) { // **************************************************************************** bool ResultDatabase::IsFileEmpty(string fileName) { - bool fileEmpty; ifstream file(fileName.c_str()); @@ -401,6 +400,7 @@ bool ResultDatabase::IsFileEmpty(string fileName) { if (!file.good()) { return true; } else { + bool fileEmpty; fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof()); file.close(); diff --git a/samples/1_Utils/hipCommander/hipCommander.cpp b/samples/1_Utils/hipCommander/hipCommander.cpp index 37eb0845b1..e95ecd82a1 100644 --- a/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/samples/1_Utils/hipCommander/hipCommander.cpp @@ -282,7 +282,7 @@ class Command { // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: class ModuleKernelCommand : public Command { public: - ModuleKernelCommand(CommandStream* cmdStream, const std::vector args) + ModuleKernelCommand(CommandStream* cmdStream, const std::vector& args) : Command(cmdStream, args), _stream(cmdStream->currentStream()) { hipModule_t module; HIPCHECK(hipModuleLoad(&module, FILENAME)); @@ -316,7 +316,7 @@ class ModuleKernelCommand : public Command { class KernelCommand : public Command { public: enum Type { Null, VectorAdd }; - KernelCommand(CommandStream* cmdStream, const std::vector args, Type kind) + KernelCommand(CommandStream* cmdStream, const std::vector& args, Type kind) : Command(cmdStream, args), _kind(kind), _stream(cmdStream->currentStream()){}; ~KernelCommand(){}; @@ -390,7 +390,7 @@ class CopyCommand : public Command { }; - void dealloc(void* p, MemType memType) { + static void dealloc(void* p, MemType memType) { if (memType == Device) { HIPCHECK(hipFree(p)); } else if (memType == PinnedHost) { @@ -433,7 +433,7 @@ class StreamSyncCommand : public Command { StreamSyncCommand(CommandStream* cmdStream, const std::vector& args) : Command(cmdStream, args), _stream(cmdStream->currentStream()){}; - const char* help() { return "synchronizes the current stream"; }; + static const char* help() { return "synchronizes the current stream"; }; void run() override { HIPCHECK(hipStreamSynchronize(_stream)); }; @@ -537,8 +537,8 @@ CopyCommand::CopyCommand(CommandStream* cmdStream, const std::vectorcurrentStream()) { + _stream(cmdStream->currentStream()), + _kind(kind) { switch (kind) { case hipMemcpyDeviceToHost: _srcType = Device; diff --git a/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp b/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp index ca312c0146..b4c8487b67 100644 --- a/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp +++ b/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp @@ -48,17 +48,14 @@ __global__ void incrementKernel(int32_t* in, int32_t* out, int32_t value, size_t int main() { int32_t incrementValue = 10; - // Host pointers - int32_t* hInput = nullptr; - int32_t* hOutput = nullptr; // Device pointers int32_t* dInput = nullptr; int32_t* dOutput = nullptr; size_t NBytes = SIZE * sizeof(int32_t); - - hInput = static_cast(malloc(NBytes)); - hOutput = static_cast(malloc(NBytes)); + // Host pointers + int32_t* hInput = static_cast(malloc(NBytes)); + int32_t* hOutput = static_cast(malloc(NBytes)); HIP_STATUS_CHECK(hipMalloc(&dInput, NBytes)); HIP_STATUS_CHECK(hipMalloc(&dOutput, NBytes)); @@ -95,4 +92,4 @@ int main() { std::cout << "success\n"; } return 0; -} \ No newline at end of file +} diff --git a/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp b/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp index a93eb141ac..a3c3f8f164 100644 --- a/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp +++ b/samples/2_Cookbook/15_static_library/device_functions/hipMain2.cpp @@ -58,8 +58,8 @@ void run_test2() { HIP_ASSERT(hipFree(A_d)); HIP_ASSERT(hipFree(B_d)); - free(A_h); - free(B_h); + delete [] A_h; + delete [] B_h; std::cout << "Test Passed!\n"; } From f86a8da09e20be70add015f0ba81b8890c822f96 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 25 Feb 2022 14:04:23 +0530 Subject: [PATCH 25/30] SWDEV-306122 - [catch2][dtest] Tests for hipGraphDestroyNode API (#2509) Added negative and functional tests for hipGraphDestroyNode API Change-Id: Ifc4183b94e8fd428c44f2bbc3016539f972b8069 --- tests/catch/unit/graph/CMakeLists.txt | 1 + tests/catch/unit/graph/hipGraphDestroyNode.cc | 139 ++++++++++++++++++ 2 files changed, 140 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphDestroyNode.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index aa23776294..8f0857c2c9 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -35,6 +35,7 @@ set(TEST_SRC hipGraphAddMemcpyNodeToSymbol.cc hipGraphExecMemsetNodeSetParams.cc hipGraphMemcpyNodeSetParamsToSymbol.cc + hipGraphDestroyNode.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphDestroyNode.cc b/tests/catch/unit/graph/hipGraphDestroyNode.cc new file mode 100644 index 0000000000..4c544a3f68 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphDestroyNode.cc @@ -0,0 +1,139 @@ +/* +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. +*/ + +/* +Testcase Scenarios of hipGraphDestroyNode API: + +Negative :: + +1) Pass nullptr to graph node + +Functional :: + +1) Create Node and destroy the node +2) Create graph with dependencies and destroy one of the dependency node + before executing the graph. +*/ + +#include +#include +#include + + +/* This test covers the negative scenarios of + hipGraphDestroyNode API */ +TEST_CASE("Unit_hipGraphDestroyNode_Negative") { + SECTION("Passing nullptr to graph Node") { + REQUIRE(hipGraphDestroyNode(nullptr) == hipErrorInvalidValue); + } +} + +/* This test covers the basic functionality of + hipGraphDestroyNode API where we create and destroy + the node +*/ +TEST_CASE("Unit_hipGraphDestroyNode_BasicFunctionality") { + char *pOutBuff_d{}; + constexpr size_t size = 1024; + hipGraph_t graph{}; + hipGraphNode_t memsetNode{}; + + HIP_CHECK(hipMalloc(&pOutBuff_d, size)); + hipMemsetParams memsetParams{}; + memsetParams.dst = reinterpret_cast(pOutBuff_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = size * sizeof(char); + memsetParams.height = 1; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams)); + REQUIRE(hipGraphDestroyNode(memsetNode) == hipSuccess); + HIP_CHECK(hipFree(pOutBuff_d)); +} + +/* +This testcase verifies the following scenario where +graph is created with dependencies and one of the dependency is +destroyed before execute the graph +*/ +TEST_CASE("Unit_hipGraphDestroyNode_DestroyDependencyNode") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + hipGraph_t graph; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyH2D_B2Copies, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + hipGraphExec_t graphExec; + size_t NElem{N}; + hipStream_t streamForGraph; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B2Copies, graph, nullptr, + 0, B_d, C_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, B_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B2Copies, &kernel_vecAdd, + 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + // Destroy one of the dependency node + HIP_CHECK(hipGraphDestroyNode(memcpyH2D_B)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, C_h, B_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); +} From daf3f429cdef512a983c5bd58b2eaf37e16fe4b0 Mon Sep 17 00:00:00 2001 From: lthakur Date: Fri, 25 Feb 2022 14:05:06 +0530 Subject: [PATCH 26/30] SWDEV-314080 - Tested All hipMemcpy**() apis with hipStreamPerThread stream obj (#2484) Change-Id: I8f429eb0cc3be2e4d62c76ccb8c1510c56a1e143 From adaa9ba27e6094447cdda26e7488520df59927db Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Fri, 25 Feb 2022 00:35:37 -0800 Subject: [PATCH 27/30] SWDEV-314665 - hipGraphInstantiateFlags enum (#2481) Change-Id: If07e415ee18f08bad85ae5e491e01f97b16a6b99 --- include/hip/hip_runtime_api.h | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index b8b511d816..f2a24e6dbe 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -4304,6 +4304,11 @@ typedef enum hipStreamUpdateCaptureDependenciesFlags { hipStreamSetCaptureDependencies, ///< Replace the dependency set with the new nodes } hipStreamUpdateCaptureDependenciesFlags; +typedef enum hipGraphInstantiateFlags { + hipGraphInstantiateFlagAutoFreeOnLaunch = + 1, ///< Automatically free memory allocated in a graph before relaunching. +} hipGraphInstantiateFlags; + /** * @brief Begins graph capture on a stream. * From 279e050a07ca065b8777380e691a1bb36853df76 Mon Sep 17 00:00:00 2001 From: sumanthtg <90063301+sumanthtg@users.noreply.github.com> Date: Fri, 25 Feb 2022 14:07:38 +0530 Subject: [PATCH 28/30] SWDEV-306122 - [catch2][dtest] hipGraph functional tests for GetNodes/GetRootNotes (#2441) Change-Id: I9b31252ddb5259b2e7c2bc7502e0494a8041ed9e --- tests/catch/unit/graph/CMakeLists.txt | 2 + tests/catch/unit/graph/hipGraphGetNodes.cc | 133 ++++++++++++++++++ .../catch/unit/graph/hipGraphGetRootNodes.cc | 133 ++++++++++++++++++ 3 files changed, 268 insertions(+) create mode 100644 tests/catch/unit/graph/hipGraphGetNodes.cc create mode 100644 tests/catch/unit/graph/hipGraphGetRootNodes.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 8f0857c2c9..afcfb2028a 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -36,6 +36,8 @@ set(TEST_SRC hipGraphExecMemsetNodeSetParams.cc hipGraphMemcpyNodeSetParamsToSymbol.cc hipGraphDestroyNode.cc + hipGraphGetNodes.cc + hipGraphGetRootNodes.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphGetNodes.cc b/tests/catch/unit/graph/hipGraphGetNodes.cc new file mode 100644 index 0000000000..c99863f037 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphGetNodes.cc @@ -0,0 +1,133 @@ +/* +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, 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. +*/ + +/** +Testcase Scenarios : +Functional :: +1) Add nodes to graph and get nodes. Verify the added nodes are present in returned list. +2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph. +3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes +will be set to NULL, and the number of nodes actually obtained will be returned in numNodes. + +*/ + +#include +#include +#include + +/** + * Functional Test for API fetching node list + */ +TEST_CASE("Unit_hipGraphGetNodes_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + constexpr auto addlEntries = 4; + hipGraph_t graph; + hipGraphNode_t memcpyNode, kernelNode; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies, nodelist; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + nodelist.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + nodelist.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, dependencies.data(), + dependencies.size(), &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(kernelNode); + nodelist.push_back(kernelNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + nodelist.push_back(memcpyNode); + + // Get numNodes by passing nodes as nullptr. + // verify : numNodes is set to actual number of nodes added + size_t numNodes{}; + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + INFO("Num of nodes returned by GetNodes : " << numNodes); + REQUIRE(numNodes == nodelist.size()); + + // Request for extra/additional nodes. + // verify : totNodes is reset to actual number of nodes + // verify : additional entries in nodes are set to nullptr + size_t totNodes = numNodes + addlEntries; + int numBytes = sizeof(hipGraphNode_t) * totNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + HIP_CHECK(hipGraphGetNodes(graph, nodes, &totNodes)); + REQUIRE(totNodes == nodelist.size()); + for (auto i = numNodes; i < numNodes + addlEntries; i++) { + REQUIRE(nodes[i] == nullptr); + } + + // Verify added nodes are present in the node entries returned + for (auto Node : nodelist) { + bool found = false; + for (size_t i = 0; i < numNodes; i++) { + if (Node == nodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Added node " << Node << " not present in returned list"); + REQUIRE(false); + } + } + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + free(nodes); +} diff --git a/tests/catch/unit/graph/hipGraphGetRootNodes.cc b/tests/catch/unit/graph/hipGraphGetRootNodes.cc new file mode 100644 index 0000000000..7238dd186b --- /dev/null +++ b/tests/catch/unit/graph/hipGraphGetRootNodes.cc @@ -0,0 +1,133 @@ +/* +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, 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. +*/ + +/** +Testcase Scenarios : +Functional :: + 1) Add nodes to graph with and without dependencies, verify the api returns list of + root nodes (i.e., nodes without dependencies). + 2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph. + 3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in + nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes. +*/ + +#include +#include +#include + +/** + * Functional Test for API fetching root node list + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + constexpr auto addlEntries = 5; + hipGraph_t graph; + hipGraphNode_t memcpyNode, kernelNode; + hipKernelNodeParams kernelNodeParams{}; + hipStream_t streamForGraph; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + std::vector dependencies, rootnodelist; + hipGraphExec_t graphExec; + size_t NElem{N}; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + rootnodelist.push_back(memcpyNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, NULL, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + dependencies.push_back(memcpyNode); + rootnodelist.push_back(memcpyNode); + + void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelNode, graph, dependencies.data(), + dependencies.size(), &kernelNodeParams)); + dependencies.clear(); + dependencies.push_back(kernelNode); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), + dependencies.size(), C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + // Get numRootNodes by passing rootnodes list as nullptr. + // verify : numRootNodes is set to actual number of root nodes added + size_t numRootNodes{}; + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + INFO("Num of nodes returned by GetRootNodes : " << numRootNodes); + REQUIRE(numRootNodes == rootnodelist.size()); + + // Request for extra/additional nodes. + // verify : totNodes is reset to actual number of root nodes present + // verify : additional entries in rootnodes list are set to nullptr + size_t totNodes = numRootNodes + addlEntries; + int numBytes = sizeof(hipGraphNode_t) * totNodes; + hipGraphNode_t* rootnodes = + reinterpret_cast(malloc(numBytes)); + REQUIRE(rootnodes != nullptr); + HIP_CHECK(hipGraphGetRootNodes(graph, rootnodes, &totNodes)); + REQUIRE(totNodes == rootnodelist.size()); + for (auto i = numRootNodes; i < numRootNodes + addlEntries; i++) { + REQUIRE(rootnodes[i] == nullptr); + } + + // Verify added nodes(without dependencies) are present + // in the root nodes fetched. + for (auto Node : rootnodelist) { + bool found = false; + for (size_t i = 0; i < numRootNodes; i++) { + if (Node == rootnodes[i]) { + found = true; + break; + } + } + + if (!found) { + INFO("Returned root node " << Node << " not present in added list"); + REQUIRE(false); + } + } + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + free(rootnodes); +} From 4ecad956c9df57c42b58859a1c5e4da5da0635e4 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 25 Feb 2022 16:25:51 +0530 Subject: [PATCH 29/30] SWDEV-306122 - [catch2][dtest] hipGraph tests for GetNodes/GetRootNodes (#2497) Change-Id: I5e9fafd0802585b6b33c391d7281151010ffc1c3 --- tests/catch/unit/graph/hipGraphGetNodes.cc | 92 ++++++++++++++- .../catch/unit/graph/hipGraphGetRootNodes.cc | 110 +++++++++++++++++- 2 files changed, 197 insertions(+), 5 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphGetNodes.cc b/tests/catch/unit/graph/hipGraphGetNodes.cc index c99863f037..18b3df1722 100644 --- a/tests/catch/unit/graph/hipGraphGetNodes.cc +++ b/tests/catch/unit/graph/hipGraphGetNodes.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -18,13 +18,20 @@ THE SOFTWARE. */ /** -Testcase Scenarios : +Testcase Scenarios +------------------ Functional :: 1) Add nodes to graph and get nodes. Verify the added nodes are present in returned list. 2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph. 3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes will be set to NULL, and the number of nodes actually obtained will be returned in numNodes. +Argument Validation :: +1) Pass graph as nullptr and verify api returns error code. +2) Pass numNodes as nullptr and other params as valid values. Expect api to return error code. +3) When there are no nodes in graph, expect numNodes to be set to zero. +4) Pass numNodes less than actual number of nodes. Expect api to populate requested number of node entries +and does update numNodes. */ #include @@ -32,7 +39,7 @@ will be set to NULL, and the number of nodes actually obtained will be returned #include /** - * Functional Test for API fetching node list + * Functional Test for hipGraphGetNodes API fetching node list */ TEST_CASE("Unit_hipGraphGetNodes_Functional") { constexpr size_t N = 1024; @@ -131,3 +138,82 @@ TEST_CASE("Unit_hipGraphGetNodes_Functional") { HIP_CHECK(hipStreamDestroy(streamForGraph)); free(nodes); } + +/** + * Test performs api parameter validation by passing various values + * as input and output parameters and validates the behavior. + * Test will include both negative and positive scenarios. + */ +TEST_CASE("Unit_hipGraphGetNodes_ParamValidation") { + hipStream_t stream{nullptr}; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + size_t Nbytes = N * sizeof(float), numNodes{}; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + INFO("Num of nodes returned by GetNodes : " << numNodes); + + int numBytes = sizeof(hipGraphNode_t) * numNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + SECTION("graph as nullptr") { + hipError_t ret = hipGraphGetNodes(nullptr, nodes, &numNodes); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("numNodes as nullptr") { + hipError_t ret = hipGraphGetNodes(graph, nodes, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("no nodes in graph") { + hipGraph_t emptyGraph{}; + HIP_CHECK(hipGraphCreate(&emptyGraph, 0)); + HIP_CHECK(hipGraphGetNodes(emptyGraph, nullptr, &numNodes)); + REQUIRE(numNodes == 0); + } + + SECTION("numNodes less than actual number of nodes") { + size_t numPartNodes = numNodes - 1; + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphGetNodes(graph, nodes, &numPartNodes)); + + // verify numPartNodes is unchanged + REQUIRE(numPartNodes == numNodes - 1); + // verify partial node list returned has valid nodes + for (size_t i = 0; i < numPartNodes; i++) { + HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType)); + REQUIRE(nodeType >= 0); + REQUIRE(nodeType < hipGraphNodeTypeCount); + } + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} diff --git a/tests/catch/unit/graph/hipGraphGetRootNodes.cc b/tests/catch/unit/graph/hipGraphGetRootNodes.cc index 7238dd186b..c2484fc8b5 100644 --- a/tests/catch/unit/graph/hipGraphGetRootNodes.cc +++ b/tests/catch/unit/graph/hipGraphGetRootNodes.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -18,13 +18,21 @@ THE SOFTWARE. */ /** -Testcase Scenarios : +Testcase Scenarios +------------------ Functional :: 1) Add nodes to graph with and without dependencies, verify the api returns list of root nodes (i.e., nodes without dependencies). 2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph. 3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes. + +Argument Validation :: + 1) Pass graph as nullptr and verify api returns error code. + 2) Pass numRootNodes as nullptr and other params as valid values. Expect api to return error code. + 3) When there are no nodes in graph, expect numRootNodes to be set to zero. + 4) Pass numRootNodes less than actual number of nodes. Expect api to populate requested number of node entries + and does update numRootNodes. */ #include @@ -41,6 +49,7 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { constexpr auto threadsPerBlock = 256; constexpr auto addlEntries = 5; hipGraph_t graph; + hipGraphNode_t memcpyNode, kernelNode; hipKernelNodeParams kernelNodeParams{}; hipStream_t streamForGraph; @@ -131,3 +140,100 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { HIP_CHECK(hipStreamDestroy(streamForGraph)); free(rootnodes); } + +/** + * Test performs api parameter validation by passing various values + * as input and output parameters and validates the behavior. + * Test will include both negative and positive scenarios. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") { + hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr}; + hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + size_t Nbytes = N * sizeof(float), numRootNodes{}; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&mstream)); + HIP_CHECK(hipEventCreate(&memsetEvent1)); + HIP_CHECK(hipEventCreate(&memsetEvent2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, mstream)); + HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1)); + HIP_CHECK(hipEventRecord(memsetEvent1, stream1)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream2)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, mstream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + INFO("Num of nodes returned by GetRootNodes : " << numRootNodes); + int numBytes = sizeof(hipGraphNode_t) * numRootNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + SECTION("graph as nullptr") { + hipError_t ret = hipGraphGetRootNodes(nullptr, nodes, &numRootNodes); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("numRootNodes as nullptr") { + hipError_t ret = hipGraphGetRootNodes(graph, nodes, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("no nodes in graph") { + hipGraph_t emptyGraph{}; + HIP_CHECK(hipGraphCreate(&emptyGraph, 0)); + HIP_CHECK(hipGraphGetRootNodes(emptyGraph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == 0); + } + + SECTION("numRootNodes less than actual number of nodes") { + size_t numPartNodes = numRootNodes - 1; + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numPartNodes)); + + // verify numPartNodes is unchanged + REQUIRE(numPartNodes == numRootNodes - 1); + // verify partial node list returned has valid nodes + for (size_t i = 0; i < numPartNodes; i++) { + HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType)); + REQUIRE(nodeType >= 0); + REQUIRE(nodeType < hipGraphNodeTypeCount); + } + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(mstream)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memsetEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} From 5eade5cfa104e0bc112bbd854aa96bbafe220753 Mon Sep 17 00:00:00 2001 From: Sourabh U Betigeri Date: Fri, 25 Feb 2022 02:56:20 -0800 Subject: [PATCH 30/30] SWDEV-294375 - Adding coalesced_groups HIP tests (#2489) Change-Id: I000a402e32d6057bab11829af2b78931fcd67fb5 --- .../coalesced_groups_shfl_down.cpp | 276 +++++++++ .../coalesced_groups_shfl_up.cpp | 260 ++++++++ .../simple_coalesced_groups.cpp | 583 ++++++++++++++++++ 3 files changed, 1119 insertions(+) create mode 100644 tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp create mode 100644 tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp create mode 100644 tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp diff --git a/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp new file mode 100644 index 0000000000..aa7b00e290 --- /dev/null +++ b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_down.cpp @@ -0,0 +1,276 @@ +/* +Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test Description: +/* This test implements sum reduction kernel, first with each threads own rank + as input and comparing the sum with expected sum output derieved from n(n-1)/2 + formula. + This sample tests functionality of intrinsics provided by thread_block_tile type, + shfl_down and shfl_xor. +*/ + +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) +#define WAVE_SIZE 32 + +__device__ int reduction_kernel_shfl_down(coalesced_group const& g, int val) { + int sz = g.size(); + + for (int i = sz / 2; i > 0; i >>= 1) { + val += g.shfl_down(val, i); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +__global__ void kernel_shfl_down (int * dPtr, int *dResults, int lane_delta, int cg_sizes) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + + if (id % cg_sizes == 0) { + coalesced_group const& g = coalesced_threads(); + int rank = g.thread_rank(); + int val = dPtr[rank]; + dResults[rank] = g.shfl_down(val, lane_delta); + return; + } +} + +__global__ void kernel_cg_group_partition(int* result, unsigned int tileSz, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + int input, outputSum, expectedSum; + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + int threadRank = tiledPartition.thread_rank(); + + input = tiledPartition.thread_rank(); + + // (n-1)(n)/2 + expectedSum = ((tileSz - 1) * tileSz / 2); + + outputSum = reduction_kernel_shfl_down(tiledPartition, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group using shfl_down is %d (expected " + "%d)\n", + tiledPartition.size() - 1, outputSum, expectedSum); + result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum; + } + return; + } +} + +void verifyResults(int* ptr, int expectedResult, int numTiles) { + for (int i = 0; i < numTiles; i++) { + if (ptr[i] != expectedResult) { + printf(" Results do not match! "); + } + } +} + +void compareResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + printf(" results do not match."); + } + } +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +static void test_group_partition(unsigned int tileSz) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = 32; + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz; + int expectedSum = ((tileSz - 1) * tileSz / 2); + int* expectedResult = new int[numTiles]; + + // numTiles = 0 when partitioning is possible. The below statement is to avoid + // out-of-bounds error and still evaluate failure case. + numTiles = (numTiles == 0) ? 1 : numTiles; + + for (int i = 0; i < numTiles; i++) { + expectedResult[i] = expectedSum; + } + + int* dResult = NULL; + int* hResult = NULL; + + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + hipMalloc(&dResult, numTiles * sizeof(int)); + + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dResult, tileSz, i); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + + + hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost); + + verifyResults(hResult, expectedSum, numTiles); + + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + delete[] expectedResult; + + printf("\n...PASSED.\n\n"); + } +} + +static void test_shfl_down() { + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int totalThreads = blockSize * threadsPerBlock; + int group_size = totalThreads / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dResults = NULL; + int lane_delta = rand() % group_size; + std::cout << "Testing coalesced_groups shfl_down with lane_delta " << lane_delta << "and group size " + << WAVE_SIZE << '\n' << std::endl; + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + // Fill up the array + for (int i = 0; i < WAVE_SIZE; i++) { + hPtr[i] = rand() % 1000; + } + + int* cpuResultsArr = (int*)malloc(group_size_in_bytes); + for (int i = 0; i < group_size; i++) { + cpuResultsArr[i] = (i + lane_delta >= group_size) ? hPtr[i] : hPtr[i + lane_delta]; + } + //printf("Array passed to GPU for computation\n"); + //printResults(hPtr, WAVE_SIZE); + hipMalloc(&dPtr, group_size_in_bytes); + hipMalloc(&dResults, group_size_in_bytes); + + hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice); + // Launch Kernel + hipLaunchKernelGGL(kernel_shfl_down, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i); + hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + //printf("GPU results: \n"); + //printResults(hPtr, WAVE_SIZE); + //printf("Printing cpu to be verified array\n"); + //printResults(cpuResultsArr, WAVE_SIZE); + + compareResults(hPtr, cpuResultsArr, group_size_in_bytes); + std::cout << "Results verified!\n"; + + hipFree(hPtr); + hipFree(dPtr); + free(cpuResultsArr); + } +} + + + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + return 0; + } + + // Test shfl_down with random group sizes + for (int i = 0; i < 100; i++) { + test_shfl_down(); + } + + std::cout << "Testing static tiled_partition for different tile sizes using shfl_down" + << std::endl; + + int testNo = 1; + std::vector tileSizes = {2, 4, 8, 16, 32}; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i); + testNo++; + } + + passed(); +} diff --git a/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp new file mode 100644 index 0000000000..62a05ec244 --- /dev/null +++ b/tests/src/runtimeApi/cooperativeGrps/coalesced_groups_shfl_up.cpp @@ -0,0 +1,260 @@ +/* +Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test Description: +/* This test implements prefix sum(scan) kernel, first with each threads own rank + as input and comparing the sum with expected serial summation output on CPU. + + This sample tests functionality of intrinsics provided by coalesced_group, + shfl_up. +*/ +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) +#define WAVE_SIZE 32 +__device__ int prefix_sum_kernel(coalesced_group const& g, int val) { + int sz = g.size(); + for (int i = 1; i < sz; i <<= 1) { + int temp = g.shfl_up(val, i); + + if (g.thread_rank() >= i) { + val += temp; + } + } + return val; +} + +__global__ void kernel_shfl_up (int * dPtr, int *dResults, int lane_delta, int cg_sizes) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + + if (id % cg_sizes == 0) { + coalesced_group g = coalesced_threads(); + int rank = g.thread_rank(); + int val = dPtr[rank]; + dResults[rank] = g.shfl_up(val, lane_delta); + return; + } + +} + +__global__ void kernel_cg_group_partition(int* dPtr, unsigned int tileSz, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + int input, outputSum; + + // we pass its own thread rank as inputs + input = threadBlockCGTy.thread_rank(); + + // Choose a leader thread to print the results + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + + input = tiledPartition.thread_rank(); + + outputSum = prefix_sum_kernel(tiledPartition, input); + + // Update the result array with the corresponsing prefix sum + dPtr[threadBlockCGTy.thread_rank()] = outputSum; + return; + } +} + +void serialScan(int* ptr, int size) { + // Fill up the array + for (int i = 0; i < size; i++) { + ptr[i] = i; + } + + int acc = 0; + for (int i = 0; i < size; i++) { + acc = acc + ptr[i]; + ptr[i] = acc; + } +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +void verifyResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + printf(" Results do not match."); + } + } +} + +static void test_group_partition(unsigned tileSz) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int* hPtr = NULL; + int* dPtr = NULL; + int* cpuPrefixSum = NULL; + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + hipMalloc(&dPtr, arrSize); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr, tileSz, i); + hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + + cpuPrefixSum = new int[tileSz]; + serialScan(cpuPrefixSum, tileSz); + //std::cout << "\nPrefix sum results on CPU\n"; + //printResults(cpuPrefixSum, tileSz); + + //std::cout << "\nPrefix sum results on GPU\n"; + //printResults(hPtr, tileSz); + std::cout << "\n"; + verifyResults(hPtr, cpuPrefixSum, tileSz); + std::cout << "Results verified!\n"; + + delete[] cpuPrefixSum; + hipFree(hPtr); + hipFree(dPtr); + } +} + +static void test_shfl_up() { + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + + int threadsPerBlock = WAVE_SIZE; + int totalThreads = blockSize * threadsPerBlock; + int group_size = totalThreads / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dResults = NULL; + int lane_delta = (rand() % group_size); + + std::cout << "Testing coalesced_groups shfl_up with lane_delta " << lane_delta + << " and group size " << WAVE_SIZE << '\n' << std::endl; + + int arrSize = blockSize * threadsPerBlock * sizeof(int); + + hipHostMalloc(&hPtr, arrSize); + // Fill up the array + for (int i = 0; i < WAVE_SIZE; i++) { + hPtr[i] = rand() % 1000; + } + //printResults(hPtr, WAVE_SIZE); + + int* cpuResultsArr = (int*)malloc(group_size_in_bytes); + for (int i = 0; i < group_size; i++) { + cpuResultsArr[i] = (i <= (lane_delta - 1)) ? hPtr[i] : hPtr[i - lane_delta]; + } + + //printf("Printing cpu results arr\n"); + //printResults(cpuResultsArr, WAVE_SIZE); + + hipMalloc(&dPtr, group_size_in_bytes); + hipMalloc(&dResults, group_size_in_bytes); + + hipMemcpy(dPtr, hPtr, group_size_in_bytes, hipMemcpyHostToDevice); + // Launch Kernel + hipLaunchKernelGGL(kernel_shfl_up, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, dPtr, dResults, lane_delta, i); + hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost); + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + //printf("GPU computation array :\n"); + //printResults(hPtr, WAVE_SIZE); + + verifyResults(hPtr, cpuResultsArr, group_size_in_bytes); + std::cout << "Results verified!\n"; + + hipFree(hPtr); + hipFree(dPtr); + free(cpuResultsArr); + } +} + +int main() { + // Use default device for validating the test + int deviceId; + ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); + hipDeviceProp_t deviceProperties; + ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + + if (!deviceProperties.cooperativeLaunch) { + std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; + if (hip_skip_tests_enabled()) { + return hip_skip_retcode(); + } else { + passed(); + } + return 0; + } + + for (int i = 0; i < 100; i++) { + test_shfl_up(); + } + + std::cout << "Testing coalesced_groups partitioning and shfl_up" << '\n' << std::endl; + + int testNo = 1; + std::vector tileSizes = {2, 4, 8, 16, 32}; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i); + testNo++; + } + passed(); +} + +/* Kogge-Stone algorithm */ \ No newline at end of file diff --git a/tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp b/tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp new file mode 100644 index 0000000000..9f2775b5dd --- /dev/null +++ b/tests/src/runtimeApi/cooperativeGrps/simple_coalesced_groups.cpp @@ -0,0 +1,583 @@ +/* +Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// Test Description: +/* This test implements sum reduction kernel, first with each threads own rank + as input and comparing the sum with expected sum output derieved from n(n-1)/2 + formula. The second part, partitions this parent group into child subgroups + a.k.a tiles using using tiled_partition() collective operation. This can be called + with a static tile size, passed in templated non-type variable-tiled_partition, + or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these + cases. + This test tests functionality of cg group partitioning, (static and dynamic) and its respective + API's size(), thread_rank(), and sync(). +*/ + +#include "test_common.h" +#include +#include +#include + +using namespace cooperative_groups; + +#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) + +#define NUM_ELEMS 10000000 +#define NUM_THREADS_PER_BLOCK 512 +#define WAVE_SIZE 32 + +/* Test coalesced group's functionality. + * + */ + +__device__ int atomicAggInc(int *ptr) { + coalesced_group g = coalesced_threads(); + int prev; + // elect the first active thread to perform atomic add + if (g.thread_rank() == 0) { + prev = atomicAdd(ptr, g.size()); + } + // broadcast previous value within the warp + // and add each active thread’s rank to it + prev = g.thread_rank() + g.shfl(prev, 0); + return prev; +} + +__global__ void kernel_shfl (int * dPtr, int *dResults, int srcLane, int cg_sizes) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group const& g = coalesced_threads(); + int rank = g.thread_rank(); + int val = dPtr[rank]; + dResults[rank] = g.shfl(val, srcLane); + return; + } +} + +__global__ void kernel_shfl_any_to_any (int *randVal, int *dsrcArr, int *dResults, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + + if (id % cg_sizes == 0) { + coalesced_group const& g = coalesced_threads(); + int rank = g.thread_rank(); + int val = randVal[rank]; + dResults[rank] = g.shfl(val, dsrcArr[rank]); + return; + } + +} + +__global__ void filter_arr(int *dst, int *nres, const int *src, int n) { + int id = threadIdx.x + blockIdx.x * blockDim.x; + + for (int i = id; i < n; i += gridDim.x * blockDim.x) { + if (src[i] > 0) dst[atomicAggInc(nres)] = src[i]; + } +} + +/* Parallel reduce kernel. + * + * Step complexity: O(log n) + * Work complexity: O(n) + * + * Note: This kernel works only with power of 2 input arrays. + */ +__device__ int reduction_kernel(coalesced_group g, int* x, int val) { + int lane = g.thread_rank(); + int sz = g.size(); + + for (int i = g.size() / 2; i > 0; i /= 2) { + // use lds to store the temporary result + x[lane] = val; + // Ensure all the stores are completed. + g.sync(); + + if (lane < i) { + val += x[lane + i]; + } + // It must work on one tiled thread group at a time, + // and it must make sure all memory operations are + // completed before moving to the next stride. + // sync() here just does that. + g.sync(); + } + + // Choose the 0'th indexed thread that holds the reduction value to return + if (g.thread_rank() == 0) { + return val; + } + // Rest of the threads return no useful values + else { + return -1; + } +} + +__global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* result, + bool isGlobalMem, int* globalMem, int cg_sizes) { + + int id = threadIdx.x + blockIdx.x * blockDim.x; + if (id % cg_sizes == 0) { + coalesced_group threadBlockCGTy = coalesced_threads(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + int* workspace = NULL; + + if (isGlobalMem) { + workspace = globalMem; + } else { + // Declare a shared memory + extern __shared__ int sharedMem[]; + workspace = sharedMem; + } + + int input, outputSum, expectedOutput; + + // input to reduction, for each thread, is its' rank in the group + input = threadBlockCGTy.thread_rank(); + + expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2; + + outputSum = reduction_kernel(threadBlockCGTy, workspace, input); + + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Sum of all ranks 0..%d in coalesced_group is %d\n\n", + (int)threadBlockCGTy.size() - 1, outputSum); + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)threadBlockCGTy.size() / tileSz, tileSz); + } + + threadBlockCGTy.sync(); + + coalesced_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz); + + // This offset allows each group to have its own unique area in the workspace array + int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank(); + + outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input); + + if (tiledPartition.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread " + "rank: %d\n", + tiledPartition.size() - 1, outputSum, input); + + result[input / (tileSz)] = outputSum; + } + return; + } +} + +__global__ void kernel_coalesced_active_groups() { + thread_block threadBlockCGTy = this_thread_block(); + int threadBlockGroupSize = threadBlockCGTy.size(); + + // input to reduction, for each thread, is its' rank in the group + int input = threadBlockCGTy.thread_rank(); + + if (threadBlockCGTy.thread_rank() == 0) { + printf(" Creating odd and even set of active thread groups based on branch divergence\n\n"); + } + + threadBlockCGTy.sync(); + + // Group all active odd threads + if (threadBlockCGTy.thread_rank() % 2) { + coalesced_group activeOdd = coalesced_threads(); + + if (activeOdd.thread_rank() == 0) { + printf(" ODD: Size of odd set of active threads is %d." + " Corresponding parent thread_rank is %d.\n\n", + activeOdd.size(), threadBlockCGTy.thread_rank()); + } + } + else { // Group all active even threads + coalesced_group activeEven = coalesced_threads(); + + if (activeEven.thread_rank() == 0) { + printf(" EVEN: Size of even set of active threads is %d." + " Corresponding parent thread_rank is %d.", + activeEven.size(), threadBlockCGTy.thread_rank()); + } + } + return; +} + +void printResults(int* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +void compareResults(int* cpu, int* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(int); i++) { + if (cpu[i] != gpu[i]) { + printf(" results do not match."); + } + } +} + +static void test_active_threads_grouping() { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + // Launch Kernel + hipLaunchKernelGGL(kernel_coalesced_active_groups, blockSize, threadsPerBlock, 0, 0); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + printf("\n...PASSED.\n\n"); +} + +// Search if the sum exists in the expected results array +void verifyResults(int* hPtr, int* dPtr, int size) { + int i = 0, j = 0; + for (i = 0; i < size; i++) { + for (j = 0; j < size; j++) { + if (hPtr[i] == dPtr[j]) { + break; + } + } + if (j == size) { + printf(" Result verification failed!"); + } + } +} + + +static void test_group_partition(unsigned int tileSz, bool useGlobalMem) { + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + int numTiles = ((blockSize * threadsPerBlock) / i) / tileSz; + + // numTiles = 0 when partitioning is possible. The below statement is to avoid + // out-of-bounds error and still evaluate failure case. + numTiles = (numTiles == 0) ? 1 : numTiles; + + // Build an array of expected reduction sum output on the host + // based on the sum of their respective thread ranks to use for verification + int* expectedSum = new int[numTiles]; + int temp = 0, sum = 0; + for (int i = 1; i <= numTiles; i++) { + sum = temp; + temp = (((tileSz * i) - 1) * (tileSz * i)) / 2; + expectedSum[i-1] = temp - sum; + } + + int* dResult = NULL; + hipMalloc(&dResult, sizeof(int) * numTiles); + + int* globalMem = NULL; + if (useGlobalMem) { + hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int)); + } + + int* hResult = NULL; + hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault); + memset(hResult, 0, numTiles * sizeof(int)); + + // Launch Kernel + if (useGlobalMem) { + hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, 0, 0, tileSz, + dResult, useGlobalMem, globalMem, i); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } else { + hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, + threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem, i); + + err = hipDeviceSynchronize(); + if (err != hipSuccess) { + fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err)); + } + } + + hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost); + verifyResults(expectedSum, hResult, numTiles); + // Free all allocated memory on host and device + hipFree(dResult); + hipFree(hResult); + if (useGlobalMem) { + hipFree(globalMem); + } + delete[] expectedSum; + + printf("\n...PASSED.\n\n"); + } +} +static void test_shfl_any_to_any() { + + std::vector cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int totalThreads = blockSize * threadsPerBlock; + int group_size = (totalThreads + i - 1) / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dsrcArr = NULL; + int* dResults = NULL; + int* srcArr = (int*)malloc(group_size_in_bytes); + int* srcArrCpu = (int*)malloc(group_size_in_bytes); + + std::cout << "Testing coalesced_groups shfl any-to-any\n" < cg_sizes = {1, 2, 3}; + for (auto i : cg_sizes) { + + hipError_t err; + int blockSize = 1; + int threadsPerBlock = WAVE_SIZE; + + int totalThreads = blockSize * threadsPerBlock; + int group_size = (totalThreads + i - 1) / i; + int group_size_in_bytes = group_size * sizeof(int); + + int* hPtr = NULL; + int* dPtr = NULL; + int* dResults = NULL; + int srcLane = rand() % 1000; + int srcLaneCpu = 0; + std::cout << "Testing coalesced_groups shfl with srcLane " << srcLane << '\n' + << " and group size " << i <(malloc(sizeof(int) * NUM_ELEMS)); + + // Generate input data. + for (int i = 0; i < NUM_ELEMS; i++) { + data_to_filter[i] = rand() % numOfBuckets; + } + + + hipMalloc(&d_data_to_filter, sizeof(int) * NUM_ELEMS); + hipMalloc(&d_filtered_data, sizeof(int) * NUM_ELEMS); + hipMalloc(&d_nres, sizeof(int)); + + hipMemcpy(d_data_to_filter, data_to_filter, + sizeof(int) * NUM_ELEMS, hipMemcpyHostToDevice); + hipMemset(d_nres, 0, sizeof(int)); + + dim3 dimBlock(NUM_THREADS_PER_BLOCK, 1, 1); + dim3 dimGrid((NUM_ELEMS / NUM_THREADS_PER_BLOCK) + 1, 1, 1); + + filter_arr<<>>(d_filtered_data, d_nres, d_data_to_filter, + NUM_ELEMS); + + + hipMemcpy(&nres, d_nres, sizeof(int), hipMemcpyDeviceToHost); + + filtered_data = reinterpret_cast(malloc(sizeof(int) * nres)); + + hipMemcpy(filtered_data, d_filtered_data, sizeof(int) * nres, + hipMemcpyDeviceToHost); + + int *host_filtered_data = + reinterpret_cast(malloc(sizeof(int) * NUM_ELEMS)); + + // Generate host output with host filtering code. + int host_flt_count = 0; + for (int i = 0; i < NUM_ELEMS; i++) { + if (data_to_filter[i] > 0) { + host_filtered_data[host_flt_count++] = data_to_filter[i]; + } + } + + printf("\nWarp Aggregated Atomics %s \n", + (host_flt_count == nres) ? "PASSED" : "FAILED"); + + // Now, testing shfl collective + std::cout << "Now testing shfl collective as a broadcast" << '\n' << std::endl; + + for (int i = 0; i < 100; i++) { + test_shfl_broadcast(); + } + + + // Now, testing shfl collective + std::cout << "Now testing shfl operations any-to-any member lanes" << '\n' << std::endl; + + for (int i = 0; i < 100; i++) { + test_shfl_any_to_any(); + } + + // Now, pass a already coalesced_group that was partitioned + /* Test coalesced group partitioning */ + std::cout << "Now testing coalesced_groups partitioning" << '\n' << std::endl; + + int testNo = 1; + for (int memTy = 0; memTy < 2; memTy++) { + std::vector tileSizes = {2, 4, 8, 16, 32}; + for (auto i : tileSizes) { + std::cout << "TEST " << testNo << ":" << '\n' << std::endl; + test_group_partition(i, memTy); + testNo++; + } + } + + std::cout << "Now grouping active threads based on branch divergence" << '\n' << std::endl; + test_active_threads_grouping(); + + passed(); + return 0; +}