From 4c513305a12bc9857dc83fc0e966286164a3ea9c Mon Sep 17 00:00:00 2001 From: SrinivasaRao Date: Fri, 8 Nov 2024 09:49:40 +0530 Subject: [PATCH] SWDEV-493299-Test Env variable to avoid silent errors with hipGetLastError feature Change-Id: Iec5720021f626f2021595db5bb2d2c99c0deeefd [ROCm/hip-tests commit: 08b2adef479d6cf1aa322a704341a558b2402ac5] --- .../catch/unit/errorHandling/CMakeLists.txt | 5 +- .../unit/errorHandling/hipGetLastError.cc | 945 +++++++++++------- .../errorHandling/hipGetLastErrorEnv_Exe.cc | 27 + .../unit/errorHandling/hipPeekAtLastError.cc | 310 +++++- .../hipPeekAtLastErrorEnv_Exe.cc | 27 + 5 files changed, 962 insertions(+), 352 deletions(-) create mode 100644 projects/hip-tests/catch/unit/errorHandling/hipGetLastErrorEnv_Exe.cc create mode 100644 projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastErrorEnv_Exe.cc diff --git a/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt b/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt index b9a6de0afa..660740e2ad 100644 --- a/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/errorHandling/CMakeLists.txt @@ -8,8 +8,11 @@ set(TEST_SRC hipGetLastError.cc hipPeekAtLastError.cc ) - +add_executable(hipGetLastErrorEnv_Exe EXCLUDE_FROM_ALL hipGetLastErrorEnv_Exe.cc) +add_executable(hipPeekAtLastErrorEnv_Exe EXCLUDE_FROM_ALL hipPeekAtLastErrorEnv_Exe.cc) hip_add_exe_to_target(NAME ErrorHandlingTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMPILE_OPTIONS -std=c++17) + +add_dependencies(build_tests hipGetLastErrorEnv_Exe hipPeekAtLastErrorEnv_Exe) diff --git a/projects/hip-tests/catch/unit/errorHandling/hipGetLastError.cc b/projects/hip-tests/catch/unit/errorHandling/hipGetLastError.cc index c6424688ee..525593939a 100644 --- a/projects/hip-tests/catch/unit/errorHandling/hipGetLastError.cc +++ b/projects/hip-tests/catch/unit/errorHandling/hipGetLastError.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 - 2023 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 - 2024 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 @@ -22,33 +22,40 @@ THE SOFTWARE. #include #include #include +#include #include //NOLINT #include //NOLINT +#ifdef _WIN64 +#include +#define setenv(x, y, z) _putenv_s(x, y) +#define unsetenv(x) _putenv(x) +#endif + static constexpr size_t WIDTH = 1024; static constexpr size_t HEIGHT = 1024; static constexpr size_t N = 1024 * 1024; static constexpr size_t Nbytes = N * sizeof(int); /** -* @addtogroup hipGetLastError hipGetLastError -* @{ -* @ingroup ErrorHandlingTest -* `hipError_t hipGetLastError ( void )` - -* Returns the last error from a runtime call. -*/ + * @addtogroup hipGetLastError hipGetLastError + * @{ + * @ingroup ErrorHandlingTest + * `hipError_t hipGetLastError ( void )` - + * Returns the last error from a runtime call. + */ /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMalloc api invalid arg call. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.2 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMalloc api invalid arg call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGetLastError_Positive_Basic") { HIP_CHECK(hipGetLastError()); @@ -58,29 +65,27 @@ TEST_CASE("Unit_hipGetLastError_Positive_Basic") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with ThreadedZigZagTest api call. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.2 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with ThreadedZigZagTest api call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipGetLastError_Positive_Threaded") { - class HipGetLastErrorThreadedTest : public - ThreadedZigZagTest { + class HipGetLastErrorThreadedTest + : public ThreadedZigZagTest { public: - void TestPart2() { - REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); - } - void TestPart3() { - HIP_CHECK(hipGetLastError()); - } - void TestPart4() { - REQUIRE_THREAD(hipGetLastError() == hipErrorInvalidValue); + void TestPart2() { + REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); + } + void TestPart3() { HIP_CHECK(hipGetLastError()); } + void TestPart4() { + REQUIRE_THREAD(hipGetLastError() == hipErrorInvalidValue); } }; @@ -89,16 +94,16 @@ TEST_CASE("Unit_hipGetLastError_Positive_Threaded") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpyPeerAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpyPeerAsync api invalid arg call + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemcpyPeerAsync") { const auto device_count = HipTest::getDeviceCount(); @@ -142,70 +147,69 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpyPeerAsync") { HIP_CHECK(hipSetDevice(dst_device)); HIP_CHECK(hipFree(B_d)); } else { - INFO("Peer access cannot be enabled between devices " << src_device << - " and devices " << dst_device); + INFO("Peer access cannot be enabled between devices " + << src_device << " and devices " << dst_device); } } #if HT_AMD /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpyDtoHAsync api invalid arg call -*  Verify hipGetLastError status with hipMemcpyDtoDAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpyDtoHAsync api invalid arg call + *  Verify hipGetLastError status with hipMemcpyDtoDAsync api invalid arg call + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemcpyDtoHAsync") { int *A_d, *B_d, *A_h; hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - HipTest::initArrays(&A_d, &B_d, nullptr, - &A_h, nullptr, nullptr, N, false); + HipTest::initArrays(&A_d, &B_d, nullptr, &A_h, nullptr, nullptr, N, + false); SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") { HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemcpyDtoHAsync(A_h, (hipDeviceptr_t)A_d, - Nbytes * 2, stream), - hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyDtoHAsync(A_h, (hipDeviceptr_t)A_d, Nbytes * 2, stream), + hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); } SECTION("Verify with hipMemcpyDtoDAsync api invalid arg call") { HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemcpyDtoDAsync((hipDeviceptr_t)A_d, - (hipDeviceptr_t)B_d, Nbytes * 2, stream), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpyDtoDAsync((hipDeviceptr_t)A_d, (hipDeviceptr_t)B_d, + Nbytes * 2, stream), + hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); } HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(A_d, B_d, nullptr, - A_h, nullptr, nullptr, false); + HipTest::freeArrays(A_d, B_d, nullptr, A_h, nullptr, nullptr, false); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpyParam2DAsync api invalid arg -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpyParam2DAsync api invalid arg + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemcpyParam2DAsync") { CHECK_IMAGE_SUPPORT - float* A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, * A_d{nullptr}; + float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr}, *A_d{nullptr}; size_t pitch_A; size_t width{WIDTH * sizeof(float)}; constexpr auto memsetval{100}; @@ -213,11 +217,11 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpyParam2DAsync") { HIP_CHECK(hipStreamCreate(&stream)); // Allocating and Initializing the data - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, width, HEIGHT)); - HipTest::initArrays(nullptr, nullptr, nullptr, - &A_h, &B_h, &C_h, width*HEIGHT, false); - HipTest::setDefaultData(WIDTH*HEIGHT, A_h, B_h, C_h); + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, HEIGHT)); + HipTest::initArrays(nullptr, nullptr, nullptr, &A_h, &B_h, &C_h, + width * HEIGHT, false); + HipTest::setDefaultData(WIDTH * HEIGHT, A_h, B_h, C_h); HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, WIDTH, HEIGHT)); // Device to Host @@ -234,30 +238,28 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpyParam2DAsync") { desc.Height = HEIGHT; HIP_CHECK(hipGetLastError()); - desc.WidthInBytes = pitch_A+1; - HIP_CHECK_ERROR(hipMemcpyParam2DAsync(&desc, stream), - hipErrorInvalidValue); + desc.WidthInBytes = pitch_A + 1; + HIP_CHECK_ERROR(hipMemcpyParam2DAsync(&desc, stream), hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); // DeAllocating the memory HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipStreamDestroy(stream)); - HipTest::freeArrays(nullptr, nullptr, nullptr, - A_h, B_h, C_h, false); + HipTest::freeArrays(nullptr, nullptr, nullptr, A_h, B_h, C_h, false); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipDrvMemcpy3DAsync api invalid arg -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipDrvMemcpy3DAsync api invalid arg + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipDrvMemcpy3DAsync") { hipStream_t stream; @@ -273,16 +275,16 @@ TEST_CASE("Unit_hipGetLastError_with_hipDrvMemcpy3DAsync") { #endif /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpy3DAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpy3DAsync api invalid arg call + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy3DAsync") { CHECK_IMAGE_SUPPORT @@ -294,31 +296,31 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy3DAsync") { hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); - int *hData = reinterpret_cast(malloc(size)); + int *hData = reinterpret_cast(malloc(size)); REQUIRE(hData != nullptr); memset(hData, 0, size); // Initialize host buffer - HipTest::setDefaultData(width*height*depth, hData, nullptr, nullptr); + HipTest::setDefaultData(width * height * depth, hData, nullptr, nullptr); - hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, - 0, 0, 0, hipChannelFormatKindSigned); + hipChannelFormatDesc channelDesc = hipCreateChannelDesc( + sizeof(int) * 8, 0, 0, 0, hipChannelFormatKindSigned); HIP_CHECK(hipMalloc3DArray(&devArray, &channelDesc, - make_hipExtent(width, height, 2), hipArrayDefault)); + make_hipExtent(width, height, 2), + hipArrayDefault)); hipMemcpy3DParms myparams; memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); myparams.srcPos = make_hipPos(0, 0, 0); myparams.dstPos = make_hipPos(0, 0, 0); myparams.extent = make_hipExtent(width, height, depth); - myparams.srcPtr = make_hipPitchedPtr(hData, width * sizeof(int), - width, height); + myparams.srcPtr = + make_hipPitchedPtr(hData, width * sizeof(int), width, height); myparams.dstArray = devArray; myparams.kind = hipMemcpyHostToDevice; HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemcpy3DAsync(&myparams, stream), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemcpy3DAsync(&myparams, stream), hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); @@ -329,20 +331,20 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy3DAsync") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpy2DToArrayAsync api invalid arg -*  Verify hipGetLastError status with hipMemcpy2DFromArrayAsync api invalid arg -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpy2DToArrayAsync api invalid arg + *  Verify hipGetLastError status with hipMemcpy2DFromArrayAsync api invalid arg + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy2D_To_From_ArrayAsync") { - int *hData = reinterpret_cast(malloc(WIDTH)); + int *hData = reinterpret_cast(malloc(WIDTH)); REQUIRE(hData != nullptr); memset(hData, 0, WIDTH); hipStream_t stream; @@ -350,8 +352,9 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy2D_To_From_ArrayAsync") { SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") { HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, WIDTH, - WIDTH, HEIGHT, hipMemcpyHostToDevice, stream), + HIP_CHECK_ERROR(hipMemcpy2DToArrayAsync(nullptr, 0, 0, hData, WIDTH, WIDTH, + HEIGHT, hipMemcpyHostToDevice, + stream), hipErrorInvalidHandle); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidHandle); HIP_CHECK(hipGetLastError()); @@ -359,7 +362,8 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy2D_To_From_ArrayAsync") { SECTION("Verify with hipMemcpyDtoHAsync api invalid arg call") { HIP_CHECK(hipGetLastError()); HIP_CHECK_ERROR(hipMemcpy2DFromArrayAsync(hData, WIDTH, nullptr, 0, 0, - WIDTH, HEIGHT, hipMemcpyDeviceToHost, stream), + WIDTH, HEIGHT, + hipMemcpyDeviceToHost, stream), hipErrorInvalidHandle); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidHandle); HIP_CHECK(hipGetLastError()); @@ -370,139 +374,139 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy2D_To_From_ArrayAsync") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipStreamAttachMemAsync api invalid arg -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipStreamAttachMemAsync api invalid arg + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipStreamAttachMemAsync") { - void* d_memory{nullptr}; + void *d_memory{nullptr}; HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemPrefetchAsync(reinterpret_cast(d_memory), - 0, hipMemAttachHost, 0), + HIP_CHECK_ERROR(hipMemPrefetchAsync(reinterpret_cast(d_memory), 0, + hipMemAttachHost, 0), hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipWaitExternalSemaphoresAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipWaitExternalSemaphoresAsync api + *invalid arg call Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipWaitExternalSemaphoresAsync") { hipExternalSemaphoreWaitParams wait_params = {}; wait_params.params.fence.value = 1; HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(nullptr, - &wait_params, 1, nullptr), - hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipWaitExternalSemaphoresAsync(nullptr, &wait_params, 1, nullptr), + hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipSignalExternalSemaphoresAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipSignalExternalSemaphoresAsync api + *invalid arg call Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipSignalExternalSemaphoresAsync") { hipExternalSemaphoreSignalParams signal_params = {}; signal_params.params.fence.value = 1; HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(nullptr, - &signal_params, 1, nullptr), - hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipSignalExternalSemaphoresAsync(nullptr, &signal_params, 1, nullptr), + hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemPrefetchAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemPrefetchAsync api invalid arg + *call Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemPrefetchAsync") { HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemPrefetchAsync(nullptr, 1024, 0), - hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemPrefetchAsync(nullptr, 1024, 0), hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpy2DAsync api invalid arg call -*  Verify hipGetLastError status with hipMemset2DAsync api invalid arg call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpy2DAsync api invalid arg call + *  Verify hipGetLastError status with hipMemset2DAsync api invalid arg call + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy2DAsync") { CHECK_IMAGE_SUPPORT - int* A_h{nullptr}, *A_d{nullptr}; + int *A_h{nullptr}, *A_d{nullptr}; size_t pitch_A; size_t width{WIDTH * sizeof(int)}; hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); // Allocating memory - A_h = reinterpret_cast(malloc(Nbytes)); + A_h = reinterpret_cast(malloc(Nbytes)); REQUIRE(A_h != nullptr); - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, width, WIDTH)); + HIP_CHECK( + hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, WIDTH)); REQUIRE(A_d != nullptr); // Initialize the data - HipTest::setDefaultData(WIDTH*HEIGHT, A_h, nullptr, nullptr); + HipTest::setDefaultData(WIDTH * HEIGHT, A_h, nullptr, nullptr); SECTION("Verify with hipMemcpy2DAsync api invalid arg call") { HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemcpy2DAsync(A_h, WIDTH*2, A_d, pitch_A, - WIDTH*sizeof(int), WIDTH, hipMemcpyDeviceToHost, stream), + HIP_CHECK_ERROR(hipMemcpy2DAsync(A_h, WIDTH * 2, A_d, pitch_A, + WIDTH * sizeof(int), WIDTH, + hipMemcpyDeviceToHost, stream), hipErrorInvalidPitchValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidPitchValue); HIP_CHECK(hipGetLastError()); } SECTION("Verify with hipMemset2DAsync api invalid arg call") { HIP_CHECK(hipGetLastError()); - HIP_CHECK_ERROR(hipMemset2DAsync(A_d, pitch_A, 22, - WIDTH*sizeof(int), WIDTH*9, stream), + HIP_CHECK_ERROR(hipMemset2DAsync(A_d, pitch_A, 22, WIDTH * sizeof(int), + WIDTH * 9, stream), hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); @@ -515,16 +519,16 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemcpy2DAsync") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemsetAsync api invalid arg call. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemsetAsync api invalid arg call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_hipMemsetAsync") { int *A_d; @@ -545,16 +549,16 @@ TEST_CASE("Unit_hipGetLastError_with_hipMemsetAsync") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpyAsync api invalid arg call. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpyAsync api invalid arg call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_with_MemCpyAsync") { int *A_d, *B_d, *C_d; @@ -575,8 +579,9 @@ TEST_CASE("Unit_hipGetLastError_with_MemCpyAsync") { HIP_CHECK(hipStreamSynchronize(stream)); // testing to check error manually - HIP_CHECK_ERROR(hipMemcpyAsync(C_h, C_d, Nbytes+N, hipMemcpyDeviceToHost, 0), - hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyAsync(C_h, C_d, Nbytes + N, hipMemcpyDeviceToHost, 0), + hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); @@ -587,17 +592,17 @@ TEST_CASE("Unit_hipGetLastError_with_MemCpyAsync") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipMemcpyAsync api invalid arg call -* Check in other thread this error should not report by hipGetLastError() -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipMemcpyAsync api invalid arg call + * Check in other thread this error should not report by hipGetLastError() + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ // Inside thread, both hipGetLastError() api call should not return error static void thread_wait_func(int sleep_time) { @@ -628,8 +633,9 @@ TEST_CASE("Unit_hipGetLastError_with_MemCpyAsync_thread") { std::thread t(thread_wait_func, 2); // testing to check error manually - HIP_CHECK_ERROR(hipMemcpyAsync(C_h, C_d, Nbytes+N, hipMemcpyDeviceToHost, 0), - hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemcpyAsync(C_h, C_d, Nbytes + N, hipMemcpyDeviceToHost, 0), + hipErrorInvalidValue); t.join(); @@ -643,18 +649,18 @@ TEST_CASE("Unit_hipGetLastError_with_MemCpyAsync_thread") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipGraphAddMemcpyNode1D api -* Create graph with one node as error consciously so it produces an error, -* which will be used to verify the behavior of hipGetLastError api. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.3 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipGraphAddMemcpyNode1D api + * Create graph with one node as error consciously so it produces an error, + * which will be used to verify the behavior of hipGetLastError api. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.3 + */ TEST_CASE("Unit_hipGetLastError_with_hipGraphAddMemcpyNode1D") { constexpr auto blocksPerCU = 6; // to hide latency @@ -681,12 +687,12 @@ TEST_CASE("Unit_hipGetLastError_with_hipGraphAddMemcpyNode1D") { Nbytes, hipMemcpyDeviceToHost)); hipKernelNodeParams kNodeParams{}; - void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + void *kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; kNodeParams.func = reinterpret_cast(HipTest::vectorADD); kNodeParams.gridDim = dim3(blocks); kNodeParams.blockDim = dim3(threadsPerBlock); kNodeParams.sharedMemBytes = 0; - kNodeParams.kernelParams = reinterpret_cast(kernelArgs); + kNodeParams.kernelParams = reinterpret_cast(kernelArgs); kNodeParams.extra = nullptr; HIP_CHECK(hipGraphAddKernelNode(&kVecAdd, graph, nullptr, 0, &kNodeParams)); @@ -695,8 +701,9 @@ TEST_CASE("Unit_hipGetLastError_with_hipGraphAddMemcpyNode1D") { // behavior of hipGetLastError api. HIP_CHECK(hipGetLastError()); HIP_CHECK_ERROR(hipGraphAddMemcpyNode1D(&memcpy_E, graph, nullptr, 0, C_h, - C_d, Nbytes * 2, hipMemcpyDeviceToHost), - hipErrorInvalidValue); + C_d, Nbytes * 2, + hipMemcpyDeviceToHost), + hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); HIP_CHECK(hipGetLastError()); @@ -722,20 +729,20 @@ TEST_CASE("Unit_hipGetLastError_with_hipGraphAddMemcpyNode1D") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipStreamEndCapture api invalid arg call -* Create a graph1 with stream with ketnelNode as vector_ADD and -* hipStreamEndCapture on graph1 with hipGraphInstantiate to create graphExec -* Again hipStreamEndCapture on graph2 which will return hipErrorIllegalState -* now verify the behavior of hipGetLastError api with this call. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.3 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipStreamEndCapture api invalid arg + *call Create a graph1 with stream with ketnelNode as vector_ADD and + * hipStreamEndCapture on graph1 with hipGraphInstantiate to create graphExec + * Again hipStreamEndCapture on graph2 which will return hipErrorIllegalState + * now verify the behavior of hipGetLastError api with this call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.3 + */ TEST_CASE("Unit_hipGetLastError_with_hipStreamBegin_EndCapture") { int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; @@ -774,16 +781,16 @@ TEST_CASE("Unit_hipGetLastError_with_hipStreamBegin_EndCapture") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with hipGraphCreate api invalid arg call. -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with hipGraphCreate api invalid arg call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_error_check_with_hipGraphCreate") { hipGraph_t graph; @@ -797,16 +804,16 @@ TEST_CASE("Unit_hipGetLastError_error_check_with_hipGraphCreate") { } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status where a success call after hip runtime error -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status where a success call after hip runtime error + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ #if HT_NVIDIA TEST_CASE("Unit_hipGetLastError_success_before_hipGetLastError") { @@ -824,19 +831,19 @@ TEST_CASE("Unit_hipGetLastError_success_before_hipGetLastError") { #endif /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status should update with new api invalid arg call. -* Api hipGraphCreate -> return error hipErrorInvalidValue -* Api hipDeviceGetGraphMemAttribute -> return error hipErrorInvalidDevice -* Now hipGetLastError() api shoud return hipErrorInvalidDevice -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status should update with new api invalid arg call. + * Api hipGraphCreate -> return error hipErrorInvalidValue + * Api hipDeviceGetGraphMemAttribute -> return error hipErrorInvalidDevice + * Now hipGetLastError() api shoud return hipErrorInvalidDevice + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ TEST_CASE("Unit_hipGetLastError_success_before_hipGetLastError_check_again") { int value = 0; @@ -846,27 +853,26 @@ TEST_CASE("Unit_hipGetLastError_success_before_hipGetLastError_check_again") { HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); - HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(-1, - hipGraphMemAttrUsedMemCurrent, &value), hipErrorInvalidDevice); + HIP_CHECK_ERROR( + hipDeviceGetGraphMemAttribute(-1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidDevice); HIP_CHECK(hipGetLastError()); } /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with divide_by_zero kernel call -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with divide_by_zero kernel call + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ -static void __global__ devideKernl(int *i, int x, int y) { - *i = x/(x-y); -} +static void __global__ devideKernl(int *i, int x, int y) { *i = x / (x - y); } TEST_CASE("Unit_hipGetLastError_with_Kernel_divide_by_zero") { int *i_d; @@ -888,23 +894,21 @@ TEST_CASE("Unit_hipGetLastError_with_Kernel_divide_by_zero") { #if HT_NVIDIA /** -* Test Description -* ------------------------ -*  - Verify hipGetLastError status with Inalid Configuration in kernel call -* Kernel call invalid configuration- blocks=1 & threadsPerBlock=0 -* Kernel call invalid configuration- blocks=1 & threadsPerBlock=1 -* Kernel call invalid configuration- blocks=0 & threadsPerBlock=0 -* Test source -* ------------------------ -*  - unit/errorHandling/hipGetLastError.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 6.0 -*/ + * Test Description + * ------------------------ + *  - Verify hipGetLastError status with Inalid Configuration in kernel call + * Kernel call invalid configuration- blocks=1 & threadsPerBlock=0 + * Kernel call invalid configuration- blocks=1 & threadsPerBlock=1 + * Kernel call invalid configuration- blocks=0 & threadsPerBlock=0 + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.0 + */ -static void __global__ incrementKernl(int *i) { - *i += 1; -} +static void __global__ incrementKernl(int *i) { *i += 1; } TEST_CASE("Unit_hipGetLastError_with_Kernel_Invalid_Configuration") { int *i_d; @@ -957,8 +961,275 @@ TEST_CASE("Unit_hipGetLastError_with_Kernel_Invalid_Configuration") { HIP_CHECK(hipStreamDestroy(stream)); } #endif - +/* +New environment variable DEBUG_HIP_7_PREVIEW is added. +This is used for enabling the backward incompatible changes before the next +major ROCm release 7.0. This will be removed after the ROCm release 7.0. +*/ /** - * End doxygen group hipGetLastError. + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status + * with hipMalloc api invalid arg call. + * Status should be last Error reported in the thread/Runtime. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGetLastError_With_EnvVar_Positive_Basic") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status + * with an Error - Success calls + * Each time status should return the corresponding Error when it called. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGetLastError_With_EnvVar_Chk_Updated_Status") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + hipGraph_t graph; + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + int *C_d; + HIP_CHECK(hipMalloc(&C_d, 1024)); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + HIP_CHECK(hipFree(C_d)); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status + * along with hipPeekAtLastError. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGetLastError_With_EnvVar_Chk_Along_hipPeekAtLastError") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + hipGraph_t graph; + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + int *C_d; + HIP_CHECK(hipMalloc(&C_d, 1024)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status + * with different Error and Success combinations. + * Each time status should return the corresponding Error when it called. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGetLastError_With_EnvVar_Error_Combinations") { + int value = 0; + hipGraph_t graph; + SECTION("A case with Error-Error") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + HIP_CHECK(hipGetLastError()); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } + SECTION("A case with Error-Success-Error-Success") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } + SECTION("A case with Success-Error-Error-Success") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } + SECTION("A Case with Success-Error-Success-Error") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } +} + +static void thread_func() { + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status with a + * runtime api invalid arg call. + * Check in other thread this error should not report by hipGetLastError() + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipGetLastError_With_EnvVar_With_Thread") { + hipGraph_t graph; + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + std::thread t(thread_func); + t.join(); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status + * in the multiple processes. + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +// Disabled due to test failure +#if 0 +TEST_CASE("Unit_hipGetLastError_With_EnvVar_MultiProcess") { + hipGraph_t graph; + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + /*int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue);*/ + hip::SpawnProc proc("hipGetLastErrorEnv_Exe", true); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + REQUIRE(proc.run() == 1); + //HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +#endif +static void __global__ emptyKernl() { } +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipGetLastError status + * with Invalid Configuration in kernel call. + * Kernel call invalid configuration- blocks=0 & threadsPerBlock=0 + * Test source + * ------------------------ + *  - unit/errorHandling/hipGetLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ + +// Below test failed on NVIDIA due to error mismatch produced by the Invalid Kernel config. +// For more details please check the ticket SWDEV-501851 comments. +#if HT_AMD +TEST_CASE("Unit_hipGetLastError_With_EnvVar_Kernel_Invalid_Config") { + hipError_t ret; + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + hipLaunchKernelGGL(emptyKernl, dim3(0), dim3(0), 0, 0); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + ret = hipGetLastError(); + REQUIRE(ret == hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + HIP_CHECK(hipFree(A_d)); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +#endif +/** + * End doxygen group ErrorTest. * @} */ + diff --git a/projects/hip-tests/catch/unit/errorHandling/hipGetLastErrorEnv_Exe.cc b/projects/hip-tests/catch/unit/errorHandling/hipGetLastErrorEnv_Exe.cc new file mode 100644 index 0000000000..7be3cb3ec3 --- /dev/null +++ b/projects/hip-tests/catch/unit/errorHandling/hipGetLastErrorEnv_Exe.cc @@ -0,0 +1,27 @@ +/*Copyright (c) 2024 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_api.h" +#include +int main() { + hipError_t err; + err = hipGetLastError(); + if (err == hipSuccess) + return 1; + else + return 0; +} diff --git a/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastError.cc b/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastError.cc index 3b0ddfddf1..27f6f74963 100644 --- a/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastError.cc +++ b/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastError.cc @@ -1,16 +1,13 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - +Copyright (c) 2022-2024 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 @@ -21,8 +18,15 @@ THE SOFTWARE. */ #include +#include #include +#ifdef _WIN64 +#include +#define setenv(x, y, z) _putenv_s(x, y) +#define unsetenv(x) _putenv(x) +#endif + /** * @addtogroup hipPeekAtLastError hipPeekAtLastError * @{ @@ -34,7 +38,8 @@ THE SOFTWARE. /** * Test Description * ------------------------ - * - Validate that `hipErrorInvalidValue` is returned after invalid `hipMalloc` call. + * - Validate that `hipErrorInvalidValue` is returned after invalid `hipMalloc` + * call. * - Validate that `hipSuccess` is returned again for getting the last error. * Test source * ------------------------ @@ -54,9 +59,10 @@ TEST_CASE("Unit_hipPeekAtLastError_Positive_Basic") { /** * Test Description * ------------------------ - * - Validate that appropriate error is returned when working with multiple threads. - * - Validate that appropriate error is returned for getting the last error when working with - * multiple threads. + * - Validate that appropriate error is returned when working with multiple + * threads. + * - Validate that appropriate error is returned for getting the last error + * when working with multiple threads. * - Cause error on purpose within one of the threads. * Test source * ------------------------ @@ -66,21 +72,297 @@ TEST_CASE("Unit_hipPeekAtLastError_Positive_Basic") { * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipPeekAtLastError_Positive_Threaded") { - class HipPeekAtLastErrorTest : public ThreadedZigZagTest { + class HipPeekAtLastErrorTest + : public ThreadedZigZagTest { public: - void TestPart2() { REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); } + void TestPart2() { + REQUIRE_THREAD(hipMalloc(nullptr, 1) == hipErrorInvalidValue); + } void TestPart3() { HIP_CHECK(hipPeekAtLastError()); HIP_CHECK(hipGetLastError()); } - void TestPart4() { REQUIRE_THREAD(hipPeekAtLastError() == hipErrorInvalidValue); } + void TestPart4() { + REQUIRE_THREAD(hipPeekAtLastError() == hipErrorInvalidValue); + } }; HipPeekAtLastErrorTest test; test.run(); } +/* + * New environment variable DEBUG_HIP_7_PREVIEW is added. + * This is used for enabling the backward incompatible changes before the next + * major ROCm release 7.0. + * This will be removed after the ROCm release 7.0. + */ +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status + * with hipMalloc api invalid arg call. + * Status should be last Error reported in the thread/Runtime. + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_Positive_Basic") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status + * with an Error - Success calls + * Each time status should return the corresponding Error when it called. + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_Chk_Updated_Status") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + hipGraph_t graph; + int value = 0; + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + int *C_d; + HIP_CHECK(hipMalloc(&C_d, 1024)); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidDevice); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status + * along with the hipGetLastError API call. + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_Chk_Along_hipGetLastError") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + hipGraph_t graph; + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipGetLastError(), hipSuccess); + int *C_d; + HIP_CHECK_ERROR(hipPeekAtLastError(), hipSuccess); + HIP_CHECK(hipMalloc(&C_d, 1024)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipSuccess); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} /** -* End doxygen group ErrorTest. -* @} -*/ + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status + * with different Error and Success combinations. + * Each time status should return the corresponding Error when it called. + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_Error_Combinations") { + int value = 0; + hipGraph_t graph; + SECTION("A case with Error-Error-Success-Success") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + HIP_CHECK(hipPeekAtLastError()); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidDevice); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidDevice); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } + SECTION("A case with Error-Success-Error-Success") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidDevice); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } + SECTION("A case with Success-Error-Error-Success") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidDevice); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } + SECTION("A Case with Success-Error-Success-Error") { + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute( + -1, hipGraphMemAttrUsedMemCurrent, &value), + hipErrorInvalidDevice); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidDevice); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } + } +} + +static void thread_func() { + HIP_CHECK_ERROR(hipPeekAtLastError(), hipSuccess); + HIP_CHECK_ERROR(hipMalloc(nullptr, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK(hipFree(A_d)); +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status with a + * runtime api invalid arg call. + * Check in other thread this error should not report by hipPeekAtLastError() + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_With_Thread") { + hipGraph_t graph; + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + std::thread t(thread_func); + t.join(); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status + * in the multiple processes. + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +// Disabled due to test failure +#if 0 +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_MultiProcess") { + hipGraph_t graph; + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + /*int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue);*/ + hip::SpawnProc proc("hipPeekAtLastErrorEnv_Exe", true); + HIP_CHECK_ERROR(hipGraphCreate(&graph, 1), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + REQUIRE(proc.run() == 1); + //HIP_CHECK(hipFree(A_d)); + HIP_CHECK_ERROR(hipPeekAtLastError(), hipErrorInvalidValue); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } else { + INFO("Env Not set"); + } +} +#endif +static void __global__ emptyKernl() {} +/** + * Test Description + * ------------------------ + *  - With Env Var DEBUG_HIP_7_PREVIEW, Verify hipPeekAtLastError status + * with Invalid Configuration in kernel call. + * Kernel call invalid configuration- blocks=0 & threadsPerBlock=0 + * Test source + * ------------------------ + *  - unit/errorHandling/hipPeekAtLastError.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 6.4 + */ +// Below test failed on NVIDIA due to error mismatch produced by the Invalid Kernel config. +// For more details please check the ticket SWDEV-501896 comments. +#if HT_AMD +TEST_CASE("Unit_hipPeekAtLastError_With_EnvVar_Kernel_Invalid_Config") { + hipError_t ret; + if (setenv("DEBUG_HIP_7_PREVIEW", "1", 1) == 0) { + hipLaunchKernelGGL(emptyKernl, dim3(0), dim3(0), 0, 0); + int *A_d; + HIP_CHECK(hipMalloc(&A_d, 1024)); + ret = hipPeekAtLastError(); + REQUIRE(ret == hipErrorInvalidValue); + HIP_CHECK(hipFree(A_d)); + unsetenv("DEBUG_HIP_7_PREVIEW"); + } +} +#endif +/** + * End doxygen group ErrorTest. + * @} + */ + diff --git a/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastErrorEnv_Exe.cc b/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastErrorEnv_Exe.cc new file mode 100644 index 0000000000..11064b4822 --- /dev/null +++ b/projects/hip-tests/catch/unit/errorHandling/hipPeekAtLastErrorEnv_Exe.cc @@ -0,0 +1,27 @@ +/*Copyright (c) 2024 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_api.h" +#include +int main() { + hipError_t err; + err = hipPeekAtLastError(); + if (err == hipSuccess) + return 1; + else + return 0; +}