From 76ee85ff82864389e04e4d419bcb179b235ad629 Mon Sep 17 00:00:00 2001 From: agodavar Date: Thu, 27 Feb 2020 20:06:10 +0530 Subject: [PATCH] Fix hipExtLaunchMultiKernelMultiDevice compilation issue Fix compilation error on hip-hcc+clang , hip-vdi+clang Enabled hipExtLaunchMultiKernelMultiDevice test on hip-vdi path hipExtLaunchMultiKernelMultiDevice common declaration for all paths Change-Id: I76031840614fce8e12a8e845548fa43a389a741a [ROCm/hip commit: 6a5d04209c79f9f9aed4c799ddf2d44e5931befe] --- .../include/hip/hcc_detail/functional_grid_launch.hpp | 10 ---------- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 2 -- projects/hip/src/hip_module.cpp | 8 ++++++++ .../module/hipExtLaunchMultiKernelMultiDevice.cpp | 2 +- 4 files changed, 9 insertions(+), 13 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index 76a04fa355..cf4422070f 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/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 inline __attribute__((visibility("hidden"))) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 0712db17f9..936170b1fd 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/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/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 18b869cee7..8901f2c4aa 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/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/projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp index cc7a06d66b..fe9b06713c 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp +++ b/projects/hip/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 */