Support defaultDriver test with Clang+Hcc Runtime

Fix the following issues:
1.Ignore hidden arguments of kernel functions.
2.Look up both origial function name and function name with .kd postfix
when argments are retrived from module.
3.Addition, fix compiling issue of LaunchKernel test app.

Change-Id: I9400943f2f02433cb4409b19c0cac3626c2bc454


[ROCm/hip commit: 9a005eacf3]

projects/hip/src/hip_module.cpp

@@ -1030,19 +1030,29 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t
     if (!*func) return hipErrorInvalidValue;
 
     std::string name_str(name);
+    std::string namekd_str(name_str + ".kd");
+
     auto kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent);
 
     if (kernel.handle == 0u) {
-        name_str.append(".kd");
-        kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent);
+        kernel = find_kernel_by_name(hmod->executable, namekd_str.c_str(), agent);
     }
 
     if (kernel.handle == 0u) return hipErrorNotFound;
 
+    auto it = hmod->kernargs.find(name_str); //Look up args from the original name
+    if (it == hmod->kernargs.end()) {
+        it = hmod->kernargs.find(namekd_str); //Look up args from .kd name
+        if (it == hmod->kernargs.end()) {
+            return hipErrorNotFound;
+        }
+    }
+
     // TODO: refactor the whole ihipThisThat, which is a mess and yields the
     //       below, due to hipFunction_t being a pointer to ihipModuleSymbol_t.
+
     func[0][0] = *static_cast<hipFunction_t>(
-        Kernel_descriptor{kernel_object(kernel), name_str, hmod->kernargs[name_str]});
+        Kernel_descriptor{kernel_object(kernel), name_str, it->second});
 
     return hipSuccess;
 }

projects/hip/src/program_state.inl

@@ -714,6 +714,27 @@ public:
                 != AMD_COMGR_STATUS_SUCCESS)
                 return;
 
+            //Look up “.value_kind” to decide whether to ignore it
+            //See http://llvm.org/docs/AMDGPUUsage.html#code-object-v3-metadata-mattr-code-object-v3
+            amd_comgr_metadata_node_t arg_value_kind_md;
+            if (amd_comgr_metadata_lookup(arg_md, ".value_kind", &arg_value_kind_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
+
+            std::string arg_value_kind{ metadata_to_string(arg_value_kind_md) };
+
+            if (amd_comgr_destroy_metadata(arg_value_kind_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
+
+            if (arg_value_kind.find("hidden_") == 0) {
+                if (amd_comgr_destroy_metadata(arg_md)
+                    != AMD_COMGR_STATUS_SUCCESS)
+                    return;
+
+                continue; //Ignore hidden arg
+            }
+
             amd_comgr_metadata_node_t arg_size_md;
             if (amd_comgr_metadata_lookup(arg_md, ".size", &arg_size_md)
                 != AMD_COMGR_STATUS_SUCCESS)

projects/hip/tests/src/gcc/LaunchKernel.c

@@ -20,7 +20,7 @@
 
 /* HIT_START
  * BUILD_CMD: gpu.o %hc -I%hip-path/include -g -c %S/gpu.cpp -o %T/gpu.o EXCLUDE_HIP_PLATFORM nvcc vdi
- * 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 vdi
+ * BUILD_CMD: launchkernel.o %hc -D__HIP_PLATFORM_HCC__ -g -I%hip-path/include -c %S/LaunchKernel.c -o %T/launchkernel.o EXCLUDE_HIP_PLATFORM nvcc vdi
  * 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 vdi
  * TEST: %t EXCLUDE_HIP_PLATFORM nvcc vdi
  * HIT_END
@@ -36,7 +36,7 @@ bool LaunchKernelArg()
   dim3 blocks 	    = {1,1,1};
   dim3 threads      = {1,1,1};
 
-  HIPCHECK(hipLaunchKernel(kernel, blocks, threads, NULL, 0, 0));
+  HIPCHECK(hipLaunchKernel((const void *)kernel, 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(kernel1, blocks, threads, Args, 0, 0));
+  HIPCHECK(hipLaunchKernel((const void *)kernel1, 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(kernel2, blocks, threads, Args,0,0));
+  HIPCHECK(hipLaunchKernel((const void *)kernel2, 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(kernel3, blocks, threads, Args,0,0));
+  HIPCHECK(hipLaunchKernel((const void *)kernel3, 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(kernel4, blocks, threads, Args, 0, 0));
+  HIPCHECK(hipLaunchKernel((const void *)kernel4, blocks, threads, Args, 0, 0));
 
   // Get the result back to host memory
   HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));