diff --git a/catch/unit/executionControl/hipExtLaunchKernel.cc b/catch/unit/executionControl/hipExtLaunchKernel.cc index db2118c71b..21b7719360 100644 --- a/catch/unit/executionControl/hipExtLaunchKernel.cc +++ b/catch/unit/executionControl/hipExtLaunchKernel.cc @@ -77,58 +77,58 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") { SECTION("gridDim.x == 0") { HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{0, 1, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("gridDim.y == 0") { HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 0, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("gridDim.z == 0") { HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 0}, dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.x == 0") { HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{0, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.y == 0") { HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 0, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.z == 0") { HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 0}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.x > maxBlockDimX") { const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.y > maxBlockDimY") { const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, y, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.z > maxBlockDimZ") { const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, z}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { @@ -137,7 +137,7 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") { HIP_CHECK_ERROR( hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim}, nullptr, 0, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { diff --git a/catch/unit/module/hipHccModuleLaunchKernel.cc b/catch/unit/module/hipHccModuleLaunchKernel.cc index 5124115b7a..c9dbca8138 100644 --- a/catch/unit/module/hipHccModuleLaunchKernel.cc +++ b/catch/unit/module/hipHccModuleLaunchKernel.cc @@ -118,17 +118,17 @@ TEST_CASE("Unit_hipHccModuleLaunchKernel_NegTst") { SECTION("-1 to localWorkSizeX(fifth argument)") { HIP_CHECK_ERROR(hipHccModuleLaunchKernel(kernelFunc, width, 1, 1, -1, 1, 1, 0, 0, kernelArgs, nullptr, nullptr, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("-1 to localWorkSizeY(sixth argument)") { HIP_CHECK_ERROR(hipHccModuleLaunchKernel(kernelFunc, width, 1, 1, width, -1, 1, 0, 0, kernelArgs, nullptr, nullptr, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("-1 to localWorkSizeZ(seventh argument)") { HIP_CHECK_ERROR(hipHccModuleLaunchKernel(kernelFunc, width, 1, 1, width, 1, -1, 0, 0, kernelArgs, nullptr, nullptr, nullptr), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("-1 to sharedMemBytes(eighth argument)") { HIP_CHECK_ERROR(hipHccModuleLaunchKernel(kernelFunc, width, 1, 1, width, 1, diff --git a/catch/unit/module/hip_module_launch_kernel_common.hh b/catch/unit/module/hip_module_launch_kernel_common.hh index cba8329cc5..eeb064f903 100644 --- a/catch/unit/module/hip_module_launch_kernel_common.hh +++ b/catch/unit/module/hip_module_launch_kernel_common.hh @@ -121,8 +121,7 @@ template void ModuleLaunchKernelPositiveParamet template void ModuleLaunchKernelNegativeParameters( bool extLaunch = false) { hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); - hipError_t expectedErrorZeroBlockDim = (extLaunch == true) ? hipErrorInvalidConfiguration - : hipErrorInvalidValue; + hipError_t expectedErrorZeroBlockDim = hipErrorInvalidConfiguration; hipError_t expectedErrorOverCapacityGridDim = (extLaunch == true) ? hipSuccess : hipErrorInvalidValue; @@ -134,17 +133,17 @@ template void ModuleLaunchKernelNegativeParamet SECTION("gridDimX == 0") { HIP_CHECK_ERROR(func(f, 0, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("gridDimY == 0") { HIP_CHECK_ERROR(func(f, 1, 0, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("gridDimZ == 0") { HIP_CHECK_ERROR(func(f, 1, 1, 0, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDimX == 0") { @@ -183,19 +182,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), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } 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), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } 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), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("blockDimX * blockDimY * blockDimZ > MaxThreadsPerBlock") { @@ -203,7 +202,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), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("sharedMemBytes > max shared memory per block") { @@ -226,7 +225,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), - hipErrorInvalidValue); + hipErrorInvalidConfiguration); } SECTION("Invalid extra") {