diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 16606e8016..b8952f8a68 100644
--- a/src/hip_module.cpp
+++ b/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;
 }
diff --git a/src/program_state.inl b/src/program_state.inl
index 272addd053..2e959a9928 100644
--- a/src/program_state.inl
+++ b/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)
diff --git a/tests/src/gcc/LaunchKernel.c b/tests/src/gcc/LaunchKernel.c
index 042089a71a..d2fc854510 100644
--- a/tests/src/gcc/LaunchKernel.c
+++ b/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));