From fbdda38cf730bcc8d97b777c23ae868188331ee5 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 10 Jun 2022 16:46:23 +0100 Subject: [PATCH 01/35] Add missing tests for hipHostMalloc (#2566) --- tests/catch/include/hip_test_helper.hh | 3 +- tests/catch/stress/memory/CMakeLists.txt | 1 + tests/catch/stress/memory/hipHostMalloc.cc | 52 +++++++++++++++ .../stress/memory/hipMallocManagedStress.cc | 1 - tests/catch/unit/memory/hipHostMallocTests.cc | 27 +++++--- tests/catch/unit/memory/hipMemCoherencyTst.cc | 66 ++++++++----------- 6 files changed, 100 insertions(+), 50 deletions(-) create mode 100644 tests/catch/stress/memory/hipHostMalloc.cc diff --git a/tests/catch/include/hip_test_helper.hh b/tests/catch/include/hip_test_helper.hh index 04aec9318b..fbbcb6cb06 100644 --- a/tests/catch/include/hip_test_helper.hh +++ b/tests/catch/include/hip_test_helper.hh @@ -51,8 +51,7 @@ static size_t getMemoryAmount() { #endif } -static size_t getHostThreadCount(const size_t memPerThread, - const size_t maxThreads) { +static inline size_t getHostThreadCount(const size_t memPerThread, const size_t maxThreads) { if (memPerThread == 0) return 0; auto memAmount = getMemoryAmount(); const auto processor_count = std::thread::hardware_concurrency(); diff --git a/tests/catch/stress/memory/CMakeLists.txt b/tests/catch/stress/memory/CMakeLists.txt index da5dfa8182..e55869f8e3 100644 --- a/tests/catch/stress/memory/CMakeLists.txt +++ b/tests/catch/stress/memory/CMakeLists.txt @@ -4,6 +4,7 @@ set(TEST_SRC hipMemcpyMThreadMSize.cc hipMallocManagedStress.cc hipMemPrftchAsyncStressTst.cc + hipHostMalloc.cc ) hip_add_exe_to_target(NAME memory diff --git a/tests/catch/stress/memory/hipHostMalloc.cc b/tests/catch/stress/memory/hipHostMalloc.cc new file mode 100644 index 0000000000..da9fa977ff --- /dev/null +++ b/tests/catch/stress/memory/hipHostMalloc.cc @@ -0,0 +1,52 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip_test_common.hh" +#include "hip_test_helper.hh" + +// Stress allocation tests +// Try to allocate as much memory as possible +// But since max allocation can fail, we need to try the next value + +TEST_CASE("Stress_hipHostMalloc_MaxAllocation") { + size_t devMemAvail{0}, devMemFree{0}; + HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemAvail)); + auto hostMemFree = HipTest::getMemoryAmount() /* In MB */ * 1024 * 1024; // In bytes + REQUIRE(devMemFree > 0); + REQUIRE(devMemAvail > 0); + REQUIRE(hostMemFree > 0); + + size_t memFree = std::min(devMemFree, hostMemFree); // which is the limiter cpu or gpu + char* d_ptr{nullptr}; + size_t counter{0}; + + INFO("Max Allocation of " << memFree << " bytes!"); + while (hipHostMalloc(&d_ptr, memFree) != hipSuccess && memFree > 1) { + counter++; + INFO("Attempt to allocate " << memFree << " bytes out of " << devMemFree << "bytes Failed!"); + memFree >>= 1; // reduce the memory to be allocated by half + REQUIRE(counter <= 2); // Make sure that we are atleast able to allocate 1/4th of max memory + } + + HIP_CHECK(hipMemset(d_ptr, 1, memFree)); + HIP_CHECK(hipDeviceSynchronize()); // Flush caches + REQUIRE(std::all_of(d_ptr, d_ptr + memFree, [](unsigned char n) { return n == 1; })); + HIP_CHECK(hipHostFree(d_ptr)); +} + diff --git a/tests/catch/stress/memory/hipMallocManagedStress.cc b/tests/catch/stress/memory/hipMallocManagedStress.cc index e582056685..8a11ab35ef 100644 --- a/tests/catch/stress/memory/hipMallocManagedStress.cc +++ b/tests/catch/stress/memory/hipMallocManagedStress.cc @@ -67,7 +67,6 @@ __global__ void KernelMul_MngdMem(int *Hmm, int *Dptr, size_t n) { 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; diff --git a/tests/catch/unit/memory/hipHostMallocTests.cc b/tests/catch/unit/memory/hipHostMallocTests.cc index de445e2e24..b0f3d4dbfa 100644 --- a/tests/catch/unit/memory/hipHostMallocTests.cc +++ b/tests/catch/unit/memory/hipHostMallocTests.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 @@ -29,33 +29,40 @@ Testcase Scenarios : */ #include +#include /** * Performs argument validation of hipHostMalloc api. */ TEST_CASE("Unit_hipHostMalloc_ArgValidation") { - hipError_t ret; +#if HT_NVIDIA + HipTest::HIP_SKIP_TEST("TODO: Need to debug"); +#endif constexpr size_t allocSize = 1000; - char *ptr; + char* ptr; SECTION("Pass ptr as nullptr") { - ret = hipHostMalloc(static_cast(nullptr), allocSize); - REQUIRE(ret != hipSuccess); + HIP_CHECK_ERROR(hipHostMalloc(static_cast(nullptr), allocSize), hipErrorInvalidValue); } SECTION("Size as max(size_t)") { - ret = hipHostMalloc(&ptr, std::numeric_limits::max()); - REQUIRE(ret != hipSuccess); + HIP_CHECK_ERROR(hipHostMalloc(&ptr, std::numeric_limits::max()), + hipErrorMemoryAllocation); } SECTION("Flags as max(uint)") { - ret = hipHostMalloc(&ptr, allocSize, - std::numeric_limits::max()); - REQUIRE(ret != hipSuccess); + HIP_CHECK_ERROR(hipHostMalloc(&ptr, allocSize, std::numeric_limits::max()), + hipErrorInvalidValue); } SECTION("Pass size as zero and check ptr reset") { HIP_CHECK(hipHostMalloc(&ptr, 0)); REQUIRE(ptr == nullptr); } + + SECTION("Pass hipHostMallocCoherent and hipHostMallocNonCoherent simultaneously") { + HIP_CHECK_ERROR( + hipHostMalloc(&ptr, allocSize, hipHostMallocCoherent | hipHostMallocNonCoherent), + hipErrorInvalidValue); + } } diff --git a/tests/catch/unit/memory/hipMemCoherencyTst.cc b/tests/catch/unit/memory/hipMemCoherencyTst.cc index 73bf4aeb09..67f455de29 100644 --- a/tests/catch/unit/memory/hipMemCoherencyTst.cc +++ b/tests/catch/unit/memory/hipMemCoherencyTst.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 @@ -34,17 +34,9 @@ #include #include -__global__ void CoherentTst(int *ptr, int PeakClk) { - // Incrementing the value by 1 - int64_t GpuFrq = (PeakClk * 1000); - int64_t StrtTck = clock64(); - atomicAdd(ptr, 1); - // The following while loop checks the value in ptr for around 3-4 seconds - while ((clock64() - StrtTck) <= (3 * GpuFrq)) { - if (*ptr == 3) { - atomicAdd(ptr, 1); - return; - } +__global__ void CoherentTst(int* ptr) { // ptr was set to 1 + atomicAdd(ptr, 1); // now ptr is 2 + while (atomicCAS(ptr, 3, 4) != 3) { // wait till ptr is 3, then change it to 4 } } @@ -59,34 +51,34 @@ __global__ void SquareKrnl(int *ptr) { static bool YES_COHERENT = false; // The function tests the coherency of allocated memory -static void TstCoherency(int *Ptr, bool HmmMem) { - int *Dptr = nullptr, peak_clk; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); +// If this test hangs, means there is issue in coherency +static void TstCoherency(int* ptr, bool hmmMem) { + int* dptr = nullptr; + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + // storing value 1 in the memory created above - *Ptr = 1; - // Getting gpu frequency - HIP_CHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0)); - if (!HmmMem) { - HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&Dptr), Ptr, - 0)); - CoherentTst<<<1, 1, 0, strm>>>(Dptr, peak_clk); + *ptr = 1; + + if (!hmmMem) { + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&dptr), ptr, 0)); + CoherentTst<<<1, 1, 0, stream>>>(dptr); } else { - CoherentTst<<<1, 1, 0, strm>>>(Ptr, peak_clk); + CoherentTst<<<1, 1, 0, stream>>>(ptr); } - // looping until the value is 2 for 3 seconds - std::chrono::steady_clock::time_point start = - std::chrono::steady_clock::now(); - while (std::chrono::duration_cast( - std::chrono::steady_clock::now() - start).count() < 3) { - if (*Ptr == 2) { - *Ptr += 1; - break; - } - } - HIP_CHECK(hipStreamSynchronize(strm)); - HIP_CHECK(hipStreamDestroy(strm)); - if (*Ptr == 4) { + + std::chrono::steady_clock::time_point start = std::chrono::steady_clock::now(); + while (std::chrono::duration_cast(std::chrono::steady_clock::now() - start) + .count() < 3 && + *ptr == 2) { + } // wait till ptr is 2 from kernel or 3 seconds + + *ptr += 1; // increment it to 3 + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + + if (*ptr == 4) { YES_COHERENT = true; } } From db05dc361bce8eb6a22a5d30ba3b31f1086961df Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 10 Jun 2022 16:46:54 +0100 Subject: [PATCH 02/35] Add test for hipHostGetDevicePointer (#2611) --- tests/catch/unit/memory/CMakeLists.txt | 2 + .../unit/memory/hipHostGetDevicePointer.cc | 78 +++++++++++++++++++ 2 files changed, 80 insertions(+) create mode 100644 tests/catch/unit/memory/hipHostGetDevicePointer.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 3ce80dd780..d36f0a1b85 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -41,6 +41,7 @@ set(TEST_SRC hipMemPtrGetInfo.cc hipPointerGetAttributes.cc hipHostGetFlags.cc + hipHostGetDevicePointer.cc hipMallocManaged_MultiScenario.cc hipMemsetNegative.cc hipMemset.cc @@ -106,6 +107,7 @@ set(TEST_SRC hipMemcpy_MultiThread.cc hipHostRegister.cc hipHostGetFlags.cc + hipHostGetDevicePointer.cc hipMallocManaged_MultiScenario.cc hipMemsetNegative.cc hipMemset.cc diff --git a/tests/catch/unit/memory/hipHostGetDevicePointer.cc b/tests/catch/unit/memory/hipHostGetDevicePointer.cc new file mode 100644 index 0000000000..7c3e689e05 --- /dev/null +++ b/tests/catch/unit/memory/hipHostGetDevicePointer.cc @@ -0,0 +1,78 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +TEST_CASE("Unit_hipHostGetDevicePointer_Negative") { + int* hPtr{nullptr}; + HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); + + SECTION("Nullptr as device") { + HIP_CHECK_ERROR(hipHostGetDevicePointer(nullptr, hPtr, 0), hipErrorInvalidValue); + } + + SECTION("Nullptr as host") { + int* dPtr{nullptr}; + HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast(&dPtr), nullptr, 0), + hipErrorInvalidValue); + } + + // Not adding check for flags since CUDA spec states that there might be more values added to it + HIP_CHECK(hipHostFree(hPtr)); +} + +template __global__ void set(T* ptr, T val) { *ptr = val; } + +TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") { + int* hPtr{nullptr}; + HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); + + auto kernel = set; + constexpr int value = 10; + + SECTION("Set the value on device - Get device ptr") { + int* dPtr{nullptr}; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&dPtr), hPtr, 0)); + REQUIRE(dPtr != nullptr); + + kernel<<<1, 1>>>(dPtr, value); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(*hPtr == value); + } + + SECTION("Set the value on device - by hipHostRegister") { + int res{0}; // Stuff on stack + HIP_CHECK(hipHostRegister(&res, sizeof(int), 0)); // Lets map stack memory :) + + int* dPtr{nullptr}; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&dPtr), &res, 0)) + + kernel<<<1, 1>>>(dPtr, value); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipHostUnregister(&res)); + + REQUIRE(value == 10); + } + + HIP_CHECK(hipHostFree(hPtr)); +} From 03ca05e443cf0996daa9269ba05484e3dbb7845b Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Fri, 10 Jun 2022 21:17:10 +0530 Subject: [PATCH 03/35] SWDEV-317716 - [catch2][dtest] Add test for hipDeviceGetUuid (#2614) Change-Id: Ic2940cd45cacc32212df7c76ee8f4ca5e7e3740c --- tests/catch/unit/device/CMakeLists.txt | 1 + tests/catch/unit/device/hipDeviceGetUuid.cc | 51 +++++++++++++++++++++ 2 files changed, 52 insertions(+) create mode 100644 tests/catch/unit/device/hipDeviceGetUuid.cc diff --git a/tests/catch/unit/device/CMakeLists.txt b/tests/catch/unit/device/CMakeLists.txt index 778cd199dc..e6d45594ea 100644 --- a/tests/catch/unit/device/CMakeLists.txt +++ b/tests/catch/unit/device/CMakeLists.txt @@ -15,6 +15,7 @@ set(TEST_SRC hipRuntimeGetVersion.cc hipSetDeviceFlags.cc hipSetGetDevice.cc + hipDeviceGetUuid.cc ) hip_add_exe_to_target(NAME DeviceTest diff --git a/tests/catch/unit/device/hipDeviceGetUuid.cc b/tests/catch/unit/device/hipDeviceGetUuid.cc new file mode 100644 index 0000000000..512ddb6a2b --- /dev/null +++ b/tests/catch/unit/device/hipDeviceGetUuid.cc @@ -0,0 +1,51 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* + * Conformance test for checking functionality of + * hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device); + */ +#include +#include +#include + +/** + * hipDeviceGetUuid tests + * Scenario1: Validates the returned UUID + * Scenario2: Validates returned error code for UUID = nullptr + * Scenario3 & 4: Validates returned error code for invalid device + */ +TEST_CASE("Unit_hipDeviceGetUuid") { + int numDevices = 0; + hipDevice_t device; + hipUUID uuid; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + for (int i = 0; i < numDevices; i++) { + HIP_CHECK(hipDeviceGet(&device, i)); + // Scenario 1 + HIP_CHECK(hipDeviceGetUuid(&uuid, device)); + REQUIRE_FALSE(!strcmp(uuid.bytes, "")); + // Scenario 2 + REQUIRE_FALSE(hipSuccess == hipDeviceGetUuid(nullptr, device)); + } + // Scenario 3 + REQUIRE_FALSE(hipSuccess == hipDeviceGetUuid(&uuid, -1)); + // Scenario 4 + REQUIRE_FALSE(hipSuccess == hipDeviceGetUuid(&uuid, numDevices)); +} From 54b1822167e7058ca194a3e70faf13c0dcd96665 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 10 Jun 2022 16:47:34 +0100 Subject: [PATCH 04/35] Adding/Moving tests for hipMemcpy[From/To]Symbol[Async] (#2638) --- tests/catch/unit/memory/CMakeLists.txt | 6 - .../catch/unit/memory/hipMemcpyFromSymbol.cc | 179 ++++++++++++++++-- .../unit/memory/hipMemcpyFromSymbolAsync.cc | 47 ----- tests/catch/unit/memory/hipMemcpyToSymbol.cc | 43 ----- .../unit/memory/hipMemcpyToSymbolAsync.cc | 47 ----- 5 files changed, 163 insertions(+), 159 deletions(-) delete mode 100644 tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc delete mode 100644 tests/catch/unit/memory/hipMemcpyToSymbol.cc delete mode 100644 tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index d36f0a1b85..70298c246e 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -58,9 +58,6 @@ set(TEST_SRC hipMallocManaged.cc hipMemRangeGetAttribute.cc hipMemcpyFromSymbol.cc - hipMemcpyFromSymbolAsync.cc - hipMemcpyToSymbol.cc - hipMemcpyToSymbolAsync.cc hipPtrGetAttribute.cc hipMemcpyPeer.cc hipMemcpyPeerAsync.cc @@ -123,9 +120,6 @@ set(TEST_SRC hipMallocManaged.cc hipMemRangeGetAttribute.cc hipMemcpyFromSymbol.cc - hipMemcpyFromSymbolAsync.cc - hipMemcpyToSymbol.cc - hipMemcpyToSymbolAsync.cc hipPtrGetAttribute.cc hipMemcpyPeer.cc hipMemcpyPeerAsync.cc diff --git a/tests/catch/unit/memory/hipMemcpyFromSymbol.cc b/tests/catch/unit/memory/hipMemcpyFromSymbol.cc index 347b4589a1..ff38ddd060 100644 --- a/tests/catch/unit/memory/hipMemcpyFromSymbol.cc +++ b/tests/catch/unit/memory/hipMemcpyFromSymbol.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,26 +18,173 @@ THE SOFTWARE. */ #include -#define SIZE 1024 -/* Test verifies hipMemcpyFromSymbol API Negative scenarios. +__device__ int devSymbol[10]; + +/* Test verifies hipMemcpy[From/To]Symbol[Async] API Negative scenarios. */ -TEST_CASE("Unit_hipMemcpyFromSymbol_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbol(S, HIP_SYMBOL(Sd), - SIZE, 0, hipMemcpyDeviceToHost)); +TEST_CASE("Unit_hipMemcpyFromToSymbol_Negative") { + SECTION("Invalid Src Ptr") { + int result{0}; + HIP_CHECK_ERROR( + hipMemcpyFromSymbol(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbol(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidSymbol); + HIP_CHECK_ERROR( + hipMemcpyToSymbolAsync(nullptr, &result, sizeof(int), 0, hipMemcpyHostToDevice, nullptr), + hipErrorInvalidSymbol); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(nullptr, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); } - SECTION("Passing NULL Pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbol(S, nullptr, - SIZE, 0, hipMemcpyDeviceToHost)); + SECTION("Invalid Dst Ptr") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidSymbol); + HIP_CHECK_ERROR( + hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), nullptr, sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyFromSymbolAsync(&result, nullptr, sizeof(int), 0, hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidSymbol); } - HIP_CHECK(hipFree(Sd)); + SECTION("Invalid Size") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int) * 100, 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * 100, 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + + SECTION("Invalid Offset") { + int result{0}; + HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, hipMemcpyHostToDevice), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int), 300, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 300, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidValue); + } + + SECTION("Invalid Direction") { + int result{0}; + HIP_CHECK_ERROR( + hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, hipMemcpyHostToDevice), + hipErrorInvalidMemcpyDirection); + HIP_CHECK_ERROR( + hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, hipMemcpyDeviceToHost), + hipErrorInvalidMemcpyDirection); + HIP_CHECK_ERROR(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &result, sizeof(int), 0, + hipMemcpyDeviceToHost, nullptr), + hipErrorInvalidMemcpyDirection); + HIP_CHECK_ERROR(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyHostToDevice, nullptr), + hipErrorInvalidMemcpyDirection); + } +} + +/* + * Test Verifies hipMemcpyToSymbol/hipMemcpyFromSymbol and Async Variants for simple use case + * For single valuea To and From Symbol + * For Array Values To and From Symbol + * For Array Values with offset To and From Symbol + * For Sync and Async Variants*/ +TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") { + enum StreamTestType { NullStream = 0, StreamPerThread, CreatedStream, NoStream }; + + /* Test type NoStream - Use Sync variants, else use async variants */ + auto streamType = GENERATE(StreamTestType::NoStream, StreamTestType::NullStream, + StreamTestType::StreamPerThread, StreamTestType::CreatedStream); + + hipStream_t stream{nullptr}; + + if (streamType == StreamTestType::StreamPerThread) { + stream = hipStreamPerThread; + } else if (streamType == StreamTestType::CreatedStream) { + HIP_CHECK(hipStreamCreate(&stream)); + } + INFO("Stream :: " << streamType); + + SECTION("Singular Value") { + int set{42}; + int result{0}; + if (streamType == StreamTestType::NoStream) { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &set, sizeof(int))); + HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int))); + } else { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &set, sizeof(int), 0, + hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int), 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + REQUIRE(result == set); + } + + SECTION("Array Values") { + constexpr size_t size{10}; + int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2}; + int result[size] = {0}; + if (streamType == StreamTestType::NoStream) { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, sizeof(int) * size)); + HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + } else { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, sizeof(int) * size, 0, + hipMemcpyHostToDevice, stream)); + + HIP_CHECK(hipMemcpyFromSymbolAsync(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } + + SECTION("Offset'ed Values") { + constexpr size_t size{10}; + constexpr size_t offset = 5 * sizeof(int); + int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2}; + int result[size] = {0}; + if (streamType == StreamTestType::NoStream) { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, offset)); + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set + 5, offset, offset)); + HIP_CHECK(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), sizeof(int) * size)); + } else { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, offset, 0, hipMemcpyHostToDevice, + stream)); + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set + 5, offset, offset, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(result, HIP_SYMBOL(devSymbol), offset, 0, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(result + 5, HIP_SYMBOL(devSymbol), offset, offset, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + } + for (size_t i = 0; i < size; i++) { + REQUIRE(result[i] == set[i]); + } + } } diff --git a/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc b/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc deleted file mode 100644 index ae6ebc3e82..0000000000 --- a/tests/catch/unit/memory/hipMemcpyFromSymbolAsync.cc +++ /dev/null @@ -1,47 +0,0 @@ -/* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#include -#define SIZE 1024 - -/* Test verifies hipMemcpyFromSymbolAsync API Negative scenarios. - */ - -TEST_CASE("Unit_hipMemcpyFromSymbolAsync_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbolAsync(S, HIP_SYMBOL(Sd), - SIZE, 0, hipMemcpyDeviceToHost, stream)); - } - - SECTION("Passing NULL pointer") { - REQUIRE(hipSuccess != hipMemcpyFromSymbolAsync(S, nullptr, - SIZE, 0, hipMemcpyDeviceToHost, stream)); - } - - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipFree(Sd)); -} diff --git a/tests/catch/unit/memory/hipMemcpyToSymbol.cc b/tests/catch/unit/memory/hipMemcpyToSymbol.cc deleted file mode 100644 index 23e5768c3c..0000000000 --- a/tests/catch/unit/memory/hipMemcpyToSymbol.cc +++ /dev/null @@ -1,43 +0,0 @@ -/* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#include -#define SIZE 1024 - -/* Test verifies hipMemcpyToSymbol API Negative scenarios. - */ - -TEST_CASE("Unit_hipMemcpyToSymbol_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbol(HIP_SYMBOL(Sd), S, - SIZE, 0, hipMemcpyDeviceToHost)); - } - - SECTION("Passing NULL Pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbol(nullptr, S, - SIZE, 0, hipMemcpyDeviceToHost)); - } - - HIP_CHECK(hipFree(Sd)); -} diff --git a/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc b/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc deleted file mode 100644 index 42b82128f9..0000000000 --- a/tests/catch/unit/memory/hipMemcpyToSymbolAsync.cc +++ /dev/null @@ -1,47 +0,0 @@ -/* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#include -#define SIZE 1024 - -/* Test verifies hipMemcpyToSymbolAsync API Negative scenarios. - */ - -TEST_CASE("Unit_hipMemcpyToSymbolAsync_Negative") { - void *Sd; - char S[SIZE]="This is not a device symbol"; - - HIP_CHECK(hipMalloc(&Sd, SIZE)); - - hipStream_t stream; - HIP_CHECK(hipStreamCreate(&stream)); - - SECTION("Passing void pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbolAsync(HIP_SYMBOL(Sd), S, - SIZE, 0, hipMemcpyHostToDevice, stream)); - } - - SECTION("Passing NULL pointer") { - REQUIRE(hipSuccess != hipMemcpyToSymbolAsync(nullptr, S, - SIZE, 0, hipMemcpyHostToDevice, stream)); - } - - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipFree(Sd)); -} From f356d2c33daee03118f37355d02a6b0a1f214abd Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 13 Jun 2022 08:23:17 +0530 Subject: [PATCH 05/35] Update hipHostMallocTests.cc std::numeric_limits::max() conflicts with max() defined in windows headers. Wrap it in parenthesis to prevent macro expansion. --- tests/catch/unit/memory/hipHostMallocTests.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/catch/unit/memory/hipHostMallocTests.cc b/tests/catch/unit/memory/hipHostMallocTests.cc index b0f3d4dbfa..9b25f613ac 100644 --- a/tests/catch/unit/memory/hipHostMallocTests.cc +++ b/tests/catch/unit/memory/hipHostMallocTests.cc @@ -46,12 +46,12 @@ TEST_CASE("Unit_hipHostMalloc_ArgValidation") { } SECTION("Size as max(size_t)") { - HIP_CHECK_ERROR(hipHostMalloc(&ptr, std::numeric_limits::max()), + HIP_CHECK_ERROR(hipHostMalloc(&ptr, (std::numeric_limits::max)()), hipErrorMemoryAllocation); } SECTION("Flags as max(uint)") { - HIP_CHECK_ERROR(hipHostMalloc(&ptr, allocSize, std::numeric_limits::max()), + HIP_CHECK_ERROR(hipHostMalloc(&ptr, allocSize, (std::numeric_limits::max)()), hipErrorInvalidValue); } From d12d0ebc578601de138765ee4b1ddd2dcbc79edf Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 15:33:27 +0530 Subject: [PATCH 06/35] SWDEV-327030 - Fix incorrect checks to find components in CMake (#2680) Change-Id: I11529018f2b5d878ca2b728ef7084163ccad317f --- hip-lang-config.cmake.in | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hip-lang-config.cmake.in b/hip-lang-config.cmake.in index 1a72643a1f..dbc8e5348d 100644 --- a/hip-lang-config.cmake.in +++ b/hip-lang-config.cmake.in @@ -94,7 +94,7 @@ find_path(HSA_HEADER hsa/hsa.h /opt/rocm/include ) -if (HSA_HEADER-NOTFOUND) +if (NOT HSA_HEADER) message (FATAL_ERROR "HSA header not found! ROCM_PATH environment not set") endif() @@ -136,7 +136,7 @@ set_property(TARGET hip-lang::device APPEND PROPERTY ) # Add support for __fp16 and _Float16, explicitly link with compiler-rt -if(CLANGRT_BUILTINS-NOTFOUND) +if(NOT CLANGRT_BUILTINS) message(FATAL_ERROR "clangrt builtins lib not found") else() set_property(TARGET hip-lang::device APPEND PROPERTY From 69fcc3a9cdaaf3f92cd1f07c0ed703885ffd8b03 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 15:33:49 +0530 Subject: [PATCH 07/35] SWDEV-306122 - [catch2][dtest] hipGraph tests for hipStreamGetCaptureInfo api (#2696) Change-Id: I73d42a87797cf70c5ac6093bc308203eae5cca00 --- .../config/config_amd_windows.json | 4 + tests/catch/unit/graph/CMakeLists.txt | 1 + .../unit/graph/hipStreamGetCaptureInfo.cc | 231 ++++++++++++++++++ 3 files changed, 236 insertions(+) create mode 100644 tests/catch/unit/graph/hipStreamGetCaptureInfo.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 44dc69cc36..8602662d49 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -80,6 +80,10 @@ "Unit_hipStreamBeginCapture_hipStreamPerThread", "Unit_hiprtc_functional", "Unit_hipStreamValue_Write", + "Unit_hipStreamGetCaptureInfo_BasicFunctional", + "Unit_hipStreamGetCaptureInfo_hipStreamPerThread", + "Unit_hipStreamGetCaptureInfo_UniqueID", + "Unit_hipStreamGetCaptureInfo_ArgValidation", "# Following test is related to ticket EXSWCPHIPT-41", "Unit_hipStreamGetPriority_happy" ] diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 704bf58cfb..430832e03e 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -58,6 +58,7 @@ set(TEST_SRC hipGraphExecMemcpyNodeSetParams.cc hipStreamBeginCapture.cc hipGraphAddMemcpyNode1D.cc + hipStreamGetCaptureInfo.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipStreamGetCaptureInfo.cc b/tests/catch/unit/graph/hipStreamGetCaptureInfo.cc new file mode 100644 index 0000000000..13b5270926 --- /dev/null +++ b/tests/catch/unit/graph/hipStreamGetCaptureInfo.cc @@ -0,0 +1,231 @@ +/* +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) Start stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusActive + and identifier returned is valid/non-zero. +2) End stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone + and identifier is not returned/updated by api. +3) Begin capture on hipStreamPerThread and get capture info. Verify api is success, capture status is hipStreamCaptureStatusActive + and identifier returned is valid/non-zero. +4) End capture on hipStreamPerThread, get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone + and identifier is not returned/updated by api. +5) Perform multiple captures and verify the identifier returned is unique. + +Argument Validation/Negative: +1) Pass pId as nullptr and verify api doesn’t crash and returns success. +2) Pass pCaptureStatus as nullptr and verify api doesn’t crash and returns error code. + +*/ + +#include +#include +#include + +constexpr size_t N = 1000000; +constexpr int LAUNCH_ITERS = 1; + +/** + * Validates stream capture info, launches graph and verify results + */ +void validateStreamCaptureInfo(hipStream_t mstream) { + hipStream_t stream1{nullptr}, stream2{nullptr}, streamForLaunch{nullptr}; + hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + size_t Nbytes = N * sizeof(float); + 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(&streamForLaunch)); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + // Create cross stream dependencies. + // memset operations are done on stream1 and stream2 + // and they are joined back to mainstream + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + 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); + + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + unsigned long long capSequenceID = 0; // NOLINT + HIP_CHECK(hipStreamGetCaptureInfo(mstream, &captureStatus, &capSequenceID)); + + // verify capture status is active and sequence id is valid + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + REQUIRE(capSequenceID > 0); + + // End capture and verify graph is returned + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + REQUIRE(graph != nullptr); + + // verify capture status is inactive and sequence id is not updated + capSequenceID = 0; + HIP_CHECK(hipStreamGetCaptureInfo(mstream, &captureStatus, &capSequenceID)); + REQUIRE(captureStatus == hipStreamCaptureStatusNone); + REQUIRE(capSequenceID == 0); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(graphExec != nullptr); + + // Replay the recorded sequence multiple times + for (int i = 0; i < LAUNCH_ITERS; i++) { + HIP_CHECK(hipGraphLaunch(graphExec, streamForLaunch)); + } + + HIP_CHECK(hipStreamSynchronize(streamForLaunch)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForLaunch)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memsetEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + INFO("A and C not matching at " << i << " C_h[i] " << C_h[i] + << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + free(A_h); + free(C_h); +} + +/** + * Basic Functional Test for stream capture and getting capture info. + * Regular/custom stream is used for stream capture. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_BasicFunctional") { + hipStream_t streamForCapture; + + HIP_CHECK(hipStreamCreate(&streamForCapture)); + validateStreamCaptureInfo(streamForCapture); + HIP_CHECK(hipStreamDestroy(streamForCapture)); +} + +/** + * Test performs stream capture on hipStreamPerThread and validates + * capture info. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_hipStreamPerThread") { + validateStreamCaptureInfo(hipStreamPerThread); +} + +/** + * Test starts stream capture on multiple streams and verifies uniqueness of + * identifiers returned. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_UniqueID") { + constexpr int numStreams = 100; + hipStream_t streams[numStreams]{}; + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + std::vector idlist; + unsigned long long capSequenceID{}; // NOLINT + hipGraph_t graph{nullptr}; + + for (int i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + HIP_CHECK(hipStreamBeginCapture(streams[i], hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamGetCaptureInfo(streams[i], &captureStatus, + &capSequenceID)); + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + REQUIRE(capSequenceID > 0); + idlist.push_back(capSequenceID); + } + + for (int i = 0; i < numStreams; i++) { + for (int j = i+1; j < numStreams; j++) { + if (idlist[i] == idlist[j]) { + INFO("Same identifier returned for stream " + << i << " and stream " << j); + REQUIRE(false); + } + } + } + + for (int i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamEndCapture(streams[i], &graph)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streams[i])); + } +} + +/** + * Argument validation/Negative tests for api + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_ArgValidation") { + hipError_t ret; + hipStream_t stream; + hipStreamCaptureStatus captureStatus; + unsigned long long capSequenceID; // NOLINT + HIP_CHECK(hipStreamCreate(&stream)); + + SECTION("Capture ID location as nullptr") { + ret = hipStreamGetCaptureInfo(stream, &captureStatus, nullptr); + // Capture ID is optional + REQUIRE(ret == hipSuccess); + } + + SECTION("Capture Status location as nullptr") { + ret = hipStreamGetCaptureInfo(stream, nullptr, &capSequenceID); + REQUIRE(ret == hipErrorInvalidValue); + } + + HIP_CHECK(hipStreamDestroy(stream)); +} From 5133509730c21238d43b0e65bb70b6973a4999f4 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 15:34:29 +0530 Subject: [PATCH 08/35] SWDEV-306122 - [catch2][dtest] hipGraph added negative test for hipStreamEndCapture (#2708) Change-Id: I6b617c2c52b6960b25db533d5bc50e238c83af7d --- tests/catch/unit/graph/CMakeLists.txt | 1 + tests/catch/unit/graph/hipStreamEndCapture.cc | 176 ++++++++++++++++++ 2 files changed, 177 insertions(+) create mode 100644 tests/catch/unit/graph/hipStreamEndCapture.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 430832e03e..2937060930 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -59,6 +59,7 @@ set(TEST_SRC hipStreamBeginCapture.cc hipGraphAddMemcpyNode1D.cc hipStreamGetCaptureInfo.cc + hipStreamEndCapture.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipStreamEndCapture.cc b/tests/catch/unit/graph/hipStreamEndCapture.cc new file mode 100644 index 0000000000..dd6db743a7 --- /dev/null +++ b/tests/catch/unit/graph/hipStreamEndCapture.cc @@ -0,0 +1,176 @@ +/* +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. +*/ + +/** +Negative Testcase Scenarios : +1) Pass stream as nullptr and verify there is no crash, api returns error code. +2) Pass graph as nullptr and verify there is no crash, api returns error code. +3) Pass graph as nullptr and and stream as hipStreamPerThread verify there + is no crash, api returns error code. +4) End capture on stream where capture has not yet started and verify + error code is returned. +5) Destroy stream and try to end capture. +6) Destroy Graph and try to end capture. +7) Begin capture on a thread with mode other than hipStreamCaptureModeRelaxed + and try to end capture from different thread. Expect to return + hipErrorStreamCaptureWrongThread. +*/ + +#include +#include + +TEST_CASE("Unit_hipStreamEndCapture_Negative") { + hipError_t ret; + SECTION("Pass stream as nullptr") { + hipGraph_t graph; + ret = hipStreamEndCapture(nullptr, &graph); + REQUIRE(hipErrorIllegalState == ret); + } +#if HT_NVIDIA + SECTION("Pass graph as nullptr") { + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + ret = hipStreamEndCapture(stream, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Pass graph as nullptr and stream as hipStreamPerThread") { + ret = hipStreamEndCapture(hipStreamPerThread, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + SECTION("End capture on stream where capture has not yet started") { + hipStream_t stream; + hipGraph_t graph; + HIP_CHECK(hipStreamCreate(&stream)); + ret = hipStreamEndCapture(stream, &graph); + REQUIRE(hipErrorIllegalState == ret); + HIP_CHECK(hipStreamDestroy(stream)); + } + SECTION("Destroy stream and try to end capture") { + hipStream_t stream; + hipGraph_t graph; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamDestroy(stream)); + ret = hipStreamEndCapture(stream, &graph); + REQUIRE(hipErrorContextIsDestroyed == ret); + } + SECTION("Destroy graph and try to end capture in between") { + hipStream_t stream{nullptr}; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 100000; + size_t Nbytes = N * sizeof(float); + 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); + + // 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(hipStreamCreate(&stream)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + 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(hipGraphDestroy(graph)); + ret = hipStreamEndCapture(stream, &graph); + REQUIRE(hipSuccess == ret); + + free(A_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipStreamDestroy(stream)); + } +} + +static void thread_func(hipStream_t stream, hipGraph_t graph) { + HIP_ASSERT(hipErrorStreamCaptureWrongThread == + hipStreamEndCapture(stream, &graph)); +} + +TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") { + hipStream_t stream{nullptr}; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 100000; + size_t Nbytes = N * sizeof(float); + 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); + + // 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(hipStreamCreate(&stream)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + 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)); + + std::thread t(thread_func, stream, graph); + t.join(); + +#if HT_AMD + HIP_CHECK(hipStreamEndCapture(stream, &graph)); +#endif + + free(A_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipGraphDestroy(graph)); +} + From 604de0f8170e05d134fee2b8d0aa9a7fc2e5fedb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Mon, 13 Jun 2022 11:05:02 +0100 Subject: [PATCH 09/35] EXSWCPHIPT-42 - Add HIP RTC support to the test framework (#2719) * EXSWCPHIPT-42 - Add HIP RTC support to the test framework * Removed ifdef from hipTestContext class --- tests/README.md | 16 + tests/catch/CMakeLists.txt | 8 + .../external/Catch2/cmake/Catch2/Catch.cmake | 9 + tests/catch/hipTestMain/hip_test_context.cc | 53 +++- tests/catch/include/hip_test_common.hh | 55 +++- tests/catch/include/hip_test_context.hh | 57 +++- tests/catch/include/hip_test_rtc.hh | 279 ++++++++++++++++++ tests/catch/include/kernel_mapping.hh | 27 ++ tests/catch/include/kernels.hh | 55 ++++ tests/catch/kernels/CMakeLists.txt | 8 + tests/catch/kernels/Set.cpp | 6 + tests/catch/kernels/vectorADD.inl | 10 + tests/catch/unit/event/Unit_hipEventRecord.cc | 10 +- tests/catch/unit/memory/hipHostMalloc.cc | 18 +- 14 files changed, 569 insertions(+), 42 deletions(-) create mode 100644 tests/catch/include/hip_test_rtc.hh create mode 100644 tests/catch/include/kernel_mapping.hh create mode 100644 tests/catch/include/kernels.hh create mode 100644 tests/catch/kernels/CMakeLists.txt create mode 100644 tests/catch/kernels/Set.cpp create mode 100644 tests/catch/kernels/vectorADD.inl diff --git a/tests/README.md b/tests/README.md index 33912a5726..1b1e54a2e3 100644 --- a/tests/README.md +++ b/tests/README.md @@ -147,6 +147,22 @@ For example, Here "-C performance" indicate the "performance" configuration of ctest. ``` +### RTC Testing + +To enable RTC testing, cmake needs to be passed the DRTC_TESTING=1 options. + +When this option is passed, all tests that support this functionality will be run using HIP RTC to compile and run. + +To enable HIP RTC support for a specific test: + 1 - Move all its kernels to tests/catch/kernels (one file per kernel) + 2 - Update tests/catch/kernels/CMakeLists.txt + 3 - Update tests/catch/include/kernels.hh + 4 - Update tests/catch/include/kernel_mapping.hh + 5 - Include kernels.hh + 6 - Call hipTest::launchKernel() function instead of hipLaunchKernelGGL() + +Note: HIP RTC does not do implicit casting of kernel parameters. This requires the test writer to explicitly do all the casting before running the kernel. The code will not compile otherwise. + ### If a test fails - how to debug a test Find the test and commandline that fail: diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index 40a61fd6f7..1a272d03aa 100644 --- a/tests/catch/CMakeLists.txt +++ b/tests/catch/CMakeLists.txt @@ -104,10 +104,17 @@ set(CATCH2_INCLUDE ${CATCH2_PATH}/cmake/Catch2/catch_include.cmake.in) include_directories( ${CATCH2_PATH} "./include" + "./kernels" ${HIP_PATH}/include ${JSON_PARSER} ) +option(RTC_TESTING "Run tests using HIP RTC to compile the kernels" OFF) +if (RTC_TESTING) + add_definitions(-DRTC_TESTING=ON) +endif() +add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/") + file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/hipTestMain) file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/script) set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake) @@ -177,6 +184,7 @@ add_custom_target(build_tests) # Tests folder add_subdirectory(unit) add_subdirectory(ABM) +add_subdirectory(kernels) add_subdirectory(hipTestMain) add_subdirectory(stress) add_subdirectory(TypeQualifiers) diff --git a/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake b/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake index 5287241143..62e4984afc 100644 --- a/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake @@ -203,7 +203,16 @@ function(hip_add_exe_to_target) "${list_args}" ) # Create shared lib of all tests + if(NOT RTC_TESTING) + add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $ $) + else () add_executable(${_NAME} EXCLUDE_FROM_ALL ${_TEST_SRC} $) + if(HIP_PLATFORM STREQUAL "amd") + target_link_libraries(${_NAME} hiprtc) + else() + target_link_libraries(${_NAME} nvrtc) + endif() + endif() catch_discover_tests(${_NAME} PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") if(UNIX) set(_LINKER_LIBS ${_LINKER_LIBS} stdc++fs) diff --git a/tests/catch/hipTestMain/hip_test_context.cc b/tests/catch/hipTestMain/hip_test_context.cc index ec55e2395e..a6e3a08609 100644 --- a/tests/catch/hipTestMain/hip_test_context.cc +++ b/tests/catch/hipTestMain/hip_test_context.cc @@ -23,10 +23,10 @@ void TestContext::detectPlatform() { std::string TestContext::substringFound(std::vector list, std::string filename) { std::string match = ""; - for(unsigned int i = 0; i < list.size() ; i++) { + for (unsigned int i = 0; i < list.size(); i++) { if (filename.find(list.at(i)) != std::string::npos) { - match = list.at(i); - break; + match = list.at(i); + break; } } return match; @@ -35,20 +35,19 @@ std::string TestContext::substringFound(std::vector list, std::stri std::string TestContext::getMatchingConfigFile(std::string config_dir) { std::string configFileToUse; - for(auto& p: fs::recursive_directory_iterator(config_dir)) { + for (auto& p : fs::recursive_directory_iterator(config_dir)) { fs::path filename = p.path(); std::string cur_arch = "TODO"; - std::string arch = substringFound(amd_arch_list_,filename.filename().string()); - std::string platform = substringFound(platform_list_,filename.filename().string()); - std::string os = substringFound(os_list_,filename.filename().string()); + std::string arch = substringFound(amd_arch_list_, filename.filename().string()); + std::string platform = substringFound(platform_list_, filename.filename().string()); + std::string os = substringFound(os_list_, filename.filename().string()); // if arch found then use that exit from loop - if(arch == cur_arch) { + if (arch == cur_arch) { configFileToUse = filename.string(); break; - // match the platform/os and continue to look - } else if((platform == config_.platform) && - (os == config_.os || os == "all")) { - configFileToUse = filename.string(); + // match the platform/os and continue to look + } else if ((platform == config_.platform) && (os == config_.os || os == "all")) { + configFileToUse = filename.string(); } } return configFileToUse; @@ -60,10 +59,10 @@ std::string& TestContext::getJsonFile() { config_dir = config_dir.parent_path(); int levels = 0; bool configFolderFound = false; - std::vector configList; + std::vector configList; std::string configFile; // check a max of 5 levels down the executable path - while(levels < 5) { + while (levels < 5) { fs::path temp_path = config_dir; temp_path /= "hipTestMain"; temp_path /= "config"; @@ -185,7 +184,7 @@ bool TestContext::parseJsonFile() { return false; } - const picojson::object &o = v.get(); + const picojson::object& o = v.get(); for (picojson::object::const_iterator i = o.begin(); i != o.end(); ++i) { // Processing for DisabledTests if (i->first == "DisabledTests") { @@ -196,7 +195,7 @@ bool TestContext::parseJsonFile() { for (auto ai = val.begin(); ai != val.end(); ai++) { std::string tmp = ai->get(); std::string newRegexName; - for(const auto &c : tmp) { + for (const auto& c : tmp) { if (c == '*') newRegexName += ".*"; else @@ -209,3 +208,25 @@ bool TestContext::parseJsonFile() { return true; } + +void TestContext::cleanContext() { + for (auto& pair : compiledKernels) { + REQUIRE(hipSuccess == hipModuleUnload(pair.second.module)); + } +} + +void TestContext::trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule, + hipFunction_t kernelFunction) { + rtcState state{loadedModule, kernelFunction}; + compiledKernels[kernelNameExpression] = state; +} + +hipFunction_t TestContext::getFunction(const std::string kernelNameExpression) { + auto it{compiledKernels.find(kernelNameExpression)}; + + if (it != compiledKernels.end()) { + return it->second.kernelFunction; + } else { + return nullptr; + } +} \ No newline at end of file diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index 496f202b8c..63b8354421 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021 - 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 @@ -22,6 +22,7 @@ THE SOFTWARE. #pragma once #include "hip_test_context.hh" +#include #include #include @@ -171,8 +172,60 @@ static inline void HIP_SKIP_TEST(char const* const reason) noexcept { // ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test. std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl; } + +/** + * @brief Helper template that returns the expected arguments of a kernel. + * + * @return constexpr std::tuple the expected arguments of the kernel. + */ +template std::tuple getExpectedArgs(void(FArgs...)){}; + +/** + * @brief Asserts that the types of the arguments of a function match exactly with the types in the + * function signature. + * This is necessary because HIP RTC does not do implicit casting of the kernel + * parameters. + * In order to get the kernel function signature, this function should only called when + * RTC is disabled. + * + * @tparam F the kernel function + * @tparam Args the parameters that will be passed to the kernel. + */ +template void validateArguments(F f, Args...) { + using expectedArgsTuple = decltype(getExpectedArgs(f)); + static_assert(std::is_same>::value, + "Kernel arguments types must match exactly!"); } +/** + * @brief Launch a kernel using either HIP or HIP RTC. + * + * @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a + * template). + * @tparam K The kernel type. Expects a function or template when RTC is disabled. Expects a + * function pointer instead when RTC is enabled. + * @tparam Dim Can be either dim3 or int. + * @tparam Args A list of kernel arguments to be forwarded. + * @param kernel The kernel to be launched (defined in kernels.hh) + * @param numBlocks + * @param numThreads + * @param memPerBlock + * @param stream + * @param packedArgs A list of kernel arguments to be forwarded. + */ +template +void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, + hipStream_t stream, Args&&... packedArgs) { +#ifndef RTC_TESTING + validateArguments(kernel, packedArgs...); + kernel<<>>(std::forward(packedArgs)...); +#else + launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, + std::forward(packedArgs)...); +#endif +} +} // namespace HipTest + // This must be called in the beginning of image test app's main() to indicate whether image // is supported. diff --git a/tests/catch/include/hip_test_context.hh b/tests/catch/include/hip_test_context.hh index e6429454ae..2d630b70f7 100644 --- a/tests/catch/include/hip_test_context.hh +++ b/tests/catch/include/hip_test_context.hh @@ -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 @@ -22,9 +22,11 @@ THE SOFTWARE. #pragma once #include +#include #include #include #include +#include #if defined(_WIN32) #define HT_WIN 1 @@ -57,9 +59,9 @@ static int _log_enable = (std::getenv("HT_LOG_ENABLE") ? 1 : 0); } typedef struct Config_ { - std::string json_file; // Json file - std::string platform; // amd/nvidia - std::string os; // windows/linux + std::string json_file; // Json file + std::string platform; // amd/nvidia + std::string os; // windows/linux } Config; class TestContext { @@ -69,14 +71,20 @@ class TestContext { std::string current_test; std::set skip_test; std::string json_file_; - std::vector platform_list_ = {"amd" , "nvidia"}; - std::vector os_list_ = {"windows", "linux", "all"}; - std::vector amd_arch_list_ = {}; + std::vector platform_list_ = {"amd", "nvidia"}; + std::vector os_list_ = {"windows", "linux", "all"}; + std::vector amd_arch_list_ = {}; + + struct rtcState { + hipModule_t module; + hipFunction_t kernelFunction; + }; + + std::unordered_map compiledKernels{}; Config config_; std::string& getJsonFile(); - std::string substringFound( std::vector list, - std::string filename); + std::string substringFound(std::vector list, std::string filename); void detectOS(); void detectPlatform(); void fillConfig(); @@ -86,10 +94,11 @@ class TestContext { std::string getMatchingConfigFile(std::string config_dir); const Config& getConfig() const { return config_; } + TestContext(int argc, char** argv); public: - static const TestContext& get(int argc = 0, char** argv = nullptr) { + static TestContext& get(int argc = 0, char** argv = nullptr) { static TestContext instance(argc, argv); return instance; } @@ -103,6 +112,34 @@ class TestContext { const std::string& getCurrentTest() const { return current_test; } std::string currentPath() const; + /** + * @brief Unload all loaded modules. + * Note: This function needs to be called at the end of each test that uses RTC. + * It is not possible to unload the loaded modules without adding explicit code to the end + * of each test. This function exists only to provide a clean way to exit a test when using RTC. + * However, not unloading a module explicitly shouldn't have any effect on the outcome of + * the test. + */ + void cleanContext(); + + /** + * @brief Keeps track of all the already compiled rtc kernels. + * + * @param kernelNameExpression The name expression (e.g. hipTest::vectorADD). + * @param loadedModule The loaded module. + * @param kernelFunction The hipFunction that will be used to run the kernel in the future. + */ + void trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule, + hipFunction_t kernelFunction); + + /** + * @brief Get the already compiled hip rtc kernel function if it exists. + * + * @param kernelNameExpression The name expression (e.g. hipTest::vectorADD). + * @return the hipFunction if it exists. nullptr otherwise + */ + hipFunction_t getFunction(const std::string kernelNameExpression); + TestContext(const TestContext&) = delete; void operator=(const TestContext&) = delete; }; diff --git a/tests/catch/include/hip_test_rtc.hh b/tests/catch/include/hip_test_rtc.hh new file mode 100644 index 0000000000..0860ae51db --- /dev/null +++ b/tests/catch/include/hip_test_rtc.hh @@ -0,0 +1,279 @@ +/* +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 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. +*/ + +#pragma once +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "hip/hip_runtime_api.h" +#include "hip_test_context.hh" + +namespace HipTest { + +struct KernelArgument { + const void* ptr; + size_t sizeRequirement; + size_t alignmentRequirement; +}; + +/** + * @brief Reconstructs the name expression for the kernel. + * + * @param kernelName the name of the kernel (e.g. "HipTest::VectorADD") + * @param typenames the typenames used by this kernel (e.g. "float"). + * @return std::string the reconstructed expression (e.g. "VectorADD""). Returns kernelName + * instead if the kernel is not a template. + */ +inline std::string reconstructExpression(std::string& kernelName, + std::vector& typenames) { + std::string kernelExpression = kernelName; + if (typenames.size() > 0) { + kernelExpression += "<" + typenames[0]; + for (size_t i = 1; i < typenames.size(); ++i) { + kernelExpression += "," + typenames[i]; + } + kernelExpression += ">"; + } + + return kernelExpression; +} + +/** + * @brief Packs the kernel arguments into the format expected by hipModuleLaunchKernel + * + * @param args list of arguments for the kernel and their alignemnt requirements. + * @return std::vector the packed arguments ready to be passed on to hipModuleLaunchKernel + */ +inline std::vector alignArguments(std::vector& args) { + std::vector alignedArguments{}; + int count = 0; + for (auto& arg : args) { + const char* argPtr{reinterpret_cast(arg.ptr)}; + + /* + * Details about the padding formula can be found at: + * https://en.wikipedia.org/wiki/Data_structure_alignment#Data_structure_padding + */ + int paddingNeeded = -count & (arg.alignmentRequirement - 1); + alignedArguments.insert(std::end(alignedArguments), paddingNeeded, 0); + count += paddingNeeded; + + alignedArguments.insert(std::end(alignedArguments), argPtr, argPtr + arg.sizeRequirement); + count += arg.sizeRequirement; + } + return alignedArguments; +} + +inline std::vector getKernelCode(hiprtcProgram& rtcProgram) { + size_t codeSize; + REQUIRE(HIPRTC_SUCCESS == hiprtcGetCodeSize(rtcProgram, &codeSize)); + + std::vector code(codeSize); + REQUIRE(HIPRTC_SUCCESS == hiprtcGetCode(rtcProgram, code.data())); + + return code; +} + +/** + * @brief Compiles a kernel using HIP RTC + * + * @param rtcKernel the name of the kernel to compile. + * @param kernelNameExpression the name expression to be added to the RTC program (e.g. + * HipTest::VectorADD) + * @return hiprtcProgram the compiled rtc program. + */ +inline hiprtcProgram compileRTC(std::string& rtcKernel, std::string& kernelNameExpression) { + std::string fileName = mapKernelToFileName.at(rtcKernel); + std::string filePath{KERNELS_PATH + fileName}; + + INFO("Opening Kernel File: " << filePath); + std::ifstream kernelFile{filePath}; + REQUIRE(kernelFile.is_open()); + + std::stringstream stringStream; + std::string line; + while (getline(kernelFile, line)) { + /* Skip the include directive since it is not part of the kernel */ + if (line.find("#include") != std::string::npos) { + continue; + } + stringStream << line << '\n'; + } + kernelFile.close(); + + std::string kernelCode{stringStream.str()}; + INFO("RTC Kernel Code:\n" << kernelCode) + + hiprtcProgram rtcProgram; + hiprtcCreateProgram(&rtcProgram, kernelCode.c_str(), (fileName + ".cu").c_str(), 0, nullptr, + nullptr); + + std::vector options{}; +#ifdef __HIP_PLATFORM_AMD__ + + int deviceCount; + REQUIRE(hipSuccess == hipGetDeviceCount(&deviceCount)); + + std::set architectures{}; + for (int i = 0; i < deviceCount; ++i) { + hipDeviceProp_t props; + REQUIRE(hipSuccess == hipGetDeviceProperties(&props, i)); + architectures.insert(std::string{"--gpu-architecture="} + props.gcnArchName); + } + + for (auto& architecture : architectures) { + options.push_back(architecture.c_str()); + } +#else + options.push_back("--fmad=false"); +#endif + + REQUIRE(HIPRTC_SUCCESS == hiprtcAddNameExpression(rtcProgram, kernelNameExpression.c_str())); + REQUIRE(HIPRTC_SUCCESS == hiprtcCompileProgram(rtcProgram, 1, options.data())); + + return rtcProgram; +} + +/** + * @brief Get a typename as a string + * + * @tparam T The typename + * @return std::string the string representation of T + */ +template std::string getTypeName() { + std::string name, prefix, suffix; + + +#ifdef __clang__ + name = __PRETTY_FUNCTION__; + prefix = "std::string HipTest::getTypeName() [T = "; + suffix = "]"; +#elif defined(__GNUC__) + name = __PRETTY_FUNCTION__; + prefix = "std::string HipTest::getTypeName() [with T = "; + suffix = "; std::string = std::__cxx11::basic_string]"; +#elif defined(_MSC_VER) + name = __FUNCSIG__; + prefix = "std::string __cdecl HipTest::getTypeName<"; + suffix = ">(void)"; +#endif + + return name.substr(prefix.size(), name.rfind(suffix) - prefix.size()); +} + +/** + * @brief Tells the user that the kernels are using HIP RTC. Prints only once per test. + * + */ +static inline void printInfo() { + static bool alreadyPrinted{false}; + + if (!alreadyPrinted) { + std::cout << "INFO: This test is running using HIP RTC to compile and run the kernels." + << std::endl; + alreadyPrinted = true; + } +} + +/** + * @brief Compiles and launches a kernel using HIP RTC + * + * @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a + * template). + * @tparam Args A list of kernel arguments to be forwarded. + * @param getKernelName A function wrapper that returns the name of the kernel to launch (check + * kernels.hh for more info) + * @param numBlocks + * @param numThreads + * @param memPerBlock + * @param stream + * @param packedArgs A list of kernel arguments to be forwarded. + */ +template +void launchRTCKernel(std::string (*getKernelName)(), dim3 numBlocks, dim3 numThreads, + std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) { + + printInfo(); + TestContext& testContext = TestContext::get(); + std::string kernelName = (*getKernelName)(); + + std::vector kernelTypenames{std::string(HipTest::getTypeName())...}; + std::string kernelExpression = reconstructExpression(kernelName, kernelTypenames); + + if (testContext.getFunction(kernelExpression) == nullptr) { + hiprtcProgram rtcProgram{compileRTC(kernelName, kernelExpression)}; + std::vector compiledCode{getKernelCode(rtcProgram)}; + + hipModule_t module; + + REQUIRE(hipSuccess == hipModuleLoadData(&module, compiledCode.data())); + + hipFunction_t kernelFunction; + + const char* loweredName; + REQUIRE(HIPRTC_SUCCESS == + hiprtcGetLoweredName(rtcProgram, kernelExpression.c_str(), &loweredName)); + REQUIRE(hipSuccess == hipModuleGetFunction(&kernelFunction, module, loweredName)); + + /* After obtaining the kernelFunction, the program is no longer needed. So it can be destroyed */ + REQUIRE(HIPRTC_SUCCESS == hiprtcDestroyProgram(&rtcProgram)); + + testContext.trackRtcState(kernelExpression, module, kernelFunction); + } + + hipFunction_t kernelFunction = testContext.getFunction(kernelExpression); + + std::vector args = { + {reinterpret_cast(&packedArgs), sizeof(Args), alignof(Args)}...}; + + std::vector alignedArguments{alignArguments(args)}; + size_t argumentsSize{alignedArguments.size()}; + + void* config_array[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, alignedArguments.data(), + HIP_LAUNCH_PARAM_BUFFER_SIZE, reinterpret_cast(&argumentsSize), + HIP_LAUNCH_PARAM_END}; + + REQUIRE(hipSuccess == + hipModuleLaunchKernel(kernelFunction, numBlocks.x, numBlocks.y, numBlocks.z, numThreads.x, + numThreads.y, numThreads.z, memPerBlock, stream, nullptr, + config_array)); +} + +/** + * @brief Template overload for when numBlocks and numThreads is an integer. + * + */ +template +void launchRTCKernel(std::string kernelName, int numBlocks, int numThreads, + std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) { + launchRTCKernel(kernelName, dim3(numBlocks), dim3(numThreads), memPerBlock, stream, + std::forward(packedArgs)...); +} + +} // namespace HipTest diff --git a/tests/catch/include/kernel_mapping.hh b/tests/catch/include/kernel_mapping.hh new file mode 100644 index 0000000000..32ca93b5e6 --- /dev/null +++ b/tests/catch/include/kernel_mapping.hh @@ -0,0 +1,27 @@ +/* +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 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. +*/ + +#pragma once + +#include + +const std::map mapKernelToFileName{ + {"Set", "Set.cpp"}, + {"HipTest::vectorADD", "vectorADD.inl"}, +}; \ No newline at end of file diff --git a/tests/catch/include/kernels.hh b/tests/catch/include/kernels.hh new file mode 100644 index 0000000000..ab2a92bfdb --- /dev/null +++ b/tests/catch/include/kernels.hh @@ -0,0 +1,55 @@ +/* +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 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. +*/ + +#pragma once + +#include +#include + +#ifndef RTC_TESTING + +__global__ void Set(int* Ad, int val); + +/* Kernel Templates */ +#include "vectorADD.inl" + +#else + +/* + * Wrapper Macros that create a string representation of the kernel name. + * In the case of kernel templates, a variadic template is used to ensure compatibility with + * the launchKernel template when RTC is not enabled. If the kernel is inside a namespace, use the + * "_NS" version of the Macro. + */ +#define FUNCTION_WRAPPER(param) \ + std::string param() { return #param; } +#define TEMPLATE_WRAPPER(param) \ + template std::string param() { return #param; } +#define FUNCTION_WRAPPER_NS(param, namespace) \ + std::string param() { return #namespace "::" #param; } +#define TEMPLATE_WRAPPER_NS(param, namespace) \ + template std::string param() { return #namespace "::" #param; } + +FUNCTION_WRAPPER(Set); + +namespace HipTest { +TEMPLATE_WRAPPER_NS(vectorADD, HipTest); +} + +#endif diff --git a/tests/catch/kernels/CMakeLists.txt b/tests/catch/kernels/CMakeLists.txt new file mode 100644 index 0000000000..91e1ab69ff --- /dev/null +++ b/tests/catch/kernels/CMakeLists.txt @@ -0,0 +1,8 @@ +if(NOT RTC_TESTING) + set(TEST_SRC + Set.cpp + ) + + add_library(KERNELS EXCLUDE_FROM_ALL OBJECT ${TEST_SRC}) + target_compile_options(KERNELS PUBLIC -std=c++17) +endif() diff --git a/tests/catch/kernels/Set.cpp b/tests/catch/kernels/Set.cpp new file mode 100644 index 0000000000..72ad74dde2 --- /dev/null +++ b/tests/catch/kernels/Set.cpp @@ -0,0 +1,6 @@ +#include + +__global__ void Set(int* Ad, int val) { + int tx = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tx] = val; +} \ No newline at end of file diff --git a/tests/catch/kernels/vectorADD.inl b/tests/catch/kernels/vectorADD.inl new file mode 100644 index 0000000000..4be89b6e89 --- /dev/null +++ b/tests/catch/kernels/vectorADD.inl @@ -0,0 +1,10 @@ +namespace HipTest { +template __global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < NELEM; i += stride) { + C_d[i] = A_d[i] + B_d[i]; + } +} +} \ No newline at end of file diff --git a/tests/catch/unit/event/Unit_hipEventRecord.cc b/tests/catch/unit/event/Unit_hipEventRecord.cc index c3dc1a6086..d2f13d6e12 100644 --- a/tests/catch/unit/event/Unit_hipEventRecord.cc +++ b/tests/catch/unit/event/Unit_hipEventRecord.cc @@ -23,8 +23,8 @@ THE SOFTWARE. // Through manual inspection of the reported timestamps, can determine if recording a NULL event // forces synchronization : set #include -#include - +#include +#include #include TEST_CASE("Unit_hipEventRecord") { @@ -61,8 +61,8 @@ TEST_CASE("Unit_hipEventRecord") { // Record the start event HIP_CHECK(hipEventRecord(start, NULL)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, - static_cast(A_d), static_cast(B_d), C_d, N); + HipTest::launchKernel(HipTest::vectorADD, blocks, threadsPerBlock, 0, 0, +static_cast(A_d), static_cast(B_d), C_d, N); HIP_CHECK(hipEventRecord(stop, NULL)); HIP_CHECK(hipEventSynchronize(stop)); @@ -87,5 +87,5 @@ TEST_CASE("Unit_hipEventRecord") { HIP_CHECK(hipEventDestroy(stop)); HipTest::checkVectorADD(A_h, B_h, C_h, N, true); - + TestContext::get().cleanContext(); } diff --git a/tests/catch/unit/memory/hipHostMalloc.cc b/tests/catch/unit/memory/hipHostMalloc.cc index 70459cf9d2..94f943c6bf 100644 --- a/tests/catch/unit/memory/hipHostMalloc.cc +++ b/tests/catch/unit/memory/hipHostMalloc.cc @@ -30,8 +30,9 @@ This testfile verifies the following scenarios of hipHostMalloc API */ #include -#include +#include #include +#include #define SYNC_EVENT 0 #define SYNC_STREAM 1 @@ -41,11 +42,6 @@ std::vector syncMsg = {"event", "stream", "device"}; static constexpr int numElements{1024 * 16}; static constexpr size_t sizeBytes{numElements * sizeof(int)}; -__global__ void Set(int* Ad, int val) { - int tx = threadIdx.x + blockIdx.x * blockDim.x; - Ad[tx] = val; -} - void CheckHostPointer(int numElements, int* ptr, unsigned eventFlags, int syncMethod, std::string msg) { std::cerr << "test: CheckHostPointer " @@ -70,10 +66,10 @@ void CheckHostPointer(int numElements, int* ptr, unsigned eventFlags, const int expected = 13; // Init array to know state: - hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42); + HipTest::launchKernel(Set, dimGrid, dimBlock, 0, 0x0, ptr, -42); HIP_CHECK(hipDeviceSynchronize()); - hipLaunchKernelGGL(Set, dimGrid, dimBlock, 0, s, ptr, expected); + HipTest::launchKernel(Set, dimGrid, dimBlock, 0, s, ptr, expected); HIP_CHECK(hipEventRecord(e, s)); // Host waits for event : @@ -137,9 +133,9 @@ TEST_CASE("Unit_hipHostMalloc_Basic") { dim3 dimGrid(LEN / 512, 1, 1); dim3 dimBlock(512, 1, 1); - hipLaunchKernelGGL(HipTest::vectorADD, dimGrid, dimBlock, + HipTest::launchKernel(HipTest::vectorADD, dimGrid, dimBlock, 0, 0, static_cast(A_d), - static_cast(B_d), C_d, LEN); + static_cast(B_d), C_d, static_cast(LEN)); HIP_CHECK(hipMemcpy(C_h, C_d, LEN*sizeof(float), hipMemcpyDeviceToHost)); HIP_CHECK(hipDeviceSynchronize()); @@ -148,6 +144,7 @@ TEST_CASE("Unit_hipHostMalloc_Basic") { HIP_CHECK(hipHostFree(A_h)); HIP_CHECK(hipHostFree(B_h)); HIP_CHECK(hipHostFree(C_h)); + TestContext::get().cleanContext(); } } /* @@ -185,6 +182,7 @@ TEST_CASE("Unit_hipHostMalloc_NonCoherent") { SYNC_STREAM, ptrType); CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_EVENT, ptrType); + TestContext::get().cleanContext(); } /* From ab6fc6dfdeca64352a72c6c1241774e036d203f2 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 15:35:15 +0530 Subject: [PATCH 10/35] SWDEV-334659 - Fixes build error because of an unused variable in the test (#2724) Change-Id: Ib1b39ee57efc49217e58332abf843b44b575d880 --- tests/catch/unit/graph/hipGraphAddHostNode.cc | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphAddHostNode.cc b/tests/catch/unit/graph/hipGraphAddHostNode.cc index 596a57d755..f66e813352 100644 --- a/tests/catch/unit/graph/hipGraphAddHostNode.cc +++ b/tests/catch/unit/graph/hipGraphAddHostNode.cc @@ -56,14 +56,16 @@ static void __global__ vector_square(int *B_d, int *D_d) { } } 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. + // The callback func is hipHostFn_t which is + // of type void (*)(void*). This test is designed to + // work with global variables, hence the workaround to + // print this *ptr value to avoid type mismatch errors. 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]); + INFO("Ignore this garbage value" << *A); REQUIRE(false); } } From ab1b5f9d375abf78e959fcce485d844cf8943927 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 15:35:33 +0530 Subject: [PATCH 11/35] SWDEV-336378 - cmake support for vulkan example (#2725) Change-Id: I5ac28d2f57c1732e243674aa904cabe23be34797 --- .../2_Cookbook/20_hip_vulkan/CMakeLists.txt | 97 ++++++++++++ .../20_hip_vulkan/SineWaveSimulation.cpp | 147 ++++++++++++++++++ samples/2_Cookbook/20_hip_vulkan/buildcmd.txt | 7 +- samples/2_Cookbook/20_hip_vulkan/main.cpp | 16 +- 4 files changed, 264 insertions(+), 3 deletions(-) create mode 100644 samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt create mode 100644 samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp diff --git a/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt b/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt new file mode 100644 index 0000000000..288c47c3d2 --- /dev/null +++ b/samples/2_Cookbook/20_hip_vulkan/CMakeLists.txt @@ -0,0 +1,97 @@ + +# Copyright (c) 2020 - 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 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. + +# hipcc.bat fails to qualify as a valid compiler for CMAKE_CXX_COMPILER_ID = ROCMClang +# so the simple compiler test is skipped and forced to use hipcc.bat as compiler +set(CMAKE_C_COMPILER_WORKS 1) +set(CMAKE_CXX_COMPILER_WORKS 1) +set(CMAKE_CXX_STANDARD 14) +project(hipVulkan) + +cmake_minimum_required(VERSION 3.10) +set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake;${CMAKE_MODULE_PATH}") + +if (NOT DEFINED ROCM_PATH ) + set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." ) +endif () + +# Search for rocm in common locations +list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) + +# need to set rocm_path for windows +# since clang and hip are two different folders during build/install step +if (WIN32 AND HIPINFO_INTERNAL_BUILD) + set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --rocm-path=${HIP_PATH}") +endif() + + + +# Find hip +find_package(hip REQUIRED) +if (WIN32) +find_package(GLFW3) +if(NOT GLFW_FOUND) + if(EXISTS "${GLFW_PATH}") + message(STATUS "FOUND GLFW SDK: ${GLFW_PATH}") + elseif (EXISTS "$ENV{GLFW_PATH}") + message(STATUS "FOUND GLFW SDK: $ENV{GLFW_PATH}") + set(GLFW_PATH $ENV{GLFW_PATH}) + else() + message("Error: Unable to locate GLFW SDK. please specify GLFW_PATH") + return() + endif() +endif() +endif(WIN32) +find_package(Vulkan) +if(NOT Vulkan_FOUND) + if(EXISTS "${VULKAN_PATH}") + message(STATUS "Vulkan SDK: ${VULKAN_PATH}") + elseif (EXISTS "$ENV{VULKAN_SDK}") + message(STATUS "FOUND VULKAN SDK: $ENV{VULKAN_SDK}") + set(VULKAN_PATH $ENV{VULKAN_SDK}) + else() + message("Error: Unable to locate Vulkan SDK. please specify VULKAN_PATH") + return() + endif() +endif() +set(VULKAN_PATH ${Vulkan_INCLUDE_DIRS}) +STRING(REGEX REPLACE "/[Ii]nclude" "" VULKAN_PATH ${VULKAN_PATH}) +# Include Vulkan header files from Vulkan SDK +include_directories(AFTER ${VULKAN_PATH}/include) +link_directories(${VULKAN_PATH}/bin;${VULKAN_PATH}/lib;) +link_directories(${GLFW_PATH}/lib-vc2019) +# Set compiler and linker +set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE}) +set(CMAKE_BUILD_TYPE Release) + +# Create the excutable +add_executable(hipVulkan VulkanBaseApp.cpp VulkanBaseApp.h main.cpp SineWaveSimulation.cpp SineWaveSimulation.h linmath.h) +include_directories(${HIP_PATH}/include) +include_directories(${GLFW_PATH}/include) + +# Link with HIP +if (WIN32) + target_link_libraries(hipVulkan advapi32 hip::host vulkan-1 glfw3dll) +else (WIN32) + target_link_libraries(hipVulkan hip::host vulkan glfw) +endif (WIN32) + diff --git a/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp b/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp new file mode 100644 index 0000000000..41adae8065 --- /dev/null +++ b/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.cpp @@ -0,0 +1,147 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#include "SineWaveSimulation.h" +#include +//#include +#include "hip/hip_runtime.h" + + +__global__ void sinewave(float *heightMap, unsigned int width, unsigned int height, float time) +{ + const float freq = 4.0f; + const size_t stride = gridDim.x * blockDim.x; + + // Iterate through the entire array in a way that is + // independent of the grid configuration + for (size_t tid = blockIdx.x * blockDim.x + threadIdx.x; tid < width * height; tid += stride) { + // Calculate the x, y coordinates + const size_t y = tid / width; + const size_t x = tid - y * width; + // Normalize x, y to [0,1] + const float u = ((2.0f * x) / width) - 1.0f; + const float v = ((2.0f * y) / height) - 1.0f; + // Calculate the new height value + const float w = 0.5f * sinf(u * freq + time) * cosf(v * freq + time); + // Store this new height value + heightMap[tid] = w; + } +} + +SineWaveSimulation::SineWaveSimulation(size_t width, size_t height) + : m_heightMap(nullptr), m_width(width), m_height(height) +{ +} + +void SineWaveSimulation::initCudaLaunchConfig(int device) +{ + hipDeviceProp_t prop = {}; + checkHIPErrors(hipSetDevice(device)); + checkHIPErrors(hipGetDeviceProperties(&prop, device)); + + // We don't need large block sizes, since there's not much inter-thread communication + m_threads = prop.warpSize; + + // Use the occupancy calculator and fill the gpu as best as we can + checkHIPErrors(hipOccupancyMaxActiveBlocksPerMultiprocessor(&m_blocks, sinewave, prop.warpSize, 0)); + m_blocks *= prop.multiProcessorCount; + + // Go ahead and the clamp the blocks to the minimum needed for this height/width + m_blocks = std::min(m_blocks, (int)((m_width * m_height + m_threads - 1) / m_threads)); +} + +int SineWaveSimulation::initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE) +{ + int current_device = 0; + int device_count = 0; + int devices_prohibited = 0; + + hipDeviceProp_t deviceProp; + checkHIPErrors(hipGetDeviceCount(&device_count)); + + if (device_count == 0) { + fprintf(stderr, "CUDA error: no devices supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + + // Find the GPU which is selected by Vulkan + while (current_device < device_count) { + hipGetDeviceProperties(&deviceProp, current_device); + + if ((deviceProp.computeMode != hipComputeModeProhibited)) { + // Compare the cuda device UUID with vulkan UUID + // FIXME + int ret = 0; // memcmp((void*)&deviceProp.uuid, vkDeviceUUID, UUID_SIZE); + if (ret == 0) + { + checkHIPErrors(hipSetDevice(current_device)); + checkHIPErrors(hipGetDeviceProperties(&deviceProp, current_device)); + printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n", + current_device, deviceProp.name, deviceProp.major, + deviceProp.minor); + + return current_device; + } + + } else { + devices_prohibited++; + } + + current_device++; + } + + if (devices_prohibited == device_count) { + fprintf(stderr, + "HIP error:" + " No Vulkan-HIP Interop capable GPU found.\n"); + exit(EXIT_FAILURE); + } + + return -1; +} + +SineWaveSimulation::~SineWaveSimulation() +{ + m_heightMap = NULL; +} + +void SineWaveSimulation::initSimulation(float *heights) +{ + m_heightMap = heights; +} + +void SineWaveSimulation::stepSimulation(float time, hipStream_t stream) +{ + hipLaunchKernelGGL(sinewave, dim3(m_blocks), dim3(m_threads), 0, stream , m_heightMap, m_width, m_height, time); + getLastHIPError("Failed to launch CUDA simulation"); + //hipStreamSynchronize(stream); +} diff --git a/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt b/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt index b847596510..791466ec42 100644 --- a/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt +++ b/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt @@ -5,7 +5,12 @@ o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.vert -V -o vert.spv o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.frag -V -o frag.spv +to build without cmake: • set HCC_AMDGPU_TARGET=gfx906:sramecc-:xnack- (for your graphic card, you can get the name from hipinfo ) -$• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14 +• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14 • run a.exe, you should see a 3D sinewave simulation +to build with cmake on windows: +• mkdir build; cd build +• cmake.exe -GNinja -DCMAKE_CXX_COMPILER_ID=ROCMClang -DCMAKE_C_COMPILER_ID=ROCMClang -DCMAKE_PREFIX_PATH=d:\driver2\drivers\drivers\compute\hip_sdk + diff --git a/samples/2_Cookbook/20_hip_vulkan/main.cpp b/samples/2_Cookbook/20_hip_vulkan/main.cpp index 0788243e52..93c8db75db 100644 --- a/samples/2_Cookbook/20_hip_vulkan/main.cpp +++ b/samples/2_Cookbook/20_hip_vulkan/main.cpp @@ -37,7 +37,7 @@ #include #include #include "linmath.h" -#include "hip_runtime.h" +#include "hip/hip_runtime.h" #include "SineWaveSimulation.h" @@ -50,6 +50,18 @@ std::string execution_path; #define ENABLE_VALIDATION (true) #endif +#ifndef _WIN64 +#define MAX_PATH 260 + +int GetModuleFileName(void* hndl, char* name, int size) +{ + FILE* stream = fopen("/proc/self/cmdline", "r"); + fgets(name, size, stream); + fclose(stream); + return strlen(name); +} +#endif + class VulkanCudaSineWave : public VulkanBaseApp { @@ -93,7 +105,7 @@ public: throw std::runtime_error("Requested height and width is too large for this sample!"); } // Add our compiled vulkan shader files - TCHAR buffer[MAX_PATH] = { 0 }; + char buffer[MAX_PATH] = { 0 }; //assuming none unicode GetModuleFileName(NULL, buffer, MAX_PATH); std::string str3 = std::string(buffer); std::string str1 = "vert.spv" ; //sdkFindFilePath("sinewave.vert", execution_path.c_str()); From ba893e877968b94428a02384ed390454e669d36d Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 19:26:48 +0530 Subject: [PATCH 12/35] SWDEV-311271 - Add basic tests for memory pool (#2618) Change-Id: Ib33fa0b867482c40c6d66fb699cdd462ff43eed6 --- tests/catch/unit/memory/CMakeLists.txt | 1 + tests/catch/unit/memory/hipMemPoolApi.cc | 545 +++++++++++++++++++++++ 2 files changed, 546 insertions(+) create mode 100644 tests/catch/unit/memory/hipMemPoolApi.cc diff --git a/tests/catch/unit/memory/CMakeLists.txt b/tests/catch/unit/memory/CMakeLists.txt index 70298c246e..fb5c5677e6 100644 --- a/tests/catch/unit/memory/CMakeLists.txt +++ b/tests/catch/unit/memory/CMakeLists.txt @@ -59,6 +59,7 @@ set(TEST_SRC hipMemRangeGetAttribute.cc hipMemcpyFromSymbol.cc hipPtrGetAttribute.cc + hipMemPoolApi.cc hipMemcpyPeer.cc hipMemcpyPeerAsync.cc hipMemcpyWithStream.cc diff --git a/tests/catch/unit/memory/hipMemPoolApi.cc b/tests/catch/unit/memory/hipMemPoolApi.cc new file mode 100644 index 0000000000..e799c4feec --- /dev/null +++ b/tests/catch/unit/memory/hipMemPoolApi.cc @@ -0,0 +1,545 @@ +/* + Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +/* Test Case Description: + 1) This testcase verifies the basic scenario - supported on + all devices +*/ + +#include +#include +#include +#include +#include +#include +#include +#include + +constexpr hipMemPoolProps kPoolProps = { + hipMemAllocationTypePinned, + hipMemHandleTypeNone, + { + hipMemLocationTypeDevice, + 0 + }, + nullptr, + {0} +}; + +/* + This testcase verifies HIP Mem Pool API basic scenario - supported on all devices + */ + +TEST_CASE("Unit_hipMemPoolApi_Basic") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + int numElements = 64 * 1024 * 1024; + float *A = nullptr, *B = nullptr; + + hipMemPool_t mem_pool = nullptr; + int device = 0; + HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, device)); + HIP_CHECK(hipDeviceSetMemPool(device, mem_pool)); + HIP_CHECK(hipDeviceGetMemPool(&mem_pool, device)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipMallocAsync(&A, numElements * sizeof(float), stream)); + INFO("hipMallocAsync result: " << A); + + HIP_CHECK(hipFreeAsync(A, stream)); + // Reset the default memory pool usage for the following tests + hipMemPoolAttr attr = hipMemPoolAttrUsedMemHigh; + std::uint64_t value64 = 0; + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value64)); + + size_t min_bytes_to_hold = 1024 * 1024; + HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold)); + + attr = hipMemPoolReuseFollowEventDependencies; + int value = 0; + + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); + + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value)); + + hipMemAccessDesc desc_list = {}; + int count = 1; + HIP_CHECK(hipMemPoolSetAccess(mem_pool, &desc_list, count)); + + hipMemAccessFlags flags = hipMemAccessFlagsProtNone; + hipMemLocation location = {}; + HIP_CHECK(hipMemPoolGetAccess(&flags, mem_pool, &location)); + + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + HIP_CHECK(hipMallocFromPoolAsync(&B, numElements * sizeof(float), mem_pool, stream)); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + + HIP_CHECK(hipStreamDestroy(stream)); +} + +constexpr auto wait_ms = 500; + +__global__ void kernel500ms(float* hostRes, int clkRate) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + hostRes[tid] = tid + 1; + __threadfence_system(); + // expecting that the data is getting flushed to host here! + uint64_t start = clock64()/clkRate, cur; + if (clkRate > 1) { + do { cur = clock64()/clkRate-start;}while (cur < wait_ms); + } else { + do { cur = clock64()/start;}while (cur < wait_ms); + } +} + +TEST_CASE("Unit_hipMemPoolApi_BasicAlloc") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + + float* B, *C; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&B), numElements * sizeof(float), mem_pool, stream)); + + numElements = 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); + + int blocks = 1024; + int clkRate; + hipMemPoolAttr attr; + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + + kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); + + attr = hipMemPoolAttrReservedMemCurrent; + std::uint64_t res_before_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_before_sync)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + std::uint64_t res_after_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_sync)); + // Sync must releaae memory to OS + REQUIRE(res_after_sync < res_before_sync); + + int value = 0; + attr = hipMemPoolReuseFollowEventDependencies; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value)); + // Default enabled + REQUIRE(1 == value); + + attr = hipMemPoolReuseAllowOpportunistic; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value)); + // Default enabled + REQUIRE(1 == value); + + attr = hipMemPoolReuseAllowInternalDependencies; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value)); + // Default enabled + REQUIRE(1 == value); + + attr = hipMemPoolAttrReleaseThreshold; + std::uint64_t value64 = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Default is 0 + REQUIRE(0 == value64); + + attr = hipMemPoolAttrReservedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Must be bigger than current + REQUIRE(value64 > res_after_sync); + + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current usage query works - just small buffer left + REQUIRE(sizeof(float) * 1024 == value64); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the high watermark usage works - the both buffers must be reported + REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipMemPoolApi_BasicTrim") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + + float* B, *C; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&B), numElements * sizeof(float), mem_pool, stream)); + + numElements = 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); + + int blocks = 2; + int clkRate; + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + + kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); + + hipMemPoolAttr attr; + attr = hipMemPoolAttrReleaseThreshold; + // The pool must hold 128MB + std::uint64_t threshold = 128 * 1024 * 1024; + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &threshold)); + + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); + + // Get reserved memory before trim + attr = hipMemPoolAttrReservedMemCurrent; + std::uint64_t res_before_trim = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_before_trim)); + + size_t min_bytes_to_hold = sizeof(float) * 1024; + HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold)); + + std::uint64_t res_after_trim = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_trim)); + // Trim must be a nop because execution isn't done + REQUIRE(res_before_trim == res_after_trim); + + HIP_CHECK(hipStreamSynchronize(stream)); + + std::uint64_t res_after_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_sync)); + // Since hipMemPoolAttrReleaseThreshold is 128 MB sync does nothing to the freed memory + REQUIRE(res_after_trim == res_after_sync); + + HIP_CHECK(hipMemPoolTrimTo(mem_pool, min_bytes_to_hold)); + + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &res_after_trim)); + // Validate memory after real trim. The pool must hold less memory than before + REQUIRE(res_after_trim < res_after_sync); + + attr = hipMemPoolAttrReleaseThreshold; + std::uint64_t value64 = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the threshold query works + REQUIRE(threshold == value64); + + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current usage query works - just small buffer left + REQUIRE(sizeof(float) * 1024 == value64); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the high watermark usage works - the both buffers must be reported + REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipMemPoolApi_BasicReuse") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + + float *A, *B, *C; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + + numElements = 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); + + int blocks = 2; + int clkRate; + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + + kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); + + hipMemPoolAttr attr; + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + + numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&B), numElements * sizeof(float), mem_pool, stream)); + // Runtime must reuse the pointer + REQUIRE(A == B); + + // Make a sync before the second kernel launch to make sure memory B isn't gone + HIP_CHECK(hipStreamSynchronize(stream)); + + // Second kernel launch with new memory + kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); + + HIP_CHECK(hipStreamSynchronize(stream)); + + attr = hipMemPoolAttrUsedMemCurrent; + std::uint64_t value64 = 0; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current usage reports the both buffers + REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the high watermark usage works - the both buffers must be reported + REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current usage reports just one buffer, because the above free doesn't hold memory + REQUIRE(sizeof(float) * 1024 == value64); + + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipMemPoolApi_Opportunistic") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + + hipMemPoolAttr attr; + int blocks = 2; + int clkRate; + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + + float *A, *B, *C; + hipStream_t stream, stream2; + + // Create 2 async non-blocking streams + HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + HIP_CHECK(hipStreamCreateWithFlags(&stream2, hipStreamNonBlocking)); + + size_t numElements = 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&C), numElements * sizeof(float), mem_pool, stream)); + int value = 0; + + SECTION("Disallow Opportunistic - No Reuse") { + numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + + // Disable all default pool states + attr = hipMemPoolReuseFollowEventDependencies; + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); + attr = hipMemPoolReuseAllowOpportunistic; + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); + attr = hipMemPoolReuseAllowInternalDependencies; + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); + + // Run kernel for 500 ms in the first stream + kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); + + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + + // Sleep for 1 second GPU should be idle by now + std::this_thread::sleep_for(std::chrono::milliseconds(1000)); + + numElements = 8 * 1024 * 1024; + // Allocate memory for the second stream + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&B), numElements * sizeof(float), mem_pool, stream2)); + // Without Opportunistic state runtime must allocate another buffer + REQUIRE(A != B); + + // Run kernel with the new memory in the second stream + kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate); + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamSynchronize(stream2)); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream2)); + } + + SECTION("Allow Opportunistic - Reuse") { + numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + + value = 1; + attr = hipMemPoolReuseAllowOpportunistic; + // Enable Opportunistic + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); + + // Run kernel for 500 ms in the first stream + kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); + + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + + // Sleep for 1 second GPU should be idle by now + std::this_thread::sleep_for(std::chrono::milliseconds(1000)); + + numElements = 8 * 1024 * 1024; + // Allocate memory for the second stream + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&B), numElements * sizeof(float), mem_pool, stream2)); + // With Opportunistic state runtime will reuse freed buffer A + REQUIRE(A == B); + + // Run kernel with the new memory in the second stream + kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate); + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamSynchronize(stream2)); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream2)); + } + + SECTION("Allow Opportunistic - No Reuse") { + numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A), numElements * sizeof(float), mem_pool, stream)); + + value = 1; + attr = hipMemPoolReuseAllowOpportunistic; + // Enable Opportunistic + HIP_CHECK(hipMemPoolSetAttribute(mem_pool, attr, &value)); + + // Run kernel for 500 ms in the first stream + kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); + + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + + numElements = 8 * 1024 * 1024; + // Allocate memory for the second stream + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&B), numElements * sizeof(float), mem_pool, stream2)); + // With Opportunistic state runtime can't reuse freed buffer A, because it's still busy with the kernel + REQUIRE(A != B); + + // Run kernel with the new memory in the second stream + kernel500ms<<<32, blocks, 0, stream2>>>(B, clkRate); + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamSynchronize(stream2)); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream2)); + } + + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamDestroy(stream2)); +} + +TEST_CASE("Unit_hipMemPoolApi_Default") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool; + HIP_CHECK(hipDeviceGetDefaultMemPool(&mem_pool, 0)); + + float *A, *B, *C; + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + size_t numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&A), numElements * sizeof(float), stream)); + + numElements = 1024; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&C), numElements * sizeof(float), stream)); + + int blocks = 2; + int clkRate; + HIP_CHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + + kernel500ms<<<32, blocks, 0, stream>>>(A, clkRate); + + hipMemPoolAttr attr; + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(A), stream)); + + numElements = 8 * 1024 * 1024; + HIP_CHECK(hipMallocAsync(reinterpret_cast(&B), numElements * sizeof(float), stream)); + // Runtime must reuse the pointer + REQUIRE(A == B); + + // Make a sync before the second kernel launch to make sure memory B isn't gone + HIP_CHECK(hipStreamSynchronize(stream)); + + // Second kernel launch with new memory + kernel500ms<<<32, blocks, 0, stream>>>(B, clkRate); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(B), stream)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + std::uint64_t value64 = 0; + attr = hipMemPoolAttrReservedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current reserved just 4KB alloc + REQUIRE(sizeof(float) * 1024 == value64); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the high watermark usage works - the both buffers must be reported + REQUIRE(sizeof(float) * (8 * 1024 * 1024 + 1024) == value64); + + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mem_pool, attr, &value64)); + // Make sure the current usage reports just one buffer, because the above free doesn't hold memory + REQUIRE(sizeof(float) * 1024 == value64); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(C), stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} From 1c7d42064b36154b59bae6a89bfdc480cdb14282 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 13 Jun 2022 19:27:08 +0530 Subject: [PATCH 13/35] SWDEV-332522 - array test cases for streamOpsWrite & streamOpsWait (#2659) Change-Id: I746c3bae592e0251b8c9482d28bd21e1bd094256 --- .../streamOperations/hipstream_operations.cpp | 32 +++++++++++++++++-- 1 file changed, 30 insertions(+), 2 deletions(-) diff --git a/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp b/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp index f47328d573..cd4d9fd425 100644 --- a/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp +++ b/tests/src/runtimeApi/streamOperations/hipstream_operations.cpp @@ -58,6 +58,7 @@ constexpr unsigned int writeFlag = 0; constexpr float SLEEP_MS = 100; constexpr uint32_t DATA_INIT = 0x1234; constexpr uint32_t DATA_UPDATE = 0X4321; +constexpr unsigned int ARR_SIZE = 5; struct TEST_WAIT { int compareOp; @@ -203,31 +204,49 @@ void testWrite() { uint64_t* host_ptr64 = (uint64_t *) malloc(sizeof(uint64_t)); uint32_t* host_ptr32 = (uint32_t *) malloc(sizeof(uint32_t)); + uint64_t h_mem64[ARR_SIZE]; + uint32_t h_mem32[ARR_SIZE]; + uint64_t* h_mem64ptr = h_mem64; + uint32_t* h_mem32ptr = h_mem32; std::cout << " hipStreamWriteValue: testing ... \n"; HIPCHECK(hipExtMallocWithFlags((void **)&signalPtr, 8, hipMallocSignalMemory)); void* device_ptr64; void* device_ptr32; - + uint64_t* d_mem64ptr; + uint32_t* d_mem32ptr; *host_ptr64 = 0x0; *host_ptr32 = 0x0; *signalPtr = 0x0; hipHostRegister(host_ptr64, sizeof(uint64_t), 0); hipHostRegister(host_ptr32, sizeof(uint32_t), 0); + hipHostRegister(h_mem64ptr, sizeof(uint64_t) * ARR_SIZE, 0); + hipHostRegister(h_mem32ptr, sizeof(uint32_t) * ARR_SIZE, 0); // Test writting registered host pointer HIPCHECK(hipStreamWriteValue64(stream, host_ptr64, value64, writeFlag)); HIPCHECK(hipStreamWriteValue32(stream, host_ptr32, value32, writeFlag)); + //test writting to an array + for (int indx = 0; indx < ARR_SIZE; indx++) { + HIPCHECK(hipStreamWriteValue64(stream, h_mem64ptr + indx, ARR_SIZE - indx, writeFlag)); + HIPCHECK(hipStreamWriteValue32(stream, h_mem32ptr + indx, ARR_SIZE - indx, writeFlag)); + } + hipStreamSynchronize(stream); HIPASSERT(*host_ptr64 == value64); HIPASSERT(*host_ptr32 == value32); - + for (int indx = 0; indx < ARR_SIZE ; indx++) { + HIPASSERT(*(h_mem64ptr + indx) == (ARR_SIZE - indx)); + HIPASSERT(*(h_mem32ptr + indx) == (ARR_SIZE - indx)); + } // Test writting device pointer hipHostGetDevicePointer((void**)&device_ptr64, host_ptr64, 0); hipHostGetDevicePointer((void**)&device_ptr32, host_ptr32, 0); + hipHostGetDevicePointer((void**)&d_mem64ptr, h_mem64ptr, 0); + hipHostGetDevicePointer((void**)&d_mem32ptr, h_mem32ptr, 0); // Reset values *host_ptr64 = 0x0; @@ -235,10 +254,19 @@ void testWrite() { HIPCHECK(hipStreamWriteValue64(stream, device_ptr64, value64, writeFlag)); HIPCHECK(hipStreamWriteValue32(stream, device_ptr32, value32, writeFlag)); + for (int indx = 0; indx < ARR_SIZE; indx++) { + HIPCHECK(hipStreamWriteValue64(stream, d_mem64ptr + indx, indx, writeFlag)); + HIPCHECK(hipStreamWriteValue32(stream, d_mem32ptr + indx, indx, writeFlag)); + } + hipStreamSynchronize(stream); HIPASSERT(*host_ptr64 == value64); HIPASSERT(*host_ptr32 == value32); + for (int indx = 0; indx < ARR_SIZE ; indx++) { + HIPASSERT(*(h_mem64ptr + indx) == indx); + HIPASSERT(*(h_mem32ptr + indx) == indx); + } // Test Writing to Signal Memory HIPCHECK(hipStreamWriteValue64(stream, signalPtr, value64, writeFlag)); From 70d43137e08f190bad38077538105805b6741137 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 01:52:10 +0530 Subject: [PATCH 14/35] SWDEV-306122 - [catch2][dtest] Added est for hipStreamIsCapturing API. (#2715) Change-Id: I9e85d393c51962f97c5e5399556ef823b3e33782 --- tests/catch/unit/graph/CMakeLists.txt | 1 + .../catch/unit/graph/hipStreamIsCapturing.cc | 215 ++++++++++++++++++ 2 files changed, 216 insertions(+) create mode 100644 tests/catch/unit/graph/hipStreamIsCapturing.cc diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 2937060930..036e9477f7 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -58,6 +58,7 @@ set(TEST_SRC hipGraphExecMemcpyNodeSetParams.cc hipStreamBeginCapture.cc hipGraphAddMemcpyNode1D.cc + hipStreamIsCapturing.cc hipStreamGetCaptureInfo.cc hipStreamEndCapture.cc ) diff --git a/tests/catch/unit/graph/hipStreamIsCapturing.cc b/tests/catch/unit/graph/hipStreamIsCapturing.cc new file mode 100644 index 0000000000..090348a25b --- /dev/null +++ b/tests/catch/unit/graph/hipStreamIsCapturing.cc @@ -0,0 +1,215 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +constexpr unsigned blocks = 512; +constexpr unsigned threadsPerBlock = 256; +constexpr size_t N = 100000; +constexpr size_t Nbytes = N * sizeof(float); + +/** +API - hipStreamIsCapturing +Negative Testcase Scenarios : Negative + 1) Check capture status with null pCaptureStatus. + 2) Check capture status with hipStreamPerThread and null pCaptureStatus. +Functional Testcase Scenarios : + 1) Check capture status with null stream. + 2) Check capture status with hipStreamPerThread. + 3) Functional : Create a stream, call api and check + capture status is hipStreamCaptureStatusNone. + 4) Functional : Start capturing a stream and check + capture status returned as hipStreamCaptureStatusActive. + 5) Functional : Stop capturing a stream and check + status is returned as hipStreamCaptureStatusNone. + 6) Functional : Use hipStreamPerThread, call api and check + capture status is hipStreamCaptureStatusNone. + 7) Functional : Start capturing using hipStreamPerThread and check + capture status returned as hipStreamCaptureStatusActive. + 8) Functional : Stop capturing using hipStreamPerThread and check + status is returned as hipStreamCaptureStatusNone. +*/ + + +TEST_CASE("Unit_hipStreamIsCapturing_Negative") { + hipError_t ret; + hipStream_t stream{}; + + SECTION("Check capture status with null pCaptureStatus.") { + ret = hipStreamIsCapturing(stream, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Check capture status with hipStreamPerThread and" + " nullptr as pCaptureStatus.") { + ret = hipStreamIsCapturing(hipStreamPerThread, nullptr); + REQUIRE(hipErrorInvalidValue == ret); + } +} + +TEST_CASE("Unit_hipStreamIsCapturing_Functional_Basic") { + hipStreamCaptureStatus cStatus; + + SECTION("Check capture status with null stream.") { + HIP_CHECK(hipStreamIsCapturing(nullptr, &cStatus)); + REQUIRE(hipStreamCaptureStatusNone == cStatus); + } + SECTION("Check capture status with hipStreamPerThread.") { + HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus)); + REQUIRE(hipStreamCaptureStatusNone == cStatus); + } +} + +/** +Testcase Scenarios : + 1) Functional : Create a stream, call api and check + capture status is hipStreamCaptureStatusNone. + 2) Functional : Start capturing a stream and check + capture status returned as hipStreamCaptureStatusActive. + 3) Functional : Stop capturing a stream and check + status is returned as hipStreamCaptureStatusNone. +*/ + +TEST_CASE("Unit_hipStreamIsCapturing_Functional") { + float *A_d, *C_d; + float *A_h, *C_h; + hipStream_t stream{nullptr}; + hipGraph_t graph{nullptr}; + hipStreamCaptureStatus cStatus; + + 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(hipStreamCreate(&stream)); + + SECTION("Check the stream capture status before start capturing.") { + HIP_CHECK(hipStreamIsCapturing(stream, &cStatus)); + REQUIRE(hipStreamCaptureStatusNone == cStatus); + } + + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + + SECTION("Start capturing a stream and check the status.") { + HIP_CHECK(hipStreamIsCapturing(stream, &cStatus)); + REQUIRE(hipStreamCaptureStatusActive == cStatus); + } + + 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)); + + SECTION("Stop capturing a stream and check the status.") { + HIP_CHECK(hipStreamIsCapturing(stream, &cStatus)); + REQUIRE(hipStreamCaptureStatusNone == cStatus); + } + + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + + free(A_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} + +/** +Testcase Scenarios : + 1) Functional : Use hipStreamPerThread, call api and check + capture status is hipStreamCaptureStatusNone. + 2) Functional : Start capturing using hipStreamPerThread and check + capture status returned as hipStreamCaptureStatusActive. + 3) Functional : Stop capturing using hipStreamPerThread and check + status is returned as hipStreamCaptureStatusNone. +*/ + +TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") { + float *A_d, *C_d; + float *A_h, *C_h; + hipGraph_t graph{nullptr}; + hipStreamCaptureStatus cStatus; + + 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); + + SECTION("Check the stream capture status before start capturing.") { + HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus)); + REQUIRE(hipStreamCaptureStatusNone == cStatus); + } + + HIP_CHECK(hipStreamBeginCapture(hipStreamPerThread, + hipStreamCaptureModeGlobal)); + + SECTION("Start capturing a stream and check the status.") { + HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus)); + REQUIRE(hipStreamCaptureStatusActive == cStatus); + } + + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, + hipStreamPerThread)); + + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, hipStreamPerThread)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, hipStreamPerThread, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, + hipStreamPerThread)); + + HIP_CHECK(hipStreamEndCapture(hipStreamPerThread, &graph)); + + SECTION("Stop capturing a stream and check the status.") { + HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus)); + REQUIRE(hipStreamCaptureStatusNone == cStatus); + } + + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipGraphDestroy(graph)); + + free(A_h); + free(C_h); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} From 78aaa848a4470eb78c5e25f615856d51462b6ed6 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 02:35:57 +0530 Subject: [PATCH 15/35] SWDEV-333962 - Update Device Side Malloc test (#2666) Change-Id: I5f872bae95e194e6ed679ea84d29383187c86cbb --- tests/src/deviceLib/hipDeviceMalloc.cpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/tests/src/deviceLib/hipDeviceMalloc.cpp b/tests/src/deviceLib/hipDeviceMalloc.cpp index 187abfcd2a..314d815a51 100644 --- a/tests/src/deviceLib/hipDeviceMalloc.cpp +++ b/tests/src/deviceLib/hipDeviceMalloc.cpp @@ -17,8 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM all - * TEST: %t EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../src/test_common.cpp + * TEST: %t * HIT_END */ #include "test_common.h" @@ -82,10 +82,10 @@ __device__ __host__ FloatT calc(FloatT A, FloatT B, enum CalcKind CK) { // Copy value from A, B to allocated memory. template __global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) { - int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x - + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x - + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x - * hipBlockDim_y; + int tx = threadIdx.x + blockDim.x * blockIdx.x + + (threadIdx.y + blockDim.y * blockIdx.y) * blockDim.x + + (threadIdx.z + blockDim.z * blockIdx.z) * blockDim.x + * blockDim.y; if (tx == 0) { *pA = (FloatT*)malloc(sizeof(FloatT) * LEN); *pB = (FloatT*)malloc(sizeof(FloatT) * LEN); @@ -100,10 +100,10 @@ __global__ void kernel_alloc(FloatT* A, FloatT* B, FloatT** pA, FloatT** pB) { // containing the address of the device-side allocated array. template __global__ void kernel_free(FloatT** pA, FloatT** pB, FloatT* C, enum CalcKind CK) { - int tx = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x - + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x - + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x - * hipBlockDim_y; + int tx = threadIdx.x + blockDim.x * blockIdx.x + + (threadIdx.y + blockDim.y * blockIdx.y) * blockDim.x + + (threadIdx.z + blockDim.z * blockIdx.z) * blockDim.x + * blockDim.y; C[tx] = calc((*pA)[tx], (*pB)[tx], CK); if (tx == 0) { free(*pA); From 14a136c3e2170654c112d7e6eb30fc77fe83d18f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 06:19:09 +0530 Subject: [PATCH 16/35] SWDEV-329687 - add instructions on building hip tests (#2737) Change-Id: I0f3567ee5f3170c5d50d39caa3ce3f40a9d152f7 --- docs/markdown/hip_build.md | 82 +++++++++++++++++++++++++++++++++++--- 1 file changed, 77 insertions(+), 5 deletions(-) diff --git a/docs/markdown/hip_build.md b/docs/markdown/hip_build.md index d61103aed2..17bc1af955 100755 --- a/docs/markdown/hip_build.md +++ b/docs/markdown/hip_build.md @@ -8,10 +8,12 @@ * [Set the environment variables](#set-the-environment-variables) * [Build HIP](#build-hip) * [Default paths and environment variables](#default-paths-and-environment-variables) + * [Build HIP Tests](#build-hip-tests) - [Build HIP on NVIDIA platform](#build-hip-on-NVIDIA-platform) * [Get HIP source code](#get-hip-source-code) * [Set the environment variables](#set-the-environment-variables) * [Build HIP](#build-hip) + * [Build HIP tests](#build-hip-tests) - [Run HIP](#run-hip) @@ -30,7 +32,6 @@ sudo apt install comgr sudo apt-get -y install rocm-dkms ``` - ## NVIDIA platform Install Nvidia driver and pre-build packages (see HIP Installation Guide at https://docs.amd.com/ for the release) @@ -61,7 +62,6 @@ ROCM_PATH is path where ROCM is installed. BY default ROCM_PATH is at /opt/rocm. git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hipamd.git git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/hip.git git clone -b $HIP_BRANCH https://github.com/ROCm-Developer-Tools/ROCclr.git -git clone -b $HIP_BRANCH https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git ``` ## Set the environment variables @@ -69,8 +69,6 @@ git clone -b $HIP_BRANCH https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtim ``` export HIPAMD_DIR="$(readlink -f hipamd)" export HIP_DIR="$(readlink -f hip)" -export ROCclr_DIR="$(readlink -f ROCclr)" -export OPENCL_DIR="$(readlink -f ROCm-OpenCL-Runtime)" ``` ROCclr is defined on AMD platform that HIP use Radeon Open Compute Common Language Runtime (ROCclr), which is a virtual device interface that HIP runtimes interact with different backends. @@ -84,7 +82,7 @@ See https://github.com/ROCm-Developer-Tools/hipamd ``` cd "$HIPAMD_DIR" mkdir -p build; cd build -cmake -DHIP_COMMON_DIR=$HIP_DIR -DAMD_OPENCL_PATH=$OPENCL_DIR -DROCCLR_PATH=$ROCCLR_DIR -DCMAKE_PREFIX_PATH="/" -DCMAKE_INSTALL_PREFIX=$PWD/install .. +cmake -DHIP_COMMON_DIR=$HIP_DIR -DCMAKE_PREFIX_PATH="/" -DCMAKE_INSTALL_PREFIX=$PWD/install .. make -j$(nproc) sudo make install ``` @@ -103,6 +101,77 @@ By default, release version of AMDHIP is built. After make install command, make sure HIP_PATH is pointed to $PWD/install/hip. +## Build HIP tests + +### Build HIP directed tests +Developers can build HIP directed tests right after build HIP commands, + +``` +sudo make install +make -j$(nproc) build_tests +``` +By default, all HIP directed tests will be built and generated under the folder $HIPAMD_DIR/build/directed_tests. +Take HIP directed device APIs tests, as an example, all available test applications will have executable files generated under, +$HIPAMD_DIR/build/directed_tests/runtimeApi/device. + +Run all HIP directed_tests, use the command, + +``` +ctest +``` +Or +``` +make test +``` + +Build and run a single directed test, use the follow command as an example, + +``` +make directed_tests.texture.hipTexObjPitch +cd $HIPAMD_DIR/build/directed_tests/texcture +./hipTexObjPitch +``` +Please note, the integrated HIP directed tests, will be deprecated in future release. + + +### Build HIP catch tests + +After build and install HIP commands, catch tests can be built via the following instructions, + +``` +cd "$HIP_DIR" +mkdir -p build; cd build +export HIP_PATH=$HIPAMD_DIR/build +cmake ../tests/catch/ -DHIP_PLATFORM=amd +make -j$(nproc) build_tests +ctest # run tests +``` + +HIP catch tests are built under the folder $HIP_DIR/build. + +To run a single catch test, the following is an example, + +``` +cd $HIP_DIR/build/unit/texture +./TextureTest +``` + +### Build HIP Catch2 standalone test + +HIP Catch2 supports build a standalone test, for example, + +``` +export PATH=$HIP_DIR/bin:$PATH +export HIP_PATH=$HIPAMD_DIR/build/install + +hipcc $HIP_DIR/tests/catch/unit/memory/hipPointerGetAttributes.cc -I ./tests/catch/include ./tests/catch/hipTestMain/standalone_main.cc -I ./tests/catch/external/Catch2 -g -o hipPointerGetAttributes +./hipPointerGetAttributes +... + +All tests passed +``` + +HIP catch tests, especially new architectured Catch2, will be official HIP tests in the repository and can be built alone as with the instructions shown above. # Build HIP on NVIDIA platform @@ -131,6 +200,9 @@ make -j$(nproc) sudo make install ``` +## Build HIP tests +Build HIP tests commands on NVIDIA platform are basically the same as AMD, except set -DHIP_PLATFORM=nvidia. + # Run HIP Compile and run the [square sample](https://github.com/ROCm-Developer-Tools/HIP/tree/rocm-5.0.x/samples/0_Intro/square). From a94de8f202dedeedb28eae5d5bc6969b42113ec3 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 11:22:27 +0530 Subject: [PATCH 17/35] SWDEV-339113 - update coordinate in HIP sample kernels (#2736) Change-Id: I8ea179b4ba8f1c0ebec830a5aa5947e843f06e42 --- samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp | 4 ++-- samples/2_Cookbook/1_hipEvent/hipEvent.cpp | 4 ++-- samples/2_Cookbook/3_shared_memory/sharedMemory.cpp | 4 ++-- samples/2_Cookbook/4_shfl/shfl.cpp | 2 +- samples/2_Cookbook/5_2dshfl/2dshfl.cpp | 4 ++-- samples/2_Cookbook/5_2dshfl/Readme.md | 2 +- 6 files changed, 10 insertions(+), 10 deletions(-) diff --git a/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp b/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp index 012debc148..fe71a03c5f 100644 --- a/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp +++ b/samples/2_Cookbook/19_cmake_lang/MatrixTranspose.cpp @@ -37,8 +37,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; out[y * width + x] = in[x * width + y]; } diff --git a/samples/2_Cookbook/1_hipEvent/hipEvent.cpp b/samples/2_Cookbook/1_hipEvent/hipEvent.cpp index a83b99beb6..81013c1dd1 100644 --- a/samples/2_Cookbook/1_hipEvent/hipEvent.cpp +++ b/samples/2_Cookbook/1_hipEvent/hipEvent.cpp @@ -35,8 +35,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; out[y * width + x] = in[x * width + y]; } diff --git a/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp b/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp index cff03f0a72..8bd489dbf3 100644 --- a/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp +++ b/samples/2_Cookbook/3_shared_memory/sharedMemory.cpp @@ -38,8 +38,8 @@ THE SOFTWARE. __global__ void matrixTranspose(float* out, float* in, const int width) { __shared__ float sharedMem[WIDTH * WIDTH]; - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; sharedMem[y * width + x] = in[x * width + y]; diff --git a/samples/2_Cookbook/4_shfl/shfl.cpp b/samples/2_Cookbook/4_shfl/shfl.cpp index 7902fd153a..de1ff7a950 100644 --- a/samples/2_Cookbook/4_shfl/shfl.cpp +++ b/samples/2_Cookbook/4_shfl/shfl.cpp @@ -36,7 +36,7 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int x = blockDim.x * blockIdx.x + threadIdx.x; float val = in[x]; diff --git a/samples/2_Cookbook/5_2dshfl/2dshfl.cpp b/samples/2_Cookbook/5_2dshfl/2dshfl.cpp index d2c26e7933..269ad58383 100644 --- a/samples/2_Cookbook/5_2dshfl/2dshfl.cpp +++ b/samples/2_Cookbook/5_2dshfl/2dshfl.cpp @@ -36,8 +36,8 @@ THE SOFTWARE. // Device (Kernel) function, it must be void __global__ void matrixTranspose(float* out, float* in, const int width) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; float val = in[y * width + x]; out[x * width + y] = __shfl(val, y * width + x); diff --git a/samples/2_Cookbook/5_2dshfl/Readme.md b/samples/2_Cookbook/5_2dshfl/Readme.md index fa10c71d6c..0419ef736c 100644 --- a/samples/2_Cookbook/5_2dshfl/Readme.md +++ b/samples/2_Cookbook/5_2dshfl/Readme.md @@ -31,7 +31,7 @@ We will be using the Simple Matrix Transpose application from the previous tutor In the same sourcecode, we used for MatrixTranspose. We'll add the following: ``` - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + int y = blockDim.y * blockIdx.y + threadIdx.y; out[x*width + y] = __shfl(val,y*width + x); ``` From 55e8bda5adbeaad405f990fdabc493149f26c1ed Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 11:29:21 +0530 Subject: [PATCH 18/35] SWDEV-330666 - Test cases for null stream and/or priority. (#2651) Change-Id: I9f5c99513fa1c754c2ab2e3cb9dcd5b6a3bb0176 --- .../catch/unit/stream/hipStreamGetPriority.cc | 44 +++++++++++++++++++ 1 file changed, 44 insertions(+) diff --git a/tests/catch/unit/stream/hipStreamGetPriority.cc b/tests/catch/unit/stream/hipStreamGetPriority.cc index ac68a9396a..e2384edd8e 100644 --- a/tests/catch/unit/stream/hipStreamGetPriority.cc +++ b/tests/catch/unit/stream/hipStreamGetPriority.cc @@ -89,6 +89,50 @@ TEST_CASE("Unit_hipStreamGetPriority_happy") { } } +/** + * both stream and priority passed as nullptr. + */ +TEST_CASE("Unit_hipStreamGetPriority_nullptr_nullptr") { + auto res = hipStreamGetPriority(nullptr,nullptr); + REQUIRE(res == hipErrorInvalidValue); +} + + +/** + * valid stream and priority passed as nullptr. + */ +TEST_CASE("Unit_hipStreamGetPriority_stream_nullptr") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + auto res = hipStreamGetPriority(stream,nullptr); + REQUIRE(res == hipErrorInvalidValue); + + HIP_CHECK(hipStreamDestroy(stream)); +} + + +/** + * nullptr stream and valid priority + */ +TEST_CASE("Unit_hipStreamGetPriority_nullptr_priority") { + int priority = -1; + HIP_CHECK(hipStreamGetPriority(nullptr,&priority)); +} + +/** + * both stream and priority passed as valid. + */ +TEST_CASE("Unit_hipStreamGetPriority_stream_priority") { + int priority = -1; + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + + HIP_CHECK(hipStreamGetPriority(stream,&priority)); + + HIP_CHECK(hipStreamDestroy(stream)); +} + #if HT_AMD /** * Create stream with CUMask and check priority is returned as expected. From ac1d594ef10c54c38d7c49f40f34adeae9c519a5 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 22:08:21 +0530 Subject: [PATCH 19/35] SWDEV-337331 - Windows initial graph tests (#2739) Enables most of graph tests on Windows. Change-Id: I32bd9a327168378776f89d7a6c4a83ba33861c01 --- .../config/config_amd_windows.json | 33 ------------------- .../graph/hipGraphAddMemcpyNodeFromSymbol.cc | 2 ++ .../graph/hipGraphAddMemcpyNodeToSymbol.cc | 2 ++ tests/catch/unit/graph/hipGraphClone.cc | 1 + ...pGraphExecMemcpyNodeSetParamsFromSymbol.cc | 1 + 5 files changed, 6 insertions(+), 33 deletions(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 8602662d49..9fc36b8e4c 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -30,35 +30,9 @@ "Unit_hipStreamPerThread_DeviceReset_1", "Unit_hipStreamPerThread_DeviceReset_2", "Unit_hipManagedKeyword_MultiGpu", - "Unit_hipGraphAddDependencies_Functional", - "Unit_hipGraph_BasicFunctional", - "Unit_hipGraphClone_Functional", - "Unit_hipGraphClone_MultiThreaded", "Unit_hipGraphAddHostNode_ClonedGraphwithHostNode", - "Unit_hipGraphAddHostNode_VectorSquare", - "Unit_hipGraphAddHostNode_BasicFunc", - "Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemory", - "Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalConstMemory", - "Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel", - "Unit_hipGraphExecHostNodeSetParams_Negative", - "Unit_hipGraphExecHostNodeSetParams_ClonedGraphwithHostNode", - "Unit_hipGraphExecHostNodeSetParams_BasicFunc", - "Unit_hipGraphAddMemcpyNodeToSymbol_GlobalMemory", - "Unit_hipGraphAddMemcpyNodeToSymbol_GlobalConstMemory", - "Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel", - "Unit_hipGraphDestroyNode_DestroyDependencyNode", - "Unit_hipGraphGetNodes_Functional", - "Unit_hipGraphGetRootNodes_Functional", - "Unit_hipGraphHostNodeSetParams_ClonedGraphwithHostNode", "Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph", "Unit_hipGraphAddChildGraphNode_SingleChildNode", - "Unit_hipGraphExecMemcpyNodeSetParams1D_Functional", - "Unit_hipGraphRemoveDependencies_ChangeComputeFunc", - "Unit_hipGraphExecUpdate_Negative_CountDiffer", - "Unit_hipGraphExecUpdate_Functional", - "Unit_hipGraphExecEventRecordNodeSetEvent_Negative", - "Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative", - "Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Functional", "Unit_hipPtrGetAttribute_Simple", "Unit_hipStreamCreateWithPriority_ValidateWithEvents", "Unit_hipEvent", @@ -67,15 +41,8 @@ "Unit_hipHostMalloc_Default", "Unit_hipStreamCreate_MultistreamBasicFunctionalities", "Unit_hipEventIpc", - "Unit_hipGraphAddDependencies_NegTest", - "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", - "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime", - "Unit_hipGraphAddEventRecordNode_Functional_WithFlags", - "Unit_hipGraphAddEventRecordNode_MultipleRun", "Unit_hipMalloc3D_Negative", "Unit_hipPointerGetAttribute_MappedMem", - "Unit_hipGraphAddMemcpyNode1D_Functional", - "Unit_hipGraphAddMemcpyNode1D_Negative", "Unit_hipStreamBeginCapture_BasicFunctional", "Unit_hipStreamBeginCapture_hipStreamPerThread", "Unit_hiprtc_functional", diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc index 7d9cbaffed..fa05c7fe66 100644 --- a/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNodeFromSymbol.cc @@ -273,6 +273,7 @@ void hipGraphAddMemcpyNodeFromSymbol_GlobalMemory(bool device_ctxchg = false, // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Validating the result for (int i = 0; i < SIZE; i++) { @@ -419,6 +420,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNodeFromSymbol_GlobalMemoryWithKernel") { // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Validating the result for (int i = 0; i < SIZE; i++) { diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc index 5a5b08f77e..a4d0bede3e 100644 --- a/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNodeToSymbol.cc @@ -250,6 +250,7 @@ void hipGraphAddMemcpyNodeToSymbol_GlobalMemory(bool device_ctxchg = false, // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Validating the result for (int i = 0; i < SIZE; i++) { @@ -384,6 +385,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNodeToSymbol_MemcpyToSymbolNodeWithKernel") { // Instantiate and launch the graph HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Validating the result for (int i = 0; i < SIZE; i++) { diff --git a/tests/catch/unit/graph/hipGraphClone.cc b/tests/catch/unit/graph/hipGraphClone.cc index 4a6cbc6ec9..75b83c82a0 100644 --- a/tests/catch/unit/graph/hipGraphClone.cc +++ b/tests/catch/unit/graph/hipGraphClone.cc @@ -295,6 +295,7 @@ TEST_CASE("Unit_hipGraphClone_MultiThreaded") { HIP_CHECK(hipGraphInstantiate(&graphExec, clonedgraph, nullptr, nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); for (size_t i = 0; i < N; i++) { if (A_h[i] != B_h[i]) { diff --git a/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol.cc b/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol.cc index a40cc4d433..b2a92f94db 100644 --- a/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol.cc +++ b/tests/catch/unit/graph/hipGraphExecMemcpyNodeSetParamsFromSymbol.cc @@ -269,6 +269,7 @@ void hipGraphExecMemcpyNodeSetParamsFromSymbol_GlobalMem(bool useConstVar) { } HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Validating the result for (int i = 0; i < SIZE; i++) { From 3e7a9bb177b68547b670f548c3b265303cdb7d00 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 22:44:37 +0530 Subject: [PATCH 20/35] SWDEV-323441 - Pahes-II : per thread default stream (#2738) Change-Id: I9f76df450207794b8c887b55fb16aea8044d69b4 --- include/hip/hip_runtime_api.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index e6a5dc2ab8..edbca6df75 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1192,7 +1192,6 @@ typedef enum hipGraphInstantiateFlags { hipGraphInstantiateFlagAutoFreeOnLaunch = 1, ///< Automatically free memory allocated in a graph before relaunching. } hipGraphInstantiateFlags; -#include // Doxygen end group GlobalDefs /** @} */ @@ -6813,6 +6812,8 @@ static inline hipError_t hipMallocManaged(T** devPtr, size_t size, #endif #endif +#include + #if USE_PROF_API #include #endif From 0cf0ead620899b649fb4cef8a959527b588e9b6d Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 15 Jun 2022 23:28:10 +0530 Subject: [PATCH 21/35] SWDEV-339657 - Update catch codes with correct coordinates (#2735) Change-Id: I8ca78d770a742bf6c2abede494759caa923fcc19 --- .../stress/memory/hipMemPrftchAsyncStressTst.cc | 4 ++-- .../stress/printf/Stress_printf_ComplexKernels.cc | 6 +++--- .../stress/printf/Stress_printf_SimpleKernels.cc | 12 ++++++------ .../graph/hipGraphExecEventRecordNodeSetEvent.cc | 2 +- tests/catch/unit/memory/hipMemPrefetchAsync.cc | 4 ++-- .../catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc | 4 ++-- tests/catch/unit/texture/hipBindTex2DPitch.cc | 4 ++-- .../catch/unit/texture/hipNormalizedFloatValueTex.cc | 2 +- 8 files changed, 19 insertions(+), 19 deletions(-) diff --git a/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc b/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc index bc2cb027da..0a159b8dbb 100644 --- a/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc +++ b/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc @@ -26,8 +26,8 @@ THE SOFTWARE. // Kernel function __global__ void MemPrftchAsyncKernel1(int* Hmm, size_t N) { - size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { Hmm[i] = Hmm[i] * Hmm[i]; } diff --git a/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc b/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc index 16796d7972..eec8e162f7 100644 --- a/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc +++ b/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc @@ -230,19 +230,19 @@ __device__ __host__ struct printInfo startPrint(uint32_t tid, // This kernel is launched only in X dimension __global__ void kernel_complex_opX(uint32_t *a, uint32_t *b, uint32_t iterCount) { - uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x; startPrint(tid, iterCount, a, b); } // This kernel is launched only in Y dimension __global__ void kernel_complex_opY(uint32_t *a, uint32_t *b, uint32_t iterCount) { - uint32_t tid = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; + uint32_t tid = threadIdx.y + blockIdx.y * blockDim.y; startPrint(tid, iterCount, a, b); } // This kernel is launched only in Z dimension __global__ void kernel_complex_opZ(uint32_t *a, uint32_t *b, uint32_t iterCount) { - uint32_t tid = hipThreadIdx_z + hipBlockIdx_z * hipBlockDim_z; + uint32_t tid = threadIdx.z + blockIdx.z * blockDim.z; startPrint(tid, iterCount, a, b); } #ifdef __linux__ diff --git a/tests/catch/stress/printf/Stress_printf_SimpleKernels.cc b/tests/catch/stress/printf/Stress_printf_SimpleKernels.cc index e116593a4a..1ec20ec873 100644 --- a/tests/catch/stress/printf/Stress_printf_SimpleKernels.cc +++ b/tests/catch/stress/printf/Stress_printf_SimpleKernels.cc @@ -85,7 +85,7 @@ __global__ void kernel_printf_conststr(uint iterCount) { // 'g' grid size such that (total bytes per iteration)*n*b*g ≈ N GB, // where N is user input. __global__ void kernel_printf_two_conditionalstr(uint iterCount) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; uint mod_tid = (tid % 2); if (0 == mod_tid) { for (uint count = 0; count < iterCount; count++) { @@ -101,7 +101,7 @@ __global__ void kernel_printf_two_conditionalstr(uint iterCount) { // iterations per thread using 'b' block size and 'g' grid size such that // (total bytes per iteration)*n*b*g ≈ N GB, where N is user input. __global__ void kernel_printf_single_conditionalstr(uint iterCount) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; uint mod_tid = (tid % 2); if (0 == mod_tid) { for (uint count = 0; count < iterCount; count++) { @@ -115,7 +115,7 @@ __global__ void kernel_printf_single_conditionalstr(uint iterCount) { // iterations per thread using 'b' block size and 'g' grid size such // that (total bytes per iteration)*n*b*g ≈ N GB, where N is user input. __global__ void kernel_printf_variablestr(uint iterCount, int *ret) { - uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint tid = threadIdx.x + blockIdx.x * blockDim.x; int retlocal = 0; const char *const_str = "Hello World from Device.Iam printing (threadID,number)="; @@ -134,7 +134,7 @@ __global__ void kernel_printf_variablestr(uint iterCount, int *ret) { // size and 'g' grid size such that // (total bytes per iteration)*n*b*g ≈ N GB, where N is user input. __global__ void kernel_dependent_calc(uint32_t iterCount, int *ret) { - uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x; int retlocal = 0; const char *const_str = "Hello World from Device.Iam printing number="; @@ -158,7 +158,7 @@ __global__ void kernel_dependent_calc(uint32_t iterCount, int *ret) { // (total bytes per iteration)*n*b*g ≈ N GB, where N is user input. __global__ void kernel_dependent_calc_atomic(uint32_t iterCount, int *ret) { - uint32_t tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x; int retlocal = 0; const char *const_str = "Hello World from Device.Iam printing number="; @@ -209,7 +209,7 @@ __global__ void kernel_shared_mem() { __shared__ uint32_t sharedMem; sharedMem = 0; __syncthreads(); - atomicAdd(&sharedMem, hipThreadIdx_x); + atomicAdd(&sharedMem, threadIdx.x); __syncthreads(); printf("%s%u\n", CONST_STR3, sharedMem); } diff --git a/tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc b/tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc index 69c60635bf..5a1cbe9997 100644 --- a/tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc +++ b/tests/catch/unit/graph/hipGraphExecEventRecordNodeSetEvent.cc @@ -57,7 +57,7 @@ Testcase Scenarios : * Kernel Functions to copy. */ static __global__ void copy_ker_func(int* a, int* b) { - int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; + int tx = blockIdx.x*blockDim.x + threadIdx.x; if (tx < LEN) b[tx] = a[tx]; } diff --git a/tests/catch/unit/memory/hipMemPrefetchAsync.cc b/tests/catch/unit/memory/hipMemPrefetchAsync.cc index 6f8fe65c74..17ef618b77 100644 --- a/tests/catch/unit/memory/hipMemPrefetchAsync.cc +++ b/tests/catch/unit/memory/hipMemPrefetchAsync.cc @@ -20,8 +20,8 @@ THE SOFTWARE. #include // Kernel function __global__ void MemPrftchAsyncKernel(int* C_d, const int* A_d, size_t N) { - size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { C_d[i] = A_d[i] * A_d[i]; } diff --git a/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc b/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc index c5a51d9488..482a73cd02 100644 --- a/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc +++ b/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc @@ -37,8 +37,8 @@ THE SOFTWARE. // Kernel function __global__ void MemPrftchAsyncKernel1(int* Hmm, size_t N) { - size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; for (size_t i = offset; i < N; i += stride) { Hmm[i] = Hmm[i] * Hmm[i]; } diff --git a/tests/catch/unit/texture/hipBindTex2DPitch.cc b/tests/catch/unit/texture/hipBindTex2DPitch.cc index eff6a9f939..c1d82ebfbe 100644 --- a/tests/catch/unit/texture/hipBindTex2DPitch.cc +++ b/tests/catch/unit/texture/hipBindTex2DPitch.cc @@ -28,8 +28,8 @@ texture tex; // texture object is a kernel argument static __global__ void texture2dCopyKernel(TYPE_t* dst) { - int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; - int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; + int x = threadIdx.x + blockIdx.x * blockDim.x; + int y = threadIdx.y + blockIdx.y * blockDim.y; if ( (x < SIZE_W) && (y < SIZE_H) ) { dst[SIZE_W*y+x] = tex2D(tex, x, y); } diff --git a/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc b/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc index e9cf002eae..8113117a64 100644 --- a/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc +++ b/tests/catch/unit/texture/hipNormalizedFloatValueTex.cc @@ -43,7 +43,7 @@ template __global__ void normalizedValTextureTest(unsigned int numElements, float* pDst) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - unsigned int elementID = hipThreadIdx_x; + unsigned int elementID = threadIdx.x; if (elementID >= numElements) return; float coord = elementID/static_cast(numElements); From 96d654e3afa3b79867cade26ebceeb0bab4c56b7 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 16 Jun 2022 01:57:42 +0530 Subject: [PATCH 22/35] SWDEV-336007 - Add sample which verifies CXX language support with amdclang++ (#2733) Change-Id: Ib4dbe1ce66a7dbd4f60f0a2134d62bdf366aeac5 --- .../21_cmake_hip_cxx_clang/CMakeLists.txt | 34 +++++++ .../21_cmake_hip_cxx_clang/README.md | 19 ++++ .../21_cmake_hip_cxx_clang/square.cpp | 98 +++++++++++++++++++ 3 files changed, 151 insertions(+) create mode 100644 samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt create mode 100644 samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md create mode 100644 samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp diff --git a/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt b/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt new file mode 100644 index 0000000000..3cf4f35336 --- /dev/null +++ b/samples/2_Cookbook/21_cmake_hip_cxx_clang/CMakeLists.txt @@ -0,0 +1,34 @@ +# 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 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. + +cmake_minimum_required(VERSION 3.21.3) + +project(cmake_cxx_amdclang++_test + DESCRIPTION "Verifies CXX Language Support with amdclang++" + LANGUAGES CXX) + +# Find HIP +find_package(hip REQUIRED) + +# Create the excutable +add_executable(square square.cpp) + +# Link with HIP +target_link_libraries(square hip::device) diff --git a/samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md b/samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md new file mode 100644 index 0000000000..3d9765308c --- /dev/null +++ b/samples/2_Cookbook/21_cmake_hip_cxx_clang/README.md @@ -0,0 +1,19 @@ +### This sample tests CXX Language support with amdclang++ +I. Build +mkdir -p build; cd build +rm -rf *; +CXX=`hipconfig -l`/amdclang++ cmake .. (or) +cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ .. (or) +cmake -DCMAKE_CXX_COMPILER=/opt/rocm-X.Y.Z/llvm/bin/amdclang++ .. +make + +II. Test +$ ./square +info: running on device Vega 20 [Radeon Pro Vega 20] +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +info: copy Host2Device +info: launch 'vector_square' kernel +info: copy Device2Host +info: check result +PASSED! diff --git a/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp b/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp new file mode 100644 index 0000000000..c3ed799edc --- /dev/null +++ b/samples/2_Cookbook/21_cmake_hip_cxx_clang/square.cpp @@ -0,0 +1,98 @@ +/* +Copyright (c) 2015-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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#define CHECK(cmd) \ +{\ + hipError_t error = cmd;\ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + + +/* + * Square each element in the array A and write to array C. + */ +template +__global__ void +vector_square(T *C_d, T *A_d, size_t N) +{ + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; + + for (size_t i=offset; i Date: Thu, 16 Jun 2022 04:57:59 +0100 Subject: [PATCH 23/35] EXSWCPHIPT-109 - hipDeviceGetName tests (#2731) --- tests/catch/unit/device/hipDeviceGetName.cc | 102 ++++++++++++++++---- 1 file changed, 85 insertions(+), 17 deletions(-) diff --git a/tests/catch/unit/device/hipDeviceGetName.cc b/tests/catch/unit/device/hipDeviceGetName.cc index c0884d0093..ec4913d146 100644 --- a/tests/catch/unit/device/hipDeviceGetName.cc +++ b/tests/catch/unit/device/hipDeviceGetName.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 @@ -21,49 +21,117 @@ THE SOFTWARE. * Conformance test for checking functionality of * hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device); */ +#include #include #include #include +#include +#include +#include + +constexpr size_t LEN = 256; -#define LEN 256 /** * hipDeviceGetName tests * Scenario1: Validates the name string with hipDeviceProp_t.name[256] * Scenario2: Validates returned error code for name = nullptr * Scenario3: Validates returned error code for len = 0 * Scenario4: Validates returned error code for len < 0 + * Scenario5: Validates returned error code for an invalid device + * Scenario6: Validates partially filling the name into a char array */ -TEST_CASE("Unit_hipDeviceGetName-NegTst") { +TEST_CASE("Unit_hipDeviceGetName_NegTst") { + std::array name; + int numDevices = 0; - char name[LEN]; - hipDevice_t device; HIP_CHECK(hipGetDeviceCount(&numDevices)); + + std::vector devices(numDevices); for (int i = 0; i < numDevices; i++) { - HIP_CHECK(hipDeviceGet(&device, i)); - HIP_CHECK(hipDeviceGetName(name, LEN, device)); - // Scenario2 - CHECK_FALSE(hipSuccess == hipDeviceGetName(nullptr, LEN, device)); + HIP_CHECK(hipDeviceGet(&devices[i], i)); + } + + SECTION("Valid Device") { + const auto device = GENERATE_COPY(from_range(std::begin(devices), std::end(devices))); + + SECTION("Nullptr for name argument") { + // Scenario2 + HIP_CHECK_ERROR(hipDeviceGetName(nullptr, name.size(), device), hipErrorInvalidValue); + } #if HT_AMD // These test scenarios fail on NVIDIA. - // Scenario3 - CHECK_FALSE(hipSuccess == hipDeviceGetName(name, 0, device)); - // Scenario4 - CHECK_FALSE(hipSuccess == hipDeviceGetName(name, -1, device)); + SECTION("Zero name length") { + // Scenario3 + HIP_CHECK_ERROR(hipDeviceGetName(name.data(), 0, device), hipErrorInvalidValue); + } + + SECTION("Negative name length") { + // Scenario4 + HIP_CHECK_ERROR(hipDeviceGetName(name.data(), -1, device), hipErrorInvalidValue); + } #endif } + SECTION("Invalid Device") { + hipDevice_t badDevice = devices.back() + 1; + + constexpr size_t timeout = 100; + size_t timeoutCount = 0; + while (std::find(std::begin(devices), std::end(devices), badDevice) != std::end(devices)) { + badDevice += 1; + timeoutCount += 1; + REQUIRE(timeoutCount < timeout); // give up after a while + } + + // Scenario5 + HIP_CHECK_ERROR(hipDeviceGetName(name.data(), name.size(), badDevice), hipErrorInvalidDevice); + } } -TEST_CASE("Unit_hipDeviceGetName-CheckPropName") { +TEST_CASE("Unit_hipDeviceGetName_CheckPropName") { int numDevices = 0; - char name[LEN]; + std::array name; hipDevice_t device; hipDeviceProp_t prop; HIP_CHECK(hipGetDeviceCount(&numDevices)); for (int i = 0; i < numDevices; i++) { HIP_CHECK(hipDeviceGet(&device, i)); - HIP_CHECK(hipDeviceGetName(name, LEN, device)); + HIP_CHECK(hipDeviceGetName(name.data(), name.size(), device)); HIP_CHECK(hipGetDeviceProperties(&prop, device)); + // Scenario1 - CHECK_FALSE(0 != strcmp(name, prop.name)); + CHECK(strncmp(name.data(), prop.name, name.size()) == 0); } } + +TEST_CASE("Unit_hipDeviceGetName_PartialFill") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-108"); + return; +#endif + std::array name; + + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + auto ordinal = GENERATE_COPY(range(0, numDevices)); + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, ordinal)); + HIP_CHECK(hipDeviceGetName(name.data(), name.size(), device)); + + auto start = std::begin(name); + auto end = std::end(name); + const auto len = std::distance(start, std::find(start, end, 0)); + + // fill up only half of the length + const auto fillLen = len / 2; + constexpr char fillValue = 1; + std::fill(start, end, fillValue); + + // Scenario6 + HIP_CHECK(hipDeviceGetName(name.data(), fillLen, device)); + + const auto strEnd = start + fillLen - 1; + REQUIRE(std::all_of(start, strEnd, [](char& c) { return c != 0; })); + REQUIRE(*strEnd == 0); + REQUIRE(std::all_of(strEnd+1, end, [](char& c) { return c == fillValue; })); +} From f0e03e9d45d9ad2c15bf5ec4d44452f17216ade8 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 16 Jun 2022 09:28:30 +0530 Subject: [PATCH 24/35] SWDEV-1 - Add -ldl as linking flag for catch2 tests. (#2741) Change-Id: I81c759363f35c3822447563e51c2db670edf9cb3 --- tests/catch/external/Catch2/cmake/Catch2/Catch.cmake | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake b/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake index 62e4984afc..603dc1b994 100644 --- a/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/tests/catch/external/Catch2/cmake/Catch2/Catch.cmake @@ -216,6 +216,7 @@ function(hip_add_exe_to_target) catch_discover_tests(${_NAME} PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") if(UNIX) set(_LINKER_LIBS ${_LINKER_LIBS} stdc++fs) + set(_LINKER_LIBS ${_LINKER_LIBS} -ldl) else() # res files are built resource files using rc files. # use llvm-rc exe to build the res files From 3a84eb3a5f6a484bbc0d9b157842aff3f919db97 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 17 Jun 2022 05:39:11 +0530 Subject: [PATCH 25/35] SWDEV-306122 - [catch2][dtest] Adding following tests for hipGraphEventWaitNodeSetEvent() (#2590) - Functional Tests - Negative Tests Change-Id: I2421c3def1a7c6865dca29bde7e741948e67345d --- tests/catch/include/hip_test_kernels.hh | 9 +- tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/hipGraphEventWaitNodeSetEvent.cc | 317 ++++++++++++++++++ 3 files changed, 326 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc diff --git a/tests/catch/include/hip_test_kernels.hh b/tests/catch/include/hip_test_kernels.hh index 04b00b5ad3..d59fd41b5a 100644 --- a/tests/catch/include/hip_test_kernels.hh +++ b/tests/catch/include/hip_test_kernels.hh @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 - 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 @@ -89,4 +89,11 @@ template __global__ void vector_square(const T* A_d, T* C_d, size_t } } +template __global__ void vector_cubic(const T* A_d, T* C_d, size_t N_ELMTS) { + size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = gputhread; i < N_ELMTS; i += stride) { + C_d[i] = A_d[i] * A_d[i] * A_d[i]; + } +} } // namespace HipTest diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 036e9477f7..444daeec45 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -49,6 +49,7 @@ set(TEST_SRC hipGraphInstantiate.cc hipGraphExecUpdate.cc hipGraphExecEventRecordNodeSetEvent.cc + hipGraphEventWaitNodeSetEvent.cc hipGraphMemsetNodeGetParams.cc hipGraphMemsetNodeSetParams.cc hipGraphExecMemcpyNodeSetParamsFromSymbol.cc diff --git a/tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc b/tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc new file mode 100644 index 0000000000..8751ffe8ed --- /dev/null +++ b/tests/catch/unit/graph/hipGraphEventWaitNodeSetEvent.cc @@ -0,0 +1,317 @@ +/* +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 : + 1) Set a different type of event using hipGraphEventWaitNodeSetEvent and + validate using hipGraphEventWaitNodeGetEvent. + 2) Create a graph1 with memset (Value 1) node, event record node (event A) + and memset (Value 2) node and event record node (event B). Create a + graph2 with Event Wait (event A ) node and memcpyd2h. Instantiate + graph1 on stream1 and graph2 on stream2. Set the Event Wait node event + to B using hipGraphEventWaitNodeSetEvent. Launch graphs. Wait for the + event to complete. Verify the results. + 3) Negative Scenarios + - Input node parameter is passed as nullptr. + - Input event parameter is passed as nullptr. + - Input node is an empty node. + - Input node is a memset node. + - Input node is a event record node. + - Input node is an uninitialized node. + - Input event is an uninitialized node. +*/ + +#include +#include +#include + +#define LEN 512 + +/** + * Local Function + */ +static void validateEventWaitNodeSetEvent(unsigned flag) { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create events + hipEvent_t event1, event2, event_out; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreateWithFlags(&event2, flag)); + hipGraphNode_t eventwait; + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, + event1)); + // Set a different event + HIP_CHECK(hipGraphEventWaitNodeSetEvent(eventwait, event2)); + HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out)); + // validate set event and get event are same + REQUIRE(event2 == event_out); + // Free resources + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +} + +/** + * Local Function + */ +static void setEventRecordNode() { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create events + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventrec; + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, + event1)); + // Set a different event eventrec using hipGraphEventWaitNodeSetEvent + REQUIRE(hipErrorInvalidValue == + hipGraphEventWaitNodeSetEvent(eventrec, event2)); + // Free resources + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +} + +/** + * Scenario 2 + */ +TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetProp") { + size_t memsize = LEN * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, LEN); + size_t NElem{LEN}; + hipGraph_t graph1, graph2; + hipStream_t streamForGraph1, streamForGraph2; + hipGraphExec_t graphExec1, graphExec2; + HIP_CHECK(hipStreamCreate(&streamForGraph1)); + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphCreate(&graph2, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph2)); + + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreateWithFlags(&event1, hipEventDefault)); + HIP_CHECK(hipEventCreateWithFlags(&event2, hipEventBlockingSync)); + hipGraphNode_t event_rec_node, event_wait_node; + int *inp_h, *inp_d, *out_h_g1, *out_d_g1, *out_h_g2, *out_d_g2; + // Allocate host buffers + inp_h = reinterpret_cast(malloc(memsize)); + REQUIRE(inp_h != nullptr); + out_h_g1 = reinterpret_cast(malloc(memsize)); + REQUIRE(out_h_g1 != nullptr); + out_h_g2 = reinterpret_cast(malloc(memsize)); + REQUIRE(out_h_g2 != nullptr); + // Allocate device buffers + HIP_CHECK(hipMalloc(&inp_d, memsize)); + HIP_CHECK(hipMalloc(&out_d_g1, memsize)); + HIP_CHECK(hipMalloc(&out_d_g2, memsize)); + // Initialize host buffer + for (uint32_t i = 0; i < LEN; i++) { + inp_h[i] = i; + out_h_g1[i] = 0; + out_h_g2[i] = 0; + } + // Graph1 creation ........... + // Create event1 record node in graph1 + HIP_CHECK(hipGraphAddEventRecordNode(&event_rec_node, graph1, nullptr, 0, + event1)); + + // Create memcpy and kernel nodes for graph1 + hipGraphNode_t memcpyH2D, memcpyD2H_1, kernelnode_1; + hipKernelNodeParams kernelNodeParams1{}; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d, + inp_h, memsize, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_1, graph1, nullptr, 0, + out_h_g1, out_d_g1, memsize, hipMemcpyDeviceToHost)); + + void* kernelArgs1[] = {&inp_d, &out_d_g1, reinterpret_cast(&NElem)}; + kernelNodeParams1.func = + reinterpret_cast(HipTest::vector_square); + kernelNodeParams1.gridDim = dim3(blocks); + kernelNodeParams1.blockDim = dim3(threadsPerBlock); + kernelNodeParams1.sharedMemBytes = 0; + kernelNodeParams1.kernelParams = reinterpret_cast(kernelArgs1); + kernelNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_1, graph1, nullptr, 0, + &kernelNodeParams1)); + // Create dependencies for graph1 + HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D, + &event_rec_node, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &event_rec_node, + &kernelnode_1, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode_1, + &memcpyD2H_1, 1)); + + // Graph2 creation ........... + // Create event1 record node in graph2 + HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0, + event1)); + // Create memcpy and kernel nodes for graph2 + hipGraphNode_t memcpyD2H_2, kernelnode_2; + hipKernelNodeParams kernelNodeParams2{}; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_2, graph2, nullptr, 0, + out_h_g2, out_d_g2, memsize, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&inp_d, &out_d_g2, reinterpret_cast(&NElem)}; + kernelNodeParams2.func = + reinterpret_cast(HipTest::vector_cubic); + kernelNodeParams2.gridDim = dim3(blocks); + kernelNodeParams2.blockDim = dim3(threadsPerBlock); + kernelNodeParams2.sharedMemBytes = 0; + kernelNodeParams2.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams2.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode_2, graph2, nullptr, 0, + &kernelNodeParams2)); + // Create dependencies for graph2 + HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node, + &kernelnode_2, 1)); + HIP_CHECK(hipGraphAddDependencies(graph2, &kernelnode_2, + &memcpyD2H_2, 1)); + + // Instantiate and launch the graphs + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + // Set event + HIP_CHECK(hipGraphEventRecordNodeSetEvent(event_rec_node, event2)); + HIP_CHECK(hipGraphEventWaitNodeSetEvent(event_wait_node, event2)); + + HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1)); + HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2)); + HIP_CHECK(hipStreamSynchronize(streamForGraph1)); + HIP_CHECK(hipStreamSynchronize(streamForGraph2)); + // Validate output + bool btestPassed1 = true; + for (uint32_t i = 0; i < LEN; i++) { + if (out_h_g1[i] != (inp_h[i]*inp_h[i])) { + btestPassed1 = false; + break; + } + } + REQUIRE(btestPassed1 == true); + bool btestPassed2 = true; + for (uint32_t i = 0; i < LEN; i++) { + if (out_h_g2[i] != (inp_h[i]*inp_h[i]*inp_h[i])) { + btestPassed2 = false; + break; + } + } + REQUIRE(btestPassed2 == true); + // Destroy all resources + HIP_CHECK(hipFree(inp_d)); + HIP_CHECK(hipFree(out_d_g1)); + HIP_CHECK(hipFree(out_d_g2)); + free(inp_h); + free(out_h_g1); + free(out_h_g2); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); + HIP_CHECK(hipStreamDestroy(streamForGraph1)); + HIP_CHECK(hipStreamDestroy(streamForGraph2)); +} +/** + * Scenario 1 + */ +TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_SetGet") { + SECTION("Flag = hipEventDefault") { + validateEventWaitNodeSetEvent(hipEventDefault); + } + + SECTION("Flag = hipEventBlockingSync") { + validateEventWaitNodeSetEvent(hipEventBlockingSync); + } + + SECTION("Flag = hipEventDisableTiming") { + validateEventWaitNodeSetEvent(hipEventDisableTiming); + } +} + +/** + * Scenario 3 + */ +TEST_CASE("Unit_hipGraphEventWaitNodeSetEvent_Negative") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventwait; + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, + event1)); + SECTION("node = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent( + nullptr, event2)); + } + + SECTION("event = nullptr") { + REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent( + eventwait, nullptr)); + } + + SECTION("input node is empty node") { + hipGraphNode_t EmptyGraphNode; + HIP_CHECK(hipGraphAddEmptyNode(&EmptyGraphNode, graph, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == + hipGraphEventWaitNodeSetEvent(EmptyGraphNode, event2)); + } + + SECTION("input node is memset node") { + constexpr size_t Nbytes = 1024; + char *A_d; + hipGraphNode_t memset_A; + hipMemsetParams memsetParams{}; + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + 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)); + REQUIRE(hipErrorInvalidValue == + hipGraphEventWaitNodeSetEvent(memset_A, event2)); + HIP_CHECK(hipFree(A_d)); + } + + SECTION("input node is event record node") { + setEventRecordNode(); + } + + SECTION("input node is uninitialized node") { + hipGraphNode_t node_uninit{}; + REQUIRE(hipErrorInvalidValue == + hipGraphEventWaitNodeSetEvent(node_uninit, event2)); + } + + SECTION("input event is uninitialized") { + hipEvent_t event_uninit{}; + REQUIRE(hipErrorInvalidValue == hipGraphEventWaitNodeSetEvent( + eventwait, event_uninit)); + } + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +} From d3a062a4ebc1d79fadbcb2386dd9453218c8e940 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 17 Jun 2022 09:35:56 +0530 Subject: [PATCH 26/35] SWDEV-330661 - Add test for hipHostGetDevicePointer null check (#2745) Change-Id: I19668385003a107e21f8484ef28e114b7c6d8a42 --- tests/catch/unit/memory/hipHostMalloc.cc | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/tests/catch/unit/memory/hipHostMalloc.cc b/tests/catch/unit/memory/hipHostMalloc.cc index 94f943c6bf..39a2bb6719 100644 --- a/tests/catch/unit/memory/hipHostMalloc.cc +++ b/tests/catch/unit/memory/hipHostMalloc.cc @@ -231,4 +231,16 @@ TEST_CASE("Unit_hipHostMalloc_Default") { CheckHostPointer(numElements, A, 0, SYNC_DEVICE, ptrType); CheckHostPointer(numElements, A, 0, SYNC_STREAM, ptrType); CheckHostPointer(numElements, A, 0, SYNC_EVENT, ptrType); + } + +TEST_CASE("Unit_hipHostGetDevicePointer_NullCheck") { + int* d_a; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&d_a), sizeof(int))); + + auto res = hipHostGetDevicePointer(nullptr,d_a,0); + REQUIRE(res == hipErrorInvalidValue); + + HIP_CHECK(hipHostFree(d_a)); +} + From da0bd3cc85d02bdc4b3bf14b7bbe89ff9ff26710 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 17 Jun 2022 10:40:15 +0530 Subject: [PATCH 27/35] SWDEV-336536 - Changed doc for hipDeviceGetCacheConfig. (#2742) Change-Id: I0147cc8026788bad192258be28cde8497c0cd265 --- include/hip/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index edbca6df75..2ac304df2f 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1516,9 +1516,9 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId); */ hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig); /** - * @brief Set Cache configuration for a specific function + * @brief Get Cache configuration for a specific Device * - * @param [in] cacheConfig + * @param [out] cacheConfig * * @returns #hipSuccess, #hipErrorNotInitialized * Note: AMD devices and some Nvidia GPUS do not support reconfigurable cache. This hint is ignored From a49143593c11f7269c4b302a0bc52974e27f900c Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 17 Jun 2022 10:40:48 +0530 Subject: [PATCH 28/35] SWDEV-326631 - Adds sample which tests HIP language support with upstream CMake (#2743) Change-Id: I6b666cc0c5b1a53819515c5e343e517a2712d2ef --- .../22_cmake_hip_lang/CMakeLists.txt | 27 +++++ .../2_Cookbook/22_cmake_hip_lang/README.md | 16 +++ .../2_Cookbook/22_cmake_hip_lang/square.hip | 98 +++++++++++++++++++ 3 files changed, 141 insertions(+) create mode 100644 samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt create mode 100644 samples/2_Cookbook/22_cmake_hip_lang/README.md create mode 100644 samples/2_Cookbook/22_cmake_hip_lang/square.hip diff --git a/samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt b/samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt new file mode 100644 index 0000000000..bed43ff8ec --- /dev/null +++ b/samples/2_Cookbook/22_cmake_hip_lang/CMakeLists.txt @@ -0,0 +1,27 @@ +# 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 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. + +cmake_minimum_required(VERSION 3.21.3) + +project(cmake_hip_lang_support VERSION 1.0 + DESCRIPTION "HIP Language Support with upstream CMake" + LANGUAGES HIP) +# Create the executable +add_executable(square square.hip) diff --git a/samples/2_Cookbook/22_cmake_hip_lang/README.md b/samples/2_Cookbook/22_cmake_hip_lang/README.md new file mode 100644 index 0000000000..c770d75f63 --- /dev/null +++ b/samples/2_Cookbook/22_cmake_hip_lang/README.md @@ -0,0 +1,16 @@ +### This will test HIP language support in upstream CMake +I. Build +mkdir -p build; cd build +rm -rf *; cmake .. +make + +II. Test +$ ./square +info: running on device +info: allocate host mem ( 7.63 MB) +info: allocate device mem ( 7.63 MB) +info: copy Host2Device +info: launch 'vector_square' kernel +info: copy Device2Host +info: check result +PASSED! diff --git a/samples/2_Cookbook/22_cmake_hip_lang/square.hip b/samples/2_Cookbook/22_cmake_hip_lang/square.hip new file mode 100644 index 0000000000..7cee90818a --- /dev/null +++ b/samples/2_Cookbook/22_cmake_hip_lang/square.hip @@ -0,0 +1,98 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +#define CHECK(cmd) \ +{\ + hipError_t error = cmd;\ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + + +/* + * Square each element in the array A and write to array C. + */ +template +__global__ void +vector_square(T *C_d, T *A_d, size_t N) +{ + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; + + for (size_t i=offset; i Date: Fri, 17 Jun 2022 10:41:14 +0530 Subject: [PATCH 29/35] SWDEV-339995 - Adding hiprtcGetBitcode and hiprtcGetBitcodeSize. (#2744) Change-Id: I90ef8c90418ec3a1bcb596d513c1e77031fe77eb --- include/hip/hiprtc.h | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/include/hip/hiprtc.h b/include/hip/hiprtc.h index 9d4a598e01..61f7649305 100644 --- a/include/hip/hiprtc.h +++ b/include/hip/hiprtc.h @@ -269,6 +269,29 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code); */ hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet); +/** + * @brief Gets the pointer of compiled bitcode by the runtime compilation program instance. + * + * @param [in] prog runtime compilation program instance. + * @param [out] code char pointer to bitcode. + * @return HIPRTC_SUCCESS + * + * @see hiprtcResult + */ +hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode); + +/** + * @brief Gets the size of compiled bitcode by the runtime compilation program instance. + * + * + * @param [in] prog runtime compilation program instance. + * @param [out] code the size of bitcode. + * @return HIPRTC_SUCCESS + * + * @see hiprtcResult + */ +hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size); + /** * @brief Creates the link instance via hiprtc APIs. * From 4900ce120ef18ce2ebed0953542141a3d54a5b29 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 17 Jun 2022 11:02:17 +0530 Subject: [PATCH 30/35] SWDEV-341955 - Windows hipcc perl to handle quoted args (#2746) example - hipcc.pl -DGREETING=\"Hello\" hello.cpp Change-Id: Ib7192d46226563a44c87f3fbe2c22ef2ab84ac3e --- bin/hipcc | 8 +++++++- bin/hipcc.pl | 5 ++++- 2 files changed, 11 insertions(+), 2 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 32494dd345..86622c7a36 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -22,7 +22,6 @@ # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1; -use strict; use warnings; use File::Basename; @@ -30,6 +29,13 @@ use File::Spec::Functions 'catfile'; # TODO: By default select perl script until change incorporated in HIP build script. my $HIPCC_USE_PERL_SCRIPT = 1; +my $isWindows = ($^O eq 'MSWin32' or $^O eq 'msys'); +# escapes args with quotes SWDEV-341955 +foreach $arg (@ARGV) { + if ($isWindows) { + $arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\]/\\$&/g; + } +} my $SCRIPT_DIR=dirname(__FILE__); if ($HIPCC_USE_PERL_SCRIPT) { diff --git a/bin/hipcc.pl b/bin/hipcc.pl index 24edf7c59c..645ae62d03 100755 --- a/bin/hipcc.pl +++ b/bin/hipcc.pl @@ -595,9 +595,12 @@ foreach $arg (@ARGV) # common characters such as alphanumerics. # Do the quoting here because sometimes the $arg is changed in the loop # Important to have all of '-Xlinker' in the set of unquoted characters. - if (not $isWindows and $escapeArg) { # Windows needs different quoting, ignore for now + if (not $isWindows and $escapeArg) { $arg =~ s/[^-a-zA-Z0-9_=+,.\/]/\\$&/g; } + if ($isWindows and $escapeArg) { + $arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\]/\\$&/g; + } $toolArgs .= " $arg" unless $swallowArg; $prevArg = $arg; } From 14ed06e2887ddfe7052e7bfc3c52e2991fcc6f82 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 17 Jun 2022 11:38:51 +0530 Subject: [PATCH 31/35] Disable Unit_hipStreamPerThread_DeviceReset_2 Workaround flaky CI node --- .jenkins/Jenkinsfile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.jenkins/Jenkinsfile b/.jenkins/Jenkinsfile index 3cd6a96649..d64775c026 100644 --- a/.jenkins/Jenkinsfile +++ b/.jenkins/Jenkinsfile @@ -101,7 +101,7 @@ def hipBuildTest(String backendLabel) { # Check if backend label contains string "amd" or backend host is a server with amd gpu if [[ $backendLabel =~ amd ]]; then export HT_CONFIG_FILE="$HIP_DIR/tests/catch/hipTestMain/config/config_amd_linux.json" - LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple' + LLVM_PATH=/opt/rocm/llvm ctest -E 'Unit_hipGraphChildGraphNodeGetGraph_Functional|Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative|Unit_hipPtrGetAttribute_Simple|Unit_hipStreamPerThread_DeviceReset_2' sleep 120 else make test From 902eae64e32cf3c30b7e794b73cfb1c3d5012ff5 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Sun, 19 Jun 2022 20:49:15 -0700 Subject: [PATCH 32/35] SWDEV-327563 - Win:skip mempool and getUuid failing tests (#2748) --- tests/catch/hipTestMain/config/config_amd_windows.json | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 9fc36b8e4c..5ae00f6dc1 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -52,6 +52,13 @@ "Unit_hipStreamGetCaptureInfo_UniqueID", "Unit_hipStreamGetCaptureInfo_ArgValidation", "# Following test is related to ticket EXSWCPHIPT-41", - "Unit_hipStreamGetPriority_happy" + "Unit_hipStreamGetPriority_happy", + "Unit_hipMemPoolApi_Basic", + "Unit_hipMemPoolApi_BasicAlloc", + "Unit_hipMemPoolApi_BasicTrim", + "Unit_hipMemPoolApi_BasicReuse", + "Unit_hipMemPoolApi_Opportunistic", + "Unit_hipMemPoolApi_Default", + "Unit_hipDeviceGetUuid" ] } From 5afcd13390c6c2ada328db4e6769ba6c770d2ba1 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 20 Jun 2022 10:37:13 +0100 Subject: [PATCH 33/35] Add HIP_CHECK_THREAD and REQUIRE_THREAD macro for multi threaded HIP API tests (#2664) --- tests/catch/README.md | 125 +++++++++++-- tests/catch/hipTestMain/hip_test_context.cc | 38 ++++ tests/catch/include/hip_test_checkers.hh | 184 +++++++++++++++----- tests/catch/include/hip_test_common.hh | 133 +++++++++----- tests/catch/include/hip_test_context.hh | 28 +++ tests/catch/unit/stream/streamCommon.cc | 6 +- tests/catch/unit/stream/streamCommon.hh | 1 + 7 files changed, 414 insertions(+), 101 deletions(-) diff --git a/tests/catch/README.md b/tests/catch/README.md index ade70de8e9..7d1ff6265b 100644 --- a/tests/catch/README.md +++ b/tests/catch/README.md @@ -34,6 +34,22 @@ Some useful functions are: This information can be accessed in any test via using: `TestContext::get().isAmd()`. +## Adding test for a specific platform +There might be some functionality which is not present on some platforms. Those tests can be hidden inside following macros. + +- ```HT_AMD``` is 1 when tests are running on AMD platform and 0 on NVIDIA. +- ```HT_NVIDIA``` is 1 when tests are running on NVIDIA platform and 0 on AMD + +Usage: + +```cpp +#if HT_AMD +TEST_CASE("hipExtAPIs") { + // ... +} +#endif +``` + ## Config file schema Some tests can be skipped using a config file placed in hipTestMain/config folder. Multiple config files can be defined for different configurations. The naming convention for the file needs to be "config_platform_os_archname.json" @@ -56,16 +72,108 @@ The schema of the json file is as follows: } ``` -## Env Variables +## Environment Variables - `HT_CONFIG_FILE` : This variable can be set to the config file name or full path. Disabled tests will be read from this. - `HT_LOG_ENABLE` : This is for debugging the HIP Test Framework itself. Setting it to 1, all `LogPrintf` will be printed on screen +## Test Macros +### Single Thread Macros +These macros are to be used when your test is calling HIP APIs via the main thread. + +- `HIP_CHECK` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. + + - Usage: ```HIP_CHECK(hipMalloc(&dPtr, 10));``` + +- ```HIP_CHECK_ERROR``` : This macro takes in a HIP API and tests its result against a provided result. This can be used when the API is expected to fail with a particular result. + + - Usage: ```HIP_CHECK_ERROR(hipMalloc(&dPtr, 0), hipErrorInvalidValue);``` + +- ```HIPRTC_CHECK``` : This macro takes in a HIPRTC API and tests its result against HIPRTC_SUCCESS. + + - Usage: ```HIPRTC_CHECK(hiprtcCompileProgram(prog, count, options));``` + +- ```HIP_ASSERT``` : This macro takes in a bool condition as input and does a ```REQUIRE``` on the condition. + + - Usage: ```HIP_ASSERT(result == 10);``` + +### Multi Thread Macros +These macros are to be used when you call HIP APIs in a multi threaded way. They exist because Catch2 ```REQUIRE``` and ```CHECK``` macros can not handle multi threaded calls. To solve this problem, two macros are added```HIP_CHECK_THREAD``` and ```REQUIRE_THREAD``` which can be used to check result of HIP APIs and test assertions respectively. The results can be validate after the threads join via ```HIP_CHECK_THREAD_FINALIZE```. + +Note: These should used in ```std::thread``` only. For multi proc guidelines look at [MultiProc Macros](#multi-process-macros) and [SpawnProc Class](#multiproc-management-class) + +- ```HIP_CHECK_THREAD``` : This macro takes in a HIP API and tests for its result to be either ```hipSuccess``` or ```hipErrorPeerAccessAlreadyEnabled```. It can also tell other threads if an error has occured in one of the HIP API and can prematurely stop the threads. + +- ```REQUIRE_THREAD``` : This macro takes in a bool condition and tests for its result to be true. If this check fails, it can signal other threads to terminate early. + +- ```HIP_CHECK_THREAD_FINALIZE``` : This macro checks for the results logged by ```HIP_CHECK_THREAD```. This needs to be called after the threads have joined. + +Please also note that you can not return values in functions calling ```HIP_CHECK_THREAD``` or ```REQUIRE_THREAD``` macro. + + Usage: + + ```cpp + auto threadFunc = []() { + int *dPtr{nullptr}; + HIP_CHECK_THREAD(hipMalloc(&dPtr, 10)); + REQUIRE_THREAD(dPtr != nullptr); + // Some other work + }; + + // Launch threads + std::vector threadPool; + for(...) { + threadPool.emplace_back(std::thread(threadFunc)); + } + + // Join threads + for(auto &i : threadPool) { + i.join(); + } + + // Validate all results + HIP_CHECK_THREAD_FINALIZE(); + ``` + +### Skipping Tests if certain criteria is not met +If there arises a condition where certain flag is disabled and due to which a test can not run at that time, the following macro can be of use. It will highlight the test in ctest report as well. + +- ```HIP_SKIP_TEST``` : The api takes in an input of the reason as well and prints out the line HIP_SKIP_THIS_TEST. This causes ctest to mark the test as skipped and the test shows up in the report as skipped prompting proper response from the team. + + Usage: + + ```cpp + TEST_CASE("TestOnlyOnXnack") { + if(!XNACKEnabled) { + HIP_SKIP_TEST("Test only runs on system with XNACK enabled"); + return; + } + // Rest of test functionality + } + ``` + +### Multi Process Macros +These macros are to be called in multi process tests, inside a process which gets spawned. The reasoning is the same, Catch2 does not support multi process checks. + +- ```HIPCHECK``` : Same as ```HIP_CHECK``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process. + +- ```HIPASSERT``` : Same as ```HIP_ASSERT``` but will not call Catch2's ```REQUIRE``` on the HIP API. It will print if there is a mismatch and exit the process. + +## MultiProc Management Class +There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn a process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf output. +Sample Usage: +```cpp +hip::SpawnProc proc(, ); +REQUIRE(0 == proc.run()); // Test of return value of the proc +REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process +``` +The process can be a standalone exe (see tests/catch/unit/printfExe for more information). + ## Enabling New Tests -Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=ON```. After porting existing tests, this will be turned on by default. +Initially, the new tests can be enabled via using ```-DHIP_CATCH_TEST=1```. After porting existing tests, this will be turned on by default. ## Building a single test ```bash -hipcc -I/tests/newTests/include /tests/newTests/hipTestMain/standalone_main.cc -I/tests/newTests/external/Catch2 -g -o +hipcc -I/tests/catch/include /tests/catch/hipTestMain/standalone_main.cc -I/tests/catch/external/Catch2 -g -o ``` ## Debugging support @@ -87,16 +195,7 @@ Tests fall in 5 categories and its file name prefix are as follows: - Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment - Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc y): Performance tests are used to get results of HIP APIs. -There is a special interface available for process isolation. ```hip::SpawnProc``` in ```hip_test_process.hh```. Using this interface test can spawn of process and place passing conditions on its return value or its output to stdout. This can be useful for testing printf tests. -Sample Usage: -```cpp -hip::SpawnProc proc(, ); -REQUIRE(0 == proc.run()); // Test of return value of the proc -REQUIRE(exepctedOutput == proc.getOutput()); // Test on expected output of the process -``` -The process can be a standalone exe (see tests/catch/unit/printfExe for more information). - -General Guidelines: +# General Guidelines: - Do not use the catch2 tags. Tags wont be used for filtering - Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails) - Check return of each HIP API and fail whenever there is a misma tch with hipSuccess or hiprtcSuccess. diff --git a/tests/catch/hipTestMain/hip_test_context.cc b/tests/catch/hipTestMain/hip_test_context.cc index a6e3a08609..597813b4c8 100644 --- a/tests/catch/hipTestMain/hip_test_context.cc +++ b/tests/catch/hipTestMain/hip_test_context.cc @@ -229,4 +229,42 @@ hipFunction_t TestContext::getFunction(const std::string kernelNameExpression) { } else { return nullptr; } +} + +void TestContext::addResults(HCResult r) { + std::unique_lock lock(resultMutex); + results.push_back(r); + if ((!r.conditionsResult) || + ((r.result != hipSuccess) && (r.result != hipErrorPeerAccessAlreadyEnabled))) { + hasErrorOccured_.store(true); + } +} + +void TestContext::finalizeResults() { + std::unique_lock lock(resultMutex); + // clear the results whatever happens + std::shared_ptr emptyVec(nullptr, [this](auto) { results.clear(); }); + + for (const auto& i : results) { + INFO("HIP API Result check\n File:: " + << i.file << "\n Line:: " << i.line << "\n API:: " << i.call + << "\n Result:: " << i.result << "\n Result Str:: " << hipGetErrorString(i.result)); + REQUIRE(((i.result == hipSuccess) || (i.result == hipErrorPeerAccessAlreadyEnabled))); + REQUIRE(i.conditionsResult); + } + hasErrorOccured_.store(false); // Clear the flag +} + +bool TestContext::hasErrorOccured() { return hasErrorOccured_.load(); } + +TestContext::~TestContext() { + // Show this message when there are unchecked results + if (results.size() != 0) { + std::cerr << "HIP_CHECK_THREAD_FINALIZE() has not been called after HIP_CHECK_THREAD\n" + << "Please call HIP_CHECK_THREAD_FINALIZE after joining threads\n" + << "There is/are " << results.size() << " unchecked results from threads." + << std::endl; + std::abort(); // Crash to bring users attention to this message and avoid accidental passing of + // tests without checking for errors + } } \ No newline at end of file diff --git a/tests/catch/include/hip_test_checkers.hh b/tests/catch/include/hip_test_checkers.hh index 77ac0f4d74..fce61fe677 100644 --- a/tests/catch/include/hip_test_checkers.hh +++ b/tests/catch/include/hip_test_checkers.hh @@ -23,17 +23,17 @@ THE SOFTWARE. #pragma once #include "hip_test_common.hh" #include -#include -#include +#include +#include #include -#define guarantee(cond, str) \ - { \ - if (!(cond)) { \ - INFO("guarantee failed: " << str); \ - abort(); \ - } \ - } +#define guarantee(cond, str) \ + { \ + if (!(cond)) { \ + INFO("guarantee failed: " << str); \ + abort(); \ + } \ + } namespace HipTest { @@ -73,15 +73,15 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM return mismatchCount; } -template // pointer type -bool checkArray(T* hData, T* hOutputData, size_t width, size_t height,size_t depth = 1) { +template // pointer type +bool checkArray(T* hData, T* hOutputData, size_t width, size_t height, size_t depth = 1) { for (size_t i = 0; i < depth; i++) { for (size_t j = 0; j < height; j++) { for (size_t k = 0; k < width; k++) { - int offset = i*width*height + j*width + k; + int offset = i * width * height + j * width + k; if (hData[offset] != hOutputData[offset]) { - INFO("Mismatch at [" << i << "," << j << "," << k << "]:" - << hData[offset] << "----" << hOutputData[offset]); + INFO("Mismatch at [" << i << "," << j << "," << k << "]:" << hData[offset] << "----" + << hOutputData[offset]); CHECK(false); return false; } @@ -120,7 +120,7 @@ template void setDefaultData(size_t numElements, T* A_h, T* B_h, T* if (A_h) A_h[i] = 3; if (B_h) B_h[i] = 4; if (C_h) C_h[i] = 5; - } else if(std::is_same::value || std::is_same::value) { + } else if (std::is_same::value || std::is_same::value) { if (A_h) A_h[i] = 'a'; if (B_h) B_h[i] = 'b'; if (C_h) C_h[i] = 'c'; @@ -185,6 +185,110 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N, return initArraysForHost(A_h, B_h, C_h, N, usePinnedHost); } +// Threaded version of setDefaultData to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template void setDefaultDataT(size_t numElements, T* A_h, T* B_h, T* C_h) { + // Initialize the host data: + + for (size_t i = 0; i < numElements; i++) { + if (std::is_same::value || std::is_same::value) { + if (A_h) A_h[i] = 3; + if (B_h) B_h[i] = 4; + if (C_h) C_h[i] = 5; + } else if (std::is_same::value || std::is_same::value) { + if (A_h) A_h[i] = 'a'; + if (B_h) B_h[i] = 'b'; + if (C_h) C_h[i] = 'c'; + } else { + if (A_h) A_h[i] = 3.146f + i; + if (B_h) B_h[i] = 1.618f + i; + if (C_h) C_h[i] = 1.4f + i; + } + } +} + +// Threaded version of initArraysForHost to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template +void initArraysForHostT(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost = false) { + size_t Nbytes = N * sizeof(T); + + if (usePinnedHost) { + if (A_h) { + HIP_CHECK_THREAD(hipHostMalloc((void**)A_h, Nbytes)); + } + if (B_h) { + HIP_CHECK_THREAD(hipHostMalloc((void**)B_h, Nbytes)); + } + if (C_h) { + HIP_CHECK_THREAD(hipHostMalloc((void**)C_h, Nbytes)); + } + } else { + if (A_h) { + *A_h = (T*)malloc(Nbytes); + REQUIRE_THREAD(*A_h != nullptr); + } + + if (B_h) { + *B_h = (T*)malloc(Nbytes); + REQUIRE_THREAD(*B_h != nullptr); + } + + if (C_h) { + *C_h = (T*)malloc(Nbytes); + REQUIRE_THREAD(*C_h != nullptr); + } + } + + setDefaultDataT(N, A_h ? *A_h : nullptr, B_h ? *B_h : nullptr, C_h ? *C_h : nullptr); +} + +// Threaded version of initArrays to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template +void initArraysT(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N, + bool usePinnedHost = false) { + size_t Nbytes = N * sizeof(T); + + if (A_d) { + HIP_CHECK_THREAD(hipMalloc(A_d, Nbytes)); + } + if (B_d) { + HIP_CHECK_THREAD(hipMalloc(B_d, Nbytes)); + } + if (C_d) { + HIP_CHECK_THREAD(hipMalloc(C_d, Nbytes)); + } + + initArraysForHostT(A_h, B_h, C_h, N, usePinnedHost); +} + +// Threaded version of freeArraysForHost to be called from multi thread tests +// Call HIP_CHECK_THREAD_FINALIZE after joining +template void freeArraysForHostT(T* A_h, T* B_h, T* C_h, bool usePinnedHost) { + if (usePinnedHost) { + if (A_h) { + HIP_CHECK_THREAD(hipHostFree(A_h)); + } + if (B_h) { + HIP_CHECK_THREAD(hipHostFree(B_h)); + } + if (C_h) { + HIP_CHECK_THREAD(hipHostFree(C_h)); + } + } else { + if (A_h) { + free(A_h); + } + if (B_h) { + free(B_h); + } + if (C_h) { + free(C_h); + } + } +} + template bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) { if (usePinnedHost) { if (A_h) { @@ -210,6 +314,21 @@ template bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePin return true; } +template +void freeArraysT(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) { + if (A_d) { + HIP_CHECK_THREAD(hipFree(A_d)); + } + if (B_d) { + HIP_CHECK_THREAD(hipFree(B_d)); + } + if (C_d) { + HIP_CHECK_THREAD(hipFree(C_d)); + } + + freeArraysForHostT(A_h, B_h, C_h, usePinnedHost); +} + template bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) { if (A_d) { @@ -226,20 +345,6 @@ bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHo } template -unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock, - size_t N) { - int device; - HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t props; - HIP_CHECK(hipGetDeviceProperties(&props, device)); - - unsigned blocks = props.multiProcessorCount * blocksPerCU; - if (blocks * threadsPerBlock > N) { - blocks = (N + threadsPerBlock - 1) / threadsPerBlock; - } - return blocks; -} -template static bool assemblyFile_Verification(std::string assemfilename, std::string inst) { std::string filePath = "./catch/unit/deviceLib/"; bool result = false; @@ -254,34 +359,27 @@ static bool assemblyFile_Verification(std::string assemfilename, std::string ins while (getline(file, line)) { line_pos++; if ((std::is_same::value)) { - if (!start_pos && - std::regex_search(line, - std::regex("Begin function (.*)AtomicCheck"))) { + if (!start_pos && std::regex_search(line, std::regex("Begin function (.*)AtomicCheck"))) { start_pos = line_pos; } - if (!last_pos && - std::regex_search(line, - std::regex(".Lfunc_end0-(.*)AtomicCheck"))) { + if (!last_pos && std::regex_search(line, std::regex(".Lfunc_end0-(.*)AtomicCheck"))) { last_pos = line_pos; break; } } else { - if ((start_match != 2) && std::regex_search(line, - std::regex("Begin function (.*)AtomicCheck"))) { + if ((start_match != 2) && + std::regex_search(line, std::regex("Begin function (.*)AtomicCheck"))) { start_match++; - if (start_match == 2) - start_pos = line_pos; + if (start_match == 2) start_pos = line_pos; } - if (!last_pos && std::regex_search(line, - std::regex("func_end1-(.*)AtomicCheck"))) { + if (!last_pos && std::regex_search(line, std::regex("func_end1-(.*)AtomicCheck"))) { last_pos = line_pos; break; } } if (start_pos) { result = std::regex_search(line, std::regex(inst)); - if (result) - break; + if (result) break; } } } else { diff --git a/tests/catch/include/hip_test_common.hh b/tests/catch/include/hip_test_common.hh index 63b8354421..fe272f7ace 100644 --- a/tests/catch/include/hip_test_common.hh +++ b/tests/catch/include/hip_test_common.hh @@ -22,9 +22,14 @@ THE SOFTWARE. #pragma once #include "hip_test_context.hh" -#include #include +#include +#include #include +#include +#include +#include +#include #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__); @@ -33,22 +38,51 @@ THE SOFTWARE. { \ hipError_t localError = error; \ if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ - INFO("Error: " << hipGetErrorString(localError) << " Code: " << localError << " Str: " \ - << #error << " In File: " << __FILE__ << " At line: " << __LINE__); \ + INFO("Error: " << hipGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ REQUIRE(false); \ } \ } +// Threaded HIP_CHECKs +#define HIP_CHECK_THREAD(error) \ + { \ + /*To see if error has occured in previous threads, stop execution */ \ + if (TestContext::get().hasErrorOccured() == true) { \ + return; /*This will only work with std::thread and not with std::async*/ \ + } \ + auto localError = error; \ + HCResult result(__LINE__, __FILE__, localError, #error); \ + TestContext::get().addResults(result); \ + } + +#define REQUIRE_THREAD(condition) \ + { \ + /*To see if error has occured in previous threads, stop execution */ \ + if (TestContext::get().hasErrorOccured() == true) { \ + return; /*This will only work with std::thread and not with std::async*/ \ + } \ + auto localResult = (condition); \ + HCResult result(__LINE__, __FILE__, hipSuccess, #condition, localResult); \ + TestContext::get().addResults(result); \ + } + +// Do not call before all threads have joined +#define HIP_CHECK_THREAD_FINALIZE() \ + { TestContext::get().finalizeResults(); } + + // Check that an expression, errorExpr, evaluates to the expected error_t, expectedError. #define HIP_CHECK_ERROR(errorExpr, expectedError) \ { \ hipError_t localError = errorExpr; \ INFO("Matching Errors: " \ - << " Expected Error: " << hipGetErrorString(expectedError) \ - << " Expected Code: " << expectedError << '\n' \ + << "\n Expected Error: " << hipGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << '\n' \ << " Actual Error: " << hipGetErrorString(localError) \ - << " Actual Code: " << localError << "\nStr: " << #errorExpr \ - << "\nIn File: " << __FILE__ << " At line: " << __LINE__); \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ REQUIRE(localError == expectedError); \ } @@ -57,8 +91,9 @@ THE SOFTWARE. { \ auto localError = error; \ if (localError != HIPRTC_SUCCESS) { \ - INFO("Error: " << hiprtcGetErrorString(localError) << " Code: " << localError << " Str: " \ - << #error << " In File: " << __FILE__ << " At line: " << __LINE__); \ + INFO("Error: " << hiprtcGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ REQUIRE(false); \ } \ } @@ -67,12 +102,6 @@ THE SOFTWARE. #define HIP_ASSERT(x) \ { REQUIRE((x)); } -#ifdef __cplusplus -#include -#include -#include -#endif - #define HIPCHECK(error) \ { \ hipError_t localError = error; \ @@ -84,19 +113,20 @@ THE SOFTWARE. } #define HIPASSERT(condition) \ - if (!(condition)) { \ - printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ - abort(); \ - } + if (!(condition)) { \ + printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ + abort(); \ + } + #if HT_NVIDIA -#define CTX_CREATE() \ - hipCtx_t context;\ +#define CTX_CREATE() \ + hipCtx_t context; \ initHipCtx(&context); #define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context)); #define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array)); #define HIP_TEX_REFERENCE hipTexRef #define HIP_ARRAY hiparray -static void initHipCtx(hipCtx_t *pcontext) { +static void initHipCtx(hipCtx_t* pcontext) { HIPCHECK(hipInit(0)); hipDevice_t device; HIPCHECK(hipDeviceGet(&device, 0)); @@ -130,9 +160,9 @@ static inline double elapsed_time(long long startTimeUs, long long stopTimeUs) { } static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) { - int device; + int device{0}; HIP_CHECK(hipGetDevice(&device)); - hipDeviceProp_t props; + hipDeviceProp_t props{}; HIP_CHECK(hipGetDeviceProperties(&props, device)); unsigned blocks = props.multiProcessorCount * blocksPerCU; @@ -143,23 +173,40 @@ static inline unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlo return blocks; } -static inline int RAND_R(unsigned* rand_seed) -{ - #if defined(_WIN32) || defined(_WIN64) - srand(*rand_seed); - return rand(); - #else - return rand_r(rand_seed); - #endif +// Threaded version of setNumBlocks - to be used in multi threaded test +// Why? because catch2 does not support multithreaded macro calls +// Make sure you call HIP_CHECK_THREAD_FINALIZE after your threads join +// Also you can not return in threaded functions, due to how HIP_CHECK_THREAD works +static inline void setNumBlocksThread(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N, + unsigned& blocks) { + int device{0}; + blocks = 0; // incase error has occured in some other thread and the next call might not execute, + // we set the blocks size to 0 + HIP_CHECK_THREAD(hipGetDevice(&device)); + hipDeviceProp_t props{}; + HIP_CHECK_THREAD(hipGetDeviceProperties(&props, device)); + + blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } +} + +static inline int RAND_R(unsigned* rand_seed) { +#if defined(_WIN32) || defined(_WIN64) + srand(*rand_seed); + return rand(); +#else + return rand_r(rand_seed); +#endif } inline bool isImageSupported() { - int imageSupport = 1; + int imageSupport = 1; #ifdef __HIP_PLATFORM_AMD__ - int device; - HIP_CHECK(hipGetDevice(&device)); - HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, - device)); + int device; + HIP_CHECK(hipGetDevice(&device)); + HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device)); #endif return imageSupport != 0; } @@ -217,8 +264,8 @@ template void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, hipStream_t stream, Args&&... packedArgs) { #ifndef RTC_TESTING - validateArguments(kernel, packedArgs...); - kernel<<>>(std::forward(packedArgs)...); + validateArguments(kernel, packedArgs...); + kernel<<>>(std::forward(packedArgs)...); #else launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, std::forward(packedArgs)...); @@ -229,6 +276,8 @@ void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerB // This must be called in the beginning of image test app's main() to indicate whether image // is supported. -#define checkImageSupport() \ - if (!HipTest::isImageSupported()) \ - { printf("Texture is not support on the device. Skipped.\n"); return; } +#define checkImageSupport() \ + if (!HipTest::isImageSupported()) { \ + printf("Texture is not support on the device. Skipped.\n"); \ + return; \ + } diff --git a/tests/catch/include/hip_test_context.hh b/tests/catch/include/hip_test_context.hh index 2d630b70f7..adcaf25aad 100644 --- a/tests/catch/include/hip_test_context.hh +++ b/tests/catch/include/hip_test_context.hh @@ -23,7 +23,11 @@ THE SOFTWARE. #pragma once #include #include + +#include +#include #include +#include #include #include #include @@ -64,6 +68,18 @@ typedef struct Config_ { std::string os; // windows/linux } Config; +// Store Multi threaded results +struct HCResult { + size_t line; // Line of check (HIP_CHECK_THREAD or REQUIRE_THREAD) + std::string file; // File name of the check + hipError_t result; // hipResult for HIP_CHECK_THREAD, for conditions its hipSuccess + std::string call; // Call of HIP API or a bool condition + bool conditionsResult; // If bool condition, result of call. For HIP Calls its true + HCResult(size_t l, std::string f, hipError_t r, std::string c, bool b = true) + : line(l), file(f), result(r), call(c), conditionsResult(b) {} +}; + + class TestContext { bool p_windows = false, p_linux = false; // OS bool amd = false, nvidia = false; // HIP Platform @@ -97,6 +113,11 @@ class TestContext { TestContext(int argc, char** argv); + // Multi threaded checks helpers + std::mutex resultMutex; + std::vector results; // Multi threaded test results buffer + std::atomic hasErrorOccured_{false}; + public: static TestContext& get(int argc = 0, char** argv = nullptr) { static TestContext instance(argc, argv); @@ -112,6 +133,11 @@ class TestContext { const std::string& getCurrentTest() const { return current_test; } std::string currentPath() const; + // Multi threaded results helpers + void addResults(HCResult r); // Add multi threaded results + void finalizeResults(); // Validate on all results + bool hasErrorOccured(); // Query if error has occured + /** * @brief Unload all loaded modules. * Note: This function needs to be called at the end of each test that uses RTC. @@ -142,4 +168,6 @@ class TestContext { TestContext(const TestContext&) = delete; void operator=(const TestContext&) = delete; + + ~TestContext(); }; diff --git a/tests/catch/unit/stream/streamCommon.cc b/tests/catch/unit/stream/streamCommon.cc index 265142e23b..14ac4000eb 100644 --- a/tests/catch/unit/stream/streamCommon.cc +++ b/tests/catch/unit/stream/streamCommon.cc @@ -95,11 +95,11 @@ __global__ void waiting_kernel(int* semaphore) { std::thread startSignalingThread(int* semaphore) { std::thread signalingThread([semaphore]() { hipStream_t signalingStream; - HIP_CHECK(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking)); + HIP_CHECK_THREAD(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking)); signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore); - HIP_CHECK(hipStreamSynchronize(signalingStream)); - HIP_CHECK(hipStreamDestroy(signalingStream)); + HIP_CHECK_THREAD(hipStreamSynchronize(signalingStream)); + HIP_CHECK_THREAD(hipStreamDestroy(signalingStream)); }); return signalingThread; diff --git a/tests/catch/unit/stream/streamCommon.hh b/tests/catch/unit/stream/streamCommon.hh index db73f1f668..1d5a1ea958 100644 --- a/tests/catch/unit/stream/streamCommon.hh +++ b/tests/catch/unit/stream/streamCommon.hh @@ -44,6 +44,7 @@ __global__ void waiting_kernel(int* semaphore = nullptr); /** * @brief Creates a thread that runs a signaling_kernel on a non-blocking stream. * hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream. + * You must call HIP_CHECK_THREAD_FINALIZE after joining this thread. * * @param semaphore memory location to signal * @return std::thread thread that has to be joined after the testing is done. From 50829a3b00566604f165594b5bc90a864030e38b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Mon, 20 Jun 2022 11:51:01 +0100 Subject: [PATCH 34/35] EXSWCPHIPT-107: Added testing for hipEventCreate (#2729) --- .../unit/event/Unit_hipEvent_Negative.cc | 33 ++++++++++++++++++- 1 file changed, 32 insertions(+), 1 deletion(-) diff --git a/tests/catch/unit/event/Unit_hipEvent_Negative.cc b/tests/catch/unit/event/Unit_hipEvent_Negative.cc index 1d8b2b29a4..76b75d6c06 100644 --- a/tests/catch/unit/event/Unit_hipEvent_Negative.cc +++ b/tests/catch/unit/event/Unit_hipEvent_Negative.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 @@ -46,3 +46,34 @@ TEST_CASE("Unit_hipEventDestroy_NullCheck") { auto res = hipEventDestroy(nullptr); REQUIRE(res != hipSuccess); } + +TEST_CASE("Unit_hipEventCreate_IncompatibleFlags") { + hipEvent_t event; + +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-106"); + return; +#endif + + HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, hipEventInterprocess), hipErrorInvalidValue); + +#if HT_AMD + HIP_CHECK_ERROR( + hipEventCreateWithFlags(&event, hipEventReleaseToDevice | hipEventReleaseToSystem), + hipErrorInvalidValue); +#endif + + unsigned allFlags{hipEventReleaseToDevice | hipEventReleaseToSystem | hipEventBlockingSync | + hipEventDisableTiming | hipEventDefault | hipEventInterprocess}; + +#if HT_AMD + HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, allFlags), hipErrorInvalidValue); +#else + /* Works on Non-AMD because hipEventReleaseToDevice / hipEventReleaseToSystem have no meaning in + * that case */ + HIP_CHECK(hipEventCreateWithFlags(&event, allFlags)); +#endif + + unsigned invalidFlag{0x08000000}; + HIP_CHECK_ERROR(hipEventCreateWithFlags(&event, invalidFlag), hipErrorInvalidValue); +} \ No newline at end of file From cc96fdec0292de186c4c79bbb69ffe99ebc468b3 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 20 Jun 2022 11:51:13 +0100 Subject: [PATCH 35/35] Add tests for multi stream usage scenarios (#2608) --- tests/catch/stress/stream/CMakeLists.txt | 4 +- tests/catch/stress/stream/streamEnqueue.cc | 231 +++++++++++++++++++++ 2 files changed, 234 insertions(+), 1 deletion(-) create mode 100644 tests/catch/stress/stream/streamEnqueue.cc diff --git a/tests/catch/stress/stream/CMakeLists.txt b/tests/catch/stress/stream/CMakeLists.txt index 357c3e5e36..ae6685da58 100644 --- a/tests/catch/stress/stream/CMakeLists.txt +++ b/tests/catch/stress/stream/CMakeLists.txt @@ -1,8 +1,10 @@ # Common Tests - Test independent of all platforms set(TEST_SRC Stress_hipStreamCreate.cc + streamEnqueue.cc ) hip_add_exe_to_target(NAME stream TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME stress_test) + TEST_TARGET_NAME stress_test + COMPILE_OPTIONS -std=c++14) diff --git a/tests/catch/stress/stream/streamEnqueue.cc b/tests/catch/stress/stream/streamEnqueue.cc new file mode 100644 index 0000000000..c44ff96afc --- /dev/null +++ b/tests/catch/stress/stream/streamEnqueue.cc @@ -0,0 +1,231 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +#include +#include +#include +#include +#include +#include + +__global__ void addVal(unsigned long long* ptr, size_t index, unsigned long long val) { + atomicAdd(ptr + index, val); +} + +// Create a copy constructible AtomicWrap around std::atomic so that we can put it in a vector +template struct AtomicWrap { + std::atomic data; + + AtomicWrap() : data() {} + + AtomicWrap(T i) : data(i) {} + + AtomicWrap(const std::atomic& a) : data(a.load()) {} + + AtomicWrap(const AtomicWrap& other) : data(other.data.load()) {} + + AtomicWrap& operator=(const AtomicWrap& other) { + data.store(other.data.load()); + return *this; + } +}; + +// Have multiple threads and enqueue commands from them on a single stream +// Validate at the end that all commands have completed successfully +TEST_CASE("Stress_StreamEnqueue_DifferentThreads") { + auto hwThreads = std::thread::hardware_concurrency(); + hwThreads = (hwThreads >= 2) ? hwThreads : 2; // Run atleast 2 threads + + std::vector> hostData(hwThreads, 0); + + unsigned long long* dPtr{nullptr}; + HIP_CHECK(hipMalloc(&dPtr, sizeof(unsigned long long) * hwThreads)); + REQUIRE(dPtr != nullptr); + + HIP_CHECK(hipMemset(dPtr, 0, sizeof(unsigned long long) * hwThreads)); + + std::random_device device; + std::mt19937 engine(device()); + + constexpr size_t maxWork = 10000; + constexpr size_t maxVal = 10; + + std::uniform_int_distribution genIndex(0, hwThreads - 1); + std::uniform_int_distribution genWork(0, maxWork); + std::uniform_int_distribution genVal(0, maxVal); + + auto enqueueKernelThread = [&](hipStream_t stream) { + auto iter = genWork(engine); // Generate work to be done via thread + for (auto i = 0; i < iter; i++) { + auto index = genIndex(engine); // Generate Index to add to + auto val = genVal(engine); // Generate value to add to the destination + hostData[index].data += val; // Replicate it on host + addVal<<<1, 1, 0, stream>>>(dPtr, static_cast(index), + static_cast(val)); // And on device + } + }; + + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + std::vector threadPool{}; + threadPool.reserve(hwThreads); + + // Launch work + for (size_t i = 0; i < hwThreads; i++) { + threadPool.emplace_back(std::thread(enqueueKernelThread, stream)); + } + + // Wait for work to finish + for (auto& i : threadPool) { + i.join(); + } + + HIP_CHECK(hipStreamDestroy(stream)); + + auto hPtr = std::make_unique(hwThreads); + HIP_CHECK( + hipMemcpy(hPtr.get(), dPtr, sizeof(unsigned long long) * hwThreads, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipFree(dPtr)); + + // Validate that CPU and GPU has the same results + for (size_t i = 0; i < hwThreads; i++) { + INFO("Check for Index " << i); + REQUIRE(hostData[i].data.load() == hPtr[i]); + } +} + +__global__ void doOperation(int* dPtr, size_t size, int val) { + auto i = threadIdx.x; + atomicAdd(dPtr + i, val); +} + +// Allocate mulitple stream for same device. +// Same device stream operate on same memory +TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { + int deviceCount{0}; + HIP_CHECK(hipGetDeviceCount(&deviceCount)); + REQUIRE(deviceCount > 0); + + // Skip the test if devices less than 2 + if (deviceCount <= 1) { + HipTest::HIP_SKIP_TEST("Skipping because devices <= 1"); + return; + } + + constexpr size_t streamPerGPU{3}; // Stream per gpu + + std::vector streamPool{}; + streamPool.reserve(deviceCount * streamPerGPU); + + std::map streamToDeviceMemory; // Map of stream and device memory + std::map> streamToHostMemory; // Map of stream and host result + std::map streamToDeviceIndex; // Map of stream and device it was created on + + constexpr size_t size = 1024; + + for (size_t i = 0; i < deviceCount; i++) { + HIP_CHECK(hipSetDevice(i)); + + for (size_t j = 0; j < streamPerGPU; j++) { + hipStream_t stream{nullptr}; + HIP_CHECK(hipStreamCreate(&stream)); + REQUIRE(stream != nullptr); + streamPool.push_back(stream); + + int* dPtr{nullptr}; + HIP_CHECK(hipMalloc(&dPtr, sizeof(int) * size)); + REQUIRE(dPtr != nullptr); + HIP_CHECK(hipMemset(dPtr, 0, sizeof(int) * size)); + + streamToDeviceMemory[stream] = dPtr; // All streams work on exclusive memory + + streamToHostMemory[stream] = AtomicWrap(0); // CPU result + + streamToDeviceIndex[stream] = i; // Capture device id for stream + } + } + + constexpr size_t maxVal = 5; + constexpr size_t maxWorkPerThread = 10000; + + // Boiler plate code to generate a random number + std::random_device device; + std::mt19937 engine(device()); + + std::uniform_int_distribution genVal(-maxVal, maxVal); + std::uniform_int_distribution genStream(0, streamPool.size() - 1); + +#if HT_NVIDIA + std::mutex ness; // On nvidia, current device needs to match stream's device +#endif + + auto enqueueKernelThread = [&]() { + for (size_t i = 0; i < maxWorkPerThread; i++) { +#if HT_NVIDIA + std::unique_lock lock(ness); // Lock on creation +#endif + hipStream_t stream = streamPool[genStream(engine)]; // Get a random stream + + // TODO use HIP_CHECK_THREAD when PR#2664 is merged + if (hipSuccess != hipSetDevice(streamToDeviceIndex[stream])) { + return; + } + + int val = genVal(engine); // Generate Value to add/sub to + + streamToHostMemory[stream].data.fetch_add(val); // Replicate result on CPU + auto dPtr = streamToDeviceMemory[stream]; + doOperation<<<1, 1024, 0, stream>>>(dPtr, size, + val); // On GPU + } + }; + + auto maxThreads = std::thread::hardware_concurrency(); + maxThreads = (maxThreads >= 2) ? maxThreads : 2; // Run atleast 2 threads + + std::vector threadPool{}; + threadPool.reserve(maxThreads); + + // Launch Threads + for (size_t i = 0; i < maxThreads; i++) { + threadPool.emplace_back(std::thread(enqueueKernelThread)); + } + + // Wait for them to stop + for (auto& i : threadPool) { + i.join(); + } + + // Sync and check results + for (auto& i : streamPool) { + HIP_CHECK(hipStreamSynchronize(i)); + auto dResult = std::make_unique(size); + HIP_CHECK(hipMemcpy(dResult.get(), streamToDeviceMemory[i], sizeof(int) * size, + hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(streamToDeviceMemory[i])); + HIP_CHECK(hipStreamDestroy(i)); + auto res = streamToHostMemory[i].data.load(); + INFO("Matching CPU: " << res << " GPU: " << dResult[0] << " Dev Ptr: " + << streamToDeviceMemory[i] << " on Device: " << streamToDeviceIndex[i]); + REQUIRE(std::all_of(dResult.get(), dResult.get() + size, [=](int r) { return r == res; })); + } +}