clang/test/OpenMP/target_parallel_generic_loop_depend_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 host codegen.
   3 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
   4 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
   5 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK
   6 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
   7 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
   8 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK
   9
  10 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  11 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  12 // RUN: %clang_cc1  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  13 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  14 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  15 // RUN: %clang_cc1  -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  16
  17 // Test target codegen - host bc file has to be created first.
  18 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  19 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK
  20 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  21 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK
  22 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  23 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK
  24 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  25 // RUN: %clang_cc1  -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK
  26
  27 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  28 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  29 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  30 // RUN: %clang_cc1  -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  31 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  32 // RUN: %clang_cc1  -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  33 // RUN: %clang_cc1  -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  34 // RUN: %clang_cc1  -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  35
  36 // expected-no-diagnostics
  37 #ifndef HEADER
  38 #define HEADER
  39
  40
  41
  42
  43
  44 // Check target registration is registered as a Ctor.
  45
  46
  47 template<typename tx, typename ty>
  48 struct TT{
  49   tx X;
  50   ty Y;
  51 };
  52
  53 int global;
  54 extern int global;
  55
  56 int foo(int n) {
  57   int a = 0;
  58   short aa = 0;
  59   float b[10];
  60   float bn[n];
  61   double c[5][10];
  62   double cn[5][n];
  63   TT<long long, char> d;
  64   static long *plocal;
  65
  66   #pragma omp target parallel loop device(global + a) depend(in: global) depend(out: a, b, cn[4])
  67   for (int i = 0; i < 10; ++i) {
  68   }
  69
  70
  71
  72
  73
  74
  75
  76   #pragma omp target parallel loop device(global + a) nowait depend(inout: global, a, bn) if(target:a)
  77   for (int i = 0; i < *plocal; ++i) {
  78     static int local1;
  79     *plocal = global;
  80     local1 = global;
  81   }
  82
  83   #pragma omp target parallel loop if(0) firstprivate(global) depend(out:global)
  84   for (int i = 0; i < global; ++i) {
  85     global += 1;
  86   }
  87
  88   return a;
  89 }
  90
  91 // Check that the offloading functions are emitted and that the arguments are
  92 // correct and loaded correctly for the target regions in foo().
  93
  94
  95
  96
  97
  98 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
  99 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 100 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
 101
 102 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 103 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 104 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
 105
 106 // Create stack storage and store argument in there.
 107 // CHECK-64:    [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i32*
 108 // CHECK-64:    load i32, i32* [[AA_CADDR]], align
 109 // CHECK-32:    load i32, i32* [[AA_ADDR]], align
 110
 111 // CHECK-64:    [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32*
 112 // CHECK-64:    store i32 [[BP1_I32]], i32* [[BP1_CAST]],
 113 // CHECK-32:    store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]],
 114
 115
 116 #endif
 117 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l66
 118 // CHECK-SAME: () #[[ATTR2:[0-9]+]] {
 119 // CHECK-NEXT:  entry:
 120 // CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB2:[0-9]+]], i32 0, ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3fooi_l66.omp_outlined)
 121 // CHECK-NEXT:    ret void
 122 //
 123 //
 124 //
 125 //
 126 //
 127 //
 128 //
 129 //
 130 //
 131 //
 132 //
 133 //
 134 //
 135 //
 136 //
 137 //
 138 //
 139 //
 140 //
 141 //
 142 //
 143 //
 144 //
 145 //
 146 //
 147 //
 148 //
 149 //
 150 //
 151 //
 152 //
 153 //
 154 //
 155 //
 156 //
 157 //
 158 //
 159 //
 160 //
 161 //
 162 //
 163 //
 164 //
 165 //
 166 //
 167 //
 168 //
 169 //
 170 //
 171 //
 172 //
 173 //
 174 //
 175 //
 176 //
 177 //
 178 //
 179 //
 180 //
 181 //
 182 //
 183 //
 184 //
 185 //
 186 //
 187 //
 188 //
 189 //
 190 //
 191 //
 192 //
 193 //
 194 //
 195 //
 196 //
 197 //
 198 //// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
 199 // SIMD-ONLY0: {{.*}}
 200 // SIMD-ONLY1: {{.*}}
 201 // TCHECK: {{.*}}