[HIP] Reclaiming hipLaunchKernel API

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

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

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__);                 \

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,

src/hip_module.cpp

@@ -1144,3 +1144,31 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
     return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
         tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
 }
+
+hipError_t hipLaunchKernel(
+    const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args,
+    size_t sharedMemBytes, hipStream_t stream)
+{
+   HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream);
+
+   hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr,
+                                                           hip_impl::target_agent(stream));
+
+   if(kd == nullptr || kd->_header == nullptr)
+       return ihipLogStatus(hipErrorInvalidValue);
+
+   size_t szKernArg = kd->_header->kernarg_segment_byte_size;
+
+   if(args == NULL && szKernArg != 0)
+      return ihipLogStatus(hipErrorInvalidValue);
+
+   void* config[]{
+        HIP_LAUNCH_PARAM_BUFFER_POINTER,
+        args,
+        HIP_LAUNCH_PARAM_BUFFER_SIZE,
+	    &szKernArg,
+        HIP_LAUNCH_PARAM_END};
+
+   return ihipLogStatus(ihipModuleLaunchKernel(tls, kd, numBlocks.x * dimBlocks.x, numBlocks.y * dimBlocks.y, numBlocks.z * dimBlocks.z,
+                          dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, (void**)&config, nullptr, nullptr, 0));
+}

tests/src/gcc/LaunchKernel.c

@@ -0,0 +1,165 @@
+
+/* Copyright (c) 2019-Present Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+
+/* HIT_START
+ * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: launchkernel.o %cc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvcc
+ * HIT_END
+ */
+
+
+#include "hip/hip_runtime.h"
+#include <stdio.h>
+#include "LaunchKernel.h"
+
+#define HIPCHECK(error)                                                                            \
+    {                                                                                              \
+        hipError_t localError = error;                                                             \
+        if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) {      \
+            printf("%serror: '%s'(%d) from %s at %s:%d%s\n", "\x1B[31m", hipGetErrorString(localError),  \
+                   localError, #error, __FILE__, __LINE__, "\x1B[0m");                             \
+            return false;	                                                                       \
+        }                                                                                          \
+    }
+
+
+bool LaunchKernelArg()
+{
+  dim3 blocks 	    = {1,1,1};
+  dim3 threads      = {1,1,1};
+
+  HIPCHECK(hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0));
+
+  return true;
+}
+
+bool LaunchKernelArg1()
+{
+  int A = 0;
+  int *A_d = NULL;
+  dim3 blocks       = {1,1,1};
+  dim3 threads      = {1,1,1};
+
+  // Allocate Device memory
+  HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
+
+  void* Args[]={A_d};
+  HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args,0,0));
+
+  // Get the result back to host memory
+  HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
+
+  HIPCHECK(hipFree(A_d));
+
+  if(A != 333)
+	return false;
+
+  return true;
+}
+
+bool LaunchKernelArg2()
+{
+  int A = 0;
+  int B = 123;
+  int *A_d = NULL;
+  int *B_d = NULL;
+
+  dim3 blocks       = {1,1,1};
+  dim3 threads      = {1,1,1};
+
+  // Allocate Device memory
+  HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
+
+  HIPCHECK(hipMalloc((void**)&B_d, sizeof(int)));
+
+  // Copy data from host memory to device memory
+  HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));
+
+  void* Args[]={A_d,B_d};
+  HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0));
+
+  // Get the result back to host memory
+  HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
+
+  HIPCHECK(hipFree(A_d));
+  HIPCHECK(hipFree(B_d));
+
+  if(A != 123)
+    return false;
+
+  return true;
+}
+
+bool LaunchKernelArg3()
+{
+  int A = 321;
+  int B = 123;
+  int C = 0;
+  int *A_d = NULL;
+  int *B_d = NULL;
+  int *C_d = NULL;
+
+  dim3 blocks       = {1,1,1};
+  dim3 threads      = {1,1,1};
+
+  // Allocate Device memory
+  HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
+
+  HIPCHECK(hipMalloc((void**)&B_d, sizeof(int)));
+
+  HIPCHECK(hipMalloc((void**)&C_d, sizeof(int)));
+
+  // Copy data from host memory to device memory
+  HIPCHECK(hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice));
+
+  HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));
+
+  void* Args[]={A_d,B_d,C_d};
+  HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0));
+
+  // Get the result back to host memory
+  HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost));
+
+  HIPCHECK(hipFree(A_d));
+  HIPCHECK(hipFree(B_d));
+  HIPCHECK(hipFree(C_d));
+
+  if(C != 444)
+    return false;
+
+  return true;
+}
+
+
+int main()
+{
+
+  if( LaunchKernelArg()  &&
+      LaunchKernelArg1() &&
+      LaunchKernelArg2() &&
+      LaunchKernelArg3())
+    {
+      printf("PASSED!\n");
+    }
+  else
+    printf("FAILED\n");
+}

tests/src/gcc/LaunchKernel.h

@@ -0,0 +1,23 @@
+
+/* Copyright (c) 2019-Present Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+extern __global__ void kernel();
+extern __global__ void kernel1(int*);
+extern __global__ void kernel2(int*,int*);
+extern __global__ void kernel3(int*,int*,int*);

tests/src/gcc/gpu.cpp

@@ -0,0 +1,46 @@
+/*
+ * Copyright (c) 2019-Present Advanced Micro Devices, Inc. All rights reserved.
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ * */
+
+
+#include<hip/hip_runtime.h>
+
+extern "C"
+{
+
+__global__ void kernel()
+{
+   int a = 20;
+}
+
+__global__ void kernel1(int *a)
+{
+   *a = 333;
+}
+
+__global__ void kernel2(int *a, int*b)
+{
+   *a = *b;
+}
+
+__global__ void kernel3(int *a, int*b, int* c)
+{
+  *c = *a+*b;
+}
+
+}//extern "C"