[OpenMP] Update target directive codegen to use 4.5 implicit data mappings.
Summary:
This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed.
Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available.
Reviewers: hfinkel, rjmccall, ABataev
Subscribers: cfe-commits, carlo.bertolli
Differential Revision: http://reviews.llvm.org/D14940
llvm-svn: 254521
diff --git a/clang/test/OpenMP/target_codegen_global_capture.cpp b/clang/test/OpenMP/target_codegen_global_capture.cpp
index af6b9ef..29469cd 100644
--- a/clang/test/OpenMP/target_codegen_global_capture.cpp
+++ b/clang/test/OpenMP/target_codegen_global_capture.cpp
@@ -1,9 +1,9 @@
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
@@ -41,27 +41,88 @@
static float Sc = 7.0;
static float Sd = 8.0;
- // CHECK-DAG: [[REFB:%.+]] = bitcast i16* [[LB]] to i8*
- // CHECK-DAG: store i8* [[REFB]], i8** [[GEPB:%.+]], align
- // CHECK-DAG: [[REFC:%.+]] = bitcast i16* [[LC]] to i8*
- // CHECK-DAG: store i8* [[REFC]], i8** [[GEPC:%.+]], align
- // CHECK-DAG: [[REFD:%.+]] = bitcast i16* [[LD]] to i8*
- // CHECK-DAG: store i8* [[REFD]], i8** [[GEPD:%.+]], align
- // CHECK-DAG: store i8* bitcast (double* [[GB]] to i8*), i8** [[GEPGB:%.+]], align
- // CHECK-DAG: store i8* bitcast (double* [[GC]] to i8*), i8** [[GEPGC:%.+]], align
- // CHECK-DAG: store i8* bitcast (double* [[GD]] to i8*), i8** [[GEPGD:%.+]], align
- // CHECK-DAG: store i8* bitcast (float* [[FB]] to i8*), i8** [[GEPFB:%.+]], align
- // CHECK-DAG: store i8* bitcast (float* [[FC]] to i8*), i8** [[GEPFC:%.+]], align
- // CHECK-DAG: store i8* bitcast (float* [[FD]] to i8*), i8** [[GEPFD:%.+]], align
- // CHECK-DAG: [[GEPB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
+ // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LB]],
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3foossssE2Sb,
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
+ // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LC]],
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3foossssE2Sc,
+ // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LD]],
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3foossssE2Sd,
+
+ // 3 local vars being captured.
+
+ // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
+ // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
+ // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
+ // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]],
+ // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
+ // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
+ // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
+ // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]],
+ // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
+ // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
+ // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
+ // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]],
+ // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static vars being captured.
+
+ // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
+ // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
+ // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
+ // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]],
+ // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
+ // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
+ // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
+ // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]],
+ // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
+ // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
+ // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
+ // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]],
+ // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static global vars being captured.
+
+ // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
+ // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
+ // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]],
+ // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
+ // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
+ // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]],
+ // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
+ // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
+ // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]],
+ // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
// CHECK: call i32 @__tgt_target
// CHECK: call void [[OFFLOADF:@.+]](
// Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
@@ -71,7 +132,7 @@
Gb += 1.0;
Sb += 1.0;
- // CHECK: define internal void [[OFFLOADF]]({{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}})
+ // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
// The parallel region only uses 3 captures.
// CHECK: call {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
// CHECK: call void @.omp_outlined.(i32* %{{.+}}, i32* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}}, {{.+}}* %{{.+}})
@@ -106,45 +167,98 @@
// CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
// Capture a, b, c, d
+ // CHECK: [[ALLOCLA:%.+]] = alloca i16
+ // CHECK: [[ALLOCLB:%.+]] = alloca i16
+ // CHECK: [[ALLOCLC:%.+]] = alloca i16
+ // CHECK: [[ALLOCLD:%.+]] = alloca i16
+ // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
+ // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
+ // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
+ // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
#pragma omp parallel
{
- // CHECK: [[ADRA:%.+]] = alloca i16*, align
- // CHECK: [[ADRB:%.+]] = alloca i16*, align
- // CHECK: [[ADRC:%.+]] = alloca i16*, align
- // CHECK: [[ADRD:%.+]] = alloca i16*, align
- // CHECK: store i16* [[A]], i16** [[ADRA]], align
- // CHECK: store i16* [[B]], i16** [[ADRB]], align
- // CHECK: store i16* [[C]], i16** [[ADRC]], align
- // CHECK: store i16* [[D]], i16** [[ADRD]], align
- // CHECK: [[REFA:%.+]] = load i16*, i16** [[ADRA]],
- // CHECK: [[REFB:%.+]] = load i16*, i16** [[ADRB]],
- // CHECK: [[REFC:%.+]] = load i16*, i16** [[ADRC]],
- // CHECK: [[REFD:%.+]] = load i16*, i16** [[ADRD]],
+ // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3barssssE2Sb,
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
+ // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3barssssE2Sc,
+ // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3barssssE2Sd,
- // CHECK: load float, float* [[BA]]
+ // 3 local vars being captured.
- // CHECK-DAG: [[CSTB:%.+]] = bitcast i16* [[REFB]] to i8*
- // CHECK-DAG: [[CSTC:%.+]] = bitcast i16* [[REFC]] to i8*
- // CHECK-DAG: [[CSTD:%.+]] = bitcast i16* [[REFD]] to i8*
- // CHECK-DAG: store i8* [[CSTB]], i8** [[GEPB:%.+]], align
- // CHECK-DAG: store i8* [[CSTC]], i8** [[GEPC:%.+]], align
- // CHECK-DAG: store i8* [[CSTD]], i8** [[GEPD:%.+]], align
- // CHECK-DAG: store i8* bitcast (double* [[GB]] to i8*), i8** [[GEPGB:%.+]], align
- // CHECK-DAG: store i8* bitcast (double* [[GC]] to i8*), i8** [[GEPGC:%.+]], align
- // CHECK-DAG: store i8* bitcast (double* [[GD]] to i8*), i8** [[GEPGD:%.+]], align
- // CHECK-DAG: store i8* bitcast (float* [[BB]] to i8*), i8** [[GEPBB:%.+]], align
- // CHECK-DAG: store i8* bitcast (float* [[BC]] to i8*), i8** [[GEPBC:%.+]], align
- // CHECK-DAG: store i8* bitcast (float* [[BD]] to i8*), i8** [[GEPBD:%.+]], align
+ // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
+ // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
+ // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
+ // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]],
+ // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
- // CHECK-DAG: [[GEPB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPBB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPBC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
- // CHECK-DAG: [[GEPBD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{.+}}
+ // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
+ // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
+ // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
+ // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]],
+ // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
+ // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
+ // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
+ // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]],
+ // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static vars being captured.
+
+ // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
+ // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
+ // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
+ // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]],
+ // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
+ // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
+ // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
+ // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]],
+ // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
+ // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
+ // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
+ // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]],
+ // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static global vars being captured.
+
+ // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
+ // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
+ // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]],
+ // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
+ // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
+ // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]],
+ // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
+ // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
+ // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]],
+ // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
// CHECK: call i32 @__tgt_target
// CHECK: call void [[OFFLOADF:@.+]](
// Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
@@ -154,7 +268,7 @@
Gb += 1.0;
Sb += 1.0;
- // CHECK: define internal void [[OFFLOADF]]({{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}}, {{.+}}* {{.*}}%{{.+}})
+ // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
// CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
// CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})