From 341ef7fdcaa6098a4737240c280015ce40f43bb6 Mon Sep 17 00:00:00 2001 From: Joseph Greathouse Date: Wed, 25 Mar 2020 16:39:24 -0500 Subject: [PATCH] Fix cooperative launch APIs to set hipGetLastError (#1935) * Fix cooperative launch APIs to set hipGetLastError Previously, the cooperative launch APIs did not properly log their errors in the global hipGetLastError variable before returning back to the user. As such, the APIs would leave hipSuccess in the last error, which would break some use cases. This fixes that problem by making a trampoline function that does the HIP_INIT_API and ihipLogStatus. * Add missing flag to the log of multi-GPU launch --- .../hip/hcc_detail/functional_grid_launch.hpp | 31 +++++++------ hipamd/src/hip_module.cpp | 46 +++++++++++++++---- 2 files changed, 54 insertions(+), 23 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp b/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp index 76a04fa355..5abe1095df 100644 --- a/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp @@ -37,14 +37,15 @@ THE SOFTWARE. hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags, hip_impl::program_state& ps); -hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX, void** kernelParams, - unsigned int sharedMemBytes, hipStream_t stream, hip_impl::program_state& ps); - -hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, - unsigned int flags, hip_impl::program_state& ps); - - +hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, + dim3 blockDim, void** args, + size_t sharedMem, hipStream_t stream, + hip_impl::program_state& ps); +hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, + unsigned int flags, + hip_impl::program_state& ps); #pragma GCC visibility push(hidden) @@ -202,22 +203,24 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, template inline __attribute__((visibility("hidden"))) -hipError_t hipLaunchCooperativeKernel(F f, dim3 gridDim, dim3 blockDimX, void** kernelParams, - unsigned int sharedMemBytes, hipStream_t stream) { - +hipError_t hipLaunchCooperativeKernel(F f, dim3 gridDim, dim3 blockDim, + void** args, size_t sharedMem, + hipStream_t stream) { hip_impl::hip_init(); auto& ps = hip_impl::get_program_state(); - return ihipLaunchCooperativeKernel(reinterpret_cast(f), gridDim, blockDimX, kernelParams, sharedMemBytes, stream, ps); + return hipLaunchCooperativeKernel(reinterpret_cast(f), gridDim, + blockDim, args, sharedMem, stream, ps); } inline __attribute__((visibility("hidden"))) -hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, - unsigned int flags) { +hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, + unsigned int flags) { hip_impl::hip_init(); auto& ps = hip_impl::get_program_state(); - return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, ps); + return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, ps); } #pragma GCC visibility pop diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 1130dec26a..794bf868e5 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -495,9 +495,8 @@ __global__ void init_gws(uint nwm1) { } } -__attribute__((visibility("default"))) hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, - dim3 blockDimX, void** kernelParams, unsigned int sharedMemBytes, + dim3 blockDim, void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream, hip_impl::program_state& ps) { #if (__hcc_workweek__ >= 20093) @@ -515,9 +514,9 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, return hipErrorInvalidConfiguration; } - size_t globalWorkSizeX = (size_t)gridDim.x * (size_t)blockDimX.x; - size_t globalWorkSizeY = (size_t)gridDim.y * (size_t)blockDimX.y; - size_t globalWorkSizeZ = (size_t)gridDim.z * (size_t)blockDimX.z; + size_t globalWorkSizeX = (size_t)gridDim.x * (size_t)blockDim.x; + size_t globalWorkSizeY = (size_t)gridDim.y * (size_t)blockDim.y; + size_t globalWorkSizeZ = (size_t)gridDim.z * (size_t)blockDim.z; if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX) { return hipErrorInvalidConfiguration; @@ -603,10 +602,10 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, // launch the main kernel in the cooperative queue result = ihipModuleLaunchKernel(tls, kd, - gridDim.x * blockDimX.x, - gridDim.y * blockDimX.y, - gridDim.z * blockDimX.z, - blockDimX.x, blockDimX.y, blockDimX.z, + gridDim.x * blockDim.x, + gridDim.y * blockDim.y, + gridDim.z * blockDim.z, + blockDim.x, blockDim.y, blockDim.z, sharedMemBytes, stream, kernelParams, nullptr, nullptr, nullptr, 0, true, impCoopParams, &coopAV); @@ -631,6 +630,20 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim, } __attribute__((visibility("default"))) +hipError_t hipLaunchCooperativeKernel(const void* func, dim3 gridDim, + dim3 blockDim, void** args, + size_t sharedMem, hipStream_t stream, + hip_impl::program_state& ps) { + + // Skipping passing in ps, because the logging function does not like it + HIP_INIT_API(hipLaunchCooperativeKernel, func, gridDim, blockDim, args, + sharedMem, stream); + + return ihipLogStatus(ihipLaunchCooperativeKernel(func, gridDim, blockDim, + args, sharedMem, stream, ps)); +} + + hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags, hip_impl::program_state& ps) { @@ -875,6 +888,21 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL #endif } +__attribute__((visibility("default"))) +hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, + unsigned int flags, + hip_impl::program_state& ps) { + + // Skipping passing in ps, because the logging function does not like it + HIP_INIT_API(hipLaunchCooperativeKernelMultiDevice, launchParamsList, + numDevices, flags); + + return ihipLogStatus(ihipLaunchCooperativeKernelMultiDevice(launchParamsList, + numDevices, + flags, ps)); +} + namespace hip_impl { hsa_executable_t executable_for(hipModule_t hmod) { return hmod->executable;