clang/test/OpenMP/nvptx_multi_target_parallel_codegen.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.bc
   4 // 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=CHECK1
   5 // 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.bc
   6 // 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=CHECK2
   7 // 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=CHECK2
   8
   9 // expected-no-diagnostics
  10 #ifndef HEADER
  11 #define HEADER
  12
  13 void work();
  14
  15 void use() {
  16       #pragma omp parallel
  17       work();
  18 }
  19
  20 int main() {
  21       #pragma omp target parallel
  22       {  use(); }
  23         #pragma omp target
  24         {  use(); }
  25 }
  26
  27 #endif
  28 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21
  29 // CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
  30 // CHECK1-NEXT:  entry:
  31 // CHECK1-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
  32 // CHECK1-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 8
  33 // CHECK1-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
  34 // 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]], -1
  36 // 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 void
  42 // CHECK1:       worker.exit:
  43 // CHECK1-NEXT:    ret void
  44 //
  45 //
  46 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_omp_outlined
  47 // 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 8
  50 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
  51 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
  52 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
  53 // CHECK1-NEXT:    call void @_Z3usev() #[[ATTR8:[0-9]+]]
  54 // CHECK1-NEXT:    ret void
  55 //
  56 //
  57 // CHECK1-LABEL: define {{[^@]+}}@_Z3usev
  58 // CHECK1-SAME: () #[[ATTR2:[0-9]+]] {
  59 // CHECK1-NEXT:  entry:
  60 // CHECK1-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 8
  61 // 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 void
  64 //
  65 //
  66 // CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23
  67 // CHECK1-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR5:[0-9]+]] {
  68 // CHECK1-NEXT:  entry:
  69 // CHECK1-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
  70 // CHECK1-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
  71 // 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]], -1
  73 // 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 void
  78 // CHECK1:       worker.exit:
  79 // CHECK1-NEXT:    ret void
  80 //
  81 //
  82 // CHECK1-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined
  83 // 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 8
  86 // CHECK1-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
  87 // CHECK1-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
  88 // CHECK1-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
  89 // CHECK1-NEXT:    call void @_Z4workv() #[[ATTR8]]
  90 // CHECK1-NEXT:    ret void
  91 //
  92 //
  93 // CHECK1-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined_wrapper
  94 // CHECK1-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR7:[0-9]+]] {
  95 // CHECK1-NEXT:  entry:
  96 // CHECK1-NEXT:    [[DOTADDR:%.*]] = alloca i16, align 2
  97 // CHECK1-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 4
  98 // CHECK1-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
  99 // CHECK1-NEXT:    [[GLOBAL_ARGS:%.*]] = alloca ptr, align 8
 100 // CHECK1-NEXT:    store i16 [[TMP0]], ptr [[DOTADDR]], align 2
 101 // CHECK1-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
 102 // CHECK1-NEXT:    store i32 0, ptr [[DOTZERO_ADDR]], align 4
 103 // 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 void
 106 //
 107 //
 108 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21
 109 // CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
 110 // CHECK2-NEXT:  entry:
 111 // CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
 112 // CHECK2-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
 113 // CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
 114 // 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]], -1
 116 // 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 void
 122 // CHECK2:       worker.exit:
 123 // CHECK2-NEXT:    ret void
 124 //
 125 //
 126 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l21_omp_outlined
 127 // 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 4
 130 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
 131 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
 132 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
 133 // CHECK2-NEXT:    call void @_Z3usev() #[[ATTR8:[0-9]+]]
 134 // CHECK2-NEXT:    ret void
 135 //
 136 //
 137 // CHECK2-LABEL: define {{[^@]+}}@_Z3usev
 138 // CHECK2-SAME: () #[[ATTR2:[0-9]+]] {
 139 // CHECK2-NEXT:  entry:
 140 // CHECK2-NEXT:    [[CAPTURED_VARS_ADDRS:%.*]] = alloca [0 x ptr], align 4
 141 // 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 void
 144 //
 145 //
 146 // CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l23
 147 // CHECK2-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR5:[0-9]+]] {
 148 // CHECK2-NEXT:  entry:
 149 // CHECK2-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
 150 // CHECK2-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
 151 // 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]], -1
 153 // 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 void
 158 // CHECK2:       worker.exit:
 159 // CHECK2-NEXT:    ret void
 160 //
 161 //
 162 // CHECK2-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined
 163 // 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 4
 166 // CHECK2-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 4
 167 // CHECK2-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 4
 168 // CHECK2-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 4
 169 // CHECK2-NEXT:    call void @_Z4workv() #[[ATTR8]]
 170 // CHECK2-NEXT:    ret void
 171 //
 172 //
 173 // CHECK2-LABEL: define {{[^@]+}}@_Z3usev_omp_outlined_wrapper
 174 // CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR7:[0-9]+]] {
 175 // CHECK2-NEXT:  entry:
 176 // CHECK2-NEXT:    [[DOTADDR:%.*]] = alloca i16, align 2
 177 // CHECK2-NEXT:    [[DOTADDR1:%.*]] = alloca i32, align 4
 178 // CHECK2-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
 179 // CHECK2-NEXT:    [[GLOBAL_ARGS:%.*]] = alloca ptr, align 4
 180 // CHECK2-NEXT:    store i16 [[TMP0]], ptr [[DOTADDR]], align 2
 181 // CHECK2-NEXT:    store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
 182 // CHECK2-NEXT:    store i32 0, ptr [[DOTZERO_ADDR]], align 4
 183 // 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 void
 186 //