1 | // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix HOST --check-prefix CHECK |
2 | // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc |
3 | // RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix DEVICE --check-prefix CHECK |
4 | // RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t |
5 | // RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix DEVICE --check-prefix CHECK |
6 | |
7 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o - | FileCheck %s --check-prefix SIMD-ONLY |
8 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc |
9 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY |
10 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t |
11 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix SIMD-ONLY |
12 | |
13 | #ifndef HEADER |
14 | #define HEADER |
15 | |
16 | // SIMD-ONLY-NOT: {{__kmpc|__tgt}} |
17 | |
18 | // DEVICE-DAG: [[C_ADDR:.+]] = internal global i32 0, |
19 | // DEVICE-DAG: [[CD_ADDR:@.+]] = global %struct.S zeroinitializer, |
20 | // HOST-DAG: @[[C_ADDR:.+]] = internal global i32 0, |
21 | // HOST-DAG: @[[CD_ADDR:.+]] = global %struct.S zeroinitializer, |
22 | |
23 | #pragma omp declare target |
24 | int foo() { return 0; } |
25 | #pragma omp end declare target |
26 | int bar() { return 0; } |
27 | #pragma omp declare target (bar) |
28 | int baz() { return 0; } |
29 | |
30 | #pragma omp declare target |
31 | int doo() { return 0; } |
32 | #pragma omp end declare target |
33 | int car() { return 0; } |
34 | #pragma omp declare target (bar) |
35 | int caz() { return 0; } |
36 | |
37 | // DEVICE-DAG: define i32 [[FOO:@.*foo.*]]() |
38 | // DEVICE-DAG: define i32 [[BAR:@.*bar.*]]() |
39 | // DEVICE-DAG: define i32 [[BAZ:@.*baz.*]]() |
40 | // DEVICE-DAG: define i32 [[DOO:@.*doo.*]]() |
41 | // DEVICE-DAG: define i32 [[CAR:@.*car.*]]() |
42 | // DEVICE-DAG: define i32 [[CAZ:@.*caz.*]]() |
43 | |
44 | static int c = foo() + bar() + baz(); |
45 | #pragma omp declare target (c) |
46 | // HOST-DAG: @[[C_CTOR:__omp_offloading__.+_c_l44_ctor]] = private constant i8 0 |
47 | // DEVICE-DAG: define internal void [[C_CTOR:@__omp_offloading__.+_c_l44_ctor]]() |
48 | // DEVICE-DAG: call i32 [[FOO]]() |
49 | // DEVICE-DAG: call i32 [[BAR]]() |
50 | // DEVICE-DAG: call i32 [[BAZ]]() |
51 | // DEVICE-DAG: ret void |
52 | |
53 | struct S { |
54 | int a; |
55 | S() = default; |
56 | S(int a) : a(a) {} |
57 | ~S() { a = 0; } |
58 | }; |
59 | |
60 | #pragma omp declare target |
61 | S cd = doo() + car() + caz() + baz(); |
62 | #pragma omp end declare target |
63 | // HOST-DAG: @[[CD_CTOR:__omp_offloading__.+_cd_l61_ctor]] = private constant i8 0 |
64 | // DEVICE-DAG: define internal void [[CD_CTOR:@__omp_offloading__.+_cd_l61_ctor]]() |
65 | // DEVICE-DAG: call i32 [[DOO]]() |
66 | // DEVICE-DAG: call i32 [[CAR]]() |
67 | // DEVICE-DAG: call i32 [[CAZ]]() |
68 | // DEVICE-DAG: ret void |
69 | |
70 | // HOST-DAG: @[[CD_DTOR:__omp_offloading__.+_cd_l61_dtor]] = private constant i8 0 |
71 | // DEVICE-DAG: define internal void [[CD_DTOR:@__omp_offloading__.+_cd_l61_dtor]]() |
72 | // DEVICE-DAG: call void |
73 | // DEVICE-DAG: ret void |
74 | |
75 | // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_ADDR]]\00" |
76 | // HOST-DAG: @.omp_offloading.entry.[[C_ADDR]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @[[C_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 |
77 | // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_ADDR]]\00" |
78 | // HOST-DAG: @.omp_offloading.entry.[[CD_ADDR]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (%struct.S* @[[CD_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section ".omp_offloading.entries", align 1 |
79 | // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_CTOR]]\00" |
80 | // HOST-DAG: @.omp_offloading.entry.[[C_CTOR]] = weak constant %struct.__tgt_offload_entry { i8* @[[C_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section ".omp_offloading.entries", align 1 |
81 | // HOST-DAG: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_CTOR]]\00" |
82 | // HOST-DAG: @.omp_offloading.entry.[[CD_CTOR]] = weak constant %struct.__tgt_offload_entry { i8* @[[CD_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section ".omp_offloading.entries", align 1 |
83 | // HOST-DAG: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_DTOR]]\00" |
84 | // HOST-DAG: @.omp_offloading.entry.[[CD_DTOR]] = weak constant %struct.__tgt_offload_entry { i8* @[[CD_DTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 4, i32 0 }, section ".omp_offloading.entries", align 1 |
85 | int maini1() { |
86 | int a; |
87 | #pragma omp target map(tofrom : a) |
88 | { |
89 | a = c; |
90 | } |
91 | return 0; |
92 | } |
93 | |
94 | // DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}} |
95 | // DEVICE: [[C:%.+]] = load i32, i32* [[C_ADDR]], |
96 | // DEVICE: store i32 [[C]], i32* % |
97 | |
98 | // HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-11]](i32* dereferenceable{{.*}}) |
99 | // HOST: [[C:%.*]] = load i32, i32* @[[C_ADDR]], |
100 | // HOST: store i32 [[C]], i32* % |
101 | |
102 | // HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}} |
103 | // HOST-DAG: !{i32 1, !"[[C_ADDR]]", i32 0, i32 {{[0-9]+}}} |
104 | |
105 | // DEVICE: !nvvm.annotations |
106 | // DEVICE-DAG: !{void ()* [[C_CTOR]], !"kernel", i32 1} |
107 | // DEVICE-DAG: !{void ()* [[CD_CTOR]], !"kernel", i32 1} |
108 | // DEVICE-DAG: !{void ()* [[CD_DTOR]], !"kernel", i32 1} |
109 | |
110 | #endif // HEADER |
111 | |
112 | |