From 35ab01f7deec517d92f3625bb3934ef6862b1769 Mon Sep 17 00:00:00 2001
From: Rahul Manocha <rahul.manocha@amd.com>
Date: Tue, 16 Apr 2024 20:04:16 +0000
Subject: [PATCH] [SWDEV-442583] - Fix hipLaunchCooperativeKernel API error
 codes

Change-Id: Id3cb4b58b9efd3aceec4377d4d1d3a053c50333e


[ROCm/clr commit: 59051ed3af1a0b6416eac299fdccd9e5721b9661]
---
 projects/clr/hipamd/src/hip_module.cpp | 10 +++++++++-
 1 file changed, 9 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp
index 1d96c6a88d..8707486424 100644
--- a/projects/clr/hipamd/src/hip_module.cpp
+++ b/projects/clr/hipamd/src/hip_module.cpp
@@ -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,