Fix C-style hipLaunchKernel (#1835)

* Fix bug in LaunchKernel test
Instead of passing the address of the gpu buffer, pass the address
of the pointer that holds the address of the gpu buffer

* Fix hipLaunchKernel's kernarg buffer construction.
The hipLaunchKernel implementation should rely on ihipModuleLaunchKernel
to construct the kernarg buffer correctly based on kernel metadata.

* Fix a bug in get_functions where the Kernel_descriptor wasn't constructed with the correct kernarg layout information.

* Fix a bug in kernarg layout parsing dealing with kernel without any arg

* teach ihipModuleLaunchKernel to handle kernel without any arg

* Add a more interesting test

[ROCm/clr commit: 14e235378f]

projects/clr/hipamd/tests/src/gcc/gpu.cpp

@@ -19,6 +19,7 @@
 
 
 #include<hip/hip_runtime.h>
+#include "LaunchKernel.h"
 
 extern "C"
 {
@@ -43,4 +44,9 @@ __global__ void kernel3(int *a, int*b, int* c)
   *c = *a+*b;
 }
 
+__global__ void kernel4(int *a, char c, short s, int i, struct things t)
+{
+   *a = c + s + i + t.c + t.s + t.i;
+}
+
 }//extern "C"