diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp
index 76a04fa355..cf4422070f 100644
--- a/include/hip/hcc_detail/functional_grid_launch.hpp
+++ b/include/hip/hcc_detail/functional_grid_launch.hpp
@@ -189,16 +189,6 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
                                      stream, &config[0]);
 }
 
-inline
-__attribute__((visibility("hidden")))
-hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
-                                              int  numDevices, unsigned int  flags) {
-    hip_impl::hip_init();
-    auto& ps = hip_impl::get_program_state();
-    return ihipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags, ps);
-
-}
-
 template <typename F>
 inline
 __attribute__((visibility("hidden")))
diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 0712db17f9..936170b1fd 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -2962,7 +2962,6 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
 hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
    uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
 
-#if __HIP_VDI__ && !defined(__HCC__)
 /**
  * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched
  * on respective streams before enqueuing any other work on the specified streams from any other threads
@@ -2977,7 +2976,6 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
 hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
                                               int  numDevices, unsigned int  flags);
 
-#endif
 
 // doxygen end Version Management
 /**
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 18b869cee7..8901f2c4aa 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -392,6 +392,14 @@ hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList
     return result;
 }
 
+__attribute__((visibility("default")))
+hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
+                                              int  numDevices, unsigned int  flags) {
+    HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags);
+    auto& ps = hip_impl::get_program_state();
+    return ihipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags, ps);
+}
+
 namespace {
 // kernel for initializing GWS
 // nwm1 is the total number of work groups minus 1
diff --git a/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp b/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp
index cc7a06d66b..fe9b06713c 100644
--- a/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp
+++ b/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp
@@ -23,7 +23,7 @@ THE SOFTWARE.
 // single GPU or multi GPUs.
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc vdi
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */