diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index fc6facabf1..9a9655577f 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -4776,6 +4776,12 @@ typedef struct hip_api_data_s { }; // hipLaunchKernel[('const void*', 'function_address'), ('dim3', 'numBlocks'), ('dim3', 'dimBlocks'), ('void**', 'args'), ('size_t', 'sharedMemBytes'), ('hipStream_t', 'stream')] #define INIT_hipLaunchKernel_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLaunchKernel.function_address = (const void*)hostFunction; \ + cb_data.args.hipLaunchKernel.numBlocks = (dim3)gridDim; \ + cb_data.args.hipLaunchKernel.dimBlocks = (dim3)blockDim; \ + cb_data.args.hipLaunchKernel.args = (void**)args; \ + cb_data.args.hipLaunchKernel.sharedMemBytes = (size_t)sharedMemBytes; \ + cb_data.args.hipLaunchKernel.stream = (hipStream_t)stream; \ }; // hipMalloc[('void**', 'ptr'), ('size_t', 'size')] #define INIT_hipMalloc_CB_ARGS_DATA(cb_data) { \ diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index c1bf8523db..71ae362449 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -656,7 +656,7 @@ hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* } - hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim, +hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMemBytes, hipStream_t stream) { STREAM_CAPTURE(hipLaunchKernel, stream, hostFunction, gridDim, blockDim, args, sharedMemBytes); @@ -664,7 +664,7 @@ hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* nullptr, 0); } - hipError_t hipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim, +hipError_t hipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args, size_t sharedMemBytes, hipStream_t stream) { HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); HIP_RETURN(hipLaunchKernel_common(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream));