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 // CK1: [[MEM_TY:%.+]] = type { [128 x i8] } 31 // CK1-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer 32 // CK1-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null 33 // CK1-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} 4 34 // CK1-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} {{8|4}} 35 // CK1-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1 36 // CK1-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1 37 38 // only nvptx side: do not outline teams region and do not call fork_teams 39 // CK1: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[ARGC:%.+]]) 40 // CK1: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}, 41 // CK1: store {{.+}} 0, {{.+}}, 42 // CK1: store i{{[0-9]+}} [[ARGC]], i{{[0-9]+}}* [[ARGCADDR]], 43 // CK1-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[ARGCADDR]] to i{{[0-9]+}}* 44 // CK1: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED1]], 45 // CK1: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE1]], 46 // CK1: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[IS_SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**)) 47 // CK1: [[KERNEL_RD:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]], 48 // CK1: [[GLOBALSTACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0 49 // CK1-64: [[ARG:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV]] 50 // CK1-32: [[ARG:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ARGCADDR]] 51 // CK1: [[ARGCADDR:%.+]] = getelementptr inbounds %struct.{{.*}}, %struct.{{.*}}* %{{.*}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0 52 // CK1: store i{{[0-9]+}} [[ARG]], i{{[0-9]+}}* [[ARGCADDR]], 53 // CK1: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}}, i32* [[ARGCADDR]]) 54 // CK1: ret void 55 // CK1-NEXT: } 56 57 // CK1: define internal void [[OUTLINED]]( 58 // CK1: store i{{[0-9]+}} 0, i{{[0-9]+}}* % 59 // CK1-NOT: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 60 61 // target region in template 62 // CK1: define {{.*}}void @{{[^,]+}}(i{{.+}}** [[ARGC:%.+]]) 63 // CK1: [[ARGCADDR:%.+]] = alloca i{{.+}}**, 64 // CK1: store i{{.+}}** [[ARGC]], i{{.+}}*** [[ARGCADDR]] 65 // CK1: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED2]], 66 // CK1: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE2]], 67 // CK1: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[IS_SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**)) 68 // CK1: [[KERNEL_RD:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]], 69 // CK1: [[GLOBALSTACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0 70 // CK1: [[ARG:%.+]] = load i{{[0-9]+}}**, i{{[0-9]+}}*** [[ARGCADDR]] 71 // CK1: [[ARGCADDR:%.+]] = getelementptr inbounds %struct.{{.*}}, %struct.{{.*}}* %{{.*}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0 72 // CK1: store i{{[0-9]+}}** [[ARG]], i{{[0-9]+}}*** [[ARGCADDR]], 73 // CK1: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}}, i8*** [[ARGCADDR]]) 74 // CK1: ret void 75 // CK1-NEXT: } 76 77 // CK1: define internal void [[OUTLINED]]( 78 // CK1: store i{{[0-9]+}}** null, i{{[0-9]+}}*** % 79 // CK1-NOT: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 80 81 82 #endif // CK1 83 84 // Test target codegen - host bc file has to be created first. 85 // 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 86 // 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 87 // 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 88 // 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 89 // expected-no-diagnostics 90 #ifdef CK2 91 92 template <typename T> 93 int tmain(T argc) { 94 int a = 10; 95 int b = 5; 96 #pragma omp target 97 #pragma omp teams num_teams(a) thread_limit(b) 98 { 99 argc = 0; 100 } 101 return 0; 102 } 103 104 int main (int argc, char **argv) { 105 int a = 20; 106 int b = 5; 107 #pragma omp target 108 #pragma omp teams num_teams(a) thread_limit(b) 109 { 110 argc = 0; 111 } 112 return tmain(argv); 113 } 114 115 // CK2: [[MEM_TY:%.+]] = type { [128 x i8] } 116 // CK2-DAG: [[SHARED_GLOBAL_RD:@.+]] = common addrspace(3) global [[MEM_TY]] zeroinitializer 117 // CK2-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null 118 // CK2-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} 4 119 // CK2-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} {{8|4}} 120 // CK2-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1 121 // CK2-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1 122 123 // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[ARGC_IN:.+]]) 124 // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}}, 125 // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}}, 126 // CK2: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}, 127 // CK2: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[AADDR]], 128 // CK2: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[BADDR]], 129 // CK2: store i{{[0-9]+}} [[ARGC_IN]], i{{[0-9]+}}* [[ARGCADDR]], 130 // CK2-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32* 131 // CK2-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32* 132 // CK2-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32* 133 // CK2: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED1]], 134 // CK2: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE1]], 135 // CK2: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[IS_SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**)) 136 // CK2: [[KERNEL_RD:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]], 137 // CK2: [[GLOBALSTACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0 138 // CK2-64: [[ARG:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV]] 139 // CK2-32: [[ARG:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ARGCADDR]] 140 // CK2: [[ARGCADDR:%.+]] = getelementptr inbounds %struct.{{.*}}, %struct.{{.*}}* %{{.*}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0 141 // CK2: store i{{[0-9]+}} [[ARG]], i{{[0-9]+}}* [[ARGCADDR]], 142 // CK2: {{%.+}} = call i32 @__kmpc_global_thread_num( 143 // CK2: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}}, i32* [[ARGCADDR]]) 144 // CK2: ret 145 146 // CK2: define internal void [[OUTLINED]]( 147 // CK2: store i{{[0-9]+}} 0, i{{[0-9]+}}* % 148 // CK2-NOT: {{.+}} = call i32 @__kmpc_push_num_teams( 149 // CK2-NOT: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 150 151 // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[BP:%.+]], i{{[0-9]+}}** [[ARGC:%.+]]) 152 // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}}, 153 // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}}, 154 // CK2: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}**, 155 // CK2: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[AADDR]], 156 // CK2: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[BADDR]], 157 // CK2: store i{{[0-9]+}}** [[ARGC]], i{{[0-9]+}}*** [[ARGCADDR]], 158 // CK2: [[IS_SHARED:%.+]] = load i16, i16* [[KERNEL_SHARED2]], 159 // CK2: [[SIZE:%.+]] = load i{{64|32}}, i{{64|32}}* [[KERNEL_SIZE2]], 160 // CK2: call void @__kmpc_get_team_static_memory(i16 0, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY]], [[MEM_TY]] addrspace(3)* [[SHARED_GLOBAL_RD]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} [[SIZE]], i16 [[IS_SHARED]], i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR]] to i8**)) 161 // CK2: [[KERNEL_RD:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]], 162 // CK2: [[GLOBALSTACK:%.+]] = getelementptr inbounds i8, i8* [[KERNEL_RD]], i{{64|32}} 0 163 // CK2: [[ARG:%.+]] = load i{{[0-9]+}}**, i{{[0-9]+}}*** [[ARGCADDR]] 164 // CK2: [[ARGCADDR:%.+]] = getelementptr inbounds %struct.{{.*}}, %struct.{{.*}}* %{{.*}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0 165 // CK2: store i{{[0-9]+}}** [[ARG]], i{{[0-9]+}}*** [[ARGCADDR]], 166 // CK2: {{%.+}} = call i32 @__kmpc_global_thread_num( 167 // CK2: call void [[OUTLINED:@.+]](i32* %{{.+}}, i32* %{{.+}}, i8*** [[ARGCADDR]]) 168 // CK2: ret void 169 170 // CK2: define internal void [[OUTLINED]]( 171 // CK2: store i{{[0-9]+}}** null, i{{[0-9]+}}*** % 172 // CK2-NOT: {{.+}} = call i32 @__kmpc_push_num_teams( 173 // CK2-NOT: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams( 174 175 #endif // CK2 176 #endif 177