diff --git a/CHANGELOG.md b/CHANGELOG.md
index 4c8ac5e97b..b78532feb0 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -6,6 +6,10 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs
 
 ### Added
 
+* New HIP APIs
+    - `hipLaunchKernelEx`  dispatches the provided kernel with the given launch configuration and forwards the kernel arguments.
+    - `hipLaunchKernelExC`  launches a HIP kernel using a generic function pointer and the specified configuration.
+    - `hipDrvLaunchKernelEx`  dispatches the device kernel represented by a HIP function object
 * New support for Open Compute Project (OCP) floating-point `FP4`/`FP6`/`FP8` as the following. For details, see [Low precision floating point document](https://rocm.docs.amd.com/projects/HIP/en/latest/reference/low_fp_types.html).
     - Data types for `FP4`/`FP6`/`FP8`.
     - HIP APIs for `FP4`/`FP6`/`FP8`, which are compatible with corresponding CUDA APIs.
@@ -35,7 +39,6 @@ HIP runtime has the following functional improvements which greatly improve runt
 
 * Error of "unable to find modules" in HIP clean up for code object module.
 
-
 ## HIP 6.4 (For ROCm 6.4)
 
 ### Added
diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp
index b773f81dfc..7d048ec67a 100644
--- a/hipamd/src/hip_module.cpp
+++ b/hipamd/src/hip_module.cpp
@@ -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) {