This introduces LipoProteinLipase (lpl), a simple tool for creating fat binaries. It represents a direct replacement of the creaky hccgenco.sh script, which had various issues. The format it uses is that of a code object bundle, generated by the Clang Offload Bundler. The output is always suffixed with the ".adipose" extension. It is shared with HCC. The hipcc script and associated tests are modified to use lpl. Help can be obtained by invoking lpl --help. A more computer-sciency / corporate friendly name is likely to be beneficial, which is a reason for choosing easily searchable/replaceable names such as lpl or adipose.

[ROCm/hip-tests commit: 182156b12b]

projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp

@@ -35,7 +35,7 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN<<2
 
-#define fileName "vcpy_kernel.code"
+#define fileName "vcpy_kernel.code.adipose"
 #define kernel_name "hello_world"
 
 #define HIP_CHECK(status) \
@@ -72,7 +72,6 @@ int main(){
 		uint32_t one = 1;
 
     struct {
-        uint32_t _hidden[6];
         void * _Ad;
         void * _Bd;
     } args;

projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp

@@ -30,7 +30,7 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN<<2
 
-#define fileName "vcpy_kernel.code"
+#define fileName "vcpy_kernel.code.adipose"
 #define kernel_name "hello_world"
 
 #define HIP_CHECK(status) \
@@ -68,7 +68,6 @@ int main(){
 		uint32_t one = 1;
 
     struct {
-        uint32_t _hidden[6];
         void * _Ad;
         void * _Bd;
     } args;

projects/hip-tests/samples/0_Intro/module_api/vcpy_kernel.cpp

@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 
-extern "C" __global__ void hello_world(hipLaunchParm lp, float *a, float *b)
+extern "C" __global__ void hello_world(float *a, float *b)
 {
     int tx = hipThreadIdx_x;
     b[tx] = a[tx];

projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp

@@ -30,7 +30,7 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN*sizeof(float)
 
-#define fileName "vcpy_kernel.code"
+#define fileName "vcpy_kernel.code.adipose"
 float myDeviceGlobal;
 float myDeviceGlobalArray[16];
 #define HIP_CHECK(cmd) \
@@ -80,7 +80,6 @@ int main(){
 		uint32_t one = 1;
 
     struct {
-        uint32_t _hidden[6];  // genco path + wrapper-gen pass use hidden arguments.
         void * _Ad;
         void * _Bd;
     } args;
@@ -111,7 +110,7 @@ int main(){
       HIP_LAUNCH_PARAM_END
     };
 
-    { 
+    {
         hipFunction_t Function;
         HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world"));
         HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
@@ -136,7 +135,7 @@ int main(){
         };
     }
 
-    { 
+    {
         hipFunction_t Function;
         HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals"));
         HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));

projects/hip-tests/samples/0_Intro/module_api_global/vcpy_kernel.cpp

@@ -27,13 +27,13 @@ THE SOFTWARE.
 extern float myDeviceGlobal;
 extern float myDeviceGlobalArray[16];;
 
-extern "C" __global__ void hello_world(hipLaunchParm lp, const float *a, float *b)
+extern "C" __global__ void hello_world(const float *a, float *b)
 {
     int tx = hipThreadIdx_x;
     b[tx] = a[tx];
 }
 
-extern "C" __global__ void test_globals(hipLaunchParm lp, const float *a, float *b)
+extern "C" __global__ void test_globals(const float *a, float *b)
 {
     int tx = hipThreadIdx_x;
     b[tx] = a[tx] + myDeviceGlobal+ myDeviceGlobalArray[tx%ARRAY_SIZE] ;