187 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _2// Test target codegen - host bc file has to be created first.3// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc4// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK15// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc6// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK27// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK28 9// expected-no-diagnostics10#ifndef HEADER11#define HEADER12 13void work();14 15void use() {16 #pragma omp parallel17 work();18}19 20int main() {21 #pragma omp target parallel22 { use(); }23 #pragma omp target24 { use(); }25}26 27#endif28// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2129// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {30// CHECK1-NEXT: entry:31// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 832// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 833// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 834// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_kernel_environment, ptr [[DYN_PTR]])35// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -136// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]37// CHECK1: user_code.entry:38// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])39// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i64 0)40// CHECK1-NEXT: call void @__kmpc_target_deinit()41// CHECK1-NEXT: ret void42// CHECK1: worker.exit:43// CHECK1-NEXT: ret void44//45//46// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_omp_outlined47// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {48// CHECK1-NEXT: entry:49// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 850// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 851// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 852// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 853// CHECK1-NEXT: call void @_Z3usev() #[[ATTR8:[0-9]+]]54// CHECK1-NEXT: ret void55//56//57// CHECK1-LABEL: define {{[^@]+}}@_Z3usev58// CHECK1-SAME: () #[[ATTR2:[0-9]+]] {59// CHECK1-NEXT: entry:60// CHECK1-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 861// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])62// CHECK1-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, ptr @_Z3usev_omp_outlined, ptr @_Z3usev_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i64 0)63// CHECK1-NEXT: ret void64//65//66// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l2367// CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR5:[0-9]+]] {68// CHECK1-NEXT: entry:69// CHECK1-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 870// CHECK1-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 871// CHECK1-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_kernel_environment, ptr [[DYN_PTR]])72// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -173// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]74// CHECK1: user_code.entry:75// CHECK1-NEXT: call void @_Z3usev() #[[ATTR8]]76// CHECK1-NEXT: call void @__kmpc_target_deinit()77// CHECK1-NEXT: ret void78// CHECK1: worker.exit:79// CHECK1-NEXT: ret void80//81//82// CHECK1-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined83// CHECK1-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {84// CHECK1-NEXT: entry:85// CHECK1-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 886// CHECK1-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 887// CHECK1-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 888// CHECK1-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 889// CHECK1-NEXT: call void @_Z4workv() #[[ATTR8]]90// CHECK1-NEXT: ret void91//92//93// CHECK1-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined_wrapper94// CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR7:[0-9]+]] {95// CHECK1-NEXT: entry:96// CHECK1-NEXT: [[DOTADDR:%.*]] = alloca i16, align 297// CHECK1-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 498// CHECK1-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 499// CHECK1-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8100// CHECK1-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2101// CHECK1-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4102// CHECK1-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4103// CHECK1-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])104// CHECK1-NEXT: call void @_Z3usev_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3:[0-9]+]]105// CHECK1-NEXT: ret void106//107//108// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21109// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {110// CHECK2-NEXT: entry:111// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4112// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4113// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4114// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_kernel_environment, ptr [[DYN_PTR]])115// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1116// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]117// CHECK2: user_code.entry:118// CHECK2-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1:[0-9]+]])119// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP1]], i32 1, i32 -1, i32 -1, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_omp_outlined, ptr null, ptr [[CAPTURED_VARS_ADDRS]], i32 0)120// CHECK2-NEXT: call void @__kmpc_target_deinit()121// CHECK2-NEXT: ret void122// CHECK2: worker.exit:123// CHECK2-NEXT: ret void124//125//126// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_omp_outlined127// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {128// CHECK2-NEXT: entry:129// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4130// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4131// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4132// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4133// CHECK2-NEXT: call void @_Z3usev() #[[ATTR8:[0-9]+]]134// CHECK2-NEXT: ret void135//136//137// CHECK2-LABEL: define {{[^@]+}}@_Z3usev138// CHECK2-SAME: () #[[ATTR2:[0-9]+]] {139// CHECK2-NEXT: entry:140// CHECK2-NEXT: [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4141// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])142// CHECK2-NEXT: call void @__kmpc_parallel_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i32 -1, i32 -1, ptr @_Z3usev_omp_outlined, ptr @_Z3usev_omp_outlined_wrapper, ptr [[CAPTURED_VARS_ADDRS]], i32 0)143// CHECK2-NEXT: ret void144//145//146// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23147// CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR5:[0-9]+]] {148// CHECK2-NEXT: entry:149// CHECK2-NEXT: [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4150// CHECK2-NEXT: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4151// CHECK2-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23_kernel_environment, ptr [[DYN_PTR]])152// CHECK2-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1153// CHECK2-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]154// CHECK2: user_code.entry:155// CHECK2-NEXT: call void @_Z3usev() #[[ATTR8]]156// CHECK2-NEXT: call void @__kmpc_target_deinit()157// CHECK2-NEXT: ret void158// CHECK2: worker.exit:159// CHECK2-NEXT: ret void160//161//162// CHECK2-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined163// CHECK2-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {164// CHECK2-NEXT: entry:165// CHECK2-NEXT: [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 4166// CHECK2-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4167// CHECK2-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4168// CHECK2-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4169// CHECK2-NEXT: call void @_Z4workv() #[[ATTR8]]170// CHECK2-NEXT: ret void171//172//173// CHECK2-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined_wrapper174// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR7:[0-9]+]] {175// CHECK2-NEXT: entry:176// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2177// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4178// CHECK2-NEXT: [[DOTZERO_ADDR:%.*]] = alloca i32, align 4179// CHECK2-NEXT: [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4180// CHECK2-NEXT: store i16 [[TMP0]], ptr [[DOTADDR]], align 2181// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4182// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4183// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])184// CHECK2-NEXT: call void @_Z3usev_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3:[0-9]+]]185// CHECK2-NEXT: ret void186//187