[SWDEV-442583] - Fix hipLaunchCooperativeKernel API error codes

Change-Id: Id3cb4b58b9efd3aceec4377d4d1d3a053c50333e [ROCm/clr commit: 59051ed3af]
2024-04-16 20:04:16 +00:00
parent 1e22a14b1b
commit 35ab01f7de
1 changed files with 9 additions and 1 deletions
@@ -705,7 +705,11 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b
                                             void** kernelParams, uint32_t sharedMemBytes,
                                             hipStream_t hStream) {
  if (!hip::isValid(hStream)) {
-    return hipErrorInvalidValue;
+    return hipErrorContextIsDestroyed;
+  }
+
+  if (f == nullptr) {
+    return hipErrorInvalidDeviceFunction;
  }

  hipFunction_t func = nullptr;
@@ -725,6 +729,10 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b
    return hipErrorInvalidConfiguration;
  }

+  if (sharedMemBytes > device->info().localMemSizePerCU_) {
+    return hipErrorCooperativeLaunchTooLarge;
+  }
+
  return ihipModuleLaunchKernel(func, static_cast<uint32_t>(globalWorkSizeX),
                                static_cast<uint32_t>(globalWorkSizeY),
                                static_cast<uint32_t>(globalWorkSizeZ), blockDim.x, blockDim.y,