r263837 - [OPENMP] Implementation of codegen for firstprivate clause of target directive
Carlo Bertolli via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 18 14:43:33 PDT 2016
Author: cbertol
Date: Fri Mar 18 16:43:32 2016
New Revision: 263837
URL: http://llvm.org/viewvc/llvm-project?rev=263837&view=rev
Log:
[OPENMP] Implementation of codegen for firstprivate clause of target directive
This patch implements the following aspects:
It extends sema to check that a variable is not reference in both a map clause and firstprivate or private. This is needed to ensure correct functioning at codegen level, apart from being useful for the user.
It implements firstprivate for target in codegen. The implementation applies to both host and nvptx devices.
It adds regression tests for codegen of firstprivate, host and device side when using the host as device, and nvptx side.
Please note that the regression test for nvptx codegen is missing VLAs. This is because VLAs currently require saving and restoring the stack which appears not to be a supported operation by nvptx backend.
It adds a check in sema regression tests for target map, firstprivate, and private clauses.
http://reviews.llvm.org/D18203
Added:
cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
cfe/trunk/test/OpenMP/target_firstprivate_codegen.cpp
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/target_firstprivate_messages.cpp
cfe/trunk/test/OpenMP/target_map_messages.cpp
cfe/trunk/test/OpenMP/target_private_messages.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=263837&r1=263836&r2=263837&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Fri Mar 18 16:43:32 2016
@@ -8133,6 +8133,8 @@ def err_omp_schedule_nonmonotonic_ordere
"'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">;
def err_omp_ordered_simd : Error<
"'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">;
+def err_omp_variable_in_map_and_dsa : Error<
+ "%0 variable cannot be in a map clause in '#pragma omp %1' directive">;
} // end of OpenMP category
let CategoryName = "Related Result Type Issue" in {
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=263837&r1=263836&r2=263837&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Fri Mar 18 16:43:32 2016
@@ -4175,6 +4175,7 @@ void CGOpenMPRuntime::emitTargetOutlined
// Emit target region as a standalone region.
auto &&CodeGen = [&CS, &D](CodeGenFunction &CGF) {
CodeGenFunction::OMPPrivateScope PrivateScope(CGF);
+ (void)CGF.EmitOMPFirstprivateClause(D, PrivateScope);
CGF.EmitOMPPrivateClause(D, PrivateScope);
(void)PrivateScope.Privatize();
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=263837&r1=263836&r2=263837&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Fri Mar 18 16:43:32 2016
@@ -7225,6 +7225,20 @@ OMPClause *Sema::ActOnOpenMPPrivateClaus
continue;
}
+ // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
+ // A list item cannot appear in both a map clause and a data-sharing
+ // attribute clause on the same construct
+ if (DSAStack->getCurrentDirective() == OMPD_target) {
+ if(DSAStack->checkMapInfoForVar(VD, /* CurrentRegionOnly = */ true,
+ [&](Expr *RE) -> bool {return true;})) {
+ Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+ << getOpenMPClauseName(OMPC_private)
+ << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+ ReportOriginalDSA(*this, DSAStack, D, DVar);
+ continue;
+ }
+ }
+
// OpenMP [2.9.3.3, Restrictions, C/C++, p.1]
// A variable of class type (or array thereof) that appears in a private
// clause requires an accessible, unambiguous default constructor for the
@@ -7456,6 +7470,19 @@ OMPClause *Sema::ActOnOpenMPFirstprivate
continue;
}
}
+ // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
+ // A list item cannot appear in both a map clause and a data-sharing
+ // attribute clause on the same construct
+ if (CurrDir == OMPD_target) {
+ if(DSAStack->checkMapInfoForVar(VD, /* CurrentRegionOnly = */ true,
+ [&](Expr *RE) -> bool {return true;})) {
+ Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+ << getOpenMPClauseName(OMPC_firstprivate)
+ << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+ ReportOriginalDSA(*this, DSAStack, D, DVar);
+ continue;
+ }
+ }
}
// Variably modified types are not supported for tasks.
@@ -9897,6 +9924,20 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClau
continue;
}
+ // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
+ // A list item cannot appear in both a map clause and a data-sharing
+ // attribute clause on the same construct
+ if (DKind == OMPD_target && VD) {
+ auto DVar = DSAStack->getTopDSA(VD, false);
+ if (isOpenMPPrivate(DVar.CKind)) {
+ Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+ << getOpenMPClauseName(DVar.CKind)
+ << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+ ReportOriginalDSA(*this, DSAStack, D, DVar);
+ continue;
+ }
+ }
+
Vars.push_back(RE);
DSAStack->addExprToVarMapInfo(D, RE);
}
Added: cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp?rev=263837&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/nvptx_target_firstprivate_codegen.cpp Fri Mar 18 16:43:32 2016
@@ -0,0 +1,235 @@
+
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fomptargets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fomptargets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+template<typename tx, typename ty>
+struct TT{
+ tx X;
+ ty Y;
+};
+
+// TCHECK: [[TT:%.+]] = type { i64, i8 }
+// TCHECK: [[S1:%.+]] = type { double }
+
+int foo(int n, double *ptr) {
+ int a = 0;
+ short aa = 0;
+ float b[10];
+ double c[5][10];
+ TT<long long, char> d;
+
+ #pragma omp target firstprivate(a)
+ {
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A1:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+ // TCHECK-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+ // TCHECK-64: [[A_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV]],
+ // TCHECK-32: [[A_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[A_ADDR_VAL]], i{{[0-9]+}}* [[A1]],
+ // TCHECK: ret void
+
+#pragma omp target firstprivate(aa,b,c,d)
+ {
+ aa += 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the
+ // target region
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A2_IN:%.+]], [10 x float]* {{.+}} [[B_IN:%.+]], [5 x [10 x double]]* {{.+}} [[C_IN:%.+]], [[TT]]* {{.+}} [[D_IN:%.+]])
+ // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B_ADDR:%.+]] = alloca [10 x float]*,
+ // TCHECK: [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*,
+ // TCHECK: [[D_ADDR:%.+]] = alloca [[TT]]*,
+ // TCHECK: [[A2_PRIV:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B_PRIV:%.+]] = alloca [10 x float],
+ // TCHECK: [[C_PRIV:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[D_PRIV:%.+]] = alloca [[TT]],
+ // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]],
+ // TCHECK: store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]],
+ // TCHECK: store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]],
+ // TCHECK: store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]],
+ // TCHECK: [[CONV_A2ADDR:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]],
+ // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]],
+ // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]],
+
+ // firstprivate(aa): a_priv = a_in
+ // TCHECK: [[A2_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV_A2ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[A2_CONV_VAL]], i{{[0-9]+}}* [[A2_PRIV]],
+
+ // firstprivate(b): memcpy(b_priv,b_in)
+ // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x float]* [[B_PRIV]] to i8*
+ // TCHECK: [[B_ADDR_REF_BCAST:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_ADDR_REF_BCAST]], {{.+}})
+
+ // firstprivate(c)
+ // TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8*
+ // TCHECK: [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_PRIV_BCAST]], i8* [[C_IN_BCAST]],{{.+}})
+
+ // firstprivate(d)
+ // TCHECK: [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8*
+ // TCHECK: [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[D_PRIV_BCAST]], i8* [[D_IN_BCAST]],{{.+}})
+
+
+ #pragma omp target firstprivate(ptr)
+ {
+ ptr[0]++;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
+ // TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
+ // TCHECK: [[PTR_PRIV:%.+]] = alloca double*,
+ // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]],
+ // TCHECK: [[PTR_IN_REF:%.+]] = load double*, double** [[PTR_ADDR]],
+ // TCHECK: store double* [[PTR_IN_REF]], double** [[PTR_PRIV]],
+
+ return a;
+}
+
+
+template<typename tx>
+tx ftemplate(int n) {
+ tx a = 0;
+ tx b[10];
+
+#pragma omp target firstprivate(a,b)
+ {
+ a += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+static
+int fstatic(int n) {
+ int a = 0;
+ char aaa = 0;
+ int b[10];
+
+#pragma omp target firstprivate(a,aaa,b)
+ {
+ a += 1;
+ aaa += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_PRIV:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]],
+// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK-64: [[A_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[A3_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A3_ADDR]] to i8*
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+
+// firstprivate(a): a_priv = a_in
+// TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_CONV]],
+// TCHECK-32: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A_PRIV]],
+
+// firstprivate(aaa)
+// TCHECK: [[A3_IN_VAL:%.+]] = load i8, i8* [[A3_CONV]],
+// TCHECK: store i{{[0-9]+}} [[A3_IN_VAL]], i{{[0-9]+}}* [[A3_PRIV]],
+
+// firstprivate(b)
+// TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
+// TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_IN_BCAST]],{{.+}})
+
+// TCHECK: ret void
+
+struct S1 {
+ double a;
+
+ int r1(int n){
+ int b = n+1;
+
+#pragma omp target firstprivate(b)
+ {
+ this->a = (double)b + 1.5;
+ }
+
+ return (int)b;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]])
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B_PRIV:%.+]] = alloca i{{[0-9]+}},
+
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+ // TCHECK-64: [[B_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}*
+
+ // firstprivate(b)
+ // TCHECK-64: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_ADDR_CONV]],
+ // TCHECK-32: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[B_IN_VAL]], i{{[0-9]+}}* [[B_PRIV]],
+
+ // TCHECK: ret void
+};
+
+
+
+int bar(int n, double *ptr){
+ int a = 0;
+ a += foo(n, ptr);
+ S1 S;
+ a += S.r1(n);
+ a += fstatic(n);
+ a += ftemplate<int>(n);
+
+ return a;
+}
+
+// template
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK-64: [[A_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+
+// firstprivate(a)
+// TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR_CONV]]
+// TCHECK-32: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]]
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A_PRIV]],
+
+// firstprivate(b)
+// TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
+// TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_IN_BCAST]],{{.+}})
+
+// TCHECK: ret void
+
+#endif
Added: cfe/trunk/test/OpenMP/target_firstprivate_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_firstprivate_codegen.cpp?rev=263837&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/target_firstprivate_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/target_firstprivate_codegen.cpp Fri Mar 18 16:43:32 2016
@@ -0,0 +1,595 @@
+// Test host codegen.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -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 -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -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 -fomptargets=i386-pc-linux-gnu -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 -fomptargets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fomptargets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fomp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fomptargets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fomp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+template<typename tx, typename ty>
+struct TT{
+ tx X;
+ ty Y;
+};
+
+// CHECK: [[TT:%.+]] = type { i64, i8 }
+// CHECK: [[S1:%.+]] = type { double }
+
+// TCHECK: [[TT:%.+]] = type { i64, i8 }
+// TCHECK: [[S1:%.+]] = type { double }
+
+// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 4]
+// CHECK: [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 128, i32 3, i32 128, i32 3, i32 3, i32 128, i32 128, i32 3, i32 3]
+// CHECK-64-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ]] 8]
+// CHECK-32-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i32] [i[[SZ]] 4]
+// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 160]
+// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 3, i32 128, i32 128, i32 128, i32 3]
+// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40]
+// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 128, i32 128, i32 3]
+// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40]
+// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 128, i32 3]
+
+
+// CHECK: define {{.*}}[[FOO:@.+]](
+int foo(int n, double *ptr) {
+ int a = 0;
+ short aa = 0;
+ float b[10];
+ float bn[n];
+ double c[5][10];
+ double cn[5][n];
+ TT<long long, char> d;
+
+ #pragma omp target firstprivate(a)
+ {
+ }
+
+ // a is passed by value to tgt_target
+ // CHECK: [[N_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // CHECK: [[PTR_ADDR:%.+]] = alloca double*,
+ // CHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // CHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+ // CHECK: [[B:%.+]] = alloca [10 x float],
+ // CHECK: [[SSTACK:%.+]] = alloca i8*,
+ // CHECK: [[C:%.+]] = alloca [5 x [10 x double]],
+ // CHECK: [[D:%.+]] = alloca [[TT]],
+ // CHECK: [[ACAST:%.+]] = alloca i{{[0-9]+}},
+ // CHECK: {{.+}} = alloca i{{[0-9]+}},
+ // CHECK: [[BASE_PTR_ARR:%.+]] = alloca [1 x i8*],
+ // CHECK: [[PTR_ARR:%.+]] = alloca [1 x i8*],
+ // CHECK: [[A2CAST:%.+]] = alloca i{{[0-9]+}},
+ // CHECK: [[BASE_PTR_ARR2:%.+]] = alloca [9 x i8*],
+ // CHECK: [[PTR_ARR2:%.+]] = alloca [9 x i8*],
+ // CHECK: [[SIZET2:%.+]] = alloca [9 x i{{[0-9]+}}],
+ // CHECK: [[BASE_PTR_ARR3:%.+]] = alloca [1 x i8*],
+ // CHECK: [[PTR_ARR3:%.+]] = alloca [1 x i8*],
+ // CHECK: [[N_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[N_ADDR]],
+ // CHECK-64: [[N_EXT:%.+]] = zext i{{[0-9]+}} [[N_ADDR_VAL]] to i{{[0-9]+}}
+ // CHECK: [[SSAVE_RET:%.+]] = call i8* @llvm.stacksave()
+ // CHECK: store i8* [[SSAVE_RET]], i8** [[SSTACK]],
+ // CHECK-64: [[BN_VLA:%.+]] = alloca float, i{{[0-9]+}} [[N_EXT]],
+ // CHECK-32: [[BN_VLA:%.+]] = alloca float, i{{[0-9]+}} [[N_ADDR_VAL]],
+ // CHECK: [[N_ADDR_VAL2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[N_ADDR]],
+ // CHECK-64: [[N_EXT2:%.+]] = zext i{{[0-9]+}} [[N_ADDR_VAL2]] to i{{[0-9]+}}
+ // CHECK-64: [[CN_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} 5, [[N_EXT2]]
+ // CHECK-32: [[CN_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} 5, [[N_ADDR_VAL2]]
+ // CHECK: [[CN_VLA:%.+]] = alloca double, i{{[0-9]+}} [[CN_SIZE]],
+ // CHECK: [[AVAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A]],
+ // CHECK-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[ACAST]] to i{{[0-9]+}}*
+ // CHECK-64: store i{{[0-9]+}} [[AVAL]], i{{[0-9]+}}* [[CONV]],
+ // CHECK-32: store i{{[0-9]+}} [[AVAL]], i{{[0-9]+}}* [[ACAST]],
+ // CHECK: [[ACAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ACAST]],
+ // CHECK: [[ACAST_TOPTR:%.+]] = inttoptr i{{[0-9]+}} [[ACAST_VAL]] to i8*
+ // CHECK: [[BASE_PTR_GEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[ACAST_TOPTR]], i8** [[BASE_PTR_GEP]],
+ // CHECK: [[ACAST_TOPTR2:%.+]] = inttoptr i{{[0-9]+}} [[ACAST_VAL]] to i8*
+ // CHECK: [[PTR_GEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[ACAST_TOPTR2]], i8** [[PTR_GEP]],
+ // CHECK: [[BASE_PTR_GEP_ARG:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BASE_PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[PTR_GEP_ARG:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PTR_ARR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: {{.+}} = call i32 @__tgt_target(i32 -1, {{.+}}, i32 1, i8** [[BASE_PTR_GEP_ARG]], i8** [[PTR_GEP_ARG]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET]], i32 0, i32 0), i32* getelementptr inbounds ([1 x i32], [1 x i32]* [[MAPT]], i32 0, i32 0))
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A1:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+ // TCHECK-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+ // TCHECK-64: [[A_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV]],
+ // TCHECK-32: [[A_ADDR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[A_ADDR_VAL]], i{{[0-9]+}}* [[A1]],
+ // TCHECK: ret void
+
+#pragma omp target firstprivate(aa,b,bn,c,cn,d)
+ {
+ aa += 1;
+ b[2] = 1.0;
+ bn[3] = 1.0;
+ c[1][2] = 1.0;
+ cn[1][3] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // CHECK: [[A2VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2]],
+ // CHECK: [[A2CASTCONV:%.+]] = bitcast i{{[0-9]+}}* [[A2CAST]] to i{{[0-9]+}}*
+ // CHECK: store i{{[0-9]+}} [[A2VAL]], i{{[0-9]+}}* [[A2CASTCONV]],
+ // CHECK: [[A2CAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2CAST]],
+ // CHECK-64: [[BN_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} [[N_EXT]], 4
+ // CHECK-32: [[BN_SIZE:%.+]] = mul{{.+}} i{{[0-9]+}} [[N_ADDR_VAL]], 4
+ // CHECK-64: [[CN_SIZE_1:%.+]] = mul{{.+}} i{{[0-9]+}} 5, [[N_EXT2]]
+ // CHECK-32: [[CN_SIZE_1:%.+]] = mul{{.+}} i{{[0-9]+}} 5, [[N_ADDR_VAL2]]
+ // CHECK: [[CN_SIZE_2:%.+]] = mul{{.+}} i{{[0-9]+}} [[CN_SIZE_1]], 8
+
+ // firstprivate(aa) --> base_ptr = aa, ptr = aa, size = 2 (short)
+ // CHECK: [[A2CAST_TO_INT:%.+]] = inttoptr i{{[0-9]+}} [[A2CAST_VAL]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_0:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[A2CAST_TO_INT]], i8** [[BASE_PTR_GEP2_0]],
+ // CHECK: [[A2CAST_TO_INT_2:%.+]] = inttoptr i{{[0-9]+}} [[A2CAST_VAL]] to i8*
+ // CHECK: [[PTR_GEP2_0:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[A2CAST_TO_INT_2]], i8** [[PTR_GEP2_0]],
+ // CHECK: [[SIZE_GEPA2:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[SIZE_GEPA2]],
+
+ // firstprivate(b): base_ptr = &b[0], ptr = &b[0], size = 40 (sizeof(float)*10)
+ // CHECK: [[BCAST:%.+]] = bitcast [10 x float]* [[B]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_1:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i8* [[BCAST]], i8** [[BASE_PTR_GEP2_1]],
+ // CHECK: [[BCAST2:%.+]] = bitcast [10 x float]* [[B]] to i8*
+ // CHECK: [[PTR_GEP2_1:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i8* [[BCAST2]], i8** [[PTR_GEP2_1]],
+ // CHECK: [[SIZE_GEPB:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i{{[0-9]+}} 40, i{{[0-9]+}}* [[SIZE_GEPB]],
+
+ // firstprivate(bn), 2 entries, n and bn: (1) base_ptr = n, ptr = n, size = 8 ; (2) base_ptr = &c[0], ptr = &c[0], size = n*sizeof(float)
+ // CHECK-64: [[N_EXT3_1:%.+]] = inttoptr i{{[0-9]+}} [[N_EXT]] to i8*
+ // CHECK-32: [[N_EXT3_1:%.+]] = inttoptr i{{[0-9]+}} [[N_ADDR_VAL]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i8* [[N_EXT3_1]], i8** [[BASE_PTR_GEP2_2]],
+ // CHECK-64: [[N_EXT3_2:%.+]] = inttoptr i{{[0-9]+}} [[N_EXT]] to i8*
+ // CHECK-32: [[N_EXT3_2:%.+]] = inttoptr i{{[0-9]+}} [[N_ADDR_VAL]] to i8*
+ // CHECK: [[PTR_GEP2_2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i8* [[N_EXT3_2]], i8** [[PTR_GEP2_2]],
+ // CHECK: [[SIZE_GEPBN_1:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i{{[0-9]+}} {{[0-9]}}, i{{[0-9]+}}* [[SIZE_GEPBN_1]],
+ // CHECK: [[VLABN_BCAST:%.+]] = bitcast float* [[BN_VLA]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_3:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 3
+ // CHECK: store i8* [[VLABN_BCAST]], i8** [[BASE_PTR_GEP2_3]],
+ // CHECK: [[SIZE_GEPBN_3:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 3
+ // CHECK: store i{{[0-9]+}} [[BN_SIZE]], i{{[0-9]+}}* [[SIZE_GEPBN_3]]
+
+ // firstprivate(c): base_ptr = &c[0], ptr = &c[0], size = 400 (5*10*sizeof(double))
+ // CHECK: [[C_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_4:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 4
+ // CHECK: store i8* [[C_BCAST]], i8** [[BASE_PTR_GEP2_4]],
+ // CHECK: [[C_BCAST2:%.+]] = bitcast [5 x [10 x double]]* [[C]] to i8*
+ // CHECK: [[PTR_GEP2_4:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 4
+ // CHECK: store i8* [[C_BCAST2]], i8** [[PTR_GEP2_4]],
+ // CHECK: [[SIZE_GEPC_4:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 4
+ // CHECK: store i{{[0-9]+}} 400, i{{[0-9]+}}* [[SIZE_GEPC_4]],
+
+ // firstprivate(cn), 3 entries, 5, n, cn: (1) base_ptr = 5, ptr = 5, size = 8; (2) (1) base_ptr = n, ptr = n, size = 8; (3) base_ptr = &cn[0], ptr = &cn[0], size = 5*n*sizeof(double)
+ // CHECK: [[BASE_PTR_GEP2_5:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 5
+ // CHECK: store i8* inttoptr (i{{[0-9]+}} 5 to i8*), i8** [[BASE_PTR_GEP2_5]],
+ // CHECK: [[PTR_GEP2_5:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 5
+ // CHECK: store i8* inttoptr (i{{[0-9]+}} 5 to i8*), i8** [[PTR_GEP2_5]],
+ // CHECK: [[SIZE_GEPCN_5:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 5
+ // CHECK: store i{{[0-9]+}} {{[0-9]}}, i{{[0-9]+}}* [[SIZE_GEPCN_5]],
+ // CHECK-64: [[CN_SZ_2_1:%.+]] = inttoptr i{{[0-9]+}} [[N_EXT2]] to i8*
+ // CHECK-32: [[CN_SZ_2_1:%.+]] = inttoptr i{{[0-9]+}} [[N_ADDR_VAL2]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_6:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 6
+ // CHECK: store i8* [[CN_SZ_2_1]], i8** [[BASE_PTR_GEP2_6]],
+ // CHECK-64: [[CN_SZ_2_2:%.+]] = inttoptr i{{[0-9]+}} [[N_EXT2]] to i8*
+ // CHECK-32: [[CN_SZ_2_2:%.+]] = inttoptr i{{[0-9]+}} [[N_ADDR_VAL2]] to i8*
+ // CHECK: [[PTR_GEP2_6:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 6
+ // CHECK: store i8* [[CN_SZ_2_2]], i8** [[PTR_GEP2_6]],
+ // CHECK: [[SIZE_GEPCN_6:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 6
+ // CHECK: store i{{[0-9]+}} {{[0-9]}}, i{{[0-9]+}}* [[SIZE_GEPCN_6]],
+ // CHECK: [[VLA_CN_BCAST:%.+]] = bitcast double* [[CN_VLA]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_7:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 7
+ // CHECK: store i8* [[VLA_CN_BCAST]], i8** [[BASE_PTR_GEP2_7]],
+ // CHECK: [[VLA_CN_BCAST2:%.+]] = bitcast double* [[CN_VLA]] to i8*
+ // CHECK: [[PTR_GEP2_7:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 7
+ // CHECK: store i8* [[VLA_CN_BCAST2]], i8** [[PTR_GEP2_7]],
+ // CHECK: [[SIZE_GEPCN_7:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 7
+ // CHECK: store i{{[0-9]+}} [[CN_SIZE_2]], i{{[0-9]+}}* [[SIZE_GEPCN_7]],
+
+ // firstprivate(d): base_ptr = &d, ptr = &d, size = 16
+ // CHECK: [[D_REF:%.+]] = bitcast [[TT]]* [[D]] to i8*
+ // CHECK: [[BASE_PTR_GEP2_8:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 8
+ // CHECK: store i8* [[D_REF]], i8** [[BASE_PTR_GEP2_8]],
+ // CHECK: [[D_REF2:%.+]] = bitcast [[TT]]* [[D]] to i8*
+ // CHECK: [[PTR_GEP2_8:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 8
+ // CHECK: store i8* [[D_REF2]], i8** [[PTR_GEP2_8]],
+ // CHECK: [[SIZE_GEPCN_8:%.+]] = getelementptr inbounds [9 x i{{[0-9]+}}], [9 x i{{[0-9]+}}]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 8
+ // CHECK: store i{{[0-9]+}} {{[0-9]+}}, i{{[0-9]+}}* [[SIZE_GEPCN_8]],
+
+
+ // CHECK: [[BASE_PTR_GEP_ARG2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BASE_PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[PTR_GEP_ARG2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[PTR_ARR2]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[SIZES_ARG2:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[SIZET2]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: {{.+}} = call i32 @__tgt_target(i32 -1, {{.+}}, i32 9, i8** [[BASE_PTR_GEP_ARG2]], i8** [[PTR_GEP_ARG2]], i[[SZ]]* [[SIZES_ARG2]], i32* getelementptr inbounds ([9 x i32], [9 x i32]* [[MAPT2]], i32 0, i32 0))
+
+ // make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the
+ // target region
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A2_IN:%.+]], [10 x float]* {{.+}} [[B_IN:%.+]], i{{[0-9]+}} [[BN_SZ:%.+]], float* {{.+}} [[BN_IN:%.+]], [5 x [10 x double]]* {{.+}} [[C_IN:%.+]], i{{[0-9]+}} [[CN_SZ1:%.+]], i{{[0-9]+}} [[CN_SZ2:%.+]], double* {{.+}} [[CN_IN:%.+]], [[TT]]* {{.+}} [[D_IN:%.+]])
+ // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B_ADDR:%.+]] = alloca [10 x float]*,
+ // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BN_ADDR:%.+]] = alloca float*,
+ // TCHECK: [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*,
+ // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[VLA_ADDR4:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CN_ADDR:%.+]] = alloca double*,
+ // TCHECK: [[D_ADDR:%.+]] = alloca [[TT]]*,
+ // TCHECK: [[A2_PRIV:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B_PRIV:%.+]] = alloca [10 x float],
+ // TCHECK: [[SSTACK:%.+]] = alloca i8*,
+ // TCHECK: [[C_PRIV:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[D_PRIV:%.+]] = alloca [[TT]],
+ // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]],
+ // TCHECK: store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[BN_SZ]], i{{[0-9]+}}* [[VLA_ADDR]],
+ // TCHECK: store float* [[BN_IN]], float** [[BN_ADDR]],
+ // TCHECK: store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[CN_SZ1]], i{{[0-9]+}}* [[VLA_ADDR2]],
+ // TCHECK: store i{{[0-9]+}} [[CN_SZ2]], i{{[0-9]+}}* [[VLA_ADDR4]],
+ // TCHECK: store double* [[CN_IN]], double** [[CN_ADDR]],
+ // TCHECK: store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]],
+ // TCHECK: [[CONV_A2ADDR:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]],
+ // TCHECK: [[BN_SZ_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
+ // TCHECK: [[BN_ADDR_REF:%.+]] = load float*, float** [[BN_ADDR]],
+ // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]],
+ // TCHECK: [[CN_SZ1_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
+ // TCHECK: [[CN_SZ2_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR4]],
+ // TCHECK: [[CN_ADDR_REF:%.+]] = load double*, double** [[CN_ADDR]],
+ // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]],
+
+ // firstprivate(aa): a_priv = a_in
+ // TCHECK: [[A2_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV_A2ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[A2_CONV_VAL]], i{{[0-9]+}}* [[A2_PRIV]],
+
+ // firstprivate(b): memcpy(b_priv,b_in)
+ // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x float]* [[B_PRIV]] to i8*
+ // TCHECK: [[B_ADDR_REF_BCAST:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_ADDR_REF_BCAST]], {{.+}})
+
+ // TCHECK: [[RET_STACK:%.+]] = call i8* @llvm.stacksave()
+ // TCHECK: store i8* [[RET_STACK]], i8** [[SSTACK]],
+
+ // firstprivate(bn)
+ // TCHECK: [[BN_PRIV:%.+]] = alloca float, i{{[0-9]+}} [[BN_SZ_VAL]],
+ // TCHECK: [[BN_COPY_SZ:%.+]] = mul{{.+}} i{{[0-9]+}} [[BN_SZ_VAL]], 4
+ // TCHECK: [[BN_PRIV__BCAST:%.+]] = bitcast float* [[BN_PRIV]] to i8*
+ // TCHECK: [[BN_REF_IN_BCAST:%.+]] = bitcast float* [[BN_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[BN_PRIV__BCAST]], i8* [[BN_REF_IN_BCAST]], i{{[0-9]+}} [[BN_COPY_SZ]],{{.+}})
+
+ // firstprivate(c)
+ // TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8*
+ // TCHECK: [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_PRIV_BCAST]], i8* [[C_IN_BCAST]],{{.+}})
+
+ // firstprivate(cn)
+ // TCHECK: [[CN_SZ:%.+]] = mul{{.+}} i{{[0-9]+}} [[CN_SZ1_VAL]], [[CN_SZ2_VAL]]
+ // TCHECK: [[CN_PRIV:%.+]] = alloca double, i{{[0-9]+}} [[CN_SZ]],
+ // TCHECK: [[CN_SZ2:%.+]] = mul{{.+}} i{{[0-9]+}} [[CN_SZ1_VAL]], [[CN_SZ2_VAL]]
+ // TCHECK: [[CN_SZ2_CPY:%.+]] = mul{{.+}} i{{[0-9]+}} [[CN_SZ2]], 8
+ // TCHECK: [[CN_PRIV_BCAST:%.+]] = bitcast double* [[CN_PRIV]] to i8*
+ // TCHECK: [[CN_IN_BCAST:%.+]] = bitcast double* [[CN_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[CN_PRIV_BCAST]], i8* [[CN_IN_BCAST]], i{{[0-9]+}} [[CN_SZ2_CPY]],{{.+}})
+
+ // firstprivate(d)
+ // TCHECK: [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8*
+ // TCHECK: [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[D_PRIV_BCAST]], i8* [[D_IN_BCAST]],{{.+}})
+
+
+ #pragma omp target firstprivate(ptr)
+ {
+ ptr[0]++;
+ }
+ // CHECK: [[PTR_ADDR_REF:%.+]] = load double*, double** [[PTR_ADDR]],
+ // CHECK: [[PTR_ADDR_BCAST:%.+]] = bitcast double* [[PTR_ADDR_REF]] to i8*
+
+ // CHECK: [[BASE_PTR_GEP3_0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[PTR_ADDR_BCAST]], i8** [[BASE_PTR_GEP3_0]],
+ // CHECK: [[PTR_ADDR_BCAST2:%.+]] = bitcast double* [[PTR_ADDR_REF]] to i8*
+ // CHECK: [[PTR_GEP3_0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[PTR_ADDR_BCAST2]], i8** [[PTR_GEP3_0]],
+
+ // CHECK: [[BASE_PTR_GEP_ARG3:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BASE_PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: [[PTR_GEP_ARG3:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PTR_ARR3]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: {{.+}} = call i32 @__tgt_target(i32 -1, {{.+}}, i32 1, i8** [[BASE_PTR_GEP_ARG3]], i8** [[PTR_GEP_ARG3]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET3]], i32 0, i32 0), i32* getelementptr inbounds ([1 x i32], [1 x i32]* [[MAPT3]], i32 0, i32 0))
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
+ // TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
+ // TCHECK: [[PTR_PRIV:%.+]] = alloca double*,
+ // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]],
+ // TCHECK: [[PTR_IN_REF:%.+]] = load double*, double** [[PTR_ADDR]],
+ // TCHECK: store double* [[PTR_IN_REF]], double** [[PTR_PRIV]],
+
+ return a;
+}
+
+
+template<typename tx>
+tx ftemplate(int n) {
+ tx a = 0;
+ tx b[10];
+
+#pragma omp target firstprivate(a,b)
+ {
+ a += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+static
+int fstatic(int n) {
+ int a = 0;
+ char aaa = 0;
+ int b[10];
+
+#pragma omp target firstprivate(a,aaa,b)
+ {
+ a += 1;
+ aaa += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_PRIV:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]],
+// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK-64: [[A_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[A3_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A3_ADDR]] to i8*
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+
+// firstprivate(a): a_priv = a_in
+// TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_CONV]],
+// TCHECK-32: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A_PRIV]],
+
+// firstprivate(aaa)
+// TCHECK: [[A3_IN_VAL:%.+]] = load i8, i8* [[A3_CONV]],
+// TCHECK: store i{{[0-9]+}} [[A3_IN_VAL]], i{{[0-9]+}}* [[A3_PRIV]],
+
+// firstprivate(b)
+// TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
+// TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_IN_BCAST]],{{.+}})
+
+// TCHECK: ret void
+
+struct S1 {
+ double a;
+
+ int r1(int n){
+ int b = n+1;
+ short int c[2][n];
+
+#pragma omp target firstprivate(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ return c[1][1] + (int)b;
+ }
+
+ // on the host side, we first generate r1, then the static function and the template above
+ // CHECK: define{{.+}} i32 {{.+}}([[S1]]* {{.+}}, i{{[0-9]+}} {{.+}})
+ // CHECK: [[BASE_PTRS4:%.+]] = alloca [5 x i8*],
+ // CHECK: [[PTRS4:%.+]] = alloca [5 x i8*],
+ // CHECK: [[SIZET4:%.+]] = alloca [5 x i{{[0-9]+}}],
+
+ // map(this: this ptr is implicitly captured (not firstprivate matter)
+ // CHECK: {{.+}} = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store {{.+}}, {{.+}},
+ // CHECK: {{.+}} = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store {{.+}}, {{.+}},
+ // CHECK: {{.+}} getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store {{.+}}, {{.+}}
+
+ // firstprivate(b): base_ptr = b, ptr = b, size = 4 (pass by-value)
+ // CHECK: [[B_CAST_PTR:%.+]] = inttoptr i{{[0-9]+}} [[B_CAST:%.+]] to i8*
+ // CHECK: [[BASE_PTRS_GEP4_1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i8* [[B_CAST_PTR]], i8** [[BASE_PTRS_GEP4_1]],
+ // CHECK: [[B_CAST_PTR2:%.+]] = inttoptr i{{[0-9]+}} [[B_CAST:%.+]] to i8*
+ // CHECK: [[PTRS_GEP4_1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i8* [[B_CAST_PTR2]], i8** [[PTRS_GEP4_1]],
+ // CHECK: [[SIZES_GEP4_1:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIZES_GEP4_1]],
+
+ // firstprivate(c), 3 entries: 2, n, c
+ // CHECK: [[BASE_PTRS_GEP4_2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i8* inttoptr (i{{[0-9]+}} 2 to i8*), i8** [[BASE_PTRS_GEP4_2]],
+ // CHECK: [[PTRS_GEP4_2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i8* inttoptr (i{{[0-9]+}} 2 to i8*), i8** [[PTRS_GEP4_2]],
+ // CHECK: [[SIZES_GEP4_2:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK-64: store i{{[0-9]+}} 8, i{{[0-9]+}}* [[SIZES_GEP4_2]],
+ // CHECK-32: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIZES_GEP4_2]],
+ // CHECK: [[N_PTR:%.+]] = inttoptr i{{[0-9]+}} [[N:%.+]] to i8*
+ // CHECK: [[BASE_PTRS_GEP4_3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3
+ // CHECK: store i8* [[N_PTR]], i8** [[BASE_PTRS_GEP4_3]],
+ // CHECK: [[N_PTR2:%.+]] = inttoptr i{{[0-9]+}} [[N:%.+]] to i8*
+ // CHECK: [[PTRS_GEP4_3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3
+ // CHECK: store i8* [[N_PTR2]], i8** [[PTRS_GEP4_3]],
+ // CHECK: [[SIZES_GEP4_3:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3
+ // CHECK-64: store i{{[0-9]+}} 8, i{{[0-9]+}}* [[SIZES_GEP4_3]],
+ // CHECK-32: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIZES_GEP4_3]],
+ // CHECK: [[B_BCAST:%.+]] = bitcast i{{[0-9]+}}* [[B:%.+]] to i8*
+ // CHECK: [[BASE_PTRS_GEP4_4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4
+ // CHECK: store i8* [[B_BCAST]], i8** [[BASE_PTRS_GEP4_4]],
+ // CHECK: [[B_BCAST2:%.+]] = bitcast i{{[0-9]+}}* [[B:%.+]] to i8*
+ // CHECK: [[PTRS_GEP4_4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4
+ // CHECK: store i8* [[B_BCAST2]], i8** [[PTRS_GEP4_4]],
+ // CHECK: [[SIZES_GEP4_4:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4
+ // CHECK: store i{{[0-9]+}} [[B_SIZE:%.+]], i{{[0-9]+}}* [[SIZES_GEP4_4]],
+
+ // only check that we use the map types stored in the global variable
+ // CHECK: call i32 @__tgt_target(i32 -1, {{.+}}, i32 5, i8** {{.+}}, i8** {{.+}}, i{{[0-9]+}}* {{.+}}, i32* getelementptr inbounds ([5 x i32], [5 x i32]* [[MAPT4]], i32 0, i32 0))
+
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}}{{.+}} [[C_IN:%.+]])
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[C_ADDR:%.+]] = alloca i{{[0-9]+}}*,
+ // TCHECK: [[B_PRIV:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[SSTACK:%.+]] = alloca i8*,
+
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]],
+ // TCHECK: store i{{[0-9]+}}* [[C_IN]], i{{[0-9]+}}** [[C_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+ // TCHECK-64: [[B_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
+ // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
+ // TCHECK: [[C_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_ADDR]],
+
+ // firstprivate(b)
+ // TCHECK-64: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_ADDR_CONV]],
+ // TCHECK-32: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[B_IN_VAL]], i{{[0-9]+}}* [[B_PRIV]],
+
+ // TCHECK: [[RET_STACK:%.+]] = call i8* @llvm.stacksave()
+ // TCHECK: store i8* [[RET_STACK:%.+]], i8** [[SSTACK]],
+
+ // firstprivate(c)
+ // TCHECK: [[C_SZ:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF]], [[VLA_ADDR_REF2]]
+ // TCHECK: [[C_PRIV:%.+]] = alloca i{{[0-9]+}}, i{{[0-9]+}} [[C_SZ]],
+ // TCHECK: [[C_SZ2:%.+]] = mul{{.+}} i{{[0-9]+}} [[VLA_ADDR_REF]], [[VLA_ADDR_REF2]]
+ // TCHECK: [[C_SZ_CPY:%.+]] = mul{{.+}} i{{[0-9]+}} [[C_SZ2]], 2
+ // TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast i{{[0-9]+}}* [[C_PRIV]] to i8*
+ // TCHECK: [[C_IN_BCAST:%.+]] = bitcast i{{[0-9]+}}* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_PRIV_BCAST]], i8* [[C_IN_BCAST]],{{.+}})
+
+ // finish
+ // TCHECK: [[RELOAD_SSTACK:%.+]] = load i8*, i8** [[SSTACK]],
+ // TCHECK: call void @llvm.stackrestore(i8* [[RELOAD_SSTACK]])
+ // TCHECK: ret void
+
+
+ // static host function
+ // CHECK: define{{.+}} i32 {{.+}}(i{{[0-9]+}} {{.+}})
+ // CHECK: [[BASE_PTRS5:%.+]] = alloca [3 x i8*],
+ // CHECK: [[PTRS5:%.+]] = alloca [3 x i8*],
+
+ // firstprivate(a): by value
+ // CHECK: [[A_CAST_PTR:%.+]] = inttoptr i{{[0-9]+}} [[A_CAST:%.+]] to i8*
+ // CHECK: [[BASE_PTRS_GEP5_0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTRS5]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[A_CAST_PTR]], i8** [[BASE_PTRS_GEP5_0]],
+ // CHECK: [[A_CAST_PTR2:%.+]] = inttoptr i{{[0-9]+}} [[A_CAST:%.+]] to i8*
+ // CHECK: [[PTRS_GEP5_0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS5]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // CHECK: store i8* [[A_CAST_PTR2]], i8** [[PTRS_GEP5_0]],
+
+ // firstprivate(aaa): by value
+ // CHECK: [[A3_CAST_PTR:%.+]] = inttoptr i{{[0-9]+}} [[A3_CAST:%.+]] to i8*
+ // CHECK: [[BASE_PTRS_GEP5_1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTRS5]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i8* [[A3_CAST_PTR]], i8** [[BASE_PTRS_GEP5_1]],
+ // CHECK: [[A3_CAST_PTR2:%.+]] = inttoptr i{{[0-9]+}} [[A3_CAST:%.+]] to i8*
+ // CHECK: [[PTRS_GEP5_1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS5]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // CHECK: store i8* [[A3_CAST_PTR2]], i8** [[PTRS_GEP5_1]],
+
+ // firstprivate(b): base_ptr = &b[0], ptr= &b[0]
+ // CHECK: [[B_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B:%.+]] to i8*
+ // CHECK: [[BASE_PTRS_GEP5_2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASE_PTRS5]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i8* [[B_BCAST]], i8** [[BASE_PTRS_GEP5_2]],
+ // CHECK: [[B_BCAST2:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B:%.+]] to i8*
+ // CHECK: [[PTRS_GEP5_2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS5]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+ // CHECK: store i8* [[B_BCAST2]], i8** [[PTRS_GEP5_2]],
+
+ // only check that the right sizes and map types are used
+ // CHECK: call i32 @__tgt_target(i32 -1, {{.+}}, i32 3, i8** {{.+}}, i8** {{.+}}, i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET5]], i32 0, i32 0), i32* getelementptr inbounds ([3 x i32], [3 x i32]* [[MAPT5]], i32 0, i32 0))
+};
+
+
+
+int bar(int n, double *ptr){
+ int a = 0;
+ a += foo(n, ptr);
+ S1 S;
+ a += S.r1(n);
+ a += fstatic(n);
+ a += ftemplate<int>(n);
+
+ return a;
+}
+
+// template host and device
+
+// CHECK: define{{.+}} i32 {{.+}}(i{{[0-9]+}} {{.+}})
+// CHECK: [[BASE_PTRS6:%.+]] = alloca [2 x i8*],
+// CHECK: [[PTRS6:%.+]] = alloca [2 x i8*],
+
+// firstprivate(a): by value
+// CHECK: [[AT_CAST_PTR:%.+]] = inttoptr i{{[0-9]+}} [[AT_CAST:%.+]] to i8*
+// CHECK: [[BASE_PTRS_GEP6_0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTRS6]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+// CHECK: store i8* [[AT_CAST_PTR]], i8** [[BASE_PTRS_GEP6_0]],
+// CHECK: [[AT_CAST_PTR2:%.+]] = inttoptr i{{[0-9]+}} [[AT_CAST:%.+]] to i8*
+// CHECK: [[PTRS_GEP6_0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS6]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+// CHECK: store i8* [[AT_CAST_PTR2]], i8** [[PTRS_GEP6_0]],
+
+// firstprivate(b): pointer
+// CHECK: [[B_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B:%.+]] to i8*
+// CHECK: [[BASE_PTRS_GEP6_1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASE_PTRS6]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+// CHECK: store i8* [[B_BCAST]], i8** [[BASE_PTRS_GEP6_1]],
+// CHECK: [[B_BCAST2:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B:%.+]] to i8*
+// CHECK: [[PTRS_GEP6_1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS6]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+// CHECK: store i8* [[B_BCAST2]], i8** [[PTRS_GEP6_1]],
+
+// CHECK: call i32 @__tgt_target(i32 -1, {{.+}}, i32 2, i8** {{.+}}, i8** {{.+}}, i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET6]], i32 0, i32 0), i32* getelementptr inbounds ([2 x i32], [2 x i32]* [[MAPT6]], i32 0, i32 0))
+
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[A_PRIV:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK-64: [[A_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+
+// firstprivate(a)
+// TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR_CONV]]
+// TCHECK-32: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR]]
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A_PRIV]],
+
+// firstprivate(b)
+// TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
+// TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_PRIV_BCAST]], i8* [[B_IN_BCAST]],{{.+}})
+
+// TCHECK: ret void
+
+#endif
Modified: cfe/trunk/test/OpenMP/target_firstprivate_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_firstprivate_messages.cpp?rev=263837&r1=263836&r2=263837&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_firstprivate_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/target_firstprivate_messages.cpp Fri Mar 18 16:43:32 2016
@@ -189,6 +189,8 @@ int main(int argc, char **argv) {
static int si;
#pragma omp target firstprivate(si) // OK
{}
+#pragma omp target map(i) firstprivate(i) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}}
+ {}
s6 = s6_0; // expected-note {{in instantiation of member function 'S6<float>::operator=' requested here}}
s7 = s7_0; // expected-note {{in instantiation of member function 'S7<S6<float> >::operator=' requested here}}
return foomain(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<int, char>' requested here}}
Modified: cfe/trunk/test/OpenMP/target_map_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_map_messages.cpp?rev=263837&r1=263836&r2=263837&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_map_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/target_map_messages.cpp Fri Mar 18 16:43:32 2016
@@ -500,6 +500,10 @@ int main(int argc, char **argv) {
#pragma omp target data map(always, tofrom: always, tofrom, x)
#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo();
+#pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as private}}
+ {}
+#pragma omp target firstprivate(j) map(j) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as firstprivate}}
+ {}
return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
}
#endif
Modified: cfe/trunk/test/OpenMP/target_private_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_private_messages.cpp?rev=263837&r1=263836&r2=263837&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_private_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/target_private_messages.cpp Fri Mar 18 16:43:32 2016
@@ -189,6 +189,8 @@ int main(int argc, char **argv) {
static int si;
#pragma omp target private(si) // OK
{}
+#pragma omp target map(i) private(i) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}
+ {}
s6 = s6_0; // expected-note {{in instantiation of member function 'S6<float>::operator=' requested here}}
s7 = s7_0; // expected-note {{in instantiation of member function 'S7<S6<float> >::operator=' requested here}}
return foomain(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<int, char>' requested here}}
More information about the cfe-commits
mailing list