From dd60eaa15696c3d94a17fcfc3d5462befe95ea33 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 8 Dec 2017 04:22:57 +0000 Subject: [PATCH] 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: 182156b12b5a529fca9279b9f300375a6171695b] --- .../samples/0_Intro/module_api/launchKernelHcc.cpp | 3 +-- .../hip-tests/samples/0_Intro/module_api/runKernel.cpp | 3 +-- .../hip-tests/samples/0_Intro/module_api/vcpy_kernel.cpp | 2 +- .../samples/0_Intro/module_api_global/runKernel.cpp | 7 +++---- .../samples/0_Intro/module_api_global/vcpy_kernel.cpp | 4 ++-- 5 files changed, 8 insertions(+), 11 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp index 7c90198b6a..eb10f64876 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/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; diff --git a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp index fb34f80b7b..84dbc395a1 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp +++ b/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; diff --git a/projects/hip-tests/samples/0_Intro/module_api/vcpy_kernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/vcpy_kernel.cpp index b728d67c93..a65d3a2e1e 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/vcpy_kernel.cpp +++ b/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]; diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp index 3f84720edf..74f4fb84e2 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp +++ b/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)); diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/vcpy_kernel.cpp b/projects/hip-tests/samples/0_Intro/module_api_global/vcpy_kernel.cpp index f24a261585..8bf67b0fa0 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/vcpy_kernel.cpp +++ b/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] ;