Skip to content

Commit b56d4f2

Browse files
dbudanov-cmplrdbudanov-cmplr
committed
Revert "Delete obsolete Clang::Driver tests"
This reverts commit 84be9c2.
1 parent 84be9c2 commit b56d4f2

File tree

4 files changed

+263
-0
lines changed

4 files changed

+263
-0
lines changed

‎clang/test/Driver/amdgpu-openmp-system-arch.c

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
// REQUIRES: system-linux
2+
// REQUIRES: x86-registered-target
3+
// REQUIRES: amdgpu-registered-target
4+
// REQUIRES: shell
5+
6+
// RUN: mkdir -p %t
7+
// RUN: rm -f %t/amdgpu_arch_gfx906
8+
// RUN: cp %S/Inputs/amdgpu-arch/amdgpu_arch_gfx906 %t/
9+
// RUN: cp %S/Inputs/amdgpu-arch/amdgpu_arch_gfx908_gfx908 %t/
10+
// RUN: chmod +x %t/amdgpu_arch_gfx906
11+
// RUN: chmod +x %t/amdgpu_arch_gfx908_gfx908
12+
13+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -nogpulib --amdgpu-arch-tool=%t/amdgpu_arch_gfx906 %s 2>&1 \
14+
// RUN: | FileCheck %s
15+
// CHECK: "-cc1"{{.*}}"-triple" "amdgcn-amd-amdhsa"{{.*}}"-target-cpu" "[[GFX:gfx906]]"
16+
// CHECK: llvm-link{{.*}}"-o" "{{.*}}amdgpu-openmp-system-arch-{{.*}}-[[GFX]]-linked-{{.*}}.bc"
17+
// CHECK: llc{{.*}}amdgpu-openmp-system-arch-{{.*}}-[[GFX]]-linked-{{.*}}.bc" "-mtriple=amdgcn-amd-amdhsa" "-mcpu=[[GFX]]" "-filetype=obj" "-o"{{.*}}amdgpu-openmp-system-arch-{{.*}}-[[GFX]]-{{.*}}.o"
18+
19+
// case when amdgpu_arch returns multiple gpus but of same arch
20+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -nogpulib --amdgpu-arch-tool=%t/amdgpu_arch_gfx908_gfx908 %s 2>&1 \
21+
// RUN: | FileCheck %s --check-prefix=CHECK-MULTIPLE
22+
// CHECK-MULTIPLE: "-cc1"{{.*}}"-triple" "amdgcn-amd-amdhsa"{{.*}}"-target-cpu" "[[GFX:gfx908]]"
23+
// CHECK-MULTIPLE: llvm-link{{.*}}"-o" "{{.*}}amdgpu-openmp-system-arch-{{.*}}-[[GFX]]-linked-{{.*}}.bc"
24+
// CHECK-MULTIPLE: llc{{.*}}amdgpu-openmp-system-arch-{{.*}}-[[GFX]]-linked-{{.*}}.bc" "-mtriple=amdgcn-amd-amdhsa" "-mcpu=[[GFX]]" "-filetype=obj" "-o"{{.*}}amdgpu-openmp-system-arch-{{.*}}-[[GFX]]-{{.*}}.o"

‎clang/test/Driver/amdgpu-openmp-toolchain.c

Lines changed: 79 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,79 @@
1+
// REQUIRES: x86-registered-target
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 --libomptarget-amdgpu-bc-path=%S/Inputs/hip_dev_lib %s 2>&1 \
4+
// RUN: | FileCheck %s
5+
6+
// verify the tools invocations
7+
// CHECK: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-x" "c"
8+
// CHECK: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-x" "ir"
9+
// CHECK: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-emit-llvm-bc"{{.*}}"-target-cpu" "gfx906" "-fcuda-is-device"{{.*}}"-mlink-builtin-bitcode"{{.*}}libomptarget-amdgpu-gfx906.bc"{{.*}}
10+
// CHECK: llvm-link{{.*}}"-o" "{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906-linked-{{.*}}.bc"
11+
// CHECK: llc{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906-linked-{{.*}}.bc" "-mtriple=amdgcn-amd-amdhsa" "-mcpu=gfx906" "-filetype=obj" "-o"{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906-{{.*}}.o"
12+
// CHECK: lld{{.*}}"-flavor" "gnu" "--no-undefined" "-shared" "-o"{{.*}}amdgpu-openmp-toolchain-{{.*}}.out" "{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906-{{.*}}.o"
13+
// CHECK: clang-offload-wrapper{{.*}}"-host" "x86_64-unknown-linux-gnu" "-o" "{{.*}}a-{{.*}}.bc" {{.*}}amdgpu-openmp-toolchain-{{.*}}.out"
14+
// CHECK: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-o" "{{.*}}a-{{.*}}.o" "-x" "ir" "{{.*}}a-{{.*}}.bc"
15+
// CHECK: ld{{.*}}"-o" "a.out"{{.*}}"{{.*}}amdgpu-openmp-toolchain-{{.*}}.o" "{{.*}}a-{{.*}}.o" "-lomp" "-lomptarget"
16+
17+
// RUN: %clang -ccc-print-phases --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 %s 2>&1 \
18+
// RUN: | FileCheck --check-prefix=CHECK-PHASES %s
19+
// phases
20+
// CHECK-PHASES: 0: input, "{{.*}}amdgpu-openmp-toolchain.c", c, (host-openmp)
21+
// CHECK-PHASES: 1: preprocessor, {0}, cpp-output, (host-openmp)
22+
// CHECK-PHASES: 2: compiler, {1}, ir, (host-openmp)
23+
// CHECK-PHASES: 3: backend, {2}, assembler, (host-openmp)
24+
// CHECK-PHASES: 4: assembler, {3}, object, (host-openmp)
25+
// CHECK-PHASES: 5: input, "{{.*}}amdgpu-openmp-toolchain.c", c, (device-openmp)
26+
// CHECK-PHASES: 6: preprocessor, {5}, cpp-output, (device-openmp)
27+
// CHECK-PHASES: 7: compiler, {6}, ir, (device-openmp)
28+
// CHECK-PHASES: 8: offload, "host-openmp (x86_64-unknown-linux-gnu)" {2}, "device-openmp (amdgcn-amd-amdhsa)" {7}, ir
29+
// CHECK-PHASES: 9: backend, {8}, assembler, (device-openmp)
30+
// CHECK-PHASES: 10: assembler, {9}, object, (device-openmp)
31+
// CHECK-PHASES: 11: linker, {10}, image, (device-openmp)
32+
// CHECK-PHASES: 12: offload, "device-openmp (amdgcn-amd-amdhsa)" {11}, image
33+
// CHECK-PHASES: 13: clang-offload-wrapper, {12}, ir, (host-openmp)
34+
// CHECK-PHASES: 14: backend, {13}, assembler, (host-openmp)
35+
// CHECK-PHASES: 15: assembler, {14}, object, (host-openmp)
36+
// CHECK-PHASES: 16: linker, {4, 15}, image, (host-openmp)
37+
38+
// handling of --libomptarget-amdgpu-bc-path
39+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx803 --libomptarget-amdgpu-bc-path=%S/Inputs/hip_dev_lib/libomptarget-amdgpu-gfx803.bc %s 2>&1 | FileCheck %s --check-prefix=CHECK-LIBOMPTARGET
40+
// CHECK-LIBOMPTARGET: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-target-cpu" "gfx803" "-fcuda-is-device" "-mlink-builtin-bitcode"{{.*}}Inputs/hip_dev_lib/libomptarget-amdgpu-gfx803.bc"{{.*}}
41+
42+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx803 -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-NOGPULIB
43+
// CHECK-NOGPULIB-NOT: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-target-cpu" "gfx803" "-fcuda-is-device" "-mlink-builtin-bitcode"{{.*}}libomptarget-amdgpu-gfx803.bc"{{.*}}
44+
45+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -ccc-print-bindings -save-temps -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx803 -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-PRINT-BINDINGS
46+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT:.*]]"],
47+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "clang",{{.*}} output: "[[HOST_BC:.*]]"
48+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[HOST_BC]]"], output: "[[HOST_S:.*]]"
49+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "clang::as", inputs: ["[[HOST_S]]"], output: "[[HOST_O:.*]]"
50+
// CHECK-PRINT-BINDINGS: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[INPUT]]"], output: "[[DEVICE_I:.*]]"
51+
// CHECK-PRINT-BINDINGS: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[DEVICE_I]]", "[[HOST_BC]]"], output: "[[DEVICE_BC:.*]]"
52+
// CHECK-PRINT-BINDINGS: "amdgcn-amd-amdhsa" - "AMDGCN::OpenMPLinker", inputs: ["[[DEVICE_BC]]"], output: "[[DEVICE_OUT:.*]]"
53+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "offload wrapper", inputs: ["[[DEVICE_OUT]]"], output: "[[OFFLOAD_WRAPPER:.*]]"
54+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[OFFLOAD_WRAPPER]]"], output: "[[OFFLOAD_S:.*]]"
55+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "clang::as", inputs: ["[[OFFLOAD_S]]"], output: "[[OFFLOAD_O:.*]]"
56+
// CHECK-PRINT-BINDINGS: "x86_64-unknown-linux-gnu" - "GNU::Linker", inputs: ["[[HOST_O]]", "[[OFFLOAD_O]]"], output:
57+
58+
// verify the llc is invoked for textual assembly output
59+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 --libomptarget-amdgpu-bc-path=%S/Inputs/hip_dev_lib -save-temps %s 2>&1 \
60+
// RUN: | FileCheck %s --check-prefix=CHECK-SAVE-ASM
61+
// CHECK-SAVE-ASM: llc{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906-linked.bc" "-mtriple=amdgcn-amd-amdhsa" "-mcpu=gfx906" "-filetype=asm" "-o"{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906.s"
62+
// CHECK-SAVE-ASM: llc{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906-linked.bc" "-mtriple=amdgcn-amd-amdhsa" "-mcpu=gfx906" "-filetype=obj" "-o"{{.*}}amdgpu-openmp-toolchain-{{.*}}-gfx906.o"
63+
64+
// check the handling of -c
65+
// RUN: %clang -ccc-print-bindings -c --target=x86_64-unknown-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 --libomptarget-amdgpu-bc-path=%S/Inputs/hip_dev_lib -save-temps %s 2>&1 \
66+
// RUN: | FileCheck %s --check-prefix=CHECK-C
67+
// CHECK-C: "x86_64-unknown-linux-gnu" - "clang",
68+
// CHECK-C: "x86_64-unknown-linux-gnu" - "clang",{{.*}}output: "[[HOST_BC:.*]]"
69+
// CHECK-C: "amdgcn-amd-amdhsa" - "clang",{{.*}}output: "[[DEVICE_I:.*]]"
70+
// CHECK-C: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[DEVICE_I]]", "[[HOST_BC]]"]
71+
// CHECK-C: "x86_64-unknown-linux-gnu" - "clang"
72+
// CHECK-C: "x86_64-unknown-linux-gnu" - "clang::as"
73+
// CHECK-C: "x86_64-unknown-linux-gnu" - "offload bundler"
74+
75+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -emit-llvm -S -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx803 -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-EMIT-LLVM-IR
76+
// CHECK-EMIT-LLVM-IR: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-emit-llvm"
77+
78+
// RUN: %clang -### -target x86_64-pc-linux-gnu -fopenmp -fno-openmp-new-driver -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx803 -lm --rocm-device-lib-path=%S/Inputs/rocm/amdgcn/bitcode %s 2>&1 | FileCheck %s --check-prefix=CHECK-LIB-DEVICE
79+
// CHECK-LIB-DEVICE: {{.*}}llvm-link{{.*}}ocml.bc"{{.*}}ockl.bc"{{.*}}oclc_daz_opt_on.bc"{{.*}}oclc_unsafe_math_off.bc"{{.*}}oclc_finite_only_off.bc"{{.*}}oclc_correctly_rounded_sqrt_on.bc"{{.*}}oclc_wavefrontsize64_on.bc"{{.*}}oclc_isa_version_803.bc"

‎clang/test/Driver/fat_archive_amdgpu.cpp

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

‎clang/test/Driver/fat_archive_nvptx.cpp

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

0 commit comments

Comments
 (0)