From ea8a7e1e5492367d50608de30400a7d8abbe2346 Mon Sep 17 00:00:00 2001 From: "Manocha, Rahul" Date: Thu, 29 May 2025 00:50:49 -0700 Subject: [PATCH] SWDEV-460098 - Fix error codes for hipLaunchKernel APIs to match cuda (#206) Co-authored-by: Branislav Brzak --- .../hipExtLaunchMultiKernelMultiDevice.cc | 2 +- .../hipLaunchCooperativeKernelMultiDevice.cc | 2 +- .../unit/executionControl/hipLaunchKernel.cc | 24 ++++++++++-------- .../hipModuleLaunchCooperativeKernel.cc | 2 +- .../module/hip_module_launch_kernel_common.hh | 25 ++++++++++--------- 5 files changed, 29 insertions(+), 26 deletions(-) diff --git a/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc b/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc index 97b1420b9c..fd1dcc8295 100644 --- a/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc +++ b/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc @@ -136,7 +136,7 @@ TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Negative_MultiKernelSameDevic } HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), 2, 0u), - hipErrorInvalidValue); + hipErrorInvalidDevice); for (const auto params : params_list) { HIP_CHECK(hipStreamDestroy(params.stream)); diff --git a/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc b/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc index c6b8503203..686adeccd7 100644 --- a/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc +++ b/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc @@ -151,7 +151,7 @@ TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Negative_MultiKernelSameDe } HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), 2, 0u), - hipErrorInvalidValue); + hipErrorInvalidDevice); for (const auto params : params_list) { HIP_CHECK(hipStreamDestroy(params.stream)); diff --git a/catch/unit/executionControl/hipLaunchKernel.cc b/catch/unit/executionControl/hipLaunchKernel.cc index 0dc072aed9..154f28f6e7 100644 --- a/catch/unit/executionControl/hipLaunchKernel.cc +++ b/catch/unit/executionControl/hipLaunchKernel.cc @@ -76,58 +76,58 @@ TEST_CASE("Unit_hipLaunchKernel_Negative_Parameters") { SECTION("gridDim.x == 0") { HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{0, 1, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("gridDim.y == 0") { HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 0, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("gridDim.z == 0") { HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 0}, dim3{1, 1, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.x == 0") { HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{0, 1, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.y == 0") { HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 0, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.z == 0") { HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 0}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.x > maxBlockDimX") { const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.y > maxBlockDimY") { const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, y, 1}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.z > maxBlockDimZ") { const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, z}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { @@ -135,7 +135,7 @@ TEST_CASE("Unit_hipLaunchKernel_Negative_Parameters") { const unsigned int dim = std::ceil(std::cbrt(max)); HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim}, nullptr, 0, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { @@ -144,7 +144,8 @@ TEST_CASE("Unit_hipLaunchKernel_Negative_Parameters") { nullptr, max, nullptr), hipErrorInvalidValue); } - + + #if HT_AMD SECTION("Invalid stream") { hipStream_t stream = nullptr; HIP_CHECK(hipStreamCreate(&stream)); @@ -153,4 +154,5 @@ TEST_CASE("Unit_hipLaunchKernel_Negative_Parameters") { nullptr, 0, stream), hipErrorInvalidValue); } + #endif } \ No newline at end of file diff --git a/catch/unit/module/hipModuleLaunchCooperativeKernel.cc b/catch/unit/module/hipModuleLaunchCooperativeKernel.cc index 415d468a82..92924d88a6 100644 --- a/catch/unit/module/hipModuleLaunchCooperativeKernel.cc +++ b/catch/unit/module/hipModuleLaunchCooperativeKernel.cc @@ -205,7 +205,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") { HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, 1, 0, stream, nullptr), - hipErrorInvalidValue); + hipErrorContextIsDestroyed); } #endif } diff --git a/catch/unit/module/hip_module_launch_kernel_common.hh b/catch/unit/module/hip_module_launch_kernel_common.hh index eeb064f903..b5dc9f7720 100644 --- a/catch/unit/module/hip_module_launch_kernel_common.hh +++ b/catch/unit/module/hip_module_launch_kernel_common.hh @@ -121,7 +121,8 @@ template void ModuleLaunchKernelPositiveParamet template void ModuleLaunchKernelNegativeParameters( bool extLaunch = false) { hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); - hipError_t expectedErrorZeroBlockDim = hipErrorInvalidConfiguration; + hipError_t expectedErrorLaunchParam = (extLaunch == true) ? hipErrorInvalidConfiguration + : hipErrorInvalidValue; hipError_t expectedErrorOverCapacityGridDim = (extLaunch == true) ? hipSuccess : hipErrorInvalidValue; @@ -133,32 +134,32 @@ template void ModuleLaunchKernelNegativeParamet SECTION("gridDimX == 0") { HIP_CHECK_ERROR(func(f, 0, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("gridDimY == 0") { HIP_CHECK_ERROR(func(f, 1, 0, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("gridDimZ == 0") { HIP_CHECK_ERROR(func(f, 1, 1, 0, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("blockDimX == 0") { HIP_CHECK_ERROR(func(f, 1, 1, 1, 0, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - expectedErrorZeroBlockDim); + expectedErrorLaunchParam); } SECTION("blockDimY == 0") { HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 0, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - expectedErrorZeroBlockDim); + expectedErrorLaunchParam); } SECTION("blockDimZ == 0") { HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 0, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - expectedErrorZeroBlockDim); + expectedErrorLaunchParam); } SECTION("gridDimX > maxGridDimX") { @@ -182,19 +183,19 @@ template void ModuleLaunchKernelNegativeParamet SECTION("blockDimX > maxBlockDimX") { const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; HIP_CHECK_ERROR(func(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("blockDimY > maxBlockDimY") { const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, y, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("blockDimZ > maxBlockDimZ") { const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, z, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("blockDimX * blockDimY * blockDimZ > MaxThreadsPerBlock") { @@ -202,7 +203,7 @@ template void ModuleLaunchKernelNegativeParamet const unsigned int dim = std::ceil(std::cbrt(max)) + 1; HIP_CHECK_ERROR( func(f, 1, 1, 1, dim, dim, dim, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + expectedErrorLaunchParam); } SECTION("sharedMemBytes > max shared memory per block") { @@ -225,7 +226,7 @@ template void ModuleLaunchKernelNegativeParamet }; // clang-format on HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args, extra, nullptr, nullptr, 0u), - hipErrorInvalidConfiguration); + hipErrorInvalidValue); } SECTION("Invalid extra") {