SWDEV-272769 - Enabled gcc based hip test

Change-Id: I48bc850df00848bf11de71e3f614273111d31b47

hipamd/tests/src/gcc/LaunchKernel.c

@@ -19,10 +19,10 @@
 
 
 /* HIT_START
- * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
- * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_AMD__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
- * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
- * TEST: %t EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
+ * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_AMD__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD_CMD: LaunchKernel %hc %T/launchkernel.o %T/gpu.o -g -Wl,--rpath=%hip-path/lib %hip-path/lib/libamdhip64.so -o %T/%t DEPENDS gpu.o launchkernel.o EXCLUDE_HIP_PLATFORM nvidia
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 
@@ -36,7 +36,7 @@ bool LaunchKernelArg()
   dim3 blocks 	    = {1,1,1};
   dim3 threads      = {1,1,1};
 
-  HIPCHECK(hipLaunchKernel((const void *)kernel, blocks, threads, NULL, 0, 0));
+  HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel), blocks, threads, NULL, 0, 0));
 
   return true;
 }
@@ -52,7 +52,7 @@ bool LaunchKernelArg1()
   HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
  
   void* Args[]={&A_d};
-  HIPCHECK(hipLaunchKernel((const void *)kernel1, blocks, threads, Args, 0, 0));
+  HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel1), blocks, threads, Args, 0, 0));
 
   // Get the result back to host memory
   HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
@@ -84,7 +84,7 @@ bool LaunchKernelArg2()
   HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice));
 
   void* Args[]={&A_d, &B_d};
-  HIPCHECK(hipLaunchKernel((const void *)kernel2, blocks, threads, Args,0,0));
+  HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel2), blocks, threads, Args,0,0));
 
   // Get the result back to host memory
   HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
@@ -123,7 +123,7 @@ bool LaunchKernelArg3()
   HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice));
 
   void* Args[]={&A_d, &B_d, &C_d};
-  HIPCHECK(hipLaunchKernel((const void *)kernel3, blocks, threads, Args,0,0));
+  HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel3), blocks, threads, Args,0,0));
 
   // Get the result back to host memory
   HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost));
@@ -154,7 +154,7 @@ bool LaunchKernelArg4()
   struct things t = {2,20,200};
   
   void* Args[]={&A_d, &c, &s, &i, &t};
-  HIPCHECK(hipLaunchKernel((const void *)kernel4, blocks, threads, Args, 0, 0));
+  HIPCHECK(hipLaunchKernel(getKernelFunc(mykernel4), blocks, threads, Args, 0, 0));
 
   // Get the result back to host memory
   HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
@@ -167,7 +167,6 @@ bool LaunchKernelArg4()
   return true;
 }
 
-
 int main()
 {
   if( LaunchKernelArg()  &&

hipamd/tests/src/gcc/LaunchKernel.h

@@ -22,17 +22,21 @@ extern "C"
 {
 #endif
 
-extern __global__ void kernel();
-extern __global__ void kernel1(int*);
-extern __global__ void kernel2(int*,int*);
-extern __global__ void kernel3(int*,int*,int*);
-
 struct things {
     char c;
     short s;
     int i;
 };
-extern __global__ void kernel4(int*, char, short, int, struct things);
+
+typedef enum func{
+   mykernel,
+   mykernel1,
+   mykernel2,
+   mykernel3,
+   mykernel4
+}func;
+
+extern const void* getKernelFunc(enum func f);
 
 #ifdef __cplusplus
 }

hipamd/tests/src/gcc/gpu.cpp

@@ -49,4 +49,15 @@ __global__ void kernel4(int *a, char c, short s, int i, struct things t)
    *a = c + s + i + t.c + t.s + t.i;
 }
 
+const void* funcTable[] = {
+   (const void*)kernel,
+   (const void*)kernel1,
+   (const void*)kernel2,
+   (const void*)kernel3,
+   (const void*)kernel4 };
+
+const void* getKernelFunc(enum func f){
+    return funcTable[f];
+ }
+
 }//extern "C"