carlo.bertolli updated this revision to Diff 51285.
carlo.bertolli added a comment.
To make clear my comment I an updating the patch following Alexey's comment and
modifying it to make it work.
You can see that now we have the following check:
CGM.getTarget().getTriple().getArch() == llvm::Triple::nvptx ||
CGM.getTarget().getTriple().getArch() == llvm::Triple::nvptx64
Repository:
rL LLVM
http://reviews.llvm.org/D18286
Files:
lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
lib/CodeGen/CGOpenMPRuntimeNVPTX.h
lib/CodeGen/CGStmtOpenMP.cpp
test/OpenMP/nvptx_teams_firstprivate_codegen.cpp
test/OpenMP/nvptx_teams_private_codegen.cpp
Index: test/OpenMP/nvptx_teams_private_codegen.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/nvptx_teams_private_codegen.cpp
@@ -0,0 +1,539 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -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++ -std=c++11 -triple i386-unknown-unknown -fomptargets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -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) {
+ 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 private(a)
+ #pragma omp teams private(a)
+ {
+ }
+
+ // generate both private variables
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK-NOT: store {{.+}}, {{.+}} [[A]],
+ // TCHECK: ret void
+
+
+ // a is implictly firstprivate
+ #pragma omp target
+ #pragma omp teams private(a)
+ {
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK-NOT: store {{.+}}, {{.+}} [[A]],
+ // TCHECK: ret void
+
+ #pragma omp target firstprivate(a)
+ #pragma omp teams private(a)
+ {
+ }
+
+ // because of firstprivate, a.addr and a parameter are
+ // created. This is fine, as long as we do not use a.addr
+ // TCHECK: define void @__omp_offloading_{{.+}}({{.+}})
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK-NOT: store {{.+}}, {{.+}} [[ATE]],
+ // TCHECK: ret void
+
+ #pragma omp target private(a)
+ #pragma omp teams private(a)
+ {
+ a = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK-NOT: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: ret void
+
+ #pragma omp target
+ #pragma omp teams private(a)
+ {
+ a = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+ // TCHECK: ret void
+
+ #pragma omp target firstprivate(a)
+ #pragma omp teams private(a)
+ {
+ a = 1;
+ }
+
+ // check that we store in a without looking at the parameter
+ // TCHECK: define void @__omp_offloading_{{.+}}({{.+}})
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: ret void
+
+ #pragma omp target private(a, aa)
+ #pragma omp teams private(a,aa)
+ {
+ a = 1;
+ aa = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+ // TCHECK: ret void
+
+ #pragma omp target
+ #pragma omp teams private(a,aa)
+ {
+ a = 1;
+ aa = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
+ // TCHECK: ret void
+
+ #pragma omp target firstprivate(a, aa)
+ #pragma omp teams private(a,aa)
+ {
+ a = 1;
+ aa = 1;
+
+ aa = a+1;
+ }
+
+ // check that we are not using the firstprivate parameter
+ // TCHECK: define void @__omp_offloading_{{.+}}({{.+}})
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+ // TCHECK: [[A_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATE]]
+ // TCHECK: [[A_INC:%.+]] = add{{.+}} i{{[0-9]+}} [[A_VAL]], 1
+ // TCHECK: [[CONV:%.+]] = trunc i{{[0-9]+}} [[A_INC]] to i{{[0-9]+}}
+ // TCHECK: store i{{[0-9]+}} [[CONV]], i{{[0-9]+}}* [[A2TE]]
+ // TCHECK: ret void
+
+ #pragma omp target private(a, b,c, d)
+ #pragma omp teams private(a, b, c, d)
+ {
+ a = 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTA:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTA:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTA:%.+]] = alloca [[TT]],
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTE:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTE:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTE:%.+]] = alloca [[TT]],
+ // TCHECK: store {{.+}} [[ATE]],
+ // TCHECK: [[B_IDX:%.+]] = getelementptr{{.+}} [[BTE]], {{.+}},
+ // TCHECK: store {{.+}} [[B_IDX]],
+ // TCHECK: [[C_IDX1:%.+]] = getelementptr {{.+}} [[CTE]], {{.+}}
+ // TCHECK: [[C_IDX2:%.+]] = getelementptr {{.+}} [[C_IDX1]], {{.+}}
+ // TCHECK: store {{.+}} [[C_IDX2]],
+ // TCHECK: [[D_X:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_X]],
+ // TCHECK: [[D_Y:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_Y]],
+
+ #pragma omp target
+ #pragma omp teams private(a, b, c, d)
+ {
+ a = 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B:%.+]] = alloca [10 x float],
+ // TCHECK: [[C:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[D:%.+]] = alloca [[TT]],
+ // TCHECK: store {{.+}} [[A]],
+ // TCHECK: [[B_IDX:%.+]] = getelementptr{{.+}} [[B]], {{.+}},
+ // TCHECK: store {{.+}} [[B_IDX]],
+ // TCHECK: [[C_IDX1:%.+]] = getelementptr {{.+}} [[C]], {{.+}}
+ // TCHECK: [[C_IDX2:%.+]] = getelementptr {{.+}} [[C_IDX1]], {{.+}}
+ // TCHECK: store {{.+}} [[C_IDX2]],
+ // TCHECK: [[D_X:%.+]] = getelementptr {{.+}} [[D]], {{.+}}
+ // TCHECK: store {{.+}} [[D_X]],
+ // TCHECK: [[D_Y:%.+]] = getelementptr {{.+}} [[D]], {{.+}}
+ // TCHECK: store {{.+}} [[D_Y]],
+
+ #pragma omp target firstprivate(a, b, c, d)
+ #pragma omp teams private(a, b, c, d)
+ {
+ a = 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}({{.+}})
+ // TCHECK: [[A_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: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTA:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTA:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTA:%.+]] = alloca [[TT]],
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTE:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTE:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTE:%.+]] = alloca [[TT]],
+ // TCHECK: store {{.+}} [[ATE]],
+ // TCHECK: [[B_IDX:%.+]] = getelementptr{{.+}} [[BTE]], {{.+}},
+ // TCHECK: store {{.+}} [[B_IDX]],
+ // TCHECK: [[C_IDX1:%.+]] = getelementptr {{.+}} [[CTE]], {{.+}}
+ // TCHECK: [[C_IDX2:%.+]] = getelementptr {{.+}} [[C_IDX1]], {{.+}}
+ // TCHECK: store {{.+}} [[C_IDX2]],
+ // TCHECK: [[D_X:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_X]],
+ // TCHECK: [[D_Y:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_Y]],
+
+ return a;
+}
+
+
+template<typename tx>
+tx ftemplate(int n) {
+ tx a = 0;
+ short aa = 0;
+ tx b[10];
+
+#pragma omp target private(a,aa,b)
+#pragma omp teams private(a,aa,b)
+ {
+ a = 1;
+ aa = 1;
+ b[2] = 1;
+ }
+
+#pragma omp target
+#pragma omp teams private(a,aa,b)
+ {
+ a = 1;
+ aa = 1;
+ b[2] = 1;
+ }
+
+#pragma omp target firstprivate(a,aa,b)
+#pragma omp teams private(a,aa,b)
+ {
+ a = 1;
+ aa = 1;
+ b[2] = 1;
+ }
+
+ return a;
+}
+
+static
+int fstatic(int n) {
+ int a = 0;
+ short aa = 0;
+ char aaa = 0;
+ int b[10];
+
+#pragma omp target private(a,aa,aaa,b)
+#pragma omp teams private(a,aa,aaa,b)
+ {
+ a = 1;
+ aa = 1;
+ aaa = 1;
+ b[2] = 1;
+ }
+
+// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+#pragma omp target
+#pragma omp teams private(a,aa,aaa,b)
+ {
+ a = 1;
+ aa = 1;
+ aaa = 1;
+ b[2] = 1;
+ }
+
+// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+#pragma omp target firstprivate(a,aa,aaa,b)
+#pragma omp teams private(a,aa,aaa,b)
+ {
+ a = 1;
+ aa = 1;
+ aaa = 1;
+ b[2] = 1;
+ }
+
+// TCHECK: define void @__omp_offloading_{{.+}}({{.+}})
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+ return a;
+}
+
+struct S1 {
+ double a;
+
+ int r1(int n){
+ int b = n+1;
+ short int c[2][5];
+
+#pragma omp target private(b,c)
+#pragma omp teams private(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]])
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[BTA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTA:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: [[BTE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTE:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+
+ // this->a = (double)b + 1.5;
+ // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[BTE]],
+ // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
+ // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
+ // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
+
+ // c[1][1] = ++a;
+ // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
+ // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
+ // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds [2 x [5 x i{{[0-9]+}}]], [2 x [5 x i{{[0-9]+}}]]* [[CTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[C_1_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
+ // TCHECK: ret void
+
+#pragma omp target
+#pragma omp teams private(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]])
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[B:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[C:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+
+ // this->a = (double)b + 1.5;
+ // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B]],
+ // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
+ // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
+ // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
+
+ // c[1][1] = ++a;
+ // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
+ // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
+ // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds [2 x [5 x i{{[0-9]+}}]], [2 x [5 x i{{[0-9]+}}]]* [[C]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[C_1_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
+ // TCHECK: ret void
+
+#pragma omp target firstprivate(b,c)
+#pragma omp teams private(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ return c[1][1] + (int)b;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}}{{.+}})
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[C_ADDR:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]]*,
+ // TCHECK: [[BTA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTA:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: [[BTE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTE:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+
+ // this->a = (double)b + 1.5;
+ // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[BTE]],
+ // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
+ // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
+ // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
+
+ // c[1][1] = ++a;
+ // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
+ // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
+ // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds [2 x [5 x i{{[0-9]+}}]], [2 x [5 x i{{[0-9]+}}]]* [[CTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[C_1_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
+ // TCHECK: ret void
+
+};
+
+
+int bar(int n){
+ int a = 0;
+ a += foo(n);
+ S1 S;
+ a += S.r1(n);
+ a += fstatic(n);
+ a += ftemplate<int>(n);
+
+ return a;
+}
+
+// template
+// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+// TCHECK: define void @__omp_offloading_{{.+}}({{.+}})
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+#endif
+
Index: test/OpenMP/nvptx_teams_firstprivate_codegen.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/nvptx_teams_firstprivate_codegen.cpp
@@ -0,0 +1,848 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fomptargets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -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++ -std=c++11 -triple i386-unknown-unknown -fomptargets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -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];
+ float bn[n];
+ double c[5][10];
+ double cn[5][n];
+ TT<long long, char> d;
+
+ #pragma omp target private(a)
+ #pragma omp teams firstprivate(a)
+ {
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: ret void
+
+
+ // a is implictly firstprivate - copy the value from parameter to private teams variable
+ #pragma omp target
+ #pragma omp teams firstprivate(a)
+ {
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A:%.+]] = 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_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]],
+ // TCHECK: ret void
+
+ #pragma omp target firstprivate(a)
+ #pragma omp teams firstprivate(a)
+ {
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = 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_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]+}}* [[ATA]],
+ // TCHECK: [[A_IN_VAL2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[A_IN_VAL2]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: ret void
+
+ #pragma omp target private(a)
+ #pragma omp teams firstprivate(a)
+ {
+ a = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: ret void
+
+ #pragma omp target
+ #pragma omp teams firstprivate(a)
+ {
+ a = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A:%.+]] = 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_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+ // TCHECK: ret void
+
+ #pragma omp target firstprivate(a)
+ #pragma omp teams firstprivate(a)
+ {
+ a = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = 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_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]+}}* [[ATA]],
+ // TCHECK: [[A_IN_VAL2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[A_IN_VAL2]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: ret void
+
+
+ #pragma omp target private(a, aa)
+ #pragma omp teams firstprivate(a,aa)
+ {
+ a = 1;
+ aa = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: [[A2TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2TA]],
+ // TCHECK: store i{{[0-9]+}} [[A2TA_VAL]], i{{[0-9]+}}* [[A2TE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+ // TCHECK: ret void
+
+ #pragma omp target
+ #pragma omp teams firstprivate(a,aa)
+ {
+ a = 1;
+ aa = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A2_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]],
+ // TCHECK-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[CONV2:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
+ // TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]],
+ // TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV2]],
+ // TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
+ // TCHECK: ret void
+
+ #pragma omp target firstprivate(a, aa)
+ #pragma omp teams firstprivate(a,aa)
+ {
+ a = 1;
+ aa = 1;
+
+ aa = a+1;
+ }
+
+ // check that we are not using the firstprivate parameter
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A2_IN:%.+]])
+ // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+ // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]],
+ // TCHECK-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[CONV2:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
+ // TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]+}}* [[ATA]],
+ // TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[CONV2]],
+ // TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2TA]],
+ // TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: [[A2TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2TA]],
+ // TCHECK: store i{{[0-9]+}} [[A2TA_VAL]], i{{[0-9]+}}* [[A2TE]],
+
+ // a = 1, aa = 1
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+
+ // aa = a+1
+ // TCHECK: [[ATE_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATE]],
+ // TCHECK: [[ATE_INC:%.+]] = add{{.+}} i{{[0-9]+}} [[ATE_VAL]], 1
+ // TCHECK: [[ATE_CONV:%.+]] = trunc {{.+}} to
+ // TCHECK: store i{{[0-9]+}} [[ATE_CONV]], i{{[0-9]+}}* [[A2TE]],
+ // TCHECK: ret void
+
+ #pragma omp target private(a, b, c, d)
+ #pragma omp teams firstprivate(a, b, c, d)
+ {
+ a = 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+// CARLO
+ // TCHECK: define void @__omp_offloading_{{.+}}()
+ // TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTA:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTA:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTA:%.+]] = alloca [[TT]],
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTE:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTE:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTE:%.+]] = alloca [[TT]],
+ // TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: [[BTE_CPY:%.+]] = bitcast [10 x float]* [[BTE]] to i8*
+ // TCHECK: [[BTA_CPY:%.+]] = bitcast [10 x float]* [[BTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[BTE_CPY]], i8* [[BTA_CPY]],{{.+}})
+ // TCHECK: [[CTE_CPY:%.+]] = bitcast [5 x [10 x double]]* [[CTE]] to i8*
+ // TCHECK: [[CTA_CPY:%.+]] = bitcast [5 x [10 x double]]* [[CTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[CTE_CPY]], i8* [[CTA_CPY]],{{.+}})
+ // TCHECK: [[DTE_CPY:%.+]] = bitcast [[TT]]* [[DTE]] to i8*
+ // TCHECK: [[DTA_CPY:%.+]] = bitcast [[TT]]* [[DTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[DTE_CPY]], i8* [[DTA_CPY]],{{.+}})
+
+ // TCHECK: store {{.+}} [[ATE]],
+ // TCHECK: [[B_IDX:%.+]] = getelementptr{{.+}} [[BTE]], {{.+}},
+ // TCHECK: store {{.+}} [[B_IDX]],
+ // TCHECK: [[C_IDX1:%.+]] = getelementptr {{.+}} [[CTE]], {{.+}}
+ // TCHECK: [[C_IDX2:%.+]] = getelementptr {{.+}} [[C_IDX1]], {{.+}}
+ // TCHECK: store {{.+}} [[C_IDX2]],
+ // TCHECK: [[D_X:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_X]],
+ // TCHECK: [[D_Y:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_Y]],
+
+ #pragma omp target
+ #pragma omp teams firstprivate(a, b, c, d)
+ {
+ a = 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x float]*{{.+}} [[B_IN:%.+]], [5 x [10 x double]]*{{.+}} [[C_IN:%.+]], [[TT]]*{{.+}} [[D_IN:%.+]])
+ // TCHECK: [[A_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: [[A:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[B:%.+]] = alloca [10 x float],
+ // TCHECK: [[C:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[D:%.+]] = alloca [[TT]],
+ // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_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-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_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 %struct.TT*, %struct.TT** [[D_ADDR]],
+
+ // TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]],
+ // TCHECK: [[B_CPY:%.+]] = bitcast [10 x float]* [[B]] to i8*
+ // TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+ // TCHECK: [[C_CPY:%.+]] = bitcast [5 x [10 x double]]* [[C]] to i8*
+ // TCHECK: [[C_IN_CPY:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_CPY]], i8* [[C_IN_CPY]],{{.+}})
+ // TCHECK: [[D_CPY:%.+]] = bitcast [[TT]]* [[D]] to i8*
+ // TCHECK: [[D_IN_CPY:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[D_CPY]], i8* [[D_IN_CPY]],{{.+}})
+
+ // TCHECK: store {{.+}} [[A]],
+ // TCHECK: [[B_IDX:%.+]] = getelementptr{{.+}} [[B]], {{.+}},
+ // TCHECK: store {{.+}} [[B_IDX]],
+ // TCHECK: [[C_IDX1:%.+]] = getelementptr {{.+}} [[C]], {{.+}}
+ // TCHECK: [[C_IDX2:%.+]] = getelementptr {{.+}} [[C_IDX1]], {{.+}}
+ // TCHECK: store {{.+}} [[C_IDX2]],
+ // TCHECK: [[D_X:%.+]] = getelementptr {{.+}} [[D]], {{.+}}
+ // TCHECK: store {{.+}} [[D_X]],
+ // TCHECK: [[D_Y:%.+]] = getelementptr {{.+}} [[D]], {{.+}}
+ // TCHECK: store {{.+}} [[D_Y]],
+
+ #pragma omp target firstprivate(a, b, c, d)
+ #pragma omp teams firstprivate(a, b, c, d)
+ {
+ a = 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x float]*{{.+}} [[B_IN:%.+]], [5 x [10 x double]]*{{.+}} [[C_IN:%.+]], [[TT]]*{{.+}} [[D_IN:%.+]])
+// TCHECK: [[A_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: [[ATA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTA:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTA:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTA:%.+]] = alloca [[TT]],
+ // TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[BTE:%.+]] = alloca [10 x float],
+ // TCHECK: [[CTE:%.+]] = alloca [5 x [10 x double]],
+ // TCHECK: [[DTE:%.+]] = alloca [[TT]],
+
+ // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_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-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[A_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 %struct.TT*, %struct.TT** [[D_ADDR]],
+
+ // TCHECK-64: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[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]+}}* [[ATA]],
+ // TCHECK: [[B_CPY:%.+]] = bitcast [10 x float]* [[BTA]] to i8*
+ // TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+ // TCHECK: [[C_CPY:%.+]] = bitcast [5 x [10 x double]]* [[CTA]] to i8*
+ // TCHECK: [[C_IN_CPY:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_CPY]], i8* [[C_IN_CPY]],{{.+}})
+ // TCHECK: [[D_CPY:%.+]] = bitcast [[TT]]* [[DTA]] to i8*
+ // TCHECK: [[D_IN_CPY:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[D_CPY]], i8* [[D_IN_CPY]],{{.+}})
+
+ // TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+ // TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+ // TCHECK: [[BTE_CPY:%.+]] = bitcast [10 x float]* [[BTE]] to i8*
+ // TCHECK: [[BTA_CPY:%.+]] = bitcast [10 x float]* [[BTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[BTE_CPY]], i8* [[BTA_CPY]],{{.+}})
+ // TCHECK: [[CTE_CPY:%.+]] = bitcast [5 x [10 x double]]* [[CTE]] to i8*
+ // TCHECK: [[CTA_CPY:%.+]] = bitcast [5 x [10 x double]]* [[CTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[CTE_CPY]], i8* [[CTA_CPY]],{{.+}})
+ // TCHECK: [[DTE_CPY:%.+]] = bitcast [[TT]]* [[DTE]] to i8*
+ // TCHECK: [[DTA_CPY:%.+]] = bitcast [[TT]]* [[DTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[DTE_CPY]], i8* [[DTA_CPY]],{{.+}})
+
+ // TCHECK: store {{.+}} [[ATE]],
+ // TCHECK: [[B_IDX:%.+]] = getelementptr{{.+}} [[BTE]], {{.+}},
+ // TCHECK: store {{.+}} [[B_IDX]],
+ // TCHECK: [[C_IDX1:%.+]] = getelementptr {{.+}} [[CTE]], {{.+}}
+ // TCHECK: [[C_IDX2:%.+]] = getelementptr {{.+}} [[C_IDX1]], {{.+}}
+ // TCHECK: store {{.+}} [[C_IDX2]],
+ // TCHECK: [[D_X:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_X]],
+ // TCHECK: [[D_Y:%.+]] = getelementptr {{.+}} [[DTE]], {{.+}}
+ // TCHECK: store {{.+}} [[D_Y]],
+
+ #pragma omp target
+ #pragma omp teams firstprivate(ptr)
+ {
+ ptr[0]++;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
+ // TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
+ // TCHECK: [[PTR:%.+]] = alloca double*,
+ // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]],
+ // TCHECK: [[PTR_ADDR_REF:%.+]] = load double*, double** [[PTR_ADDR]],
+ // TCHECK: store double* [[PTR_ADDR_REF]], double** [[PTR]],
+ // TCHECK: load double*, double** [[PTR]],
+ // TCHECK-NOT: [[PTR_ADDR]]
+
+ #pragma omp target firstprivate(ptr)
+ #pragma omp teams firstprivate(ptr)
+ {
+ ptr[0]++;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
+ // TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
+ // TCHECK: [[PTRTA:%.+]] = alloca double*,
+ // TCHECK: [[PTRTE:%.+]] = alloca double*,
+ // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]],
+ // TCHECK: [[PTR_ADDR_REF:%.+]] = load double*, double** [[PTR_ADDR]],
+ // TCHECK: store double* [[PTR_ADDR_REF]], double** [[PTRTA]],
+ // TCHECK: [[PTRTA_VAL:%.+]] = load double*, double** [[PTRTA]],
+ // TCHECK: store double* [[PTRTA_VAL]], double** [[PTRTE]],
+ // TCHECK: load double*, double** [[PTRTE]],
+ // TCHECK-NOT: [[PTRTA]]
+ // TCHECK-NOT: [[PTR_ADDR]]
+
+ return a;
+}
+
+
+template<typename tx>
+tx ftemplate(int n) {
+ tx a = 0;
+ short aa = 0;
+ tx b[10];
+
+#pragma omp target private(a,aa,b)
+#pragma omp teams firstprivate(a,aa,b)
+ {
+ a = 1;
+ aa = 1;
+ b[2] = 1;
+ }
+
+#pragma omp target
+#pragma omp teams firstprivate(a,aa,b)
+ {
+ a = 1;
+ aa = 1;
+ b[2] = 1;
+ }
+
+#pragma omp target firstprivate(a,aa,b)
+#pragma omp teams firstprivate(a,aa,b)
+ {
+ a = 1;
+ aa = 1;
+ b[2] = 1;
+ }
+
+ return a;
+}
+
+static
+int fstatic(int n) {
+ int a = 0;
+ short aa = 0;
+ char aaa = 0;
+ int b[10];
+
+#pragma omp target private(a,aa,aaa,b)
+#pragma omp teams firstprivate(a,aa,aaa,b)
+ {
+ a = 1;
+ aa = 1;
+ aaa = 1;
+ b[2] = 1;
+ }
+
+// TCHECK: define void @__omp_offloading_{{.+}}()
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+// TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+// TCHECK: [[A2TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2TA]],
+// TCHECK: store i{{[0-9]+}} [[A2TA_VAL]], i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[A3TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A3TA]],
+// TCHECK: store i{{[0-9]+}} [[A3TA_VAL]], i{{[0-9]+}}* [[A3TE]],
+// TCHECK: [[BTE_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTE]] to i8*
+// TCHECK: [[BTA_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTA]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[BTE_CPY]], i8* [[BTA_CPY]],{{.+}})
+
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+#pragma omp target
+#pragma omp teams firstprivate(a,aa,aaa,b)
+ {
+ a = 1;
+ aa = 1;
+ aaa = 1;
+ b[2] = 1;
+ }
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A2_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_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: [[A2_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[A3_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A3_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// 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]],
+// TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2_CONV]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2]],
+// TCHECK: [[A3_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A3_CONV]],
+// TCHECK: store i{{[0-9]+}} [[A3_IN_VAL]], i{{[0-9]+}}* [[A3]],
+// TCHECK: [[B_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B]] to i8*
+// TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+#pragma omp target firstprivate(a,aa,aaa,b)
+#pragma omp teams firstprivate(a,aa,aaa,b)
+ {
+ a = 1;
+ aa = 1;
+ aaa = 1;
+ b[2] = 1;
+ }
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[A2_IN:%.+]], i{{[0-9]+}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A3TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_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: [[A2_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A2_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[A3_CONV:%.+]] = bitcast i{{[0-9]+}}* [[A3_ADDR]] to i{{[0-9]+}}*
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// 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]+}}* [[ATA]],
+// TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2_CONV]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2TA]],
+// TCHECK: [[A3_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A3_CONV]],
+// TCHECK: store i{{[0-9]+}} [[A3_IN_VAL]], i{{[0-9]+}}* [[A3TA]],
+// TCHECK: [[B_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTA]] to i8*
+// TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+
+// TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+// TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+// TCHECK: [[A2TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2TA]],
+// TCHECK: store i{{[0-9]+}} [[A2TA_VAL]], i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[A3TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A3TA]],
+// TCHECK: store i{{[0-9]+}} [[A3TA_VAL]], i{{[0-9]+}}* [[A3TE]],
+// TCHECK: [[BTE_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTE]] to i8*
+// TCHECK: [[BTA_IN_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTA]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[BTE_CPY]], i8* [[BTA_IN_CPY]],{{.+}})
+
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+ return a;
+}
+
+struct S1 {
+ double a;
+
+ int r1(int n){
+ int b = n+1;
+ short int c[2][5];
+
+#pragma omp target private(b,c)
+#pragma omp teams firstprivate(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]])
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[BTA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTA:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: [[BTE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTE:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+
+ // TCHECK: [[BTA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[BTA]],
+ // TCHECK: store i{{[0-9]+}} [[BTA_VAL]], i{{[0-9]+}}* [[BTE]],
+ // TCHECK: [[CTE_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[CTE]] to i8*
+ // TCHECK: [[CTA_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[CTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[CTE_CPY]], i8* [[CTA_CPY]],{{.+}})
+
+ // this->a = (double)b + 1.5;
+ // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[BTE]],
+ // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
+ // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
+ // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
+
+ // c[1][1] = ++a;
+ // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
+ // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
+ // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds [2 x [5 x i{{[0-9]+}}]], [2 x [5 x i{{[0-9]+}}]]* [[CTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[C_1_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
+ // TCHECK: ret void
+
+#pragma omp target
+#pragma omp teams firstprivate(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[B_IN:%.+]], [2 x [5 x i{{[0-9]+}}]]*{{.+}} [[C_IN:%.+]], [[S1]]* [[TH:%.+]])
+ // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[C_ADDR:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]]*,
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[B:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[C:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]],
+ // TCHECK: store [2 x [5 x i{{[0-9]+}}]]* [[C_IN]], [2 x [5 x i{{[0-9]+}}]]** [[C_ADDR]],
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK-64: [[B_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[C_ADDR_REF:%.+]] = load [2 x [5 x i{{[0-9]+}}]]*, [2 x [5 x i{{[0-9]+}}]]** [[C_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+ // TCHECK-64: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_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]],
+ // TCHECK: [[C_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[C]] to i8*
+ // TCHECK: [[C_IN_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_CPY]], i8* [[C_IN_CPY]],{{.+}})
+
+ // this->a = (double)b + 1.5;
+ // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B]],
+ // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
+ // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
+ // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
+
+ // c[1][1] = ++a;
+ // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
+ // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
+ // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds [2 x [5 x i{{[0-9]+}}]], [2 x [5 x i{{[0-9]+}}]]* [[C]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[C_1_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
+ // TCHECK: ret void
+
+#pragma omp target firstprivate(b,c)
+#pragma omp teams firstprivate(b,c)
+ {
+ this->a = (double)b + 1.5;
+ c[1][1] = ++a;
+ }
+
+ return c[1][1] + (int)b;
+ }
+
+ // TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[B_IN:%.+]], [2 x [5 x i{{[0-9]+}}]]*{{.+}} [[C_IN:%.+]], [[S1]]* [[TH:%.+]])
+ // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[C_ADDR:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]]*,
+ // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
+ // TCHECK: [[BTA:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTA:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+ // TCHECK: [[BTE:%.+]] = alloca i{{[0-9]+}},
+ // TCHECK: [[CTE:%.+]] = alloca [2 x [5 x i{{[0-9]+}}]],
+
+ // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]],
+ // TCHECK: store [2 x [5 x i{{[0-9]+}}]]* [[C_IN]], [2 x [5 x i{{[0-9]+}}]]** [[C_ADDR]],
+ // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
+ // TCHECK-64: [[B_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}*
+ // TCHECK: [[C_ADDR_REF:%.+]] = load [2 x [5 x i{{[0-9]+}}]]*, [2 x [5 x i{{[0-9]+}}]]** [[C_ADDR]],
+ // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
+ // TCHECK-64: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_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]+}}* [[BTA]],
+ // TCHECK: [[C_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[CTA]] to i8*
+ // TCHECK: [[C_IN_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[C_ADDR_REF]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_CPY]], i8* [[C_IN_CPY]],{{.+}})
+
+ // TCHECK: [[B_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[BTA]],
+ // TCHECK: store i{{[0-9]+}} [[B_IN_VAL]], i{{[0-9]+}}* [[BTE]],
+ // TCHECK: [[C_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[CTE]] to i8*
+ // TCHECK: [[C_IN_CPY:%.+]] = bitcast [2 x [5 x i{{[0-9]+}}]]* [[CTA]] to i8*
+ // TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[C_CPY]], i8* [[C_IN_CPY]],{{.+}})
+
+ // this->a = (double)b + 1.5;
+ // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[BTE]],
+ // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
+ // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
+ // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],
+
+ // c[1][1] = ++a;
+ // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
+ // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
+ // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
+ // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
+ // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds [2 x [5 x i{{[0-9]+}}]], [2 x [5 x i{{[0-9]+}}]]* [[CTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[C_1_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
+ // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],
+ // 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_{{.+}}()
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+
+// TCHECK: [[ATA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+// TCHECK: store i{{[0-9]+}} [[ATA_VAL]], i{{[0-9]+}}* [[ATE]],
+// TCHECK: [[A2TA_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2TA]],
+// TCHECK: store i{{[0-9]+}} [[A2TA_VAL]], i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[BTE_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTE]] to i8*
+// TCHECK: [[BTA_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTA]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[BTE_CPY]], i8* [[BTA_CPY]],{{.+}})
+
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}} [[A_IN:%.+]], i{{[0-9]+}}*{{.+}} [[A2:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}*,
+// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}}*,
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}}* [[A_IN]], i{{[0-9]+}}** [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}}* [[A2_IN]], i{{[0-9]+}}** [[A2_ADDR]],
+// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK: [[A_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A_ADDR]],
+// TCHECK: [[A2_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A2_ADDR]],
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR_REF]],
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[A]],
+// TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2_ADDR_REF]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2]],
+// TCHECK: [[B_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B]] to i8*
+// TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+// TCHECK: define void @__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}} [[A_IN:%.+]], i{{[0-9]+}}*{{.+}} [[A2:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
+// TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}*,
+// TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}}*,
+// TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
+// TCHECK: [[ATA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TA:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTA:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: [[ATE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[A2TE:%.+]] = alloca i{{[0-9]+}},
+// TCHECK: [[BTE:%.+]] = alloca [10 x i{{[0-9]+}}],
+// TCHECK: store i{{[0-9]+}}* [[A_IN]], i{{[0-9]+}}** [[A_ADDR]],
+// TCHECK: store i{{[0-9]+}}* [[A2_IN]], i{{[0-9]+}}** [[A2_ADDR]],
+// TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK: [[A_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A_ADDR]],
+// TCHECK: [[A2_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[A2_ADDR]],
+// TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
+// TCHECK: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A_ADDR_REF]],
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[ATA]],
+// TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2_ADDR_REF]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2TA]],
+// TCHECK: [[B_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTA]] to i8*
+// TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+
+// TCHECK: [[A_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[ATA]],
+// TCHECK: store i{{[0-9]+}} [[A_IN_VAL]], i{{[0-9]+}}* [[ATE]],
+// TCHECK: [[A2_IN_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[A2TA]],
+// TCHECK: store i{{[0-9]+}} [[A2_IN_VAL]], i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[B_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTE]] to i8*
+// TCHECK: [[B_IN_CPY:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[BTA]] to i8*
+// TCHECK: call void @llvm.memcpy.{{.+}}(i8* [[B_CPY]], i8* [[B_IN_CPY]],{{.+}})
+
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[ATE]],
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A2TE]],
+// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[BTE]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
+// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
+// TCHECK: ret void
+
+#endif
+
Index: lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- lib/CodeGen/CGStmtOpenMP.cpp
+++ lib/CodeGen/CGStmtOpenMP.cpp
@@ -3070,7 +3070,13 @@
(void)PrivateScope.Privatize();
CGF.EmitStmt(cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());
};
- emitCommonOMPTeamsDirective(*this, S, OMPD_teams, CodeGen);
+
+ if (getLangOpts().OpenMPIsDevice && (
+ CGM.getTarget().getTriple().getArch() == llvm::Triple::nvptx ||
+ CGM.getTarget().getTriple().getArch() == llvm::Triple::nvptx64))
+ CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_teams, CodeGen);
+ else
+ emitCommonOMPTeamsDirective(*this, S, OMPD_teams, CodeGen);
}
void CodeGenFunction::EmitOMPCancellationPointDirective(
Index: lib/CodeGen/CGOpenMPRuntimeNVPTX.h
===================================================================
--- lib/CodeGen/CGOpenMPRuntimeNVPTX.h
+++ lib/CodeGen/CGOpenMPRuntimeNVPTX.h
@@ -31,33 +31,6 @@
/// \param ThreadLimit An integer expression of threads.
void emitNumTeamsClause(CodeGenFunction &CGF, const Expr *NumTeams,
const Expr *ThreadLimit, SourceLocation Loc) override;
-
- /// \brief Emits inlined function for the specified OpenMP parallel
- // directive but an inlined function for teams.
- /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID,
- /// kmp_int32 BoundID, struct context_vars*).
- /// \param D OpenMP directive.
- /// \param ThreadIDVar Variable for thread id in the current OpenMP region.
- /// \param InnermostKind Kind of innermost directive (for simple directives it
- /// is a directive itself, for combined - its innermost directive).
- /// \param CodeGen Code generation sequence for the \a D directive.
- llvm::Value *emitParallelOrTeamsOutlinedFunction(
- const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen)
- override;
-
- /// \brief Emits code for teams call of the \a OutlinedFn with
- /// variables captured in a record which address is stored in \a
- /// CapturedStruct.
- /// \param OutlinedFn Outlined function to be run by team masters. Type of
- /// this function is void(*)(kmp_int32 *, kmp_int32, struct context_vars*).
- /// \param CapturedVars A pointer to the record with the references to
- /// variables used in \a OutlinedFn function.
- ///
- void emitTeamsCall(CodeGenFunction &CGF, const OMPExecutableDirective &D,
- SourceLocation Loc, llvm::Value *OutlinedFn,
- ArrayRef<llvm::Value *> CapturedVars) override;
-
};
} // CodeGen namespace.
Index: lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
+++ lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
@@ -27,31 +27,3 @@
const Expr *NumTeams,
const Expr *ThreadLimit,
SourceLocation Loc) {}
-
-llvm::Value *CGOpenMPRuntimeNVPTX::emitParallelOrTeamsOutlinedFunction(
- const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
- OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
-
- llvm::Function *OutlinedFun = nullptr;
- if (isa<OMPTeamsDirective>(D)) {
- // no outlining happening for teams
- } else
- llvm_unreachable("parallel directive is not yet supported for nvptx "
- "backend.");
-
- return OutlinedFun;
-}
-
-void CGOpenMPRuntimeNVPTX::emitTeamsCall(CodeGenFunction &CGF,
- const OMPExecutableDirective &D,
- SourceLocation Loc,
- llvm::Value *OutlinedFn,
- ArrayRef<llvm::Value *> CapturedVars) {
-
- // just emit the statements in the teams region inlined
- auto &&CodeGen = [&D](CodeGenFunction &CGF) {
- CGF.EmitStmt(cast<CapturedStmt>(D.getAssociatedStmt())->getCapturedStmt());
- };
-
- emitInlinedDirective(CGF, OMPD_teams, CodeGen);
-}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
