From d1b75659f28d6a150bd6699dfe7de8d7c1aefb9b Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Mon, 12 Feb 2024 11:36:28 +0000 Subject: [PATCH] SWDEV-445539, SWDEV-445549 - Restore kernel launch prof data. Issue is due to INIT_hipLaunchKernel_CB_ARGS_DATA turns out to be empty when we generate new hip_prof_str.h with since extra space introduced by https://gerrit-git.amd.com/c/compute/ec/clr/+/846275/77/hipamd/src/hip_module.cpp#667 Change-Id: Id251956c1663483cc7ab11a3722769cb5b0d3664 --- hipamd/include/hip/amd_detail/hip_prof_str.h | 6 ++++++ hipamd/src/hip_module.cpp | 4 ++-- 2 files changed, 8 insertions(+), 2 deletions(-) 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));