AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-508973 - Fix the catch tests failures due to change in kernel validation error codes to match CUDA (#198)

Browse Source
This commit is contained in:
Dittakavi, Satyanvesh
2025-05-26 18:41:58 +05:30
committed by GitHub
parent 74e743dcb9
commit 1be3aa8283
3 changed files with 22 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/executionControl/hipExtLaunchKernel.cc
+10 -10
View File
@@ -77,58 +77,58 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") {
SECTION("gridDim.x == 0") {
HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast<void*>(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<void*>(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<void*>(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<void*>(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<void*>(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<void*>(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<void*>(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<void*>(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<void*>(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<void*>(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim},
nullptr, 0, nullptr, nullptr, nullptr, 0u),
hipErrorInvalidValue);
hipErrorInvalidConfiguration);
}
SECTION("sharedMemBytes > maxSharedMemoryPerBlock") {
catch/unit/module/hipHccModuleLaunchKernel.cc
+3 -3
View File
@@ -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,
catch/unit/module/hip_module_launch_kernel_common.hh
+9 -10
View File
@@ -121,8 +121,7 @@ template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveParamet
template <ExtModuleLaunchKernelSig* func> 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 <ExtModuleLaunchKernelSig* func> 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 <ExtModuleLaunchKernelSig* func> 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 <ExtModuleLaunchKernelSig* func> 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 <ExtModuleLaunchKernelSig* func> 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") {