[HIP] Reclaiming hipLaunchKernel API (#1353)

* [HIP] Reclaiming hipLaunchKernel API

* Reclaiming hipLaunchKernel : Incorporated review comments

* Incorporated review comments

* Removed hipLaunchKernel Macro from nvcc path

hipamd/include/hip/hcc_detail/functional_grid_launch.hpp

@@ -177,13 +177,4 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
                                      stream, &config[0]);
 }
 
-template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
-[[deprecated("hipLaunchKernel is deprecated and will be removed in the next "
-             "version of HIP; please upgrade to hipLaunchKernelGGL.")]]
-inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
-                            std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
-    hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
-                       std::move(args)...);
-}
-
 #pragma GCC visibility pop

hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -1520,6 +1520,34 @@ hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind,
 } // Namespace hip_impl.
 #endif
 
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/**
+ * @brief C compliant kernel launch API
+ *
+ * @param [in] function_address - kernel function pointer.
+ * @param [in] numBlocks - number of blocks
+ * @param [in] dimBlocks - dimension of a block
+ * @param [in] args - kernel arguments
+ * @param [in] sharedMemBytes - Amount of dynamic shared memory to allocate for this kernel.  The
+ *  Kernel can access this with HIP_DYNAMIC_SHARED.
+ * @param [in] stream - Stream where the kernel should be dispatched.  May be 0, in which case th
+ *  default stream is used with associated synchronization rules.
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
+ *
+ */
+
+hipError_t  hipLaunchKernel(const void* function_address,
+			    dim3 numBlocks, dim3 dimBlocks, void** args,
+				size_t sharedMemBytes, hipStream_t stream);
+
+#ifdef __cplusplus
+}
+#endif
+
 #if defined(__cplusplus)
 extern "C" {
 #endif

hipamd/include/hip/nvcc_detail/hip_runtime.h

@@ -31,11 +31,6 @@ THE SOFTWARE.
 
 typedef int hipLaunchParm;
 
-#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...)             \
-    do {                                                                                           \
-        kernelName<<<numblocks, numthreads, memperblock, streamId>>>(0, ##__VA_ARGS__);            \
-    } while (0)
-
 #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...)          \
     do {                                                                                           \
         kernelName<<<numblocks, numthreads, memperblock, streamId>>>(__VA_ARGS__);                 \

hipamd/include/hip/nvcc_detail/hip_runtime_api.h

@@ -1241,6 +1241,13 @@ inline static hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* im
         cuModuleLoadDataEx(module, image, numOptions, options, optionValues));
 }
 
+inline static hipError_t hipLaunchKernel(const void* function_address, dim3 numBlocks,
+					 dim3 dimBlocks, void** args, size_t sharedMemBytes,
+					 hipStream_t stream)
+{
+   return hipCUDAErrorTohipError(cudaLaunchKernel(function_address,numBlocks,dimBlocks,args,sharedMemBytes,stream));
+}
+
 inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX,
                                                unsigned int gridDimY, unsigned int gridDimZ,
                                                unsigned int blockDimX, unsigned int blockDimY,