From 9f6072679eff44c1d586300a1250fd866fac9955 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:26:52 +0530 Subject: [PATCH 01/28] SWDEV-1 - Add missing checks in hipMemcpy2D.cc [ROCm/hip commit: 0492446e87bd403818776221336a2c78b556a217] --- projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc index 7d9a8e1ccc..27e011089c 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2D.cc @@ -135,7 +135,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2D_multiDevice-D2D", "" size_t width{NUM_W * sizeof(TestType)}; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); From b1911b121d1635f234b7b7461e02f22281d0733c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:28:19 +0530 Subject: [PATCH 02/28] SWDEV-1 - Add missing checks in hipMemcpy2DAsync.cc [ROCm/hip commit: f44eab96039b5c70d7f435545d641aef0fe18595] --- projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc index 10924687df..76b20ac574 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DAsync.cc @@ -160,7 +160,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-Host&PinnedMem", "" hipStream_t stream; if (numDevices > 1) { - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); HIP_CHECK(hipStreamCreate(&stream)); @@ -250,7 +250,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy2DAsync_multiDevice-StreamOnDiffDevice", "" hipStream_t stream; if (numDevices > 1) { - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); From 34f099a2e5d497b60908d112556e7e645bc13134 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:29:39 +0530 Subject: [PATCH 03/28] SWDEV-1 - Add missing checks in hipMemcpy2DFromArray.cc [ROCm/hip commit: e99cd4b50a80b6602b30d8c274cf1024bbb95204] --- projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc index 0e6a7aacde..1972a5f10a 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArray.cc @@ -190,7 +190,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_multiDevicePinnedMemPeerGpu") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); hipArray *A_d{nullptr}; @@ -246,7 +246,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_multiDeviceContextChange") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); hipArray *A_d{nullptr}; From 3f2a773a4a4b46fab4eb2aebfe7718d7d3c636ff Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:30:38 +0530 Subject: [PATCH 04/28] SWDEV-1 - Add missing checks in hipMemcpy2DFromArrayAsync.cc [ROCm/hip commit: fba4cedc330192d3a3c116d2dd4945a74c39534c] --- .../hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc index 6603243585..f7c8e75b71 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy2DFromArrayAsync.cc @@ -217,7 +217,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDevicePinnedHostMem") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); hipArray *A_d{nullptr}; @@ -277,7 +277,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_multiDeviceContextChange") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(0)); hipArray *A_d{nullptr}; From e83a40029010741da17615761da0ebfe9f0cf4e7 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:50:32 +0530 Subject: [PATCH 05/28] SWDEV-1 - Add missing checks in hipMemcpyAllApiNegative.cc [ROCm/hip commit: 3d349720b763afd5ea9163cb4ef89be3713099b3] --- .../hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc index 7ea39d30be..bb54701af8 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyAllApiNegative.cc @@ -150,7 +150,7 @@ TEST_CASE("Unit_hipMemcpy_NullCheck") { &A_h, &B_h, &C_h, NUM_ELM*sizeof(float)); hipStream_t stream; - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipMemcpy(A_d, C_h, NUM_ELM*sizeof(float), hipMemcpyHostToDevice)); @@ -240,7 +240,7 @@ TEST_CASE("Unit_hipMemcpy_HalfMemCopy") { &A_h, &B_h, &C_h, NUM_ELM*sizeof(float)); hipStream_t stream; - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); SECTION("hipMemcpyHtoD half memory copy") { HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h, From 6c4dbfad4f1a129a78e6e08610c2070e5c7d9e31 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:53:04 +0530 Subject: [PATCH 06/28] SWDEV-1 - Add missing checks in hipMemcpy_MultiThread.cc [ROCm/hip commit: 27193f755b41371c9b4beaa846f0f329c586f816] --- .../hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc index 1dad378a6c..d9c08f3181 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy_MultiThread.cc @@ -105,7 +105,7 @@ void memcpyTests::Memcpy_And_verify(bool *ret_val) { for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIPCHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(T), hipMemcpyDefault)); @@ -155,7 +155,7 @@ void memcpyTests::Memcpy_And_verify(bool *ret_val) { int canAccessPeer = 0; for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]), A_h, NUM_ELM * sizeof(T))); @@ -198,7 +198,7 @@ void memcpyTests::Memcpy_And_verify(bool *ret_val) { for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIPCHECK(hipMemcpyAsync(A_d[j], A_d[i], NUM_ELM * sizeof(T), @@ -252,7 +252,7 @@ void memcpyTests::Memcpy_And_verify(bool *ret_val) { for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIPCHECK(hipSetDevice(j)); HIPCHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]), From 0f62c5bbb7eb26590975deb8e7fb470e6263f1a4 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 08:54:25 +0530 Subject: [PATCH 07/28] SWDEV-1 - Add missing checks in hipMallocManaged_MultiScenario.cc [ROCm/hip commit: 3940d3cd9c4598e0972314bd686052a8a54f9738] --- .../tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc b/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc index 971d6cf8c6..774ec6a6fd 100644 --- a/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc +++ b/projects/hip/tests/catch/unit/memory/hipMallocManaged_MultiScenario.cc @@ -93,7 +93,7 @@ TEST_CASE("Unit_hipMallocManaged_HostDeviceConcurrent") { std::thread host_thread(HostKernelDouble, Hmm, hPtr, N); KernelDouble<<>>(Hmm, dPtr, N); host_thread.join(); - hipMemcpy(resPtr, dPtr, N * sizeof(float), hipMemcpyDeviceToHost); + HIP_CHECK(hipMemcpy(resPtr, dPtr, N * sizeof(float), hipMemcpyDeviceToHost)); for (size_t i = 0; i < N; i++) { REQUIRE(hPtr[i] == resPtr[i]); From 9d302fa7bbce88e710b01d787b91f8ed9121246a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 09:14:00 +0530 Subject: [PATCH 08/28] SWDEV-1 - Add missing checks in hipMemsetNegative.cc [ROCm/hip commit: eda2e499b31612054d36957c9dfa8d0461a7a60f] --- projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc b/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc index 82a0c2879d..929ff2e3c6 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemsetNegative.cc @@ -105,7 +105,7 @@ TEST_CASE("Unit_hipMemset2D_Negative_InvalidPtr") { size_t pitch_A; HIP_CHECK(hipMallocPitch(&A_d, &pitch_A, width, height)); testHipMemset2DApis(dst, pitch_A, memsetVal, width, height); - hipFree(A_d); + HIP_CHECK(hipFree(A_d)); } TEST_CASE("Unit_hipMemset2D_Negative_InvalidSizes") { From be62a298346f19b402f1e0e7914bb4ea126949fe Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 09:16:13 +0530 Subject: [PATCH 09/28] SWDEV-1 - Add missing checks in hipMemset.cc [ROCm/hip commit: 1dc36ddc7c632fa9f263d689233648bd68a1d84c] --- .../hip/tests/catch/unit/memory/hipMemset.cc | 22 +++++++++---------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemset.cc b/projects/hip/tests/catch/unit/memory/hipMemset.cc index 600a008436..c4adb15e53 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemset.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemset.cc @@ -260,21 +260,21 @@ TEST_CASE("Unit_hipMemset_2AsyncOperations") { std::vector v; v.resize(2048); float* p2, *p3; - hipMalloc(reinterpret_cast(&p2), 4096 + 4096*2); + HIP_CHECK(hipMalloc(reinterpret_cast(&p2), 4096 + 4096*2)); p3 = p2+2048; hipStream_t s; - hipStreamCreate(&s); - hipMemsetAsync(p2, 0, 32*32*4, s); - hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s); - hipStreamSynchronize(s); + HIP_CHECK(hipStreamCreate(&s)); + HIP_CHECK(hipMemsetAsync(p2, 0, 32*32*4, s)); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s)); + HIP_CHECK(hipStreamSynchronize(s)); for (int i = 0; i < 256; ++i) { - hipMemsetAsync(p2, 0, 32*32*4, s); - hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s); + HIP_CHECK(hipMemsetAsync(p2, 0, 32*32*4, s)); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)p3, 0x3fe00000, 32*32, s)); } - hipStreamSynchronize(s); - hipDeviceSynchronize(); - hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost); - hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost); + HIP_CHECK(hipStreamSynchronize(s)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(&v[0], p2, 1024, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&v[1024], p3, 1024, hipMemcpyDeviceToHost)); REQUIRE(v[0] == 0); REQUIRE(v[1024] == 1.75f); From a71e9329cc5262646e529dce1b2d06bd87fa3784 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 09:56:00 +0530 Subject: [PATCH 10/28] SWDEV-1 - Add missing checks in hipMemset2D.cc [ROCm/hip commit: d632c8272e32cbe702d8a379e92268d94b414bcb] --- projects/hip/tests/catch/unit/memory/hipMemset2D.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemset2D.cc b/projects/hip/tests/catch/unit/memory/hipMemset2D.cc index e48b8931f0..2b2cb0463d 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemset2D.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemset2D.cc @@ -78,7 +78,7 @@ TEST_CASE("Unit_hipMemset2D_BasicFunctional") { } } - hipFree(A_d); + HIP_CHECK(hipFree(A_d)); free(A_h); } @@ -120,7 +120,7 @@ TEST_CASE("Unit_hipMemset2DAsync_BasicFunctional") { } } - hipFree(A_d); + HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipStreamDestroy(stream)); free(A_h); } @@ -169,7 +169,7 @@ TEST_CASE("Unit_hipMemset2D_UniqueWidthHeight") { } } - hipFree(A_d); + HIP_CHECK(hipFree(A_d)); free(A_h); } From dbd6338527748988b09480e208f71d704c2fda91 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 10:27:42 +0530 Subject: [PATCH 11/28] SWDEV-1 - Add missing checks in hipMemcpyPeer.cc [ROCm/hip commit: d896d6b2ee894ade781433933302cf408dd5050b] --- projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc index a8ebe9534f..48e881895d 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyPeer.cc @@ -35,7 +35,7 @@ TEST_CASE("Unit_hipMemcpyPeer_Negative") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { // Initialization of variables int *A_d{nullptr}, *B_d{nullptr}; @@ -102,7 +102,7 @@ TEST_CASE("Unit_hipMemcpyPeer_Basic") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; int *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; From bb7960222400567bb85c238e8ea05d046bb79a7c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 10:31:40 +0530 Subject: [PATCH 12/28] SWDEV-1 - Add missing checks in hipMemcpyPeerAsync.cc [ROCm/hip commit: a371c151207418f3478f55e73f4a2e951dc31812] --- projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc index 1c738f1c76..56a76cbad2 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyPeerAsync.cc @@ -39,7 +39,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Negative") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { // Initialization of variables int *A_d{nullptr}, *B_d{nullptr}; @@ -116,7 +116,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { // Initialization of Variables on GPU-0 int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; @@ -202,7 +202,7 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_StreamOnDiffDevice") { HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices > 1) { int canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, 0, 1); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); if (canAccessPeer) { int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr}; int *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr}; From 737f6cb5c9f62c84666e9802ff88adbd528fc8a4 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 10:42:36 +0530 Subject: [PATCH 13/28] SWDEV-1 - Add missing checks in hipMemsetFunctional.cc [ROCm/hip commit: 5afd5f660037d31d1fa894c48271a64cbc05992b] --- .../hip/tests/catch/unit/memory/hipMemsetFunctional.cc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc b/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc index d6f152cec9..2f2b478ae3 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemsetFunctional.cc @@ -83,7 +83,7 @@ void checkMemset(T value, size_t count, MemsetType memsetType, bool async = fals MemsetMallocType mallocType = hipDeviceMalloc_t, T* devPtr = nullptr) { hipStream_t stream{nullptr}; if (async) { - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); } // Allocate Memory @@ -254,7 +254,7 @@ template void checkMemset2D(T value, size_t width, size_t height, bool async = false, size_t pitch = 0, T* devPtr = nullptr) { hipStream_t stream{nullptr}; - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); constexpr size_t elementSize = sizeof(T); bool freeDevPtr = false; if (devPtr == nullptr) { @@ -277,7 +277,7 @@ void checkMemset2D(T value, size_t width, size_t height, bool async = false, siz if (freeDevPtr) { HIP_CHECK(hipFree(devPtr)); } - hipStreamDestroy(stream); + HIP_CHECK(hipStreamDestroy(stream)); } TEST_CASE("Unit_hipMemsetFunctional_ZeroValue_2D") { @@ -430,7 +430,7 @@ void check_device_data_3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent exten template void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool async = false) { hipStream_t stream{nullptr}; - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); if (devPitchedPtr.ptr == nullptr) { HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent)); } @@ -445,7 +445,7 @@ void checkMemset3D(hipPitchedPtr& devPitchedPtr, T value, hipExtent extent, bool if (extent.width * extent.height * extent.depth > 0) { check_device_data_3D(devPitchedPtr, value, extent); } - hipStreamDestroy(stream); + HIP_CHECK(hipStreamDestroy(stream)); } void check_memset_3D(std::string sectionStr, size_t width, size_t height, size_t depth, From e6b27bb9521bf14551a6006b4d1934f0fc5da498 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 10:44:04 +0530 Subject: [PATCH 14/28] SWDEV-1 - Add missing checks in hipMemcpyAsync.cc [ROCm/hip commit: e16ec66c0fc0baece749f0f79931f8bb0eedbf54] --- projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc index 503763aed8..52964c87bd 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc @@ -112,7 +112,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float, TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}; HIP_CHECK(hipSetDevice(0)); hipStream_t stream; - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false); @@ -147,7 +147,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int, TestType *A_Ph{nullptr}, *B_Ph{nullptr}; HIP_CHECK(hipSetDevice(0)); hipStream_t stream; - hipStreamCreate(&stream); + HIP_CHECK(hipStreamCreate(&stream)); HipTest::initArrays(&A_d, &B_d, nullptr, &A_h, &B_h, nullptr, NUM_ELM*sizeof(TestType)); From 1a250f3fe8ded269420bad07bf553c5cb1bf8a31 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 10:50:43 +0530 Subject: [PATCH 15/28] SWDEV-1 - Add missing checks in hipMemcpy.cc [ROCm/hip commit: 4a745feb5f295e3afc48143dc3f722e2860669b6] --- projects/hip/tests/catch/unit/memory/hipMemcpy.cc | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpy.cc b/projects/hip/tests/catch/unit/memory/hipMemcpy.cc index 20860d2843..84af63bea0 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpy.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpy.cc @@ -576,7 +576,7 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch", HIP_CHECK(hipSetDevice(HipTest::RAND_R(&seed) % (numDevices-1)+1)); int device; - hipGetDevice(&device); + HIP_CHECK(hipGetDevice(&device)); std::cout <<"hipMemcpy is set to happen between device 0 and device " <(&X_d, &Y_d, &Z_d, nullptr, @@ -588,10 +588,10 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch", C_h[j] = 0; } - hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost); - hipMemcpy(X_d, A_h, Nbytes, hipMemcpyHostToDevice); - hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost); - hipMemcpy(Y_d, B_h, Nbytes, hipMemcpyHostToDevice); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(X_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(Y_d, B_h, Nbytes, hipMemcpyHostToDevice)); hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, static_cast(X_d), From 57d6a09d9d2a4cbe96dd82641ca1b10def2a1e1d Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 11:13:42 +0530 Subject: [PATCH 16/28] SWDEV-1 - Add missing checks in hipMemcpyAsync.cc [ROCm/hip commit: 172c6d78a9edca384b61b664ed32f78db20735cc] --- projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc b/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc index 52964c87bd..4c4e08ec2d 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemcpyAsync.cc @@ -373,10 +373,10 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch", C_h[j] = 0; } - hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost); - hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); - hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost); - hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream)); + HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream)); hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, static_cast(X_d), From 9083aaa30add22adde6665f6e6e62139da5f835e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 11:15:26 +0530 Subject: [PATCH 17/28] SWDEV-1 - Add missing checks in hipDrvMemcpy3D.cc [ROCm/hip commit: 7fc4b7d6f45efc9df7244d89d010e5baad188ab8] --- projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc b/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc index b70c8f74b1..0b524655a9 100644 --- a/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc +++ b/projects/hip/tests/catch/unit/memory/hipDrvMemcpy3D.cc @@ -501,8 +501,8 @@ void DrvMemcpy3D::HostArray_DrvMemcpy3D(bool device_context_change) { /* DeAllocating the memory */ template void DrvMemcpy3D::DeAllocateMemory() { - hipArrayDestroy(arr); - hipArrayDestroy(arr1); + HIP_CHECK(hipArrayDestroy(arr)); + HIP_CHECK(hipArrayDestroy(arr1)); free(hData); } From 16b5c5d7a653cde6d4c6560e7e473d131cea16d4 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 11:58:41 +0530 Subject: [PATCH 18/28] SWDEV-1 - Add missing checks in MemUtils.hh [ROCm/hip commit: 47df29a9d09733bf7cd57f1024a1db63b2d28e22] --- projects/hip/tests/catch/unit/memory/MemUtils.hh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/MemUtils.hh b/projects/hip/tests/catch/unit/memory/MemUtils.hh index a05588c3b7..b80343d004 100644 --- a/projects/hip/tests/catch/unit/memory/MemUtils.hh +++ b/projects/hip/tests/catch/unit/memory/MemUtils.hh @@ -344,10 +344,10 @@ static inline void memsetCheck(T* aPtr, size_t value, memType memType, MultiDDat template static inline void freeStuff(T* aPtr, allocType type) { switch (type) { case allocType::deviceMalloc: - hipFree(aPtr); + HIP_CHECK(hipFree(aPtr)); break; case allocType::hostMalloc: - hipHostFree(aPtr); + HIP_CHECK(hipHostFree(aPtr)); break; default: // for host and device registered HIP_CHECK(hipHostUnregister(aPtr)); From 3db4c84176f62b45b6f4ca3ae05ea3af793ee0fe Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 12:00:23 +0530 Subject: [PATCH 19/28] SWDEV-1 - Add missing checks in hipMemsetSync.cc [ROCm/hip commit: 1606235b12f5dec02ea259987440c39e6e5e931d] --- projects/hip/tests/catch/unit/memory/hipMemsetSync.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc b/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc index 6770a6e16c..2a55a2a0a3 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemsetSync.cc @@ -379,10 +379,10 @@ void memsetCheck(T* aPtr, size_t value, memSetType memsetType, MultiDData& data, template void freeStuff(T* aPtr, allocType type) { switch (type) { case allocType::deviceMalloc: - hipFree(aPtr); + HIP_CHECK(hipFree(aPtr)); break; case allocType::hostMalloc: - hipHostFree(aPtr); + HIP_CHECK(hipHostFree(aPtr)); break; case allocType::hostRegisted: HIP_CHECK(hipHostUnregister(aPtr)); From a04eeb66ff9e365b362e0458680c79a37faaf3d9 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 11 Oct 2022 12:01:36 +0530 Subject: [PATCH 20/28] SWDEV-1 - Add missing checks in hipMemoryAllocateCoherent.cc [ROCm/hip commit: a0dfc6a2c80fb6bac20bce797db66bca210d909d] --- .../tests/catch/unit/memory/hipMemoryAllocateCoherent.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc b/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc index 50747e275d..75c6c87b4e 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc @@ -44,8 +44,8 @@ __global__ void Kernel(float* hostRes, int clkRate) { TEST_CASE("Unit_hipHostMalloc_CoherentAccess") { int blocks = 2; float* hostRes; - hipHostMalloc(&hostRes, blocks * sizeof(float), - hipHostMallocMapped); + HIP_CHECK(hipHostMalloc(&hostRes, blocks * sizeof(float), + hipHostMallocMapped)); hostRes[0] = 0; hostRes[1] = 0; int clkRate; @@ -60,6 +60,6 @@ TEST_CASE("Unit_hipHostMalloc_CoherentAccess") { while (hostRes[eleCounter] == 0) {printf("waiting for counter inc\n");} eleCounter++; } - hipHostFree(reinterpret_cast(hostRes)); + HIP_CHECK(hipHostFree(reinterpret_cast(hostRes))); } From e5de7bbf55eadcec22235d898027efa3814f3d87 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 12 Oct 2022 05:00:51 +0200 Subject: [PATCH 21/28] EXSWHTEC-94 - Implement helper classes and functions for memory tests (#2978) - Implement helper classes and functions for memory tests - Remove c++14 standard constraint on memory tests - Remove GenerateLinearAllocationFlagCombinations until finished [ROCm/hip commit: 5219f5707facefb8443daefd7820ff5425d3adb4] --- .../tests/catch/include/resource_guards.hh | 125 ++++++++++++++++++ projects/hip/tests/catch/include/utils.hh | 102 ++++++++++++++ .../tests/catch/unit/memory/CMakeLists.txt | 3 +- 3 files changed, 228 insertions(+), 2 deletions(-) create mode 100644 projects/hip/tests/catch/include/resource_guards.hh create mode 100644 projects/hip/tests/catch/include/utils.hh diff --git a/projects/hip/tests/catch/include/resource_guards.hh b/projects/hip/tests/catch/include/resource_guards.hh new file mode 100644 index 0000000000..7e6179c81a --- /dev/null +++ b/projects/hip/tests/catch/include/resource_guards.hh @@ -0,0 +1,125 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 + +enum class LinearAllocs { + malloc, + mallocAndRegister, + hipHostMalloc, + hipMalloc, + hipMallocManaged, +}; + +template class LinearAllocGuard { + public: + LinearAllocGuard(const LinearAllocs allocation_type, const size_t size, + const unsigned int flags = 0u) + : allocation_type_{allocation_type} { + switch (allocation_type_) { + case LinearAllocs::malloc: + ptr_ = host_ptr_ = reinterpret_cast(malloc(size)); + break; + case LinearAllocs::mallocAndRegister: + host_ptr_ = reinterpret_cast(malloc(size)); + HIP_CHECK(hipHostRegister(host_ptr_, size, flags)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&ptr_), host_ptr_, 0u)); + break; + case LinearAllocs::hipHostMalloc: + HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr_), size, flags)); + host_ptr_ = ptr_; + break; + case LinearAllocs::hipMalloc: + HIP_CHECK(hipMalloc(reinterpret_cast(&ptr_), size)); + break; + case LinearAllocs::hipMallocManaged: + HIP_CHECK(hipMallocManaged(reinterpret_cast(&ptr_), size, flags ? flags : 1u)); + host_ptr_ = ptr_; + } + } + + LinearAllocGuard(const LinearAllocGuard&) = delete; + LinearAllocGuard(LinearAllocGuard&&) = delete; + + ~LinearAllocGuard() { + // No Catch macros, don't want to possibly throw in the destructor + switch (allocation_type_) { + case LinearAllocs::malloc: + free(ptr_); + break; + case LinearAllocs::mallocAndRegister: + // Cast to void to suppress nodiscard warnings + static_cast(hipHostUnregister(host_ptr_)); + free(host_ptr_); + break; + case LinearAllocs::hipHostMalloc: + static_cast(hipHostFree(ptr_)); + break; + case LinearAllocs::hipMalloc: + case LinearAllocs::hipMallocManaged: + static_cast(hipFree(ptr_)); + } + } + + T* ptr() { return ptr_; }; + T* const ptr() const { return ptr_; }; + T* host_ptr() { return host_ptr_; } + T* const host_ptr() const { return host_ptr(); } + + private: + const LinearAllocs allocation_type_; + T* ptr_ = nullptr; + T* host_ptr_ = nullptr; +}; + +enum class Streams { nullstream, perThread, created }; + +class StreamGuard { + public: + StreamGuard(const Streams stream_type) : stream_type_{stream_type} { + switch (stream_type_) { + case Streams::nullstream: + stream_ = nullptr; + break; + case Streams::perThread: + stream_ = hipStreamPerThread; + break; + case Streams::created: + HIP_CHECK(hipStreamCreate(&stream_)); + } + } + + StreamGuard(const StreamGuard&) = delete; + StreamGuard(StreamGuard&&) = delete; + + ~StreamGuard() { + if (stream_type_ == Streams::created) { + static_cast(hipStreamDestroy(stream_)); + } + } + + hipStream_t stream() const { return stream_; } + + private: + const Streams stream_type_; + hipStream_t stream_; +}; \ No newline at end of file diff --git a/projects/hip/tests/catch/include/utils.hh b/projects/hip/tests/catch/include/utils.hh new file mode 100644 index 0000000000..9edffc6f7c --- /dev/null +++ b/projects/hip/tests/catch/include/utils.hh @@ -0,0 +1,102 @@ +/* +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 + +namespace { +inline constexpr size_t kPageSize = 4096; +} // anonymous namespace + +template +void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements) { + const auto ret = std::mismatch(expected, expected + num_elements, actual); + if (ret.first != expected + num_elements) { + const auto idx = std::distance(expected, ret.first); + INFO("Value mismatch at index: " << idx); + REQUIRE(expected[idx] == actual[idx]); + } +} + +template void ArrayFindIfNot(It begin, It end, const T expected_value) { + const auto it = std::find_if_not( + begin, end, [expected_value](const int elem) { return expected_value == elem; }); + + if (it != end) { + const auto idx = std::distance(begin, it); + INFO("Value mismatch at index " << idx); + REQUIRE(expected_value == *it); + } +} + +template +void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_elements) { + ArrayFindIfNot(array, array + num_elements, expected_value); +} + +template +__global__ void VectorIncrement(T* const vec, const T increment_value, 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 < N; i += stride) { + vec[i] += increment_value; + } +} + +template __global__ void VectorSet(T* const vec, const T value, 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 < N; i += stride) { + vec[i] = value; + } +} + +// Will execute for atleast interval milliseconds +static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { + while (interval--) { + uint64_t start = clock(); + while (clock() - start < ticks_per_ms) { + } + } +} + +inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream) { + int ticks_per_ms = 0; + // Clock rate is in kHz => number of clock ticks in a millisecond + HIP_CHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0)); + Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms); + HIP_CHECK(hipGetLastError()); +} + +template +inline bool DeviceAttributesSupport(const int device, Attributes... attributes) { + constexpr auto DeviceAttributeSupport = [](const int device, + const hipDeviceAttribute_t attribute) { + int value = 0; + HIP_CHECK(hipDeviceGetAttribute(&value, attribute, device)); + return value; + }; + return (... && DeviceAttributeSupport(device, attributes)); +} \ No newline at end of file diff --git a/projects/hip/tests/catch/unit/memory/CMakeLists.txt b/projects/hip/tests/catch/unit/memory/CMakeLists.txt index f24c63ad8c..54e7708556 100644 --- a/projects/hip/tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/memory/CMakeLists.txt @@ -183,5 +183,4 @@ endif() hip_add_exe_to_target(NAME MemoryTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests - COMPILE_OPTIONS -std=c++14) + TEST_TARGET_NAME build_tests) From 996399531cb4b335fbcb33d5fbc8e511f2234413 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 12 Oct 2022 08:31:25 +0530 Subject: [PATCH 22/28] SWDEV-348820 - Fix hipPointerAttribute_t incompatibility (#2856) - Mark hipPointerAttribute_t.memoryType deprecated - Instead add hipPointerAttribute_t.type - Update testcases & documentation to use hipPointerAttribute_t.type Change-Id: I3ba58cd2defe82102f9415bbb90cf151e08499dd [ROCm/hip commit: 96b3b9eaf58bfc9d9d7584657da2695a8b94134a] --- .../hip/docs/markdown/hip_porting_guide.md | 27 ++++++++++--------- projects/hip/include/hip/hip_runtime_api.h | 27 ++++++++++++------- .../unit/memory/hipPointerGetAttributes.cc | 18 ++++++------- .../memory/hipPointerAttributes.cpp | 14 +++++----- 4 files changed, 48 insertions(+), 38 deletions(-) diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 179ad64a2f..fe1972e3b3 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -468,31 +468,34 @@ int main() ## CU_POINTER_ATTRIBUTE_MEMORY_TYPE -To get pointer's memory type in HIP/HIP-Clang, developers should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'memoryType' as member variable. 'memoryType' indicates input pointer is allocated on device or host. +To get pointer's memory type in HIP/HIP-Clang, developers should use hipPointerGetAttributes API. First parameter of the API is hipPointerAttribute_t which has 'type' as member variable. 'type' indicates input pointer is allocated on device or host. For example: ``` double * ptr; hipMalloc(reinterpret_cast(&ptr), sizeof(double)); hipPointerAttribute_t attr; -hipPointerGetAttributes(&attr, ptr); /*attr.memoryType will have value as hipMemoryTypeDevice*/ +hipPointerGetAttributes(&attr, ptr); /*attr.type will have value as hipMemoryTypeDevice*/ double* ptrHost; hipHostMalloc(&ptrHost, sizeof(double)); hipPointerAttribute_t attr; -hipPointerGetAttributes(&attr, ptrHost); /*attr.memoryType will have value as hipMemoryTypeHost*/ +hipPointerGetAttributes(&attr, ptrHost); /*attr.type will have value as hipMemoryTypeHost*/ ``` Please note, hipMemoryType enum values are different from cudaMemoryType enum values. -For example, on AMD platform, memoryType is defined in hip_runtime_api.h, +For example, on AMD platform, hipMemoryType is defined in hip_runtime_api.h, +``` typedef enum hipMemoryType { - hipMemoryTypeHost, ///< Memory is physically located on host - hipMemoryTypeDevice, ///< Memory is physically located on device. - hipMemoryTypeArray, ///< Array memory, physically located on device. - hipMemoryTypeUnified ///< Not used currently + hipMemoryTypeHost = 0, ///< Memory is physically located on host + hipMemoryTypeDevice = 1, ///< Memory is physically located on device. (see deviceId for specific device) + hipMemoryTypeArray = 2, ///< Array memory, physically located on device. (see deviceId for specific device) + hipMemoryTypeUnified = 3, ///< Not used currently + hipMemoryTypeManaged = 4 ///< Managed memory, automaticallly managed by the unified memory system } hipMemoryType; - -Looking into CUDA toolkit, it defines memoryType as following, +``` +Looking into CUDA toolkit, it defines cudaMemoryType as following, +``` enum cudaMemoryType { cudaMemoryTypeUnregistered = 0, // Unregistered memory. @@ -500,8 +503,8 @@ enum cudaMemoryType cudaMemoryTypeDevice = 2, // Device memory. cudaMemoryTypeManaged = 3, // Managed memory } - -In this case, memoryType translation for hipPointerGetAttributes needs to be handled properly on nvidia platform to get the correct memory type in CUDA, which is done in the file nvidia_hip_runtime_api.h. +``` +In this case, memory type translation for hipPointerGetAttributes needs to be handled properly on nvidia platform to get the correct memory type in CUDA, which is done in the file nvidia_hip_runtime_api.h. So in any HIP applications which use HIP APIs involving memory types, developers should use #ifdef in order to assign the correct enum values depending on Nvidia or AMD platform. diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 1d8a6cc731..4ffb3f5241 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -153,24 +153,31 @@ typedef struct hipDeviceProp_t { } hipDeviceProp_t; -/** - * Memory type (for pointer attributes) + /* + * @brief HIP Memory type (for pointer attributes) + * @enum + * @ingroup Enumerations */ typedef enum hipMemoryType { - hipMemoryTypeHost, ///< Memory is physically located on host - hipMemoryTypeDevice, ///< Memory is physically located on device. (see deviceId for specific - ///< device) - hipMemoryTypeArray, ///< Array memory, physically located on device. (see deviceId for specific - ///< device) - hipMemoryTypeUnified, ///< Not used currently - hipMemoryTypeManaged ///< Managed memory, automaticallly managed by the unified memory system + hipMemoryTypeHost = 0, ///< Memory is physically located on host + hipMemoryTypeDevice = 1, ///< Memory is physically located on device. (see deviceId for + ///< specific device) + hipMemoryTypeArray = 2, ///< Array memory, physically located on device. (see deviceId for + ///< specific device) + hipMemoryTypeUnified = 3, ///< Not used currently + hipMemoryTypeManaged = 4 ///< Managed memory, automaticallly managed by the unified + ///< memory system } hipMemoryType; /** * Pointer attributes */ typedef struct hipPointerAttribute_t { - enum hipMemoryType memoryType; + union { + // Deprecated, use instead type + enum hipMemoryType memoryType; + enum hipMemoryType type; + }; int device; void* devicePointer; void* hostPointer; diff --git a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc index 8157c8811d..d4ba5960ce 100644 --- a/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc +++ b/projects/hip/tests/catch/unit/memory/hipPointerGetAttributes.cc @@ -43,7 +43,7 @@ constexpr size_t N{1000000}; bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) { return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) && - (lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) && + (lhs.type == rhs.type) && (lhs.device == rhs.device) && (lhs.allocationFlags == rhs.allocationFlags)); } @@ -68,7 +68,7 @@ const char* memoryTypeToString(hipMemoryType memoryType) { void resetAttribs(hipPointerAttribute_t* attribs) { attribs->hostPointer = reinterpret_cast(-1); attribs->devicePointer = reinterpret_cast(-1); - attribs->memoryType = hipMemoryTypeHost; + attribs->type = hipMemoryTypeHost; attribs->device = -2; attribs->isManaged = -1; attribs->allocationFlags = 0xffff; @@ -77,9 +77,9 @@ void resetAttribs(hipPointerAttribute_t* attribs) { void printAttribs(const hipPointerAttribute_t* attribs) { printf( - "hostPointer:%p devicePointer:%p memType:%s deviceId:%d isManaged:%d " + "hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d " "allocationFlags:%u\n", - attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType), + attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type), attribs->device, attribs->isManaged, attribs->allocationFlags); } @@ -151,14 +151,14 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { if (isDevice) { totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes; HIP_CHECK(hipMalloc(reinterpret_cast(&ptr), reference[i]._sizeBytes)); - reference[i]._attrib.memoryType = hipMemoryTypeDevice; + reference[i]._attrib.type = hipMemoryTypeDevice; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = NULL; reference[i]._attrib.allocationFlags = 0; } else { HIP_CHECK(hipHostMalloc(reinterpret_cast(&ptr), reference[i]._sizeBytes, hipHostMallocDefault)); - reference[i]._attrib.memoryType = hipMemoryTypeHost; + reference[i]._attrib.type = hipMemoryTypeHost; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = ptr; reference[i]._attrib.allocationFlags = 0; @@ -190,7 +190,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { checkPointer(ref, i, 2, (ptr + ref._sizeBytes - 1)); } - if (ref._attrib.memoryType == hipMemoryTypeDevice) { + if (ref._attrib.type == hipMemoryTypeDevice) { HIP_CHECK(hipFree(ref._pointer)); } else { HIP_CHECK(hipHostFree(ref._pointer)); @@ -369,11 +369,11 @@ TEST_CASE("Unit_hipPointerGetAttributes_GpuIter") { // Memory address and type check if (MemoryType == MemoryTypes::DeviceMemory) { - REQUIRE(attributes.memoryType == hipMemoryTypeDevice); + REQUIRE(attributes.type == hipMemoryTypeDevice); REQUIRE(attributes.devicePointer == ptr); REQUIRE(attributes.hostPointer == nullptr); } else if (MemoryType == MemoryTypes::HostMemory) { - REQUIRE(attributes.memoryType == hipMemoryTypeHost); + REQUIRE(attributes.type == hipMemoryTypeHost); REQUIRE(attributes.hostPointer == ptr); } else if (MemoryType == MemoryTypes::MappedMemory) { int* devicePtr{nullptr}; diff --git a/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp b/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp index cd9683ea5b..9666dd4693 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipPointerAttributes.cpp @@ -40,7 +40,7 @@ size_t Nbytes = 0; bool operator==(const hipPointerAttribute_t& lhs, const hipPointerAttribute_t& rhs) { return ((lhs.hostPointer == rhs.hostPointer) && (lhs.devicePointer == rhs.devicePointer) && - (lhs.memoryType == rhs.memoryType) && (lhs.device == rhs.device) && + (lhs.type == rhs.type) && (lhs.device == rhs.device) && (lhs.allocationFlags == rhs.allocationFlags)); }; @@ -65,7 +65,7 @@ const char* memoryTypeToString(hipMemoryType memoryType) { void resetAttribs(hipPointerAttribute_t* attribs) { attribs->hostPointer = (void*)(-1); attribs->devicePointer = (void*)(-1); - attribs->memoryType = hipMemoryTypeHost; + attribs->type = hipMemoryTypeHost; attribs->device = -2; attribs->isManaged = -1; attribs->allocationFlags = 0xffff; @@ -74,9 +74,9 @@ void resetAttribs(hipPointerAttribute_t* attribs) { void printAttribs(const hipPointerAttribute_t* attribs) { printf( - "hostPointer:%p devicePointer:%p memoryType:%s deviceId:%d isManaged:%d " + "hostPointer:%p devicePointer:%p type:%s deviceId:%d isManaged:%d " "allocationFlags:%u\n", - attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->memoryType), + attribs->hostPointer, attribs->devicePointer, memoryTypeToString(attribs->type), attribs->device, attribs->isManaged, attribs->allocationFlags); }; @@ -229,13 +229,13 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { if (isDevice) { totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes; HIPCHECK(hipMalloc((void**)&ptr, reference[i]._sizeBytes)); - reference[i]._attrib.memoryType = hipMemoryTypeDevice; + reference[i]._attrib.type = hipMemoryTypeDevice; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = NULL; reference[i]._attrib.allocationFlags = 0; // TODO-randomize these. } else { HIPCHECK(hipHostMalloc((void**)&ptr, reference[i]._sizeBytes, hipHostMallocDefault)); - reference[i]._attrib.memoryType = hipMemoryTypeHost; + reference[i]._attrib.type = hipMemoryTypeHost; reference[i]._attrib.devicePointer = ptr; reference[i]._attrib.hostPointer = ptr; reference[i]._attrib.allocationFlags = 0; // TODO-randomize these. @@ -265,7 +265,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) { checkPointer(ref, i, 2, (char*)ref._pointer + ref._sizeBytes - 1); } - if (ref._attrib.memoryType == hipMemoryTypeDevice) { + if (ref._attrib.type == hipMemoryTypeDevice) { hipFree(ref._pointer); } else { hipHostFree(ref._pointer); From cd091261309f96f2ea257490e2434fe73e86dc4e Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 13 Oct 2022 23:26:40 +0530 Subject: [PATCH 23/28] SWDEV-357797 - Windows: hipcc to take in args with spaces inside quotes (#2990) Avoids "--offload-arch=gfx1034 --offload-arch=gfx1030" from becoming "--offload-arch=gfx1034\ --offload-arch=gfx1030" Change-Id: I4ef0ffeaf5454c274ec4cfb0fe897e5fbc86d3c5 [ROCm/hip commit: 005fd991cc1c7bd0860f8209a000abd077923c53] --- projects/hip/bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 42fa1e8a20..dfb5a8f6b8 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -35,7 +35,7 @@ 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; + $arg =~ s/[^-a-zA-Z0-9_=+,.:\/\\ ]/\\$&/g; } } From bc393c41c7240669ac9293eaf95ec628909779d9 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 14 Oct 2022 02:31:13 +0530 Subject: [PATCH 24/28] SWDEV-327563 - Enable fixed Windows tests v4 (#2989) Change-Id: Ifa85eea8356d332f555f54516cde9c922c63f6af [ROCm/hip commit: aba1e8b921e822c24f6cb95f0c9aae1ee5a6876c] --- .../config/config_amd_windows_common.json | 34 +++++++++++-------- 1 file changed, 19 insertions(+), 15 deletions(-) diff --git a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json index 48630e0ce1..002666b345 100644 --- a/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json +++ b/projects/hip/tests/catch/hipTestMain/config/config_amd_windows_common.json @@ -6,35 +6,39 @@ "Unit_hipGraphAddHostNode_ClonedGraphwithHostNode", "Unit_hipEventIpc", "Unit_hipMalloc3D_Negative", - "Unit_hipStreamValue_Write", "Unit_hipMemPoolApi_BasicAlloc", "Unit_hipMemPoolApi_BasicTrim", "Unit_hipMemPoolApi_BasicReuse", "Unit_hipMemPoolApi_Opportunistic", "Unit_hipGraphMemcpyNodeSetParams_Functional", "Unit_hipMalloc3D_ValidatePitch", - "Unit_hipArrayCreate_happy", "Unit_hipMemAllocPitch_ValidatePitch", - "Unit_hipArrayCreate_happy - int", - "Unit_hipArrayCreate_happy - int4", - "Unit_hipArrayCreate_happy - short2", - "Unit_hipArrayCreate_happy - char", - "Unit_hipArrayCreate_happy - char4", - "Unit_hipArrayCreate_happy - float", - "Unit_hipArrayCreate_happy - float2", - "Unit_hipArrayCreate_happy - float4", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Functional", - "Unit_hipMallocManaged_MultiChunkMultiDevice", - "Unit_hipMallocManaged_TwoPointers - int", - "Unit_hipMallocManaged_TwoPointers - float", - "Unit_hipMallocManaged_TwoPointers - double", + "Unit_hipMallocManaged_OverSubscription", + "Unit_hipMallocManaged_CoherentTstWthAdvise", + "Unit_hipMallocManaged_Advanced", + "Unit_hipMemRangeGetAttribute_TstCountParam", + "Unit_hipMemRangeGetAttribute_NegativeTests", + "Unit_hipMemRangeGetAttribute_AccessedBy1", + "Unit_hipMemRangeGetAttribte_3", + "Unit_hipMemRangeGetAttribute_4", + "Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr", + "Unit_hipMemAdvise_TstFlags", + "Unit_hipMemAdvise_PrefrdLoc", + "Unit_hipMemAdvise_ReadMostly", + "Unit_hipMemAdvise_TstFlgOverrideEffect", + "Unit_hipMemAdvise_TstAccessedByFlg", + "Unit_hipMemAdvise_TstAccessedByFlg4", + "Unit_hipMemAdvise_TstMemAdvisePrefrdLoc", + "Unit_hipMemAdvise_TstMemAdviseMultiFlag", + "Unit_hipMemAdvise_ReadMosltyMgpuTst", + "Unit_hipMemAdvise_TstSetUnsetPrfrdLoc", "Unit_hipMallocManaged_DeviceContextChange - unsigned char", "Unit_hipMallocManaged_DeviceContextChange - int", "Unit_hipMallocManaged_DeviceContextChange - float", "Unit_hipMallocManaged_DeviceContextChange - double", "Unit_hipGraphNodeGetDependentNodes_Functional", "Unit_hipGraphNodeGetDependencies_Functional", - "Unit_hipMemGetInfo_Negative", "Unit_hipStreamCreateWithPriority_ValidateWithEvents", "Unit_hipStreamPerThread_StrmWaitEvt", "Unit_hipMemGetInfo_DifferentMallocSmall", From 410d2fcd459b2f923dd256ac46a485ac84634fac Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 14 Oct 2022 10:08:03 +0530 Subject: [PATCH 25/28] SWDEV-325938 - Add missing queries for occupancy (#2995) New C++ interfaces: hipOccupancyMaxPotentialBlockSizeVariableSMem() and hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags() Change-Id: Ifae9d6b94b159e59cbf616f180028abb18a96527 [ROCm/hip commit: 6d4c13adc26fa01bb033377c0323354c3a3a3cf4] --- projects/hip/include/hip/hip_runtime_api.h | 142 +++++++++++++++++++++ 1 file changed, 142 insertions(+) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 4ffb3f5241..b77f542b4b 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -6775,6 +6775,148 @@ inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( numBlocks, reinterpret_cast(f), blockSize, dynSharedMemPerBlk, flags); } +/** + * @brief Returns grid and block size that achieves maximum potential occupancy for a device function + * + * Returns in \p *min_grid_size and \p *block_size a suggested grid / + * block size pair that achieves the best potential occupancy + * (i.e. the maximum number of active warps on the current device with the smallest number + * of blocks for a particular function). + * + * @param [out] min_grid_size minimum grid size needed to achieve the best potential occupancy + * @param [out] block_size block size required for the best potential occupancy + * @param [in] func device function symbol + * @param [in] block_size_to_dynamic_smem_size - a unary function/functor that takes block size, + * and returns the size, in bytes, of dynamic shared memory needed for a block + * @param [in] block_size_limit the maximum block size \p func is designed to work with. 0 means no limit. + * @param [in] flags reserved + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidDeviceFunction, #hipErrorInvalidValue, + * #hipErrorUnknown + */ +template +static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags( + int* min_grid_size, + int* block_size, + T func, + UnaryFunction block_size_to_dynamic_smem_size, + int block_size_limit = 0, + unsigned int flags = 0) { + if (min_grid_size == nullptr || block_size == nullptr || + reinterpret_cast(func) == nullptr) { + return hipErrorInvalidValue; + } + + int dev; + hipError_t status; + if ((status = hipGetDevice(&dev)) != hipSuccess) { + return status; + } + + int max_threads_per_cu; + if ((status = hipDeviceGetAttribute(&max_threads_per_cu, + hipDeviceAttributeMaxThreadsPerMultiProcessor, dev)) != hipSuccess) { + return status; + } + + int warp_size; + if ((status = hipDeviceGetAttribute(&warp_size, + hipDeviceAttributeWarpSize, dev)) != hipSuccess) { + return status; + } + + int max_cu_count; + if ((status = hipDeviceGetAttribute(&max_cu_count, + hipDeviceAttributeMultiprocessorCount, dev)) != hipSuccess) { + return status; + } + + struct hipFuncAttributes attr; + if ((status = hipFuncGetAttributes(&attr, reinterpret_cast(func))) != hipSuccess) { + return status; + } + + // Initial limits for the execution + const int func_max_threads_per_block = attr.maxThreadsPerBlock; + if (block_size_limit == 0) { + block_size_limit = func_max_threads_per_block; + } + + if (func_max_threads_per_block < block_size_limit) { + block_size_limit = func_max_threads_per_block; + } + + const int block_size_limit_aligned = + ((block_size_limit + (warp_size - 1)) / warp_size) * warp_size; + + // For maximum search + int max_threads = 0; + int max_block_size{}; + int max_num_blocks{}; + for (int block_size_check_aligned = block_size_limit_aligned; + block_size_check_aligned > 0; + block_size_check_aligned -= warp_size) { + // Make sure the logic uses the requested limit and not aligned + int block_size_check = (block_size_limit < block_size_check_aligned) ? + block_size_limit : block_size_check_aligned; + + size_t dyn_smem_size = block_size_to_dynamic_smem_size(block_size_check); + int optimal_blocks; + if ((status = hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + &optimal_blocks, func, block_size_check, dyn_smem_size, flags)) != hipSuccess) { + return status; + } + + int total_threads = block_size_check * optimal_blocks; + if (total_threads > max_threads) { + max_block_size = block_size_check; + max_num_blocks = optimal_blocks; + max_threads = total_threads; + } + + // Break if the logic reached possible maximum + if (max_threads_per_cu == max_threads) { + break; + } + } + + // Grid size is the number of blocks per CU * CU count + *min_grid_size = max_num_blocks * max_cu_count; + *block_size = max_block_size; + + return status; +} + +/** + * @brief Returns grid and block size that achieves maximum potential occupancy for a device function + * + * Returns in \p *min_grid_size and \p *block_size a suggested grid / + * block size pair that achieves the best potential occupancy + * (i.e. the maximum number of active warps on the current device with the smallest number + * of blocks for a particular function). + * + * @param [out] min_grid_size minimum grid size needed to achieve the best potential occupancy + * @param [out] block_size block size required for the best potential occupancy + * @param [in] func device function symbol + * @param [in] block_size_to_dynamic_smem_size - a unary function/functor that takes block size, + * and returns the size, in bytes, of dynamic shared memory needed for a block + * @param [in] block_size_limit the maximum block size \p func is designed to work with. 0 means no limit. + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidDeviceFunction, #hipErrorInvalidValue, + * #hipErrorUnknown + */ +template +static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMem( + int* min_grid_size, + int* block_size, + T func, + UnaryFunction block_size_to_dynamic_smem_size, + int block_size_limit = 0) +{ + return hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(min_grid_size, block_size, func, + block_size_to_dynamic_smem_size, block_size_limit); +} + template inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) { From 0c00c63145b945afda0d2e4bbf4ae101633c251d Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 14 Oct 2022 16:02:40 +0530 Subject: [PATCH 26/28] SWDEV-355608 - Remove clang include path (#2996) Clang doesn't need to be told where to find the clang headers Change-Id: I175781731c43647676d91eedc8a915124f2479ab [ROCm/hip commit: 154009d64af014876c01f3240e25a6eb1c758dc4] --- projects/hip/bin/hipcc.pl | 6 ------ 1 file changed, 6 deletions(-) diff --git a/projects/hip/bin/hipcc.pl b/projects/hip/bin/hipcc.pl index da9559be9e..606727f4cd 100755 --- a/projects/hip/bin/hipcc.pl +++ b/projects/hip/bin/hipcc.pl @@ -184,9 +184,6 @@ if ($HIP_PLATFORM eq "amd") { $HIP_CLANG_TARGET = `$HIPCC -print-target-triple`; chomp($HIP_CLANG_TARGET); - if (! defined $HIP_CLANG_INCLUDE_PATH) { - $HIP_CLANG_INCLUDE_PATH = abs_path("$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/include"); - } if (! defined $HIP_INCLUDE_PATH) { $HIP_INCLUDE_PATH = "$HIP_PATH/include"; } @@ -199,15 +196,12 @@ if ($HIP_PLATFORM eq "amd") { print ("HIP_ROCCLR_HOME=$HIP_ROCCLR_HOME\n"); } print ("HIP_CLANG_PATH=$HIP_CLANG_PATH\n"); - print ("HIP_CLANG_INCLUDE_PATH=$HIP_CLANG_INCLUDE_PATH\n"); print ("HIP_INCLUDE_PATH=$HIP_INCLUDE_PATH\n"); print ("HIP_LIB_PATH=$HIP_LIB_PATH\n"); print ("DEVICE_LIB_PATH=$DEVICE_LIB_PATH\n"); print ("HIP_CLANG_TARGET=$HIP_CLANG_TARGET\n"); } - $HIPCXXFLAGS .= " -isystem \"$HIP_CLANG_INCLUDE_PATH/..\""; - $HIPCFLAGS .= " -isystem \"$HIP_CLANG_INCLUDE_PATH/..\""; $HIPLDFLAGS .= " -L\"$HIP_LIB_PATH\""; if ($isWindows) { $HIPLDFLAGS .= " -lamdhip64"; From 9bf1e7a709d07831a9a613f45b473a3624a9cb01 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 17 Oct 2022 13:41:29 +0530 Subject: [PATCH 27/28] SWDEV-356057 - [catch2][dtest] Adding tests for no GPU environment. (#2981) Change-Id: I1f63bd2abb1b6fe427be7264e0f5fbfa07216727 [ROCm/hip commit: 274edb77252ee663e3b51bd1f754679cbd19231e] --- .../hip/tests/catch/multiproc/CMakeLists.txt | 12 + .../tests/catch/multiproc/dummy_kernel.cpp | 26 + .../hip/tests/catch/multiproc/hipNoGpuTsts.cc | 1687 +++++++++++++++++ 3 files changed, 1725 insertions(+) create mode 100644 projects/hip/tests/catch/multiproc/dummy_kernel.cpp create mode 100644 projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc diff --git a/projects/hip/tests/catch/multiproc/CMakeLists.txt b/projects/hip/tests/catch/multiproc/CMakeLists.txt index 128384ea3a..5541067207 100644 --- a/projects/hip/tests/catch/multiproc/CMakeLists.txt +++ b/projects/hip/tests/catch/multiproc/CMakeLists.txt @@ -15,10 +15,22 @@ set(LINUX_TEST_SRC hipIpcEventHandle.cc hipIpcMemAccessTest.cc deviceAllocationMproc.cc + hipNoGpuTsts.cc ) +add_custom_target(dummy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${CMAKE_CURRENT_SOURCE_DIR}/dummy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../multiproc/dummy_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + # the last argument linker libraries is required for this test but optional to the function +if(HIP_PLATFORM MATCHES "nvidia") +hip_add_exe_to_target(NAME MultiProc + TEST_SRC ${LINUX_TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS nvrtc) +elseif(HIP_PLATFORM MATCHES "amd") hip_add_exe_to_target(NAME MultiProc TEST_SRC ${LINUX_TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) +endif() +add_dependencies(build_tests dummy_kernel.code) + diff --git a/projects/hip/tests/catch/multiproc/dummy_kernel.cpp b/projects/hip/tests/catch/multiproc/dummy_kernel.cpp new file mode 100644 index 0000000000..876fe4e83e --- /dev/null +++ b/projects/hip/tests/catch/multiproc/dummy_kernel.cpp @@ -0,0 +1,26 @@ +/* +Copyright (c) 2015 - 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 "hip/hip_runtime.h" + +extern "C" __global__ void dummy_ker() { +} diff --git a/projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc b/projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc new file mode 100644 index 0000000000..ad4b9813f4 --- /dev/null +++ b/projects/hip/tests/catch/multiproc/hipNoGpuTsts.cc @@ -0,0 +1,1687 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#include + +#ifdef __linux__ +#include +#include +#include + +// Common Definitions and Macros +#if HT_AMD + static const char* visibility_env_var = "HIP_VISIBLE_DEVICES"; +#else + static const char* visibility_env_var = "CUDA_VISIBLE_DEVICES"; +#endif + +#define UNUSED(expr) do { (void)(expr); } while (0) + +void Callback(hipStream_t stream, hipError_t status, + void* userData) { + UNUSED(stream); + HIP_CHECK(status); + REQUIRE(userData == NULL); +} + +// Dummy RTC function +static constexpr auto rtcfn { + R"( +extern "C" +__global__ void fn() +{ +} +)"}; + +// Dummy Kernel Function +static __global__ void fn() { +} + +// No GPU test case for hipGetDeviceCount +static bool NoGpuTst_hipGetDeviceCount() { + bool passed = false; + hipError_t err; + int nGpus = 0; + err = hipGetDeviceCount(&nGpus); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipGetDeviceCount: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipSetDevice +static bool NoGpuTst_hipSetDevice() { + bool passed = false; + hipError_t err; + err = hipSetDevice(0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipSetDevice: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetAttribute +static bool NoGpuTst_hipDeviceGetAttribute() { + bool passed = false; + hipError_t err; + int pi = 0; + hipDeviceAttribute_t attr = hipDeviceAttributeCudaCompatibleBegin; + err = hipDeviceGetAttribute(&pi, attr, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetAttribute: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetP2PAttribute +static bool NoGpuTst_hipDeviceGetP2PAttribute() { + bool passed = false; + hipError_t err; + int srcDevice = 0, dstDevice = 1; + int value{-1}; + err = hipDeviceGetP2PAttribute(&value, hipDevP2PAttrPerformanceRank, + srcDevice, dstDevice); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetP2PAttribute: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetPCIBusId +static bool NoGpuTst_hipDeviceGetPCIBusId() { + bool passed = false; + hipError_t err; + int device = 0; + char pciBusId[128]; + err = hipDeviceGetPCIBusId(pciBusId, 128, device); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetPCIBusId: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetByPCIBusId +static bool NoGpuTst_hipDeviceGetByPCIBusId() { + bool passed = false; + hipError_t err; + int device = 0; + char pciBusId[128]; + err = hipDeviceGetByPCIBusId(&device, pciBusId); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetByPCIBusId: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceTotalMem +static bool NoGpuTst_hipDeviceTotalMem() { + bool passed = false; + hipError_t err, err_exp; +#if HT_AMD + err_exp = hipErrorNoDevice; +#else + err_exp = hipErrorNotInitialized; +#endif + int device = 0; + size_t bytes = 0; + err = hipDeviceTotalMem(&bytes, device); + if (err == err_exp) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceTotalMem: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipFuncSetSharedMemConfig +static bool NoGpuTst_hipFuncSetSharedMemConfig() { + bool passed = false; + hipError_t err; + err = hipFuncSetSharedMemConfig(reinterpret_cast(&fn), + hipSharedMemBankSizeDefault); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncSetSharedMemConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamCreate +static bool NoGpuTst_hipStreamCreate() { + bool passed = false; + hipError_t err; + hipStream_t stream = nullptr; + err = hipStreamCreate(&stream); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamCreate: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamCreateWithPriority +static bool NoGpuTst_hipStreamCreateWithPriority() { + bool passed = false; + hipError_t err; + hipStream_t mystream = nullptr; + err = hipStreamCreateWithPriority(&mystream, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamCreateWithPriority: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamSynchronize +static bool NoGpuTst_hipStreamSynchronize() { + bool passed = false; + hipError_t err; + err = hipStreamSynchronize(0); // sync with null stream + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamSynchronize: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipStreamGetFlags +static bool NoGpuTst_hipStreamGetFlags() { + bool passed = false; + hipError_t err; + unsigned int flags = 0; + err = hipStreamGetFlags(0, &flags); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamGetFlags: " << + hipGetErrorName(err)); + } + return passed; +} +#if HT_AMD +// No GPU test case for hipExtStreamCreateWithCUMask +static bool NoGpuTst_hipExtStreamCreateWithCUMask() { + bool passed = false; + hipError_t err; + hipStream_t mystream = nullptr; + uint32_t cuMaskSize = 0; + const uint32_t cuMask = 0; + err = hipExtStreamCreateWithCUMask(&mystream, cuMaskSize, &cuMask); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipExtStreamCreateWithCUMask: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +// No GPU test case for hipStreamAddCallback +static bool NoGpuTst_hipStreamAddCallback() { + bool passed = false; + hipError_t err; + void* userData = NULL; + unsigned int flags = 0; + err = hipStreamAddCallback(0, Callback, userData, flags); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamAddCallback: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipEventCreateWithFlags +static bool NoGpuTst_hipEventCreateWithFlags() { + bool passed = false; + hipError_t err; + hipEvent_t event = 0; + err = hipEventCreateWithFlags(&event, hipEventDefault); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipEventCreateWithFlags: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceSynchronize +static bool NoGpuTst_hipDeviceSynchronize() { + bool passed = false; + hipError_t err; + err = hipDeviceSynchronize(); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceSynchronize: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipGetDevice +static bool NoGpuTst_hipGetDevice() { + bool passed = false; + hipError_t err; + int deviceId = 0; + err = hipGetDevice(&deviceId); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipGetDevice: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetCacheConfig +static bool NoGpuTst_hipDeviceGetCacheConfig() { + bool passed = false; + hipError_t err; + hipFuncCache_t cacheConfig = hipFuncCachePreferNone; + err = hipDeviceGetCacheConfig(&cacheConfig); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetCacheConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceGetSharedMemConfig +static bool NoGpuTst_hipDeviceGetSharedMemConfig() { + bool passed = false; + hipError_t err; + hipSharedMemConfig pConfig = hipSharedMemBankSizeDefault; + err = hipDeviceGetSharedMemConfig(&pConfig); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceGetSharedMemConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipDeviceSetSharedMemConfig +static bool NoGpuTst_hipDeviceSetSharedMemConfig() { + bool passed = false; + hipError_t err; + hipSharedMemConfig pConfig = hipSharedMemBankSizeDefault; + err = hipDeviceSetSharedMemConfig(pConfig); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceSetSharedMemConfig: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipPointerGetAttributes +static bool NoGpuTst_hipPointerGetAttributes() { + bool passed = false; + hipError_t err; + hipPointerAttribute_t attributes = {}; + char* A_d; + err = hipPointerGetAttributes(&attributes, A_d); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipPointerGetAttributes: " << + hipGetErrorName(err)); + } + return passed; +} +#if HT_AMD +// No GPU test case for hipExtMallocWithFlags +static bool NoGpuTst_hipExtMallocWithFlags() { + bool passed = false; + hipError_t err; + int *Ptr = nullptr; + unsigned int flags = 0; + err = hipExtMallocWithFlags(reinterpret_cast(Ptr), 128, flags); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipExtMallocWithFlags: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +// No GPU test case for hipMallocManaged +static bool NoGpuTst_hipMallocManaged() { + bool passed = false; + hipError_t err; + float *dev_ptr; + err = hipMallocManaged(&dev_ptr, 128, hipMemAttachGlobal); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocManaged: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipHostFree +static bool NoGpuTst_hipHostFree() { + bool passed = false; + hipError_t err; + void* ptr1 = nullptr; + err = hipHostFree(ptr1); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipHostFree: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test case for hipMemcpyWithStream +static bool NoGpuTst_hipMemcpyWithStream() { + bool passed = false; + hipError_t err; + int *A_h, *B_h; + A_h = reinterpret_cast(malloc(sizeof(int))); + REQUIRE(A_h != nullptr); + B_h = reinterpret_cast(malloc(sizeof(int))); + REQUIRE(B_h != nullptr); + err = hipMemcpyWithStream(A_h, B_h, 128, hipMemcpyHostToHost, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMemcpyWithStream: " << + hipGetErrorName(err)); + } + return passed; +} + +#if HT_AMD +// No GPU test case for hipMallocMipmappedArray +static bool NoGpuTst_hipMallocMipmappedArray() { + bool passed = false; + hipError_t err; + hipMipmappedArray_t mipmappedArray{}; + hipExtent extent; + hipChannelFormatDesc desc = hipCreateChannelDesc(); + err = hipMallocMipmappedArray(&mipmappedArray, &desc, extent, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocMipmappedArray: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +// No GPU test case for hipDeviceEnablePeerAccess +static bool NoGpuTst_hipDeviceEnablePeerAccess() { + bool passed = false; + hipError_t err; + int peerDeviceId = 0; + err = hipDeviceEnablePeerAccess(peerDeviceId, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipDeviceEnablePeerAccess: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipHostMalloc() API +static bool NoGpuTst_hipHostMalloc() { + bool passed = false; + hipError_t err; + int *ptr; + size_t size = 1024; + unsigned int flag = hipHostMallocDefault; + err = hipHostMalloc(&ptr, size, flag); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipHostMalloc: " <(malloc(size)); + err = hipHostRegister(hostPtr, size, reg_flag); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipHostRegister: " <(malloc(size)); + int *src = reinterpret_cast(malloc(size)); + err = hipMemcpy(dst, src, size, hipMemcpyHostToHost); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMemcpy: " <(); + hipExtent extent = make_hipExtent(64, 64, 64); + err = hipMalloc3DArray(&array, &desc, extent, hipArrayDefault); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMalloc3DArray: " <(&fn), attr, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncSetAttribute: " + <(&fn), config); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncSetCacheConfig: " + <(&fn)); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipFuncGetAttributes: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipModuleLoadData() API +static bool NoGpuTst_hipModuleLoadData() { + bool passed = false; + hipError_t err; + hipModule_t module; + std::ifstream file("dummy_kernel.code", std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) { + err = hipModuleLoadData(&module, &buffer[0]); +#if HT_AMD + if (err == hipErrorNoDevice) { + passed = true; + } +#else + if (err == hipErrorNotInitialized) { + passed = true; + } +#endif + if (passed == false) { + WARN("Error Code Returned by hipModuleLoadData: " << + hipGetErrorName(err)); + } + } else { + WARN("File Read Failed"); + } + return passed; +} + +// No GPU test for hipModuleLoadDataEx() API +static bool NoGpuTst_hipModuleLoadDataEx() { + bool passed = false; + hipError_t err; + hipModule_t module; + std::ifstream file("dummy_kernel.code", std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) { + err = hipModuleLoadDataEx(&module, &buffer[0], 0, nullptr, nullptr); +#if HT_AMD + if (err == hipErrorNoDevice) { + passed = true; + } +#else + if (err == hipErrorNotInitialized) { + passed = true; + } +#endif + if (passed == false) { + WARN("Error Code Returned by hipModuleLoadDataEx: " << + hipGetErrorName(err)); + } + } else { + WARN("File Read Failed"); + } + return true; +} + +// No GPU test for hipLaunchCooperativeKernel() API +static bool NoGpuTst_hipLaunchCooperativeKernel() { + bool passed = false; + hipError_t err; + dim3 dimBlock = dim3(1); + dim3 dimGrid = dim3(1); + err = hipLaunchCooperativeKernel(reinterpret_cast(fn), + dimGrid, dimBlock, nullptr, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipLaunchCooperativeKernel: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipLaunchCooperativeKernelMultiDevice() API +static bool NoGpuTst_hipLaunchCooperativeKernelMultiDevice() { + bool passed = false; + hipError_t err; + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + hipLaunchParams md_params[1]; + void *dev_params[1][1]; + dev_params[0][0] = nullptr; + md_params[0].func = reinterpret_cast(fn); + md_params[0].gridDim = dimGrid; + md_params[0].blockDim = dimBlock; + md_params[0].sharedMem = 0; + md_params[0].stream = hipStreamPerThread; + md_params[0].args = dev_params[0]; + + err = hipLaunchCooperativeKernelMultiDevice(md_params, 1, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipLaunchCooperativeKernelMultiDevice: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipExtLaunchMultiKernelMultiDevice() API +#if HT_AMD +static bool NoGpuTst_hipExtLaunchMultiKernelMultiDevice() { + bool passed = false; + hipError_t err; + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + hipLaunchParams md_params[1]; + md_params[0].func = reinterpret_cast(fn); + md_params[0].gridDim = dimGrid; + md_params[0].blockDim = dimBlock; + md_params[0].sharedMem = 0; + md_params[0].stream = hipStreamPerThread; + md_params[0].args = nullptr; + + err = hipExtLaunchMultiKernelMultiDevice(md_params, 1, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipExtLaunchMultiKernelMultiDevice: " << + hipGetErrorName(err)); + } + return passed; +} +#endif + +// No GPU test for hipExtLaunchMultiKernelMultiDevice() API +static bool NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessor() { + bool passed = false; + hipError_t err; + int numBlock = 0, blockSize = 256; + err = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, + fn, blockSize, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code hipOccupancyMaxActiveBlocksPerMultiprocessor: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags() +static bool NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessorFlags() { + bool passed = false; + hipError_t err; + int numBlock = 0, blockSize = 256; + err = hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(&numBlock, + fn, blockSize, 0, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags: " + << hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipOccupancyMaxPotentialBlockSize() +static bool NoGpuTst_hipOccupancyMaxPotentialBlockSize() { + bool passed = false; + hipError_t err; + int blocksizelimit = 256, blockSize = 0, gridsize = 0; + err = hipOccupancyMaxPotentialBlockSize(&gridsize, &blockSize, fn, + 0, blocksizelimit); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipOccupancyMaxPotentialBlockSize: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hiprtcVersion() +static bool NoGpuTst_hiprtcVersion() { + bool passed = false; + int major = 0, minor = 0; + hiprtcResult err = hiprtcVersion(&major, &minor); + if (err == HIPRTC_SUCCESS) { + passed = true; + } else { + WARN("Error Code Returned by hiprtcVersion: " << + hiprtcGetErrorString(err)); + } + return passed; +} + +// No GPU test for hiprtcCreateProgram() and hiprtcDestroyProgram() +static bool NoGpuTst_hiprtcCreateProgram_DestroyProg() { + bool passed = false; + hiprtcProgram prog; + hiprtcResult err = hiprtcCreateProgram(&prog, rtcfn, "rtcfn.cu", + 0, nullptr, nullptr); + if (err == HIPRTC_SUCCESS) { + passed = true; + err = hiprtcDestroyProgram(&prog); + if (err == HIPRTC_SUCCESS) { + passed = true; + } else { + passed = false; + WARN("Error Code Returned by hiprtcDestroyProgram: " << + hiprtcGetErrorString(err)); + } + } else { + WARN("Error Code Returned by hiprtcCreateProgram: " << + hiprtcGetErrorString(err)); + } + return passed; +} + +// No GPU test for hiprtcAddNameExpression() +static bool NoGpuTst_hiprtcAddNameExpression() { + bool passed = false; + hiprtcProgram prog; + hiprtcResult err = hiprtcCreateProgram(&prog, rtcfn, "rtcfn.cu", + 0, nullptr, nullptr); + if (err == HIPRTC_SUCCESS) { + err = hiprtcAddNameExpression(prog, "dummy_func"); + if (err == HIPRTC_SUCCESS) { + passed = true; + } else { + passed = false; + WARN("Error Code Returned by hiprtcAddNameExpression: " << + hiprtcGetErrorString(err)); + } + } + if (err == HIPRTC_SUCCESS) { + err = hiprtcDestroyProgram(&prog); + } + return passed; +} + +// No GPU test for hipMallocPitch() +#define SIZE_H 20 +#define SIZE_W 179 +static bool NoGpuTst_hipMallocPitch() { + bool passed = false; + hipError_t err; + int* devPtr; + size_t devPitch; + err = hipMallocPitch(reinterpret_cast(&devPtr), &devPitch, + SIZE_W*sizeof(int), SIZE_H); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocPitch: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipMallocArray() +static bool NoGpuTst_hipMallocArray() { + bool passed = false; + hipError_t err; + hipChannelFormatDesc channelDesc = + hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat); + hipArray* hipArray; + err = hipMallocArray(&hipArray, &channelDesc, 256, 256); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMallocArray: " << + hipGetErrorName(err)); + } + return passed; +} +#if HT_AMD +// No GPU test for hipMipmappedArrayCreate() +static bool NoGpuTst_hipMipmappedArrayCreate() { + bool passed = false; + hipError_t err; + unsigned int width = 256, height = 256, mipmap_level = 2; + unsigned int orig_width = width; + unsigned int orig_height = height; + width /= pow(2, mipmap_level); + height /= pow(2, mipmap_level); + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, + 0, hipChannelFormatKindFloat); + HIP_ARRAY3D_DESCRIPTOR mipmapped_array_desc; + memset(&mipmapped_array_desc, 0x00, sizeof(HIP_ARRAY3D_DESCRIPTOR)); + mipmapped_array_desc.Width = orig_width; + mipmapped_array_desc.Height = orig_height; + mipmapped_array_desc.Depth = 0; + mipmapped_array_desc.Format = HIP_AD_FORMAT_FLOAT; + mipmapped_array_desc.NumChannels = + ((channelDesc.x != 0) + (channelDesc.y != 0) +\ + (channelDesc.z != 0) + (channelDesc.w != 0)); + mipmapped_array_desc.Flags = 0; + + hipMipmappedArray* mip_array_ptr; + err = hipMipmappedArrayCreate(&mip_array_ptr, &mipmapped_array_desc, + 2*mipmap_level); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMipmappedArrayCreate: " << + hipGetErrorName(err)); + } + return passed; +} +#endif +#define SMALL_SIZE 4 +static bool NoGpuTst_hipMalloc3D() { + bool passed = false; + hipError_t err; + size_t width = SMALL_SIZE * sizeof(char); + size_t height{SMALL_SIZE}, depth{SMALL_SIZE}; + hipPitchedPtr devPitchedPtr; + hipExtent extent = make_hipExtent(width, height, depth); + err = hipMalloc3D(&devPitchedPtr, extent); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipMalloc3D: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipGraphCreate() API +static bool NoGpuTst_hipGraphCreate() { + bool passed = false; + hipError_t err; + hipGraph_t graph; + + err = hipGraphCreate(&graph, 0); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipGraphCreate: " << + hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipStreamBeginCapture() API +static bool NoGpuTst_hipStreamBeginCapture() { + bool passed = false; + hipError_t err; + + err = hipStreamBeginCapture(0, hipStreamCaptureModeGlobal); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamBeginCapture with null stream: " << + hipGetErrorName(err)); + } + + err = hipStreamBeginCapture(hipStreamPerThread, hipStreamCaptureModeGlobal); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamBeginCapture with " + "hipStreamPerThread stream: " << hipGetErrorName(err)); + } + return passed; +} + +// No GPU test for hipStreamIsCapturing() API +static bool NoGpuTst_hipStreamIsCapturing() { + bool passed = false; + hipError_t err; + hipStreamCaptureStatus cStatus; + + err = hipStreamIsCapturing(0, &cStatus); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamIsCapturing: " << + hipGetErrorName(err)); + } + err = hipStreamIsCapturing(hipStreamPerThread, &cStatus); + if (err == hipErrorNoDevice) { + passed = true; + } else { + WARN("Error Code Returned by hipStreamIsCapturing with " + "hipStreamPerThread stream: " << hipGetErrorName(err)); + } + return passed; +} +// Common function invoking individual tests +static bool NoGpuTst_Common(bool(*test_fn)()) { + bool passed = false; + int *Ptr = reinterpret_cast(mmap(NULL, sizeof(int), + PROT_READ | PROT_WRITE, + MAP_SHARED|MAP_ANONYMOUS, 0, 0)); + if (Ptr == MAP_FAILED) { + WARN("Mapping Failed\n"); + return false; + } + // forking to get number of gpus on this system + if (fork() == 0) { // child process + int TotlGpus = 0; + static_cast(hipGetDeviceCount(&TotlGpus)); + Ptr[0] = TotlGpus; + exit(0); + } else { // parent process + wait(NULL); + if (Ptr[0] > 0) { // If there are gpus active on the system + if ((setenv(visibility_env_var, "-1", 1)) != 0) { + WARN("Unable to turn on HIP_VISIBLE_DEVICES, hence exiting!"); + return false; + } + passed = test_fn(); + } else if (Ptr[0] == 0) { + WARN("GPUs not available!"); + passed = test_fn(); + } + } + return passed; +} + +// Tests +TEST_CASE("Unit_NoGpuTst_hipGetDeviceCount") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDeviceCount)); +} + +TEST_CASE("Unit_NoGpuTst_hipSetDevice") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipSetDevice)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetAttribute") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetAttribute)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetP2PAttribute") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetP2PAttribute)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetPCIBusId") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetPCIBusId)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetByPCIBusId") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetByPCIBusId)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceTotalMem") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceTotalMem)); +} + +TEST_CASE("Unit_NoGpuTst_hipFuncSetSharedMemConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncSetSharedMemConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamCreate)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamCreateWithPriority") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamCreateWithPriority)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamSynchronize") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamSynchronize)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamGetFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamGetFlags)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtStreamCreateWithCUMask") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtStreamCreateWithCUMask)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipStreamAddCallback") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamAddCallback)); +} + +TEST_CASE("Unit_NoGpuTst_hipEventCreateWithFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipEventCreateWithFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceSynchronize") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceSynchronize)); +} + +TEST_CASE("Unit_NoGpuTst_hipGetDevice") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDevice)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetCacheConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetCacheConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetSharedMemConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetSharedMemConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceSetSharedMemConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceSetSharedMemConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipPointerGetAttributes") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipPointerGetAttributes)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtMallocWithFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtMallocWithFlags)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipMallocManaged") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocManaged)); +} + +TEST_CASE("Unit_NoGpuTst_hipHostFree") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipHostFree)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemcpyWithStream") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemcpyWithStream)); +} + +TEST_CASE("Unit_NoGpuTst_hipMallocArray") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocArray)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipMallocMipmappedArray") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocMipmappedArray)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipDeviceEnablePeerAccess") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceEnablePeerAccess)); +} + +TEST_CASE("Unit_NoGpuTst_hipHostMalloc") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipHostMalloc)); +} + +TEST_CASE("Unit_NoGpuTst_hipMalloc") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMalloc)); +} + +TEST_CASE("Unit_NoGpuTst_hipHostRegister") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipHostRegister)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemcpy") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemcpy)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemAllocPitch") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemAllocPitch)); +} + +TEST_CASE("Unit_NoGpuTst_hipMemGetInfo") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMemGetInfo)); +} + +TEST_CASE("Unit_NoGpuTst_hipMalloc3DArray") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMalloc3DArray)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceCanAccessPeer") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceCanAccessPeer)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceDisablePeerAccess") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceDisablePeerAccess)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGet") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGet)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceComputeCapability") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceComputeCapability)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetName") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetName)); +} + +TEST_CASE("Unit_NoGpuTst_hipGetDeviceProperties") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDeviceProperties)); +} + +TEST_CASE("Unit_NoGpuTst_hipChooseDevice") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipChooseDevice)); +} + +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtGetLinkTypeAndHopCount") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtGetLinkTypeAndHopCount)); +} + +TEST_CASE("Unit_NoGpuTst_hipExtStreamGetCUMask") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipExtStreamGetCUMask)); +} +#endif + +TEST_CASE("Unit_NoGpuTst_hipFuncSetAttribute") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncSetAttribute)); +} + +TEST_CASE("Unit_NoGpuTst_hipFuncSetCacheConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncSetCacheConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamCreateWithFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamCreateWithFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetStreamPriorityRange") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetStreamPriorityRange)); +} + +TEST_CASE("Unit_NoGpuTst_hipEventCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipEventCreate)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamGetPriority") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamGetPriority)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceReset") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceReset)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceSetCacheConfig") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceSetCacheConfig)); +} + +TEST_CASE("Unit_NoGpuTst_hipDeviceGetLimit") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipDeviceGetLimit)); +} + +TEST_CASE("Unit_NoGpuTst_hipGetDeviceFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGetDeviceFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipSetDeviceFlags") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipSetDeviceFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipModuleLoad") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipModuleLoad)); +} + +TEST_CASE("Unit_NoGpuTst_hipFuncGetAttributes") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipFuncGetAttributes)); +} + +TEST_CASE("Unit_NoGpuTst_hipModuleLoadData") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipModuleLoadData)); +} + +TEST_CASE("Unit_NoGpuTst_hipModuleLoadDataEx") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipModuleLoadDataEx)); +} + +TEST_CASE("Unit_NoGpuTst_hipLaunchCooperativeKernel") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipLaunchCooperativeKernel)); +} + +TEST_CASE("Unit_NoGpuTst_hipLaunchCooperativeKernelMultiDevice") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipLaunchCooperativeKernelMultiDevice)); +} +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipExtLaunchMultiKernelMultiDevice") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipExtLaunchMultiKernelMultiDevice)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessor") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessor)); +} + +TEST_CASE("Unit_NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessorFlags") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipOccupancyMaxActiveBlocksPerMultiprocessorFlags)); +} + +TEST_CASE("Unit_NoGpuTst_hipOccupancyMaxPotentialBlockSize") { + REQUIRE(NoGpuTst_Common( + NoGpuTst_hipOccupancyMaxPotentialBlockSize)); +} + +TEST_CASE("Unit_NoGpuTst_hiprtcVersion") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hiprtcVersion)); +} + +TEST_CASE("Unit_NoGpuTst_hiprtcCreateProgram_DestroyProg") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hiprtcCreateProgram_DestroyProg)); +} + +TEST_CASE("Unit_NoGpuTst_hiprtcAddNameExpression") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hiprtcAddNameExpression)); +} + +TEST_CASE("Unit_NoGpuTst_hipMallocPitch") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMallocPitch)); +} + +#if HT_AMD +TEST_CASE("Unit_NoGpuTst_hipMipmappedArrayCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMipmappedArrayCreate)); +} +#endif +TEST_CASE("Unit_NoGpuTst_hipMalloc3D") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipMalloc3D)); +} + +TEST_CASE("Unit_NoGpuTst_hipGraphCreate") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipGraphCreate)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamBeginCapture") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamBeginCapture)); +} + +TEST_CASE("Unit_NoGpuTst_hipStreamIsCapturing") { + REQUIRE(NoGpuTst_Common(NoGpuTst_hipStreamIsCapturing)); +} +#endif // #ifdef __linux__ From 3fd7a5dc216e76189c795a4fa13cf1901e515a7e Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 17 Oct 2022 16:57:32 +0530 Subject: [PATCH 28/28] SWDEV-353686 - add default value for missing function (#2991) Change-Id: I76e95636ecf2c50a11d0120f3ffdfcf641841cdb [ROCm/hip commit: 24c0cb92b38688b4db8cf80198752f2e8e03e248] --- projects/hip/include/hip/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index b77f542b4b..aa36787636 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -6392,7 +6392,7 @@ hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, un * @warning : This API is marked as beta, meaning, while this is feature complete, * it is still open to changes and may have outstanding issues. */ -hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count); +hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count __dparm(1)); // doxygen end graph API /** * @}