From ce1d49b6e057aca654b4d6bf86a7c49a4a8088fd Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 08:56:44 +0530 Subject: [PATCH 1/7] SWDEV-1 - Fixed Malloc test failures (#2847) Change-Id: I27042e743dda78f34e5f6b0d548e7d507b4f0ba3 --- .../multiproc/hipMallocConcurrencyMproc.cc | 18 +++++++++++++++++- .../catch/unit/memory/hipMallocConcurrency.cc | 18 +++++++++++++++++- 2 files changed, 34 insertions(+), 2 deletions(-) diff --git a/tests/catch/multiproc/hipMallocConcurrencyMproc.cc b/tests/catch/multiproc/hipMallocConcurrencyMproc.cc index 9018f694a2..d13a31949d 100644 --- a/tests/catch/multiproc/hipMallocConcurrencyMproc.cc +++ b/tests/catch/multiproc/hipMallocConcurrencyMproc.cc @@ -103,6 +103,14 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { HIP_CHECK(hipSetDevice(gpu)); HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) { + //In concurrent calls on one GPU, we cannot verify leaking in this way + printf("%s : Memory allocation mismatch observed." + "Possible memory leak.\n", __func__); + TestPassed &= false; + } unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -122,9 +130,17 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { TestPassed = false; } + HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); + if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) { + // In concurrent calls on one GPU, we cannot verify leaking in this way + UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed." + << "Possible memory leak."); + TestPassed = false; + } + if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { // In concurrent calls on one GPU, we cannot verify leaking in this way printf( @@ -200,7 +216,7 @@ TEST_CASE("Unit_hipMalloc_ChildConcurrencyMultiGpu") { } else if (!pid) { // Child process bool TestPassedChild = false; - TestPassedChild = validateMemoryOnGPU(gpu); + TestPassedChild = validateMemoryOnGPU(gpu, true); if (TestPassedChild) { exit(resSuccess); // child exit with success status diff --git a/tests/catch/unit/memory/hipMallocConcurrency.cc b/tests/catch/unit/memory/hipMallocConcurrency.cc index 647d38d99d..1283284022 100644 --- a/tests/catch/unit/memory/hipMallocConcurrency.cc +++ b/tests/catch/unit/memory/hipMallocConcurrency.cc @@ -203,6 +203,14 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { HIPCHECK(hipSetDevice(gpu)); HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + HIPCHECK(hipMemGetInfo(&curAvl, &curTot)); + + if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) { + //In concurrent calls on one GPU, we cannot verify leaking in this way + printf("%s : Memory allocation mismatch observed." + "Possible memory leak.\n", __func__); + TestPassed &= false; + } unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -221,9 +229,17 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { TestPassed = false; } + HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIPCHECK(hipMemGetInfo(&curAvl, &curTot)); + if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) { + // In concurrent calls on one GPU, we cannot verify leaking in this way + UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed." + << "Possible memory leak."); + TestPassed = false; + } + if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { // In concurrent calls on one GPU, we cannot verify leaking in this way UNSCOPED_INFO( @@ -291,7 +307,7 @@ static bool regressAllocInLoopMthread(int gpu) { * Thread func to regress alloc and check data consistency */ static void threadFunc(int gpu) { - g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu); + g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu, true); UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << g_thTestPassed.load()); } From ec3700fe20eee7832def2fe5a8f2d530961de263 Mon Sep 17 00:00:00 2001 From: Finlay Date: Fri, 5 Aug 2022 06:29:19 +0100 Subject: [PATCH 2/7] EXSWCPHIPT-139 - added tests for hipStreamDestroy (#2548) --- tests/catch/unit/stream/CMakeLists.txt | 2 + tests/catch/unit/stream/hipStreamDestroy.cc | 97 +++++++++++++++++++++ 2 files changed, 99 insertions(+) create mode 100644 tests/catch/unit/stream/hipStreamDestroy.cc diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt index 6cdbac9a4e..9b829207cf 100644 --- a/tests/catch/unit/stream/CMakeLists.txt +++ b/tests/catch/unit/stream/CMakeLists.txt @@ -7,6 +7,7 @@ set(TEST_SRC hipStreamAddCallback.cc hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc + hipStreamDestroy.cc hipStreamGetCUMask.cc hipAPIStreamDisable.cc streamCommon.cc @@ -25,6 +26,7 @@ set(TEST_SRC hipStreamAddCallback.cc hipStreamCreateWithFlags.cc hipStreamCreateWithPriority.cc + hipStreamDestroy.cc hipAPIStreamDisable.cc # hipStreamAttachMemAsync.cc # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync # Fixing would break ABI, to be re-enabled when the fix is made. diff --git a/tests/catch/unit/stream/hipStreamDestroy.cc b/tests/catch/unit/stream/hipStreamDestroy.cc new file mode 100644 index 0000000000..54edddc7c6 --- /dev/null +++ b/tests/catch/unit/stream/hipStreamDestroy.cc @@ -0,0 +1,97 @@ +/* +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 + +namespace hipStreamDestroyTests { + +TEST_CASE("Unit_hipStreamDestroy_Default") { + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +TEST_CASE("Unit_hipStreamDestroy_Negative_DoubleDestroy") { + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipStreamDestroy(stream), hipErrorContextIsDestroyed); +} + +TEST_CASE("Unit_hipStreamDestroy_Negative_NullStream") { + HIP_CHECK_ERROR(hipStreamDestroy(nullptr), hipErrorInvalidResourceHandle); +} + +template void checkDataSet(int* deviceData) { + HIP_CHECK(hipStreamSynchronize(nullptr)); + std::array hostData{}; + HIP_CHECK( + hipMemcpy(hostData.data(), deviceData, sizeof(int) * numDataPoints, hipMemcpyDeviceToHost)); + REQUIRE(std::all_of(std::begin(hostData), std::end(hostData), [](int x) { return x == 1; })); +} + +__global__ void setToOne(int* x, size_t size) { + unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < size) { + x[idx] = 1; + } +} + +TEST_CASE("Unit_hipStreamDestroy_WithFinishedWork") { + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + + constexpr int numDataPoints = 10; + int* deviceData{}; + HIP_CHECK(hipMalloc(&deviceData, sizeof(int) * numDataPoints)); + HIP_CHECK(hipMemset(deviceData, 0, sizeof(int) * numDataPoints)); + + setToOne<<<1, numDataPoints, 0, stream>>>(deviceData, numDataPoints); + checkDataSet(deviceData); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(deviceData)); +} + +// hipStreamDestroy should return immediately then clean up the resources when the stream is empty +// of work +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ +TEST_CASE("Unit_hipStreamDestroy_WithPendingWork") { +#if HT_AMD + HipTest::HIP_SKIP_TEST( + "EXSWCPHIPT-44 - expected hipStreamDestroy to return immediately then release the resources " + "when the queued jobs are finished"); + return; +#endif + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + constexpr int numDataPoints = 10; + int* deviceData{}; + HIP_CHECK(hipMalloc(&deviceData, sizeof(int) * numDataPoints)); + HIP_CHECK(hipMemset(deviceData, 0, sizeof(int) * numDataPoints)); + + HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream); + setToOne<<<1, numDataPoints, 0, stream>>>(deviceData, numDataPoints); + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); + HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady); + HIP_CHECK(hipStreamDestroy(stream)); + checkDataSet(deviceData); + HIP_CHECK(hipFree(deviceData)); +} +#endif +} // namespace hipStreamDestroyTests From 46717c6ffe49f58c94d921cd3af693b199f6f5c1 Mon Sep 17 00:00:00 2001 From: Finlay Date: Fri, 5 Aug 2022 06:29:44 +0100 Subject: [PATCH 3/7] EXSWCPHIPT-140 - Added hipStreamCreateWithFlags tests (#2549) --- .../unit/stream/hipStreamCreateWithFlags.cc | 80 +++++++++++++++---- 1 file changed, 63 insertions(+), 17 deletions(-) diff --git a/tests/catch/unit/stream/hipStreamCreateWithFlags.cc b/tests/catch/unit/stream/hipStreamCreateWithFlags.cc index 95606f418a..ef38755028 100644 --- a/tests/catch/unit/stream/hipStreamCreateWithFlags.cc +++ b/tests/catch/unit/stream/hipStreamCreateWithFlags.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 @@ -17,23 +17,69 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -1) Validates functionality of hipStreamCreateWithFlags when stream = nullptr. -2) Validates functionality of hipStreamCreateWithFlags when flag = 0xffffffff. -*/ - +#include #include +namespace hipStreamCreateWithFlagsTests { -TEST_CASE("Unit_hipStreamCreateWithFlags_ArgValidation") { - // stream = nullptr test - SECTION("stream is nullptr") { - REQUIRE(hipStreamCreateWithFlags(nullptr, hipStreamDefault) != hipSuccess); - } - // flag value invalid test - SECTION("flag value invalid") { - hipStream_t stream; - REQUIRE(hipStreamCreateWithFlags(&stream, 0xffffffff) != hipSuccess); - } +TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_NullStream") { + HIP_CHECK_ERROR(hipStreamCreateWithFlags(nullptr, hipStreamDefault), hipErrorInvalidValue); } + +TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_InvalidFlag") { + hipStream_t stream{}; + unsigned int flag = 0xFF; + REQUIRE(flag != hipStreamDefault); + REQUIRE(flag != hipStreamNonBlocking); + HIP_CHECK_ERROR(hipStreamCreateWithFlags(&stream, flag), hipErrorInvalidValue); +} + +// create a stream and check the properties are correctly set +TEST_CASE("Unit_hipStreamCreateWithFlags_Default") { + const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking); + hipStream_t stream{}; + HIP_CHECK(hipStreamCreateWithFlags(&stream, flagUnderTest)); + + unsigned int flag{}; + HIP_CHECK(hipStreamGetFlags(stream, &flag)); + REQUIRE(flag == flagUnderTest); + + int priority{}; + HIP_CHECK(hipStreamGetPriority(stream, &priority)); + // zero is considered default priority + REQUIRE(priority == 0); + + HIP_CHECK(hipStreamDestroy(stream)); +} + +// a stream will default to blocking the null stream, but will not block the null stream when +// created with hipStreamNonBlocking +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ +TEST_CASE("Unit_hipStreamCreateWithFlags_DefaultStreamInteraction") { + const hipStream_t defaultStream = GENERATE(static_cast(nullptr), hipStreamPerThread); + const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking); + const hipError_t expectedError = (flagUnderTest == hipStreamDefault) && (defaultStream == nullptr) + ? hipErrorNotReady + : hipSuccess; + CAPTURE(defaultStream, flagUnderTest, expectedError, hipGetErrorString(expectedError)); + + hipStream_t stream{}; + HIP_CHECK(hipStreamCreateWithFlags(&stream, flagUnderTest)); + + constexpr auto delay = std::chrono::milliseconds(500); + + SECTION("default stream waiting for created stream") { + HipTest::runKernelForDuration(delay, stream); + REQUIRE(hipStreamQuery(defaultStream) == expectedError); + } + SECTION("created stream waiting for default stream") { + HipTest::runKernelForDuration(delay, defaultStream); + REQUIRE(hipStreamQuery(stream) == expectedError); + } + + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipStreamDestroy(stream)); +} +#endif + +} // namespace hipStreamCreateWithFlagsTests From 159d41531df4a9d70fde223af2302c4c609fbc15 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio?= Date: Fri, 5 Aug 2022 06:30:06 +0100 Subject: [PATCH 4/7] EXSWCPHIPT-105: Add testing for hipEventDestroy (#2727) --- tests/catch/unit/event/CMakeLists.txt | 1 + tests/catch/unit/event/hipEventDestroy.cc | 115 ++++++++++++++++++++++ 2 files changed, 116 insertions(+) create mode 100644 tests/catch/unit/event/hipEventDestroy.cc diff --git a/tests/catch/unit/event/CMakeLists.txt b/tests/catch/unit/event/CMakeLists.txt index bbcc2c9fa4..30d5a4155b 100644 --- a/tests/catch/unit/event/CMakeLists.txt +++ b/tests/catch/unit/event/CMakeLists.txt @@ -5,6 +5,7 @@ set(TEST_SRC Unit_hipEventElapsedTime.cc Unit_hipEventRecord.cc Unit_hipEventIpc.cc + hipEventDestroy.cc ) # The test used wait mechanism and doesnt play well with all arch of nvidia diff --git a/tests/catch/unit/event/hipEventDestroy.cc b/tests/catch/unit/event/hipEventDestroy.cc new file mode 100644 index 0000000000..23cb17bed5 --- /dev/null +++ b/tests/catch/unit/event/hipEventDestroy.cc @@ -0,0 +1,115 @@ +/* +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 "hip/hip_runtime_api.h" + +#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ + +static constexpr size_t vectorSize{1024}; + +/** + * @brief Launches vectorAdd kernel with a delay + */ +static inline void launchVectorAdd(float*& A_h, float*& B_h, float*& C_h, + std::chrono::milliseconds delay, hipStream_t stream = nullptr) { + float* A_d{nullptr}; + float* B_d{nullptr}; + float* C_d{nullptr}; + HipTest::initArraysForHost(&A_h, &B_h, &C_h, vectorSize, true); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&B_d), B_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&C_d), C_h, 0)); + HipTest::runKernelForDuration(delay, stream); + HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, vectorSize); +} + + +/** + * @brief Check that destroying an event before the kernel has finished running causes no errors. + * + */ +TEST_CASE("Unit_hipEventDestroy_Unfinished") { + hipEvent_t event; + + HIP_CHECK(hipEventCreate(&event)); + + float *A_h, *B_h, *C_h; + launchVectorAdd(A_h, B_h, C_h, std::chrono::milliseconds(1000)); + + HIP_CHECK(hipEventRecord(event)); + HIP_CHECK_ERROR(hipEventQuery(event), hipErrorNotReady); + HIP_CHECK(hipEventDestroy(event)); + + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, vectorSize); + HipTest::freeArraysForHost(A_h, B_h, C_h, true); +} + +/** + * @brief Check that destroying an event enqueued to a stream causes no errors. + * + */ +TEST_CASE("Unit_hipEventDestroy_WithWaitingStream") { + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + float *A_h, *B_h, *C_h; + launchVectorAdd(A_h, B_h, C_h, std::chrono::milliseconds(1000), stream); + + HIP_CHECK(hipEventRecord(event, stream)); + HIP_CHECK_ERROR(hipEventQuery(event), hipErrorNotReady); + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + + HipTest::checkVectorADD(A_h, B_h, C_h, vectorSize); + HipTest::freeArraysForHost(A_h, B_h, C_h, true); +} + +TEST_CASE("Unit_hipEventDestroy_Negative") { +#if HT_AMD + HipTest::HIP_SKIP_TEST("EXSWCPHIPT-103"); + return; +#endif + + SECTION("Invalid Event") { + hipEvent_t event{nullptr}; + HIP_CHECK_ERROR(hipEventDestroy(event), hipErrorInvalidResourceHandle); + } + + SECTION("Destroy twice") { + hipEvent_t event; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK_ERROR(hipEventDestroy(event), hipErrorContextIsDestroyed); + } +} +#endif From 8bfdde5ecb7e15d5be1d6844c3bff8fa8a44a18d Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 12:19:12 +0530 Subject: [PATCH 5/7] SWDEV-348718 - Use of GENERATE in thead function causes random failures with multithread condition. (#2842) Change-Id: Ie6722abbf5aa9936b5e06002064f7f1109066af9 --- tests/catch/unit/memory/hipArrayCreate.cc | 38 +++++++++++---------- tests/catch/unit/memory/hipMalloc3DArray.cc | 31 +++++++++-------- tests/catch/unit/memory/hipMallocArray.cc | 32 +++++++++-------- 3 files changed, 54 insertions(+), 47 deletions(-) diff --git a/tests/catch/unit/memory/hipArrayCreate.cc b/tests/catch/unit/memory/hipArrayCreate.cc index 836c0d1c45..efffc344e5 100644 --- a/tests/catch/unit/memory/hipArrayCreate.cc +++ b/tests/catch/unit/memory/hipArrayCreate.cc @@ -52,26 +52,28 @@ static constexpr auto ARRAY_LOOP{100}; static void ArrayCreate_DiffSizes(int gpu) { HIP_CHECK_THREAD(hipSetDevice(gpu)); - std::pair size = - GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); - std::array array; - size_t pavail, avail; - HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); - HIP_ARRAY_DESCRIPTOR desc; - desc.NumChannels = 1; - desc.Width = std::get<0>(size); - desc.Height = std::get<1>(size); - desc.Format = HIP_AD_FORMAT_FLOAT; + //Use of GENERATE in thead function causes random failures with multithread condition. + std::vector> runs {std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)}; + for (const auto& size : runs) { + std::array array; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + HIP_ARRAY_DESCRIPTOR desc; + desc.NumChannels = 1; + desc.Width = std::get<0>(size); + desc.Height = std::get<1>(size); + desc.Format = HIP_AD_FORMAT_FLOAT; - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK_THREAD(hipArrayDestroy(array[i])); - } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipArrayDestroy(array[i])); + } - HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); - REQUIRE_THREAD(pavail == avail); + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); + } } /* This testcase verifies hipArrayCreate API for small and big chunks data*/ diff --git a/tests/catch/unit/memory/hipMalloc3DArray.cc b/tests/catch/unit/memory/hipMalloc3DArray.cc index 1b7ee27171..7982f85048 100644 --- a/tests/catch/unit/memory/hipMalloc3DArray.cc +++ b/tests/catch/unit/memory/hipMalloc3DArray.cc @@ -49,23 +49,26 @@ static constexpr auto ARRAY_LOOP{100}; */ static void Malloc3DArray_DiffSizes(int gpu) { HIP_CHECK_THREAD(hipSetDevice(gpu)); - const int size = GENERATE(ARRAY_SIZE, BIG_ARRAY_SIZE); - int width{size}, height{size}, depth{size}; - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); - std::array arr; - size_t pavail, avail; - HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + //Use of GENERATE in thead function causes random failures with multithread condition. + std::vector runs {ARRAY_SIZE, BIG_ARRAY_SIZE}; + for (const auto& size : runs) { + size_t width{size}, height{size}, depth{size}; + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); + std::array arr; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth), + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth), hipArrayDefault)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK_THREAD(hipFreeArray(arr[i])); - } + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(arr[i])); + } - HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); - REQUIRE_THREAD(pavail == avail); + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); + } } TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") { diff --git a/tests/catch/unit/memory/hipMallocArray.cc b/tests/catch/unit/memory/hipMallocArray.cc index 24519f4b1a..8b2099a4ec 100644 --- a/tests/catch/unit/memory/hipMallocArray.cc +++ b/tests/catch/unit/memory/hipMallocArray.cc @@ -51,23 +51,25 @@ static constexpr int ARRAY_LOOP{100}; */ static void MallocArray_DiffSizes(int gpu) { HIP_CHECK_THREAD(hipSetDevice(gpu)); - std::pair size = - GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)); - hipChannelFormatDesc desc = hipCreateChannelDesc(); - std::array A_d; - size_t pavail, avail; - HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); + //Use of GENERATE in thead function causes random failures with multithread condition. + std::vector> runs {std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)}; + for (const auto& size : runs) { + hipChannelFormatDesc desc = hipCreateChannelDesc(); + std::array A_d; + size_t pavail, avail; + HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK_THREAD( - hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); - } - for (int i = 0; i < ARRAY_LOOP; i++) { - HIP_CHECK_THREAD(hipFreeArray(A_d[i])); - } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD( + hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); + } + for (int i = 0; i < ARRAY_LOOP; i++) { + HIP_CHECK_THREAD(hipFreeArray(A_d[i])); + } - HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); - REQUIRE_THREAD(pavail == avail); + HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); + REQUIRE_THREAD(pavail == avail); + } } TEST_CASE("Unit_hipMallocArray_DiffSizes") { From 768f186d69a7a0c3fef8a2aff0244d3ab16b70f6 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 12:19:24 +0530 Subject: [PATCH 6/7] SWDEV-348686 - Add extension on windows platform for executable path. (#2845) Change-Id: Id7a9b6ec49e8f210b0a74b5f5fd68d584f005b4e --- tests/catch/include/hip_test_process.hh | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/tests/catch/include/hip_test_process.hh b/tests/catch/include/hip_test_process.hh index 335e6c7bba..2113dc85e1 100644 --- a/tests/catch/include/hip_test_process.hh +++ b/tests/catch/include/hip_test_process.hh @@ -66,7 +66,12 @@ class SpawnProc { auto dir = fs::path(TestContext::get().currentPath()); dir /= exeName; exeName = dir.string(); - + // On Windows, fs::exists returns false without extension. + if (TestContext::get().isWindows()) { + if(fs::path(exeName).extension().empty()) { + exeName += ".exe"; + } + } INFO("Testing that exe exists: " << exeName); REQUIRE(fs::exists(exeName)); @@ -92,13 +97,13 @@ class SpawnProc { execCmd += " > "; execCmd += tmpFileName; } - auto res = std::system(execCmd.c_str()); if (captureOutput) { std::ifstream t(tmpFileName.c_str()); resultStr = std::string((std::istreambuf_iterator(t)), std::istreambuf_iterator()); + t.close(); } #if HT_LINUX return WEXITSTATUS(res); From 90f42bdf82b2969d189559ae2fd60d89860334c4 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 5 Aug 2022 12:19:41 +0530 Subject: [PATCH 7/7] SWDEV-339150 - [catch2][dtest] Updating test Unit_hipGetDeviceAttribute_CheckFineGrainSupport: (#2846) This PR fixes the failure of test Unit_hipGetDeviceAttribute_CheckFineGrainSupport on MI100 --- .../unit/device/hipGetDeviceAttribute.cc | 95 +++++++++++++++---- 1 file changed, 77 insertions(+), 18 deletions(-) diff --git a/tests/catch/unit/device/hipGetDeviceAttribute.cc b/tests/catch/unit/device/hipGetDeviceAttribute.cc index e21c4eac90..fb11f64d54 100644 --- a/tests/catch/unit/device/hipGetDeviceAttribute.cc +++ b/tests/catch/unit/device/hipGetDeviceAttribute.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 @@ -22,10 +22,12 @@ THE SOFTWARE. // Test the device info API extensions for HIP #include +#include +#ifdef __linux__ +#include +#endif #include - - static hipError_t test_hipDeviceGetAttribute(int deviceId, hipDeviceAttribute_t attr, int expectedValue = -1) { @@ -234,15 +236,78 @@ TEST_CASE("Unit_hipGetDeviceAttribute_CheckAttrValues") { /** * Validate the hipDeviceAttributeFineGrainSupport property in AMD. */ +#ifdef __linux__ #if HT_AMD +#define COMMAND_LEN 256 +#define BUFFER_LEN 512 + +static bool isRocmPathSet() { + FILE *fpipe; + char const *command = "echo $ROCM_PATH"; + fpipe = popen(command, "r"); + + if (fpipe == nullptr) { + printf("Unable to create command\n"); + return false; + } + char command_op[BUFFER_LEN]; + if (fgets(command_op, BUFFER_LEN, fpipe)) { + size_t len = strlen(command_op); + if (len > 1) { // This is because fgets always adds newline character + pclose(fpipe); + return true; + } + } + pclose(fpipe); + return false; +} + // This is AMD specific property test TEST_CASE("Unit_hipGetDeviceAttribute_CheckFineGrainSupport") { int deviceId; int deviceCount = 0; + FILE *fpipe; + char command[COMMAND_LEN] = ""; + const char *rocmpath = nullptr; + if (isRocmPathSet()) { + // For STG2 testing where /opt/rocm path is not present + rocmpath = "$ROCM_PATH/bin/rocminfo"; + } else { + // Check if the rocminfo tool exists + rocmpath = "/opt/rocm/bin/rocminfo"; + } + snprintf(command, COMMAND_LEN, "%s", rocmpath); + strncat(command, " | grep -i \"Segment:\\|Uuid:\"", COMMAND_LEN); + // Execute the rocminfo command and extract the segment info + fpipe = popen(command, "r"); + if (fpipe == nullptr) { + printf("Unable to create command file\n"); + return; + } HIP_CHECK(hipGetDeviceCount(&deviceCount)); - REQUIRE(deviceCount > 0); - // Check hipDeviceAttributeFineGrainSupport for all available device - // in system. + assert(deviceCount > 0); + int *fine_grained_val = new int[deviceCount]; + assert(fine_grained_val != nullptr); + bool *gpuFound = new bool[deviceCount]; + assert(gpuFound != nullptr); + for (int i = 0; i < deviceCount; i++) { + gpuFound[i] = false; + fine_grained_val[i] = 0; // Initialize to 0 + } + char command_op[BUFFER_LEN]; + int count = -1; + // Extract each segment flags + while (fgets(command_op, BUFFER_LEN, fpipe)) { + std::string rocminfo_line(command_op); + if ((std::string::npos != rocminfo_line.find("GPU-")) && + (std::string::npos != rocminfo_line.find("Uuid:"))) { + count++; + gpuFound[count] = true; + } else if (gpuFound[count] && + (std::string::npos != rocminfo_line.find("FLAGS: FINE GRAINED"))) { + fine_grained_val[count] = 1; + } + } for (int dev = 0; dev < deviceCount; dev++) { HIP_CHECK(hipSetDevice(dev)); HIP_CHECK(hipGetDevice(&deviceId)); @@ -250,21 +315,15 @@ TEST_CASE("Unit_hipGetDeviceAttribute_CheckFineGrainSupport") { HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); int value = 0; HIP_CHECK(hipDeviceGetAttribute(&value, - hipDeviceAttributeFineGrainSupport, deviceId)); - std::string gpu_arch_name(props.gcnArchName); - if (std::string::npos != gpu_arch_name.find("gfx90a")) { - // Current GPU is gfx90a architecture - REQUIRE(value == 1); - } else if (std::string::npos != gpu_arch_name.find("gfx906")) { - // Current GPU is gfx906 architecture - REQUIRE(value == 0); - } else if (std::string::npos != gpu_arch_name.find("gfx908")) { - // Current GPU is gfx908 architecture - REQUIRE(value == 0); - } + hipDeviceAttributeFineGrainSupport, deviceId)); + REQUIRE(value == fine_grained_val[dev]); } + // Validate hipDeviceAttributeFineGrainSupport + delete[] fine_grained_val; + delete[] gpuFound; } #endif +#endif /** * Validates negative scenarios for hipDeviceGetAttribute * scenario1: pi = nullptr