clang/test/OpenMP/target_reduction_codegen.cpp

   1 // Only test codegen on target side, as private clause does not require any action on the host side
   2 // Test target codegen - host bc file has to be created first.
   3 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
   4 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -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 --check-prefix TCHECK-64
   5 // RUN: %clang_cc1 -fopenmp -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
   6 // RUN: %clang_cc1 -fopenmp -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 --check-prefix TCHECK-64
   7 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
   8 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -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 --check-prefix TCHECK-32
   9 // RUN: %clang_cc1 -fopenmp -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
  10 // RUN: %clang_cc1 -fopenmp -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 --check-prefix TCHECK-32
  11
  12 // 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
  13 // 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-ONLY0 %s
  14 // 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
  15 // 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-ONLY0 %s
  16 // 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
  17 // 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-ONLY0 %s
  18 // 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
  19 // 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-ONLY0 %s
  20 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
  21
  22 // expected-no-diagnostics
  23 #ifndef HEADER
  24 #define HEADER
  25
  26 template<typename tx, typename ty>
  27 struct TT{
  28   tx X;
  29   ty Y;
  30   TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; }
  31 };
  32
  33 // TCHECK: [[S1:%.+]] = type { double }
  34
  35 int foo(int n) {
  36   int a = 0;
  37   short aa = 0;
  38   float b[10];
  39   float bn[n];
  40   double c[5][10];
  41   double cn[5][n];
  42   TT<long long, char> d;
  43
  44   #pragma omp target reduction(*:a)
  45   {
  46   }
  47
  48   // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr{{.+}} %{{.+}})
  49   // TCHECK: [[DYN_PTR:%.+]] = alloca ptr,
  50   // TCHECK: [[A:%.+]] = alloca ptr,
  51   // TCHECK: store {{.+}}, {{.+}} [[A]],
  52   // TCHECK: load ptr, ptr [[A]],
  53   // TCHECK: ret void
  54
  55 #pragma omp target reduction(+:a)
  56   {
  57     a = 1;
  58   }
  59
  60   // TCHECK:  define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr{{.+}} %{{.+}})
  61   // TCHECK: [[DYN_PTR:%.+]] = alloca ptr,
  62   // TCHECK: [[A:%.+]] = alloca ptr,
  63   // TCHECK: store {{.+}}, {{.+}} [[A]],
  64   // TCHECK: [[REF:%.+]] = load ptr, ptr [[A]],
  65   // TCHECK: store i{{[0-9]+}} 1, ptr [[REF]],
  66   // TCHECK: ret void
  67
  68   #pragma omp target reduction(-:a, aa)
  69   {
  70     a = 1;
  71     aa = 1;
  72   }
  73
  74   // TCHECK:  define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr{{.+}} [[A:%.+]], ptr{{.+}} [[AA:%.+]])
  75   // TCHECK: [[DYN_PTR:%.+]] = alloca ptr,
  76   // TCHECK:  [[A:%.+]] = alloca ptr,
  77   // TCHECK:  [[AA:%.+]] = alloca ptr,
  78   // TCHECK: store {{.+}}, {{.+}} [[A]],
  79   // TCHECK: store {{.+}}, {{.+}} [[AA]],
  80   // TCHECK: [[A_REF:%.+]] = load ptr, ptr [[A]],
  81   // TCHECK: [[AA_REF:%.+]] = load ptr, ptr [[AA]],
  82   // TCHECK:  store i{{[0-9]+}} 1, ptr [[A_REF]],
  83   // TCHECK:  store i{{[0-9]+}} 1, ptr [[AA_REF]],
  84   // TCHECK:  ret void
  85
  86   return a;
  87 }
  88
  89
  90 template<typename tx>
  91 tx ftemplate(int n) {
  92   tx a = 0;
  93   short aa = 0;
  94   tx b[10];
  95
  96 #pragma omp target reduction(+:a,aa,b)
  97   {
  98     a = 1;
  99     aa = 1;
 100     b[2] = 1;
 101   }
 102
 103   return a;
 104 }
 105
 106 static
 107 int fstatic(int n) {
 108   int a = 0;
 109   short aa = 0;
 110   char aaa = 0;
 111   int b[10];
 112
 113 #pragma omp target reduction(-:a,aa,aaa,b)
 114   {
 115     a = 1;
 116     aa = 1;
 117     aaa = 1;
 118     b[2] = 1;
 119   }
 120
 121   return a;
 122 }
 123
 124 // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr{{.+}}, ptr{{.+}}, ptr{{.+}}, ptr{{.+}})
 125 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr,
 126 // TCHECK:  [[A:%.+]] = alloca ptr,
 127 // TCHECK:  [[A2:%.+]] = alloca ptr,
 128 // TCHECK:  [[A3:%.+]] = alloca ptr,
 129 // TCHECK:  [[B:%.+]] = alloca ptr,
 130 // TCHECK: store {{.+}}, {{.+}} [[A]],
 131 // TCHECK: store {{.+}}, {{.+}} [[A2]],
 132 // TCHECK: store {{.+}}, {{.+}} [[A3]],
 133 // TCHECK: store {{.+}}, {{.+}} [[B]],
 134 // TCHECK: [[A_REF:%.+]] = load ptr, ptr [[A]],
 135 // TCHECK: [[AA_REF:%.+]] = load ptr, ptr [[AA]],
 136 // TCHECK: [[A3_REF:%.+]] = load ptr, ptr [[A3]],
 137 // TCHECK: [[B_REF:%.+]] = load ptr, ptr [[B]],
 138 // TCHECK:  store i{{[0-9]+}} 1, ptr [[A_REF]],
 139 // TCHECK:  store i{{[0-9]+}} 1, ptr [[AA_REF]],
 140 // TCHECK:  store i{{[0-9]+}} 1, ptr [[A3_REF]],
 141 // TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], ptr [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
 142 // TCHECK:  store i{{[0-9]+}} 1, ptr [[B_GEP]],
 143 // TCHECK:  ret void
 144
 145 struct S1 {
 146   double a;
 147
 148   int r1(int n){
 149     int b = n+1;
 150     short int c[2][n];
 151
 152 #pragma omp target reduction(max:b,c)
 153     {
 154       this->a = (double)b + 1.5;
 155       c[1][1] = ++a;
 156     }
 157
 158     return c[1][1] + (int)b;
 159   }
 160
 161   // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr noundef [[TH:%.+]], ptr{{.+}}, i{{[0-9]+}} noundef [[VLA:%.+]], i{{[0-9]+}} noundef [[VLA1:%.+]], ptr{{.+}})
 162   // TCHECK: [[DYN_PTR:%.+]] = alloca ptr,
 163   // TCHECK: [[TH_ADDR:%.+]] = alloca ptr,
 164   // TCHECK: [[B_ADDR:%.+]] = alloca ptr,
 165   // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
 166   // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
 167   // TCHECK: [[C_ADDR:%.+]] = alloca ptr,
 168   // TCHECK: store ptr [[TH]], ptr [[TH_ADDR]],
 169   // TCHECK: store ptr {{.+}}, ptr [[B_ADDR]],
 170   // TCHECK: store i{{[0-9]+}} [[VLA]], ptr [[VLA_ADDR]],
 171   // TCHECK: store i{{[0-9]+}} [[VLA1]], ptr [[VLA_ADDR2]],
 172   // TCHECK: store ptr {{.+}}, ptr [[C_ADDR]],
 173   // TCHECK: [[TH_ADDR_REF:%.+]] = load ptr, ptr [[TH_ADDR]],
 174   // TCHECK: [[B_REF:%.+]] = load ptr, ptr [[B_ADDR]],
 175   // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR]],
 176   // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, ptr [[VLA_ADDR2]],
 177   // TCHECK: [[C_REF:%.+]] = load ptr, ptr [[C_ADDR]],
 178
 179   // this->a = (double)b + 1.5;
 180   // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, ptr [[B_REF]],
 181   // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
 182   // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
 183   // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], ptr [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
 184   // TCHECK: store double [[NEW_A_VAL]], ptr [[A_FIELD]],
 185
 186   // c[1][1] = ++a;
 187   // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], ptr [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
 188   // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, ptr [[A_FIELD4]],
 189   // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
 190   // TCHECK: store double [[A_FIELD_INC]], ptr [[A_FIELD4]],
 191   // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
 192   // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]]
 193   // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, ptr [[C_REF]], i{{[0-9]+}} [[C_IND]]
 194   // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, ptr [[C_1_REF]], i{{[0-9]+}} 1
 195   // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], ptr [[C_1_1_REF]],
 196
 197   // finish
 198   // TCHECK: ret void
 199 };
 200
 201
 202 int bar(int n){
 203   int a = 0;
 204   a += foo(n);
 205   S1 S;
 206   a += S.r1(n);
 207   a += fstatic(n);
 208   a += ftemplate<int>(n);
 209
 210   return a;
 211 }
 212
 213 // template
 214 // TCHECK: define weak_odr protected void @__omp_offloading_{{.+}}(ptr {{[^,]+}}, ptr{{.+}}, ptr{{.+}}, ptr{{.+}})
 215 // TCHECK: [[DYN_PTR:%.+]] = alloca ptr,
 216 // TCHECK: [[A:%.+]] = alloca ptr,
 217 // TCHECK: [[A2:%.+]] = alloca ptr,
 218 // TCHECK: [[B:%.+]] = alloca ptr,
 219 // TCHECK: store {{.+}}, {{.+}} [[A]],
 220 // TCHECK: store {{.+}}, {{.+}} [[A2]],
 221 // TCHECK: store {{.+}}, {{.+}} [[B]],
 222 // TCHECK: [[A_REF:%.+]] = load ptr, ptr [[A]],
 223 // TCHECK: [[AA_REF:%.+]] = load ptr, ptr [[AA]],
 224 // TCHECK: [[B_REF:%.+]] = load ptr, ptr [[B]],
 225 // TCHECK: store i{{[0-9]+}} 1, ptr [[A_REF]],
 226 // TCHECK: store i{{[0-9]+}} 1, ptr [[AA_REF]],
 227 // TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], ptr [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
 228 // TCHECK: store i{{[0-9]+}} 1, ptr [[B_GEP]],
 229 // TCHECK: ret void
 230
 231 #endif