Index: lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp +++ lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp @@ -50,6 +50,11 @@ // just emit the statements in the teams region inlined auto &&CodeGen = [&D](CodeGenFunction &CGF) { + CodeGenFunction::OMPPrivateScope PrivateScope(CGF); + (void)CGF.EmitOMPFirstprivateClause(D, PrivateScope); + CGF.EmitOMPPrivateClause(D, PrivateScope); + (void)PrivateScope.Privatize(); + CGF.EmitStmt(cast(D.getAssociatedStmt())->getCapturedStmt()); }; 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 +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 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 +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(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: 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 +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 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 +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(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 +