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

SWDEV-523281 - CHANGELOG.md and negative test return values : hipLaunchKernelEx, hipLaunchKernelExC, hipDrvLaunchKernelEx (#155)

Browse Source 
[ROCm/clr commit: 64d6f5714a]
This commit is contained in:
GunaShekar, Ajay
2025-04-22 09:17:37 -07:00
committed by GitHub
parent f8344154a0
commit 34d3f7022b
2 changed files with 30 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_module.cpp
+26 -3
View File
@@ -824,6 +824,10 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b
return hipErrorCooperativeLaunchTooLarge;
}
if (globalWorkSizeX == 0 || globalWorkSizeY == 0 || globalWorkSizeZ == 0) {
return hipErrorInvalidConfiguration;
}
return ihipModuleLaunchKernel(func, static_cast<uint32_t>(globalWorkSizeX),
static_cast<uint32_t>(globalWorkSizeY),
static_cast<uint32_t>(globalWorkSizeZ), blockDim.x, blockDim.y,
@@ -1086,11 +1090,19 @@ hipError_t hipLinkDestroy(hipLinkState_t hip_link_state) {
hipError_t hipLaunchKernelExC(const hipLaunchConfig_t* config, const void* fPtr, void** args) {
HIP_INIT_API(hipLaunchKernelExC, config, fPtr, args);
if (fPtr == nullptr) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
if (fPtr == nullptr || config == nullptr || config->numAttrs == 0) {
if (config == nullptr) {
HIP_RETURN(hipErrorInvalidConfiguration);
}
if (config->numAttrs == 0) {
HIP_RETURN_DURATION(hipLaunchKernel_common(fPtr, config->gridDim, config->blockDim, args,
config->dynamicSmemBytes, config->stream));
}
for (size_t attr_idx = 0; attr_idx < config->numAttrs; ++attr_idx) {
hipLaunchAttribute& attr = config->attrs[attr_idx];
switch (attr.id) {
@@ -1112,9 +1124,12 @@ hipError_t hipLaunchKernelExC(const hipLaunchConfig_t* config, const void* fPtr,
hipError_t hipDrvLaunchKernelEx(const HIP_LAUNCH_CONFIG* config, hipFunction_t f,
void** kernelParams, void** extra) {
HIP_INIT_API(hipDrvLaunchKernelEx, config, f, kernelParams, extra);
if (f == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
}
if (f == nullptr || config == nullptr || config->numAttrs == 0) {
HIP_RETURN(hipErrorInvalidConfiguration);
if (config == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
size_t globalWorkSizeX = static_cast<size_t>(config->gridDimX) * config->blockDimX;
@@ -1126,6 +1141,14 @@ hipError_t hipDrvLaunchKernelEx(const HIP_LAUNCH_CONFIG* config, hipFunction_t f
HIP_RETURN(hipErrorInvalidConfiguration);
}
if (config->numAttrs == 0) {
HIP_RETURN(ihipModuleLaunchKernel(
f, static_cast<uint32_t>(globalWorkSizeX), static_cast<uint32_t>(globalWorkSizeY),
static_cast<uint32_t>(globalWorkSizeZ), config->blockDimX, config->blockDimY,
config->blockDimZ, config->sharedMemBytes, config->hStream, kernelParams, nullptr,
nullptr, nullptr, 0));
}
for (size_t attr_idx = 0; attr_idx < config->numAttrs; ++attr_idx) {
hipLaunchAttribute& attr = config->attrs[attr_idx];
switch (attr.id) {