1 // Test target codegen - host bc file has to be created first. 2 // RUN: %clang_cc1 -DCK1 -verify -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 -DCK1 -verify -fopenmp -x c++ -triple nvptx64-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 CK1 --check-prefix CK1-64 4 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 5 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 6 // expected-no-diagnostics 7 #ifndef HEADER 8 #define HEADER 9 10 #ifdef CK1 11 12 template <typename T> 13 int tmain(T argc) { 14 #pragma omp target 15 #pragma omp teams 16 argc = 0; 17 return 0; 18 } 19 20 21 int main (int argc, char **argv) { 22 #pragma omp target 23 #pragma omp teams 24 { 25 argc = 0; 26 } 27 return tmain(argv); 28 } 29 30 // only nvptx side: do not outline teams region and do not call fork_teams 31 // CK1: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[ARGC:%.+]]) 32 // CK1: {{.+}} = alloca i{{[0-9]+}}*, 33 // CK1: {{.+}} = alloca i{{[0-9]+}}*, 34 // CK1: [[ARGCADDR_PTR:%.+]] = alloca i{{[0-9]+}}*, 35 // CK1: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}, 36 // CK1: store {{.+}} 0, {{.+}}, 37 // CK1: store i{{[0-9]+}} [[ARGC]], i{{[0-9]+}}* [[ARGCADDR]], 38 // CK1-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[ARGCADDR]] to i{{[0-9]+}}* 39 // CK1-64: store i{{[0-9]+}}* [[CONV]], i{{[0-9]+}}** [[ARGCADDR_PTR]], 40 // CK1-32: store i{{[0-9]+}}* [[ARGCADDR]], i{{[0-9]+}}** [[ARGCADDR_PTR]], 41 // CK1: [[ARGCADDR_PTR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[ARGCADDR_PTR]], 42 // CK1: store i{{[0-9]+}} 0, i{{[0-9]+}}* [[ARGCADDR_PTR_REF]], 43 // CK1-NOT: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 44 // CK1: ret void 45 // CK1-NEXT: } 46 47 // target region in template 48 // CK1: define {{.*}}void @{{[^,]+}}(i{{.+}}** [[ARGC:%.+]]) 49 // CK1: [[ARGCADDR_PTR:%.+]] = alloca i{{.+}}***, 50 // CK1: [[ARGCADDR:%.+]] = alloca i{{.+}}**, 51 // CK1: store i{{.+}}** [[ARGC]], i{{.+}}*** [[ARGCADDR]] 52 // CK1: store i8*** [[ARGCADDR]], i8**** [[ARGCADDR_PTR]], 53 // CK1: [[ARGCADDR_PTR_REF:%.+]] = load i{{.+}}**, i{{.+}}*** [[ARGCADDR_PTR]], 54 // CK1: store i{{[0-9]+}}** null, i{{[0-9]+}}*** [[ARGCADDR_PTR_REF]], 55 // CK1-NOT: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 56 // CK1: ret void 57 // CK1-NEXT: } 58 59 60 #endif // CK1 61 62 // Test target codegen - host bc file has to be created first. 63 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 64 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple nvptx64-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 CK2 --check-prefix CK2-64 65 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 66 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 67 // expected-no-diagnostics 68 #ifdef CK2 69 70 template <typename T> 71 int tmain(T argc) { 72 int a = 10; 73 int b = 5; 74 #pragma omp target 75 #pragma omp teams num_teams(a) thread_limit(b) 76 { 77 argc = 0; 78 } 79 return 0; 80 } 81 82 int main (int argc, char **argv) { 83 int a = 20; 84 int b = 5; 85 #pragma omp target 86 #pragma omp teams num_teams(a) thread_limit(b) 87 { 88 argc = 0; 89 } 90 return tmain(argv); 91 } 92 93 // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[ARGC_IN:.+]]) 94 // CK2: {{.}} = alloca i{{[0-9]+}}*, 95 // CK2: {{.}} = alloca i{{[0-9]+}}*, 96 // CK2: [[ARGCADDR_PTR:%.+]] = alloca i{{[0-9]+}}*, 97 // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}}, 98 // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}}, 99 // CK2: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}, 100 // CK2-NOT: {{%.+}} = call i32 @__kmpc_global_thread_num( 101 // CK2: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[AADDR]], 102 // CK2: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[BADDR]], 103 // CK2: store i{{[0-9]+}} [[ARGC_IN]], i{{[0-9]+}}* [[ARGCADDR]], 104 // CK2-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32* 105 // CK2-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32* 106 // CK2-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32* 107 // CK2-64: store i{{[0-9]+}}* [[CONV]], i{{[0-9]+}}** [[ARGCADDR_PTR]], 108 // CK2-32: store i{{[0-9]+}}* [[ARGCADDR]], i{{[0-9]+}}** [[ARGCADDR_PTR]], 109 // CK2: [[ARGCADDR_PTR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[ARGCADDR_PTR]], 110 // CK2: store i{{[0-9]+}} 0, i{{[0-9]+}}* [[ARGCADDR_PTR_REF]], 111 // CK2-NOT: {{.+}} = call i32 @__kmpc_push_num_teams( 112 // CK2-NOT: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 113 // CK2: ret 114 115 // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[BP:%.+]], i{{[0-9]+}}** [[ARGC:%.+]]) 116 // CK2: [[ARGCADDR_PTR:%.+]] = alloca i{{[0-9]+}}***, 117 // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}}, 118 // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}}, 119 // CK2: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}**, 120 // CK2-NOT: {{%.+}} = call i32 @__kmpc_global_thread_num( 121 // CK2: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[AADDR]], 122 // CK2: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[BADDR]], 123 // CK2: store i{{[0-9]+}}** [[ARGC]], i{{[0-9]+}}*** [[ARGCADDR]], 124 // CK2: store i{{[0-9]+}}*** [[ARGCADDR]], i{{[0-9]+}}**** [[ARGCADDR_PTR]], 125 // CK2: [[ARGCADDR_PTR_REF:%.+]] = load i{{[0-9]+}}***, i{{[0-9]+}}**** [[ARGCADDR_PTR]], 126 // CK2: store i{{[0-9]+}}** null, i{{[0-9]+}}*** [[ARGCADDR_PTR_REF]], 127 // CK2-NOT: {{.+}} = call i32 @__kmpc_push_num_teams( 128 // CK2-NOT: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 129 // CK2: ret void 130 131 #endif // CK2 132 #endif 133
