From df278887feaa1c40596deaab271ca9f9116b5634 Mon Sep 17 00:00:00 2001
From: Saleel Kudchadker <Saleel.Kudchadker@amd.com>
Date: Wed, 14 Sep 2022 18:13:55 -0700
Subject: [PATCH] SWDEV-301667 - Pass all params to API Init

- Correct return for ihipLaunchKernel

Change-Id: I66e81fa635473e6e728936e4cfba9f1bc7f018ad
---
 hipamd/include/hip/amd_detail/hip_prof_str.h |  3 ---
 hipamd/src/hip_module.cpp                    |  9 +++------
 hipamd/src/hip_platform.cpp                  | 10 +++++-----
 3 files changed, 8 insertions(+), 14 deletions(-)

diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h
index 73c2fd124d..9e1aaae028 100644
--- a/hipamd/include/hip/amd_detail/hip_prof_str.h
+++ b/hipamd/include/hip/amd_detail/hip_prof_str.h
@@ -401,7 +401,6 @@ enum hip_api_id_t {
   HIP_API_ID_hipTexRefSetFilterMode = HIP_API_ID_NONE,
   HIP_API_ID_hipTexRefSetMipmapFilterMode = HIP_API_ID_NONE,
   HIP_API_ID_hipUnbindTexture = HIP_API_ID_NONE,
-  HIP_API_ID_ihipModuleLaunchKernel = HIP_API_ID_NONE,
 };
 
 // Return the HIP API string for a given callback ID
@@ -5274,8 +5273,6 @@ typedef struct hip_api_data_s {
 #define INIT_hipTexRefSetMipmapFilterMode_CB_ARGS_DATA(cb_data) {};
 // hipUnbindTexture()
 #define INIT_hipUnbindTexture_CB_ARGS_DATA(cb_data) {};
-// ihipModuleLaunchKernel()
-#define INIT_ihipModuleLaunchKernel_CB_ARGS_DATA(cb_data) {};
 
 #define INIT_NONE_CB_ARGS_DATA(cb_data) {};
 
diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp
index ade9529844..6c771900c2 100644
--- a/hipamd/src/hip_module.cpp
+++ b/hipamd/src/hip_module.cpp
@@ -338,15 +338,11 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                   uint32_t flags = 0, uint32_t params = 0, uint32_t gridId = 0,
                                   uint32_t numGrids = 0, uint64_t prevGridSum = 0,
                                   uint64_t allGridSum = 0, uint32_t firstDevice = 0) {
-  HIP_INIT_API(ihipModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
-               blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra,
-               startEvent, stopEvent, flags, params);
-
   int deviceId = hip::Stream::DeviceId(hStream);
   for (size_t dev = 0; dev < g_devices.size(); ++dev) {
     HIP_RETURN_ONFAIL(PlatformState::instance().initStatManagedVarDevicePtr(dev));
   }
-  
+
   if (f == nullptr) {
     LogPrintfError("%s", "Function passed is null");
     return hipErrorInvalidImage;
@@ -486,7 +482,8 @@ extern "C" hipError_t hipLaunchKernel_spt(const void* hostFunction, dim3 gridDim
 extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim,
                                          void** args, size_t sharedMemBytes, hipStream_t stream,
                                          hipEvent_t startEvent, hipEvent_t stopEvent, int flags) {
-  HIP_INIT_API(hipExtLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream);
+  HIP_INIT_API(hipExtLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes,
+               stream, startEvent, stopEvent, flags);
   STREAM_CAPTURE(hipExtLaunchKernel, stream, hostFunction, gridDim, blockDim, args, sharedMemBytes,
                  startEvent, stopEvent, flags);
   HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream,
diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp
index e0054d3b2b..d92737e9b4 100644
--- a/hipamd/src/hip_platform.cpp
+++ b/hipamd/src/hip_platform.cpp
@@ -616,9 +616,9 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi
   hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId);
   if ((hip_error != hipSuccess) || (func == nullptr)) {
     if (hip_error == hipErrorSharedObjectInitFailed) {
-      HIP_RETURN(hip_error);
+      return hip_error;
     } else {
-      HIP_RETURN(hipErrorInvalidDeviceFunction);
+      return hipErrorInvalidDeviceFunction;
     }
   }
   size_t globalWorkSizeX = static_cast<size_t>(gridDim.x) * blockDim.x;
@@ -627,12 +627,12 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi
   if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
-    HIP_RETURN(hipErrorInvalidConfiguration);
+    return hipErrorInvalidConfiguration;
   }
-  HIP_RETURN(ihipModuleLaunchKernel(
+  return ihipModuleLaunchKernel(
       func, static_cast<uint32_t>(globalWorkSizeX), static_cast<uint32_t>(globalWorkSizeY),
       static_cast<uint32_t>(globalWorkSizeZ), blockDim.x, blockDim.y, blockDim.z, sharedMemBytes,
-      stream, args, nullptr, startEvent, stopEvent, flags));
+      stream, args, nullptr, startEvent, stopEvent, flags);
 }
 
 // conversion routines between float and half precision