1 // REQUIRES: clang-driver 2 // REQUIRES: x86-registered-target 3 // REQUIRES: amdgpu-registered-target 4 // UNSUPPORTED: -aix 5 6 // See the steps to create a fat archive are given at the end of the file. 7 8 // Given a FatArchive, clang-offload-bundler should be called to create a 9 // device specific archive, which should be passed to llvm-link. 10 // RUN: %clang -O2 -### -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 %s -L%S/Inputs/openmp_static_device_link -lFatArchive 2>&1 | FileCheck %s 11 // CHECK: clang{{.*}}"-cc1"{{.*}}"-triple" "amdgcn-amd-amdhsa"{{.*}}"-emit-llvm-bc"{{.*}}"-target-cpu" "[[GPU:gfx[0-9]+]]"{{.*}}"-o" "[[HOSTBC:.*.bc]]" "-x" "c++"{{.*}}.cpp 12 // CHECK: clang-offload-bundler" "-unbundle" "-type=a" "-inputs={{.*}}/Inputs/openmp_static_device_link/libFatArchive.a" "-targets=openmp-amdgcn-amd-amdhsa-[[GPU]]" "-outputs=[[DEVICESPECIFICARCHIVE:.*.a]]" "-allow-missing-bundles" 13 // CHECK: llvm-link{{.*}}"[[HOSTBC]]" "[[DEVICESPECIFICARCHIVE]]" "-o" "{{.*}}-[[GPU]]-linked-{{.*}}.bc" 14 // expected-no-diagnostics 15 16 #ifndef HEADER 17 #define HEADER 18 19 #define N 10 20 21 #pragma omp declare target 22 // Functions defined in Fat Archive. 23 extern "C" void func_present(float *, float *, unsigned); 24 25 #ifdef MISSING 26 // Function not defined in the fat archive. 27 extern "C" void func_missing(float *, float *, unsigned); 28 #endif 29 30 #pragma omp end declare target 31 32 int main() { 33 float in[N], out[N], sum = 0; 34 unsigned i; 35 36 #pragma omp parallel for 37 for (i = 0; i < N; ++i) { 38 in[i] = i; 39 } 40 41 func_present(in, out, N); // Returns out[i] = a[i] * 0 42 43 #ifdef MISSING 44 func_missing(in, out, N); // Should throw an error here 45 #endif 46 47 #pragma omp parallel for reduction(+ \ 48 : sum) 49 for (i = 0; i < N; ++i) 50 sum += out[i]; 51 52 if (!sum) 53 return 0; 54 return sum; 55 } 56 57 #endif 58 59 /*********************************************** 60 Steps to create Fat Archive (libFatArchive.a) 61 ************************************************ 62 ***************** File: func_1.c *************** 63 void func_present(float* in, float* out, unsigned n){ 64 unsigned i; 65 #pragma omp target teams distribute parallel for map(to: in[0:n]) map(from: out[0:n]) 66 for(i=0; i<n; ++i){ 67 out[i] = in[i] * 0; 68 } 69 } 70 ************************************************* 71 1. Compile source file(s) to generate object file(s) 72 clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c func_1.c -o func_1_gfx906.o 73 clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 -c func_1.c -o func_1_gfx908.o 74 clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 -c func_2.c -o func_2_gfx906.o 75 clang -O2 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 -c func_2.c -o func_2_gfx908.o 76 clang -O2 -fopenmp -fopenmp-targets=nvptx64 -c func_1.c -o func_1_nvptx.o 77 clang -O2 -fopenmp -fopenmp-targets=nvptx64 -c func_2.c -o func_2_nvptx.o 78 79 2. Create a fat archive by combining all the object file(s) 80 llvm-ar cr libFatArchive.a func_1_gfx906.o func_1_gfx908.o func_2_gfx906.o func_2_gfx908.o func_1_nvptx.o func_2_nvptx.o 81 ************************************************/ 82