[HIP] Reclaiming hipLaunchKernel API (#1353)

* [HIP] Reclaiming hipLaunchKernel API * Reclaiming hipLaunchKernel : Incorporated review comments * Incorporated review comments * Removed hipLaunchKernel Macro from nvcc path [ROCm/hip commit: 5c4f78bac3]
2019-08-29 06:32:41 +05:30
@@ -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");
+}
@@ -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*);
@@ -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"