From b4343df485992a4da0b02d7d33acc71d6c74b4e9 Mon Sep 17 00:00:00 2001
From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com>
Date: Wed, 28 Jun 2023 21:52:36 +0530
Subject: [PATCH] SWDEV-400015 - Improve documentation of extra arg in
 hipModuleLaunchKernel (#3236)

Change-Id: Ie6ba5a94af6ef43c647c85842fd3561e4dc6a6ac

[ROCm/hip commit: 97ac320bd3bcdc6265b623fbd11055e1f888fcc1]
---
 projects/hip/include/hip/hip_ext.h         | 5 +++--
 projects/hip/include/hip/hip_runtime_api.h | 6 +++---
 2 files changed, 6 insertions(+), 5 deletions(-)

diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h
index 2548cbe1e5..09e9cf8b95 100644
--- a/projects/hip/include/hip/hip_ext.h
+++ b/projects/hip/include/hip/hip_ext.h
@@ -49,6 +49,9 @@ THE SOFTWARE.
  * @param [in] kernelParams  pointer to kernel parameters.
  * @param [in] extra  Pointer to kernel arguments. These are passed directly to the kernel and
  * must be in the memory layout and alignment expected by the kernel.
+ * All passed arguments must be naturally aligned according to their type. The memory address of each
+ * argument should be a multiple of its size in bytes. Please refer to hip_porting_driver_api.md
+ * for sample usage.
  * @param [in] startEvent  If non-null, specified event will be updated to track the start time of
  * the kernel launch. The event must be created before calling this API.
  * @param [in] stopEvent  If non-null, specified event will be updated to track the stop time of
@@ -57,8 +60,6 @@ THE SOFTWARE.
  * launched in any order.
  * @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
  *
- * @warning  kernellParams argument is not yet implemented in HIP, use extra instead.
- * Please refer to hip_porting_driver_api.md for sample usage.
  * HIP/ROCm actually updates the start event when the associated kernel completes.
  * Currently, timing between startEvent and stopEvent does not include the time it takes to perform
  * a system scope release/cache flush - only the time it takes to issues writes to cache.
diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h
index a1f3822366..75326d0d70 100644
--- a/projects/hip/include/hip/hip_runtime_api.h
+++ b/projects/hip/include/hip/hip_runtime_api.h
@@ -5042,15 +5042,15 @@ hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned
  * @param [in] kernelParams
  * @param [in] extra     Pointer to kernel arguments.   These are passed directly to the kernel and
  * must be in the memory layout and alignment expected by the kernel.
+ * All passed arguments must be naturally aligned according to their type. The memory address of each
+ * argument should be a multiple of its size in bytes. Please refer to hip_porting_driver_api.md
+ * for sample usage.
  *
  * Please note, HIP does not support kernel launch with total work items defined in dimension with
  * size gridDim x blockDim >= 2^32. So gridDim.x * blockDim.x, gridDim.y * blockDim.y
  * and gridDim.z * blockDim.z are always less than 2^32.
  *
  * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
- *
- * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please
- * refer to hip_porting_driver_api.md for sample usage.
  */
 hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY,
                                  unsigned int gridDimZ, unsigned int blockDimX,