711 lines · cpp
1// Test host codegen.2// 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 --check-prefix CHECK-64 --check-prefix OMP453// 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 %s4// 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 --check-prefix CHECK-64 --check-prefix OMP455// 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 --check-prefix CHECK-32 --check-prefix OMP456// 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 %s7// 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 --check-prefix CHECK-32 --check-prefix OMP458// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP519// RUN: %clang_cc1 -fopenmp -DOMP5 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s10// RUN: %clang_cc1 -fopenmp -DOMP5 -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 --check-prefix CHECK-64 --check-prefix OMP5111// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP5112// RUN: %clang_cc1 -fopenmp -DOMP5 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s13// RUN: %clang_cc1 -fopenmp -DOMP5 -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 --check-prefix CHECK-32 --check-prefix OMP5114 15// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s16// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s17// RUN: %clang_cc1 -fopenmp-simd -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 --check-prefix SIMD-ONLY0 %s18// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s19// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s20// RUN: %clang_cc1 -fopenmp-simd -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 --check-prefix SIMD-ONLY0 %s21// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s22// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s23// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s24// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s25// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s26// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s27// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}28 29// Test target codegen - host bc file has to be created first.30// 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.bc31// 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 --check-prefix TCHECK-64 --check-prefix TOMP4532// 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 %s33// 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 --check-prefix TCHECK-64 --check-prefix TOMP4534// 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.bc35// 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 --check-prefix TCHECK-32 --check-prefix TOMP4536// 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 %s37// 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 --check-prefix TCHECK-32 --check-prefix TOMP4538// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc39// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -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 --check-prefix TOMP5140// RUN: %clang_cc1 -fopenmp -DOMP5 -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 %s41// RUN: %clang_cc1 -fopenmp -DOMP5 -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 --check-prefix TOMP5142// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc43// RUN: %clang_cc1 -verify -Wno-vla -fopenmp -DOMP5 -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 --check-prefix TOMP5144// RUN: %clang_cc1 -fopenmp -DOMP5 -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 %s45// RUN: %clang_cc1 -fopenmp -DOMP5 -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 --check-prefix TOMP5146 47// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc48// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -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 --check-prefix SIMD-ONLY1 %s49// RUN: %clang_cc1 -fopenmp-simd -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 %s50// RUN: %clang_cc1 -fopenmp-simd -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 --check-prefix SIMD-ONLY1 %s51// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc52// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -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 --check-prefix SIMD-ONLY1 %s53// RUN: %clang_cc1 -fopenmp-simd -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 %s54// RUN: %clang_cc1 -fopenmp-simd -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 --check-prefix SIMD-ONLY1 %s55// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc56// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -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 %s57// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s58// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s59// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc60// RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -DOMP5 -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 %s61// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s62// RUN: %clang_cc1 -fopenmp-simd -DOMP5 -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 %s63// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}64 65// expected-no-diagnostics66#ifndef HEADER67#define HEADER68 69// CHECK-DAG: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr }70// CHECK-DAG: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%.+]] }71// CHECK-DAG: [[KMP_TASK_T]] = type { ptr, ptr, i32, %{{[^,]+}}, %{{[^,]+}} }72// CHECK-DAG: [[TT:%.+]] = type { i64, i8 }73// CHECK-DAG: [[S1:%.+]] = type { double }74// CHECK-DAG: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }75 76// TCHECK: [[ENTTY:%.+]] = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }77 78// We have 8 target regions, but only 7 that actually will generate offloading79// code, only 6 will have mapped arguments, and only 4 have all-constant map80// sizes.81 82// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [3 x i64] [i64 2, i64 4, i64 4]83// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 800]84// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2]85// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]86// CHECK-DAG: [[SIZET4:@.+]] = private unnamed_addr constant [9 x i64] [i64 4, i64 40, i64 {{8|4}}, i64 0, i64 400, i64 {{8|4}}, i64 {{8|4}}, i64 0, i64 {{16|12}}]87// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547]88// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40]89// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547]90// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40]91// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547]92// OMP45-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [5 x i64] [i64 {{8|4}}, i64 4, i64 {{8|4}}, i64 {{8|4}}, i64 0]93// OMP51-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [6 x i64] [i64 {{8|4}}, i64 4, i64 {{8|4}}, i64 {{8|4}}, i64 0, i64 1]94// OMP45-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547]95// OMP51-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547, i64 800]96// CHECK-DAG: @{{.*}} = weak constant i8 097// CHECK-DAG: @{{.*}} = weak constant i8 098// CHECK-DAG: @{{.*}} = weak constant i8 099// CHECK-DAG: @{{.*}} = weak constant i8 0100// CHECK-DAG: @{{.*}} = weak constant i8 0101// CHECK-DAG: @{{.*}} = weak constant i8 0102// CHECK-DAG: @{{.*}} = weak constant i8 0103 104// TCHECK: @{{.+}} = weak constant [[ENTTY]]105// TCHECK: @{{.+}} = weak constant [[ENTTY]]106// TCHECK: @{{.+}} = weak constant [[ENTTY]]107// TCHECK: @{{.+}} = weak constant [[ENTTY]]108// TCHECK: @{{.+}} = weak constant [[ENTTY]]109// TCHECK: @{{.+}} = weak constant [[ENTTY]]110// TCHECK: @{{.+}} = weak constant [[ENTTY]]111// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]112 113template<typename tx, typename ty>114struct TT{115 tx X;116 ty Y;117};118 119// CHECK-LABEL: get_val120long long get_val() { return 0; }121 122// CHECK: define {{.*}}[[FOO:@.+]](123int foo(int n) {124 int a = 0;125 short aa = 0;126 float b[10];127 float bn[n];128 double c[5][10];129 double cn[5][n];130 TT<long long, char> d;131 132 // CHECK-32: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i32 20, i32 1, ptr [[OMP_TASK_ENTRY:@[^,]+]], i64 -1)133 // CHECK-64: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i64 40, i64 1, ptr [[OMP_TASK_ENTRY:@[^,]+]], i64 -1)134 // CHECK: call i32 @__kmpc_omp_task(ptr @{{[^,]+}}, i32 %{{[^,]+}}, ptr [[TASK]])135 #pragma omp target simd nowait136 for (int i = 3; i < 32; i += 5) {137 }138 139 // CHECK: call void [[HVT1:@.+]](i[[SZ:32|64]] {{[^,]+}}, {{[^)]+}})140 long long k = get_val();141 #pragma omp target simd if(target: 0) linear(k : 3)142 for (int i = 10; i > 1; i--) {143 a += 1;144 }145 146 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])147 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2148 // CHECK-DAG: store ptr [[BP:%.+]], ptr [[BPARG]]149 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3150 // CHECK-DAG: store ptr [[P:%.+]], ptr [[PARG]]151 // CHECK-DAG: [[BP]] = getelementptr inbounds [3 x ptr], ptr [[BPR:%[^,]+]], i32 0, i32 0152 // CHECK-DAG: [[P]] = getelementptr inbounds [3 x ptr], ptr [[PR:%[^,]+]], i32 0, i32 0153 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BPR]], i32 0, i32 0154 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PR]], i32 0, i32 0155 // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],156 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],157 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BPR]], i32 0, i32 1158 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PR]], i32 0, i32 1159 // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],160 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],161 // CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BPR]], i32 0, i32 2162 // CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PR]], i32 0, i32 2163 // CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], ptr [[BPADDR2]],164 // CHECK-DAG: store i[[SZ]] [[VAL2]], ptr [[PADDR2]],165 166 // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0167 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]168 // CHECK: [[FAIL]]169 // CHECK: call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}}, i[[SZ]] {{[^)]+}})170 // CHECK-NEXT: br label %[[END]]171 // CHECK: [[END]]172 int lin = 12;173 #pragma omp target simd if(target: 1) linear(lin, a : get_val())174 for (unsigned long long it = 2000; it >= 600; it-=400) {175 aa += 1;176 }177 178 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10179 // CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]180 // CHECK: [[IFTHEN]]181 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])182 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2183 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]184 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3185 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]186 // CHECK-DAG: [[BPR]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0187 // CHECK-DAG: [[PR]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0188 189 // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0190 // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0191 // CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],192 // CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],193 194 // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1195 // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1196 // CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],197 // CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],198 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0199 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]200 // CHECK: [[FAIL]]201 // CHECK: call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})202 // CHECK-NEXT: br label %[[END]]203 // CHECK: [[END]]204 // CHECK-NEXT: br label %[[IFEND:.+]]205 // CHECK: [[IFELSE]]206 // CHECK: call void [[HVT3]]({{[^,]+}}, {{[^,]+}})207 // CHECK-NEXT: br label %[[IFEND]]208 // CHECK: [[IFEND]]209 210 #pragma omp target simd if(target: n>10)211 for (short it = 6; it <= 20; it-=-4) {212 a += 1;213 aa += 1;214 }215 216 // We capture 3 VLA sizes in this target region217 // CHECK-64: [[A_VAL:%.+]] = load i32, ptr %{{.+}},218 // CHECK-64: store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],219 // CHECK-64: [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],220 221 // CHECK-32: [[A_VAL:%.+]] = load i32, ptr %{{.+}},222 // CHECK-32: store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],223 // CHECK-32: [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],224 225 // CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20226 // CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]227 // CHECK: [[TRY]]228 // CHECK-64: [[BNSIZE:%.+]] = mul nuw i64 [[VLA0:%.+]], 4229 // CHECK-32: [[BNSZSIZE:%.+]] = mul nuw i32 [[VLA0:%.+]], 4230 // CHECK-32: [[BNSIZE:%.+]] = sext i32 [[BNSZSIZE]] to i64231 // CHECK: [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]]232 // CHECK-64: [[CNSIZE:%.+]] = mul nuw i64 [[CNELEMSIZE2]], 8233 // CHECK-32: [[CNSZSIZE:%.+]] = mul nuw i32 [[CNELEMSIZE2]], 8234 // CHECK-32: [[CNSIZE:%.+]] = sext i32 [[CNSZSIZE]] to i64235 236// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])237// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2238// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]239// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3240// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]241// CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4242// CHECK-DAG: store ptr [[SR:%.+]], ptr [[SARG]]243// CHECK-DAG: [[BPR]] = getelementptr inbounds [9 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0244// CHECK-DAG: [[PR]] = getelementptr inbounds [9 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0245// CHECK-DAG: [[SR]] = getelementptr inbounds [9 x i64], ptr [[S:%[^,]+]], i32 0, i32 0246 247// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX0:[0-9]+]]248// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX0]]249// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX1:[0-9]+]]250// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX1]]251// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX2:[0-9]+]]252// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX2]]253// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [9 x i64], ptr [[S]], i32 0, i32 [[IDX3:[0-9]+]]254// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX3]]255// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX3]]256// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX4:[0-9]+]]257// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX4]]258// CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX5:[0-9]+]]259// CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX5]]260// CHECK-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX6:[0-9]+]]261// CHECK-DAG: [[PADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX6]]262// CHECK-DAG: [[SADDR7:%.+]] = getelementptr inbounds [9 x i64], ptr [[S]], i32 0, i32 [[IDX7:[0-9]+]]263// CHECK-DAG: [[BPADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX7]]264// CHECK-DAG: [[PADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX7]]265// CHECK-DAG: [[BPADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX8:[0-9]+]]266// CHECK-DAG: [[PADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX8]]267 268// The names below are not necessarily consistent with the names used for the269// addresses above as some are repeated.270// CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr {{%[^,]+}},271// CHECK-DAG: store i[[SZ]] [[A_CVAL]], ptr {{%[^,]+}},272 273// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},274// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},275 276// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},277// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},278 279// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},280// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},281// CHECK-DAG: store i64 [[BNSIZE]], ptr {{%[^,]+}}282 283// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},284// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},285 286// CHECK-DAG: store i[[SZ]] 5, ptr {{%[^,]+}},287// CHECK-DAG: store i[[SZ]] 5, ptr {{%[^,]+}},288 289// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},290// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},291 292// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},293// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},294// CHECK-DAG: store i64 [[CNSIZE]], ptr {{%[^,]+}}295 296// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},297// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},298 299// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0300// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]301 302// CHECK: [[FAIL]]303// CHECK: call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})304// CHECK-NEXT: br label %[[END]]305// CHECK: [[END]]306#pragma omp target simd if (target \307 : n > 20)308 for (unsigned char it = 'z'; it >= 'a'; it+=-1) {309 a += 1;310 b[2] += 1.0;311 bn[3] += 1.0;312 c[1][2] += 1.0;313 cn[1][3] += 1.0;314 d.X += 1;315 d.Y += 1;316 }317 318 return a;319}320 321// Check that the offloading functions are emitted and that the arguments are322// correct and loaded correctly for the target regions in foo().323 324// CHECK: define internal void [[HVT0:@.+]]()325// CHECK: !llvm.loop326// CHECK: ret void327// CHECK-NEXT: }328 329// CHECK: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)330// CHECK: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])331// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0332// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]333// CHECK: [[FAIL]]334// CHECK: call void [[HVT0]]()335// CHECK-NEXT: br label %[[END]]336// CHECK: [[END]]337 338// CHECK: define internal void [[HVT1]](i[[SZ]] noundef %{{.+}}, {{.+}})339// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align340// CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align341// CHECK-64: [[AA:%.+]] = load i32, ptr [[AA_ADDR]], align342// CHECK-32: [[AA:%.+]] = load i32, ptr [[AA_ADDR]], align343// CHECK: !llvm.access.group344// CHECK: !llvm.loop345// CHECK: ret void346// CHECK-NEXT: }347 348// CHECK: define internal void [[HVT2]](i[[SZ]] noundef %{{.+}}, i[[SZ]] noundef %{{.+}}, i[[SZ]] noundef %{{.+}})349// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align350// CHECK: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align351// CHECK: [[AA:%.+]] = load i16, ptr [[AA_ADDR]], align352// CHECK: !llvm.loop353// CHECK: ret void354// CHECK-NEXT: }355 356// CHECK: define internal void [[HVT3]]357// CHECK: [[A_ADDR:%.+]] = alloca i[[SZ]], align358// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align359// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[A_ADDR]], align360// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align361// CHECK: !llvm.loop362// CHECK: ret void363// CHECK-NEXT: }364 365// CHECK: define internal void [[HVT4]]366// Create local storage for each capture.367// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]368// CHECK: [[LOCAL_B:%.+]] = alloca ptr369// CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]370// CHECK: [[LOCAL_BN:%.+]] = alloca ptr371// CHECK: [[LOCAL_C:%.+]] = alloca ptr372// CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]373// CHECK: [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]374// CHECK: [[LOCAL_CN:%.+]] = alloca ptr375// CHECK: [[LOCAL_D:%.+]] = alloca ptr376// CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]377// CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]378// CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]379// CHECK-DAG: store ptr [[ARG_BN:%.+]], ptr [[LOCAL_BN]]380// CHECK-DAG: store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]381// CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]382// CHECK-DAG: store i[[SZ]] [[ARG_VLA3:%.+]], ptr [[LOCAL_VLA3]]383// CHECK-DAG: store ptr [[ARG_CN:%.+]], ptr [[LOCAL_CN]]384// CHECK-DAG: store ptr [[ARG_D:%.+]], ptr [[LOCAL_D]]385 386// CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],387// CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],388// CHECK-DAG: [[REF_BN:%.+]] = load ptr, ptr [[LOCAL_BN]],389// CHECK-DAG: [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],390// CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],391// CHECK-DAG: [[VAL_VLA3:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA3]],392// CHECK-DAG: [[REF_CN:%.+]] = load ptr, ptr [[LOCAL_CN]],393// CHECK-DAG: [[REF_D:%.+]] = load ptr, ptr [[LOCAL_D]],394 395 396template<typename tx>397tx ftemplate(int n) {398 tx a = 0;399 short aa = 0;400 tx b[10];401 402 #pragma omp target simd if(target: n>40)403 for (long long i = -10; i < 10; i += 3) {404 a += 1;405 aa += 1;406 b[2] += 1;407 }408 409 return a;410}411 412static413int fstatic(int n) {414 int a = 0;415 short aa = 0;416 char aaa = 0;417 int b[10];418 419 #pragma omp target simd if(target: n>50)420 for (unsigned i=100; i<10; i+=10) {421 a += 1;422 aa += 1;423 aaa += 1;424 b[2] += 1;425 }426 427 return a;428}429 430struct S1 {431 double a;432 433 int r1(int n){434 int b = n+1;435 short int c[2][n];436 437#ifdef OMP5438 #pragma omp target simd if(n>60) nontemporal(a) private(a)439#else440 #pragma omp target simd if(n>60) private(a)441#endif // OMP5442 for (unsigned long long it = 2000; it >= 600; it -= 400) {443 this->a = (double)b + 1.5;444 c[1][1] = ++a;445 }446 447 return c[1][1] + (int)b;448 }449};450 451// CHECK: define {{.*}}@{{.*}}bar{{.*}}452int bar(int n){453 int a = 0;454 455 // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})456 a += foo(n);457 458 S1 S;459 // CHECK: call {{.*}}i32 [[FS1:@.+]](ptr {{.*}}, i32 {{.*}})460 a += S.r1(n);461 462 // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})463 a += fstatic(n);464 465 // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})466 a += ftemplate<int>(n);467 468 return a;469}470 471//472// CHECK: define {{.*}}[[FS1]]473//474// CHECK: ptr @llvm.stacksave.p0()475// CHECK-32: store i32 %{{.+}}, ptr %__vla_expr476// OMP51: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60477// CHECK-64: store i32 %{{.+}}, ptr [[B_ADDR:%.+]],478// CHECK-64: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_ADDR]],479 480// CHECK-32: store i32 %{{.+}}, ptr [[B_ADDR:%.+]],481// CHECK-32: [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_ADDR]],482 483// OMP45: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60484// OMP51: [[TOBOOL:%.+]] = trunc i8 %{{.+}} to i1485// OMP51: [[FROMBOOL:%.+]] = zext i1 [[TOBOOL]] to i8486// OMP51: store i8 [[FROMBOOL]], ptr [[CAP:%.+]],487// OMP51: [[SIMD_COND:%.+]] = load i[[SZ]], ptr [[CAP]],488// OMP51: [[IF:%.+]] = trunc i8 %{{.+}} to i1489// CHECK: br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]490// CHECK: [[TRY]]491// We capture 2 VLA sizes in this target region492// CHECK: [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]]493// CHECK-64: [[CSIZE:%.+]] = mul nuw i64 [[CELEMSIZE2]], 2494// CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2495// CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64496 497// OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])498// OMP45-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2499// OMP45-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]500// OMP45-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3501// OMP45-DAG: store ptr [[PR:%.+]], ptr [[PARG]]502// OMP45-DAG: [[BPR]] = getelementptr inbounds [5 x ptr], ptr [[BP:%.+]], i32 0, i32 0503// OMP45-DAG: [[PR]] = getelementptr inbounds [5 x ptr], ptr [[P:%.+]], i32 0, i32 0504// OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX0:[0-9]+]]505// OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX0]]506// OMP45-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX1:[0-9]+]]507// OMP45-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX1]]508// OMP45-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX2:[0-9]+]]509// OMP45-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX2]]510// OMP45-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX3:[0-9]+]]511// OMP45-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX3]]512// OMP45-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], ptr [[S]], i32 [[IDX4:[0-9]+]]513// OMP45-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 [[IDX4]]514// OMP45-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 [[IDX4]]515// OMP51-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])516// OMP51-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2517// OMP51-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]518// OMP51-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3519// OMP51-DAG: store ptr [[PR:%.+]], ptr [[PARG]]520// OMP51-DAG: [[BPR]] = getelementptr inbounds [6 x ptr], ptr [[BP:%.+]], i32 0, i32 0521// OMP51-DAG: [[PR]] = getelementptr inbounds [6 x ptr], ptr [[P:%.+]], i32 0, i32 0522// OMP51-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX0:[0-9]+]]523// OMP51-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX0]]524// OMP51-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX1:[0-9]+]]525// OMP51-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX1]]526// OMP51-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX2:[0-9]+]]527// OMP51-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX2]]528// OMP51-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX3:[0-9]+]]529// OMP51-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX3]]530// OMP51-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], ptr [[S]], i32 [[IDX4:[0-9]+]]531// OMP51-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX4]]532// OMP51-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX4]]533// OMP51-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BP]], i32 [[IDX5:[0-9]+]]534// OMP51-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[P]], i32 [[IDX5]]535 536// The names below are not necessarily consistent with the names used for the537// addresses above as some are repeated.538// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},539// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},540 541// CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr {{%[^,]+}},542// CHECK-DAG: store i[[SZ]] [[B_CVAL]], ptr {{%[^,]+}},543 544// CHECK-DAG: store i[[SZ]] 2, ptr {{%[^,]+}},545// CHECK-DAG: store i[[SZ]] 2, ptr {{%[^,]+}},546 547// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},548// CHECK-DAG: store i[[SZ]] %{{.+}}, ptr {{%[^,]+}},549 550// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},551// CHECK-DAG: store ptr %{{.+}}, ptr {{%[^,]+}},552// CHECK-DAG: store i64 [[CSIZE]], ptr {{%[^,]+}}553 554// OMP51-DAG: store i[[SZ]] [[SIMD_COND]], ptr {{%[^,]+}}555// OMP51-DAG: store i[[SZ]] [[SIMD_COND]], ptr {{%[^,]+}}556 557// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0558// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]559 560// CHECK: [[FAIL]]561// OMP45: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})562// OMP51: call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})563// CHECK-NEXT: br label %[[END]]564// CHECK: [[END]]565 566//567// CHECK: define {{.*}}[[FSTATIC]]568//569// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50570// CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]571// CHECK: [[IFTHEN]]572// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])573// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2574// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]575// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3576// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]577// CHECK-DAG: [[BPR]] = getelementptr inbounds [4 x ptr], ptr [[BP:%.+]], i32 0, i32 0578// CHECK-DAG: [[PR]] = getelementptr inbounds [4 x ptr], ptr [[P:%.+]], i32 0, i32 0579 580// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 0581// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 0582// CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],583// CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],584 585// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 1586// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 1587// CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],588// CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],589 590// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 2591// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 2592// CHECK-DAG: store i[[SZ]] [[VAL2:%.+]], ptr [[BPADDR2]],593// CHECK-DAG: store i[[SZ]] [[VAL2]], ptr [[PADDR2]],594 595// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 3596// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 3597// CHECK-DAG: store ptr [[VAL3:%.+]], ptr [[BPADDR3]],598// CHECK-DAG: store ptr [[VAL3]], ptr [[PADDR3]],599 600// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0601// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]602// CHECK: [[FAIL]]603// CHECK: call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})604// CHECK-NEXT: br label %[[END]]605// CHECK: [[END]]606// CHECK-NEXT: br label %[[IFEND:.+]]607// CHECK: [[IFELSE]]608// CHECK: call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})609// CHECK-NEXT: br label %[[IFEND]]610// CHECK: [[IFEND]]611 612//613// CHECK: define {{.*}}[[FTEMPLATE]]614//615// CHECK: [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40616// CHECK: br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]617// CHECK: [[IFTHEN]]618// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 1, i32 1, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])619// CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2620// CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]621// CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3622// CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]623// CHECK-DAG: [[BPR]] = getelementptr inbounds [3 x ptr], ptr [[BP:%.+]], i32 0, i32 0624// CHECK-DAG: [[PR]] = getelementptr inbounds [3 x ptr], ptr [[P:%.+]], i32 0, i32 0625 626// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 0627// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 0628// CHECK-DAG: store i[[SZ]] [[VAL0:%.+]], ptr [[BPADDR0]],629// CHECK-DAG: store i[[SZ]] [[VAL0]], ptr [[PADDR0]],630 631// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 1632// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 1633// CHECK-DAG: store i[[SZ]] [[VAL1:%.+]], ptr [[BPADDR1]],634// CHECK-DAG: store i[[SZ]] [[VAL1]], ptr [[PADDR1]],635 636// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 2637// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 2638// CHECK-DAG: store ptr [[VAL2:%.+]], ptr [[BPADDR2]],639// CHECK-DAG: store ptr [[VAL2]], ptr [[PADDR2]],640 641// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0642// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]643// CHECK: [[FAIL]]644// CHECK: call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})645// CHECK-NEXT: br label %[[END]]646// CHECK: [[END]]647// CHECK-NEXT: br label %[[IFEND:.+]]648// CHECK: [[IFELSE]]649// CHECK: call void [[HVT:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})650// CHECK-NEXT: br label %[[IFEND]]651// CHECK: [[IFEND]]652 653// Check that the offloading functions are emitted and that the arguments are654// correct and loaded correctly for the target regions of the callees of bar().655 656// CHECK: define internal void [[HVT7]]657// Create local storage for each capture.658// CHECK: [[LOCAL_THIS:%.+]] = alloca ptr659// CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]660// CHECK: [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]661// CHECK: [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]662// CHECK: [[LOCAL_C:%.+]] = alloca ptr663// OMP51: [[LOCAL_SIMD_COND_CASTED:%.+]] = alloca i[[SZ]],664// CHECK-DAG: store ptr [[ARG_THIS:%.+]], ptr [[LOCAL_THIS]]665// CHECK-DAG: store i[[SZ]] [[ARG_B:%.+]], ptr [[LOCAL_B]]666// CHECK-DAG: store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]667// CHECK-DAG: store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]668// CHECK-DAG: store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]669// Store captures in the context.670// CHECK-DAG: [[REF_THIS:%.+]] = load ptr, ptr [[LOCAL_THIS]],671// CHECK-DAG: [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],672// CHECK-DAG: [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],673// CHECK-DAG: [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],674// OMP51-DAG: [[SIMD_COND:%.+]] = load i8, ptr [[LOCAL_SIMD_COND_CASTED]],675// OMP51-DAG: trunc i8 [[SIMD_COND]] to i1676// OMP45-NOT: !nontemporal677// OMP51: store double {{.*}}!nontemporal678// OMP51: load double, {{.*}}!nontemporal679// OMP51: store double {{.*}}!nontemporal680 681// CHECK: define internal void [[HVT6]]682// Create local storage for each capture.683// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]684// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]685// CHECK: [[LOCAL_AAA:%.+]] = alloca i[[SZ]]686// CHECK: [[LOCAL_B:%.+]] = alloca ptr687// CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]688// CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]689// CHECK-DAG: store i[[SZ]] [[ARG_AAA:%.+]], ptr [[LOCAL_AAA]]690// CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]691// Store captures in the context.692// CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],693 694// CHECK: define internal void [[HVT5]]695// Create local storage for each capture.696// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]697// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]698// CHECK: [[LOCAL_B:%.+]] = alloca ptr699// CHECK-DAG: store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]700// CHECK-DAG: store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]701// CHECK-DAG: store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]702// Store captures in the context.703// CHECK-DAG: [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],704 705// OMP45-NOT: !{!"llvm.loop.vectorize.enable", i1 false}706// TOMP45-NOT: !{!"llvm.loop.vectorize.enable", i1 false}707// OMP51: !{!"llvm.loop.vectorize.enable", i1 false}708// TOMP51: !{!"llvm.loop.vectorize.enable", i1 false}709 710#endif711