Merge "Fix hipExtLaunchMultiKernelMultiDevice compilation issue" into amd-master-next

[ROCm/clr commit: 1216153c66]

projects/clr/hipamd/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")))

projects/clr/hipamd/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
 /**

projects/clr/hipamd/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

projects/clr/hipamd/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
  */