diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index bb987d87548d140c4441ba55a39ade7a328036a6..96f5a655956ce5e7da7591cda38f5607ec806c6e 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -1093,7 +1093,8 @@ void CodeGenFunction::EmitOMPReductionClauseFinal( bool WithNowait = D.getSingleClause() || isOpenMPParallelDirective(D.getDirectiveKind()) || D.getDirectiveKind() == OMPD_simd; - bool SimpleReduction = D.getDirectiveKind() == OMPD_simd; + bool SimpleReduction = D.getDirectiveKind() == OMPD_simd || + D.getDirectiveKind() == OMPD_distribute_simd; // Emit nowait reduction if nowait clause is present or directive is a // parallel directive (it always has implicit barrier). CGM.getOpenMPRuntime().emitReduction( @@ -1215,9 +1216,13 @@ void CodeGenFunction::EmitOMPLoopBody(const OMPLoopDirective &D, EmitIgnoredExpr(I); } // Update the linear variables. - for (const auto *C : D.getClausesOfKind()) { - for (auto *U : C->updates()) - EmitIgnoredExpr(U); + // In distribute directives only loop counters may be marked as linear, no + // need to generate the code for them. + if (!isOpenMPDistributeDirective(D.getDirectiveKind())) { + for (const auto *C : D.getClausesOfKind()) { + for (auto *U : C->updates()) + EmitIgnoredExpr(U); + } } // On a continue in the body, jump to the end. @@ -2054,14 +2059,11 @@ void CodeGenFunction::EmitOMPDistributeParallelForSimdDirective( void CodeGenFunction::EmitOMPDistributeSimdDirective( const OMPDistributeSimdDirective &S) { + auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) { + CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc()); + }; OMPLexicalScope Scope(*this, S, /*AsInlined=*/true); - CGM.getOpenMPRuntime().emitInlinedDirective( - *this, OMPD_distribute_simd, - [&S](CodeGenFunction &CGF, PrePostActionTy &) { - OMPLoopScope PreInitScope(CGF, S); - CGF.EmitStmt( - cast(S.getAssociatedStmt())->getCapturedStmt()); - }); + CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_simd, CodeGen); } void CodeGenFunction::EmitOMPTargetSimdDeviceFunction( @@ -3075,6 +3077,7 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, incrementProfileCounter(&S); } + emitAlignedClause(*this, S); // Emit 'then' code. { // Emit helper vars inits. @@ -3096,14 +3099,17 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, OMPPrivateScope LoopScope(*this); if (EmitOMPFirstprivateClause(S, LoopScope)) { - // Emit implicit barrier to synchronize threads and avoid data races on - // initialization of firstprivate variables and post-update of + // Emit implicit barrier to synchronize threads and avoid data races + // on initialization of firstprivate variables and post-update of // lastprivate variables. CGM.getOpenMPRuntime().emitBarrierCall( - *this, S.getLocStart(), OMPD_unknown, /*EmitChecks=*/false, - /*ForceSimpleCall=*/true); + *this, S.getLocStart(), OMPD_unknown, /*EmitChecks=*/false, + /*ForceSimpleCall=*/true); } EmitOMPPrivateClause(S, LoopScope); + if (isOpenMPSimdDirective(S.getDirectiveKind()) && + !isOpenMPParallelDirective(S.getDirectiveKind())) + EmitOMPReductionClauseInit(S, LoopScope); HasLastprivateClause = EmitOMPLastprivateClauseInit(S, LoopScope); EmitOMPPrivateLoopCounters(S, LoopScope); (void)LoopScope.Privatize(); @@ -3116,8 +3122,8 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, if (const auto *Ch = C->getChunkSize()) { Chunk = EmitScalarExpr(Ch); Chunk = EmitScalarConversion(Chunk, Ch->getType(), - S.getIterationVariable()->getType(), - S.getLocStart()); + S.getIterationVariable()->getType(), + S.getLocStart()); } } const unsigned IVSize = getContext().getTypeSize(IVExpr->getType()); @@ -3133,6 +3139,8 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, // league. The size of the chunks is unspecified in this case. if (RT.isStaticNonchunked(ScheduleKind, /* Chunked */ Chunk != nullptr)) { + if (isOpenMPSimdDirective(S.getDirectiveKind())) + EmitOMPSimdInit(S, /*IsMonotonic=*/true); CGOpenMPRuntime::StaticRTInput StaticInit( IVSize, IVSigned, /* Ordered = */ false, IL.getAddress(), LB.getAddress(), UB.getAddress(), ST.getAddress()); @@ -3174,13 +3182,38 @@ void CodeGenFunction::EmitOMPDistributeLoop(const OMPLoopDirective &S, EmitOMPDistributeOuterLoop(ScheduleKind, S, LoopScope, LoopArguments, CodeGenLoop); } - + if (isOpenMPSimdDirective(S.getDirectiveKind())) { + EmitOMPSimdFinal(S, [&](CodeGenFunction &CGF) -> llvm::Value * { + return CGF.Builder.CreateIsNotNull( + CGF.EmitLoadOfScalar(IL, S.getLocStart())); + }); + } + OpenMPDirectiveKind ReductionKind = OMPD_unknown; + if (isOpenMPParallelDirective(S.getDirectiveKind()) && + isOpenMPSimdDirective(S.getDirectiveKind())) { + ReductionKind = OMPD_parallel_for_simd; + } else if (isOpenMPParallelDirective(S.getDirectiveKind())) { + ReductionKind = OMPD_parallel_for; + } else if (isOpenMPSimdDirective(S.getDirectiveKind())) { + ReductionKind = OMPD_simd; + } else if (!isOpenMPTeamsDirective(S.getDirectiveKind()) && + S.hasClausesOfKind()) { + llvm_unreachable( + "No reduction clauses is allowed in distribute directive."); + } + EmitOMPReductionClauseFinal(S, ReductionKind); + // Emit post-update of the reduction variables if IsLastIter != 0. + emitPostUpdateForReductionClause( + *this, S, [&](CodeGenFunction &CGF) -> llvm::Value * { + return CGF.Builder.CreateIsNotNull( + CGF.EmitLoadOfScalar(IL, S.getLocStart())); + }); // Emit final copy of the lastprivate variables if IsLastIter != 0. - if (HasLastprivateClause) + if (HasLastprivateClause) { EmitOMPLastprivateClauseFinal( S, /*NoFinals=*/false, - Builder.CreateIsNotNull( - EmitLoadOfScalar(IL, S.getLocStart()))); + Builder.CreateIsNotNull(EmitLoadOfScalar(IL, S.getLocStart()))); + } } // We're now done with the loop, so jump to the continuation block. diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 2a7e2f1da49f24a4ac046452c724fc758ec34123..4c0fa7e617289d9a46e1832c9b0320bcc0782667 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -6887,14 +6887,24 @@ StmtResult Sema::ActOnOpenMPDistributeSimdDirective( // The point of exit cannot be a branch out of the structured block. // longjmp() and throw() must not violate the entry/exit criteria. CS->getCapturedDecl()->setNothrow(); + for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_distribute_simd); + ThisCaptureLevel > 1; --ThisCaptureLevel) { + CS = cast(CS->getCapturedStmt()); + // 1.2.2 OpenMP Language Terminology + // Structured block - An executable statement with a single entry at the + // top and a single exit at the bottom. + // The point of exit cannot be a branch out of the structured block. + // longjmp() and throw() must not violate the entry/exit criteria. + CS->getCapturedDecl()->setNothrow(); + } OMPLoopDirective::HelperExprs B; // In presence of clause 'collapse' with number of loops, it will // define the nested loops number. unsigned NestedLoopCount = CheckOpenMPLoop(OMPD_distribute_simd, getCollapseNumberExpr(Clauses), - nullptr /*ordered not a clause on distribute*/, AStmt, - *this, *DSAStack, VarsWithImplicitDSA, B); + nullptr /*ordered not a clause on distribute*/, CS, *this, + *DSAStack, VarsWithImplicitDSA, B); if (NestedLoopCount == 0) return StmtError(); diff --git a/clang/test/OpenMP/distribute_parallel_for_simd_lastprivate_codegen.cpp b/clang/test/OpenMP/distribute_parallel_for_simd_lastprivate_codegen.cpp index 0cf6d115f1bdca9f84dfb35924ac196f456ac6db..5b34622c5579eea7b5b3f2b26add7192b404b461 100644 --- a/clang/test/OpenMP/distribute_parallel_for_simd_lastprivate_codegen.cpp +++ b/clang/test/OpenMP/distribute_parallel_for_simd_lastprivate_codegen.cpp @@ -102,6 +102,15 @@ int main() { // LAMBDA-32: call{{.+}} @__kmpc_fork_call({{.+}}, {{.+}}, {{.+}}[[OMP_PARFOR_OUTLINED:@.+]] to void ({{.+}})*), {{.+}}, {{.+}}, {{.+}} [[G_PRIVATE]], {{.+}} [[G1_PAR]], {{.+}} [[SVAR_PRIVATE]], {{.+}} [[SFVAR_PRIVATE]]) // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( + // linear counter + // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], + // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 + // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_COUNTER_BLOCK:.+]], label %[[OMP_COUNTER_DONE:.+]] + // LAMBDA: [[OMP_COUNTER_BLOCK]]: + // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* % + // LAMBDA: br label %[[OMP_COUNTER_DONE]] + // LAMBDA: [[OMP_COUNTER_DONE]]: + // lastprivate // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 diff --git a/clang/test/OpenMP/distribute_simd_codegen.cpp b/clang/test/OpenMP/distribute_simd_codegen.cpp new file mode 100644 index 0000000000000000000000000000000000000000..d8f18c3e431b2c4e25dc0b7d1ea8f24e8dc8677f --- /dev/null +++ b/clang/test/OpenMP/distribute_simd_codegen.cpp @@ -0,0 +1,269 @@ +// Test host codegen. +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix HCHECK +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix HCHECK +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix HCHECK + +// Test target codegen - host bc file has to be created first. (no significant differences with host version of target region) +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s + +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +// CHECK-DAG: %ident_t = type { i32, i32, i32, i32, i8* } +// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" +// CHECK-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } +// CHECK-DAG: [[DEF_LOC_DISTRIBUTE_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2050, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } + +// CHECK-LABEL: define {{.*void}} @{{.*}}without_schedule_clause{{.*}}(float* {{.+}}, float* {{.+}}, float* {{.+}}, float* {{.+}}) +void without_schedule_clause(float *a, float *b, float *c, float *d) { + #pragma omp target + #pragma omp teams + #pragma omp distribute simd simdlen(8) aligned(a) + for (int i = 33; i < 32000000; i += 7) { + a[i] = b[i] * c[i] * d[i]; + } +} + +// CHECK: define {{.*}}void @{{.+}}(i32* noalias [[GBL_TIDP:%.+]], i32* noalias [[BND_TID:%.+]], float** dereferenceable({{[0-9]+}}) [[APTR:%.+]], float** dereferenceable({{[0-9]+}}) [[BPTR:%.+]], float** dereferenceable({{[0-9]+}}) [[CPTR:%.+]], float** dereferenceable({{[0-9]+}}) [[DPTR:%.+]]) +// CHECK: [[TID_ADDR:%.+]] = alloca i32* +// CHECK: [[IV:%.+iv]] = alloca i32 +// CHECK: [[LB:%.+lb]] = alloca i32 +// CHECK: [[UB:%.+ub]] = alloca i32 +// CHECK: [[ST:%.+stride]] = alloca i32 +// CHECK: [[LAST:%.+last]] = alloca i32 +// CHECK-DAG: store i32* [[GBL_TIDP]], i32** [[TID_ADDR]] +// CHECK-DAG: call void @llvm.assume( +// CHECK-DAG: store i32 0, i32* [[LB]] +// CHECK-DAG: store i32 4571423, i32* [[UB]] +// CHECK-DAG: store i32 1, i32* [[ST]] +// CHECK-DAG: store i32 0, i32* [[LAST]] +// CHECK-DAG: [[GBL_TID:%.+]] = load i32*, i32** [[TID_ADDR]] +// CHECK-DAG: [[GBL_TIDV:%.+]] = load i32, i32* [[GBL_TID]] +// CHECK: call void @__kmpc_for_static_init_{{.+}}(%ident_t* [[DEF_LOC_DISTRIBUTE_0]], i32 [[GBL_TIDV]], i32 92, i32* %.omp.is_last, i32* %.omp.lb, i32* %.omp.ub, i32* %.omp.stride, i32 1, i32 1) +// CHECK-DAG: [[UBV0:%.+]] = load i32, i32* [[UB]] +// CHECK-DAG: [[USWITCH:%.+]] = icmp sgt i32 [[UBV0]], 4571423 +// CHECK: br i1 [[USWITCH]], label %[[BBCT:.+]], label %[[BBCF:.+]] +// CHECK-DAG: [[BBCT]]: +// CHECK-DAG: br label %[[BBCE:.+]] +// CHECK-DAG: [[BBCF]]: +// CHECK-DAG: [[UBV1:%.+]] = load i32, i32* [[UB]] +// CHECK-DAG: br label %[[BBCE]] +// CHECK: [[BBCE]]: +// CHECK: [[SELUB:%.+]] = phi i32 [ 4571423, %[[BBCT]] ], [ [[UBV1]], %[[BBCF]] ] +// CHECK: store i32 [[SELUB]], i32* [[UB]] +// CHECK: [[LBV0:%.+]] = load i32, i32* [[LB]] +// CHECK: store i32 [[LBV0]], i32* [[IV]] +// CHECK: br label %[[BBINNFOR:.+]] +// CHECK: [[BBINNFOR]]: +// CHECK: [[IVVAL0:%.+]] = load i32, i32* [[IV]] +// CHECK: [[UBV2:%.+]] = load i32, i32* [[UB]] +// CHECK: [[IVLEUB:%.+]] = icmp sle i32 [[IVVAL0]], [[UBV2]] +// CHECK: br i1 [[IVLEUB]], label %[[BBINNBODY:.+]], label %[[BBINNEND:.+]] +// CHECK: [[BBINNBODY]]: +// CHECK: {{.+}} = load i32, i32* [[IV]] +// ... loop body ... +// CHECK: br label %[[BBBODYCONT:.+]] +// CHECK: [[BBBODYCONT]]: +// CHECK: br label %[[BBINNINC:.+]] +// CHECK: [[BBINNINC]]: +// CHECK: [[IVVAL1:%.+]] = load i32, i32* [[IV]] +// CHECK: [[IVINC:%.+]] = add nsw i32 [[IVVAL1]], 1 +// CHECK: store i32 [[IVINC]], i32* [[IV]] +// CHECK: br label %[[BBINNFOR]] +// CHECK: [[BBINNEND]]: +// CHECK: br label %[[LPEXIT:.+]] +// CHECK: [[LPEXIT]]: +// CHECK: call void @__kmpc_for_static_fini(%ident_t* [[DEF_LOC_DISTRIBUTE_0]], i32 [[GBL_TIDV]]) +// CHECK: ret void + + +// CHECK-LABEL: define {{.*void}} @{{.*}}static_not_chunked{{.*}}(float* {{.+}}, float* {{.+}}, float* {{.+}}, float* {{.+}}) +void static_not_chunked(float *a, float *b, float *c, float *d) { + #pragma omp target + #pragma omp teams + #pragma omp distribute simd dist_schedule(static) safelen(32) + for (int i = 32000000; i > 33; i += -7) { + a[i] = b[i] * c[i] * d[i]; + } +} + +// CHECK: define {{.*}}void @.omp_outlined.{{.*}}(i32* noalias [[GBL_TIDP:%.+]], i32* noalias [[BND_TID:%.+]], float** dereferenceable({{[0-9]+}}) [[APTR:%.+]], float** dereferenceable({{[0-9]+}}) [[BPTR:%.+]], float** dereferenceable({{[0-9]+}}) [[CPTR:%.+]], float** dereferenceable({{[0-9]+}}) [[DPTR:%.+]]) +// CHECK: [[TID_ADDR:%.+]] = alloca i32* +// CHECK: [[IV:%.+iv]] = alloca i32 +// CHECK: [[LB:%.+lb]] = alloca i32 +// CHECK: [[UB:%.+ub]] = alloca i32 +// CHECK: [[ST:%.+stride]] = alloca i32 +// CHECK: [[LAST:%.+last]] = alloca i32 +// CHECK-DAG: store i32* [[GBL_TIDP]], i32** [[TID_ADDR]] +// CHECK-DAG: store i32 0, i32* [[LB]] +// CHECK-DAG: store i32 4571423, i32* [[UB]] +// CHECK-DAG: store i32 1, i32* [[ST]] +// CHECK-DAG: store i32 0, i32* [[LAST]] +// CHECK-DAG: [[GBL_TID:%.+]] = load i32*, i32** [[TID_ADDR]] +// CHECK-DAG: [[GBL_TIDV:%.+]] = load i32, i32* [[GBL_TID]] +// CHECK: call void @__kmpc_for_static_init_{{.+}}(%ident_t* [[DEF_LOC_DISTRIBUTE_0]], i32 [[GBL_TIDV]], i32 92, i32* %.omp.is_last, i32* %.omp.lb, i32* %.omp.ub, i32* %.omp.stride, i32 1, i32 1) +// CHECK-DAG: [[UBV0:%.+]] = load i32, i32* [[UB]] +// CHECK-DAG: [[USWITCH:%.+]] = icmp sgt i32 [[UBV0]], 4571423 +// CHECK: br i1 [[USWITCH]], label %[[BBCT:.+]], label %[[BBCF:.+]] +// CHECK-DAG: [[BBCT]]: +// CHECK-DAG: br label %[[BBCE:.+]] +// CHECK-DAG: [[BBCF]]: +// CHECK-DAG: [[UBV1:%.+]] = load i32, i32* [[UB]] +// CHECK-DAG: br label %[[BBCE]] +// CHECK: [[BBCE]]: +// CHECK: [[SELUB:%.+]] = phi i32 [ 4571423, %[[BBCT]] ], [ [[UBV1]], %[[BBCF]] ] +// CHECK: store i32 [[SELUB]], i32* [[UB]] +// CHECK: [[LBV0:%.+]] = load i32, i32* [[LB]] +// CHECK: store i32 [[LBV0]], i32* [[IV]] +// CHECK: br label %[[BBINNFOR:.+]] +// CHECK: [[BBINNFOR]]: +// CHECK: [[IVVAL0:%.+]] = load i32, i32* [[IV]] +// CHECK: [[UBV2:%.+]] = load i32, i32* [[UB]] +// CHECK: [[IVLEUB:%.+]] = icmp sle i32 [[IVVAL0]], [[UBV2]] +// CHECK: br i1 [[IVLEUB]], label %[[BBINNBODY:.+]], label %[[BBINNEND:.+]] +// CHECK: [[BBINNBODY]]: +// CHECK: {{.+}} = load i32, i32* [[IV]] +// ... loop body ... +// CHECK: br label %[[BBBODYCONT:.+]] +// CHECK: [[BBBODYCONT]]: +// CHECK: br label %[[BBINNINC:.+]] +// CHECK: [[BBINNINC]]: +// CHECK: [[IVVAL1:%.+]] = load i32, i32* [[IV]] +// CHECK: [[IVINC:%.+]] = add nsw i32 [[IVVAL1]], 1 +// CHECK: store i32 [[IVINC]], i32* [[IV]] +// CHECK: br label %[[BBINNFOR]] +// CHECK: [[BBINNEND]]: +// CHECK: br label %[[LPEXIT:.+]] +// CHECK: [[LPEXIT]]: +// CHECK: call void @__kmpc_for_static_fini(%ident_t* [[DEF_LOC_DISTRIBUTE_0]], i32 [[GBL_TIDV]]) +// CHECK: ret void + + +// CHECK-LABEL: define {{.*void}} @{{.*}}static_chunked{{.*}}(float* {{.+}}, float* {{.+}}, float* {{.+}}, float* {{.+}}) +void static_chunked(float *a, float *b, float *c, float *d) { + #pragma omp target + #pragma omp teams +#pragma omp distribute simd dist_schedule(static, 5) + for (unsigned i = 131071; i <= 2147483647; i += 127) { + a[i] = b[i] * c[i] * d[i]; + } +} + +// CHECK: define {{.*}}void @.omp_outlined.{{.*}}(i32* noalias [[GBL_TIDP:%.+]], i32* noalias [[BND_TID:%.+]], float** dereferenceable({{[0-9]+}}) [[APTR:%.+]], float** dereferenceable({{[0-9]+}}) [[BPTR:%.+]], float** dereferenceable({{[0-9]+}}) [[CPTR:%.+]], float** dereferenceable({{[0-9]+}}) [[DPTR:%.+]]) +// CHECK: [[TID_ADDR:%.+]] = alloca i32* +// CHECK: [[IV:%.+iv]] = alloca i32 +// CHECK: [[LB:%.+lb]] = alloca i32 +// CHECK: [[UB:%.+ub]] = alloca i32 +// CHECK: [[ST:%.+stride]] = alloca i32 +// CHECK: [[LAST:%.+last]] = alloca i32 +// CHECK-DAG: store i32* [[GBL_TIDP]], i32** [[TID_ADDR]] +// CHECK-DAG: store i32 0, i32* [[LB]] +// CHECK-DAG: store i32 16908288, i32* [[UB]] +// CHECK-DAG: store i32 1, i32* [[ST]] +// CHECK-DAG: store i32 0, i32* [[LAST]] +// CHECK-DAG: [[GBL_TID:%.+]] = load i32*, i32** [[TID_ADDR]] +// CHECK-DAG: [[GBL_TIDV:%.+]] = load i32, i32* [[GBL_TID]] +// CHECK: call void @__kmpc_for_static_init_{{.+}}(%ident_t* [[DEF_LOC_DISTRIBUTE_0]], i32 [[GBL_TIDV]], i32 91, i32* %.omp.is_last, i32* %.omp.lb, i32* %.omp.ub, i32* %.omp.stride, i32 1, i32 5) +// CHECK-DAG: [[UBV0:%.+]] = load i32, i32* [[UB]] +// CHECK-DAG: [[USWITCH:%.+]] = icmp ugt i32 [[UBV0]], 16908288 +// CHECK: br i1 [[USWITCH]], label %[[BBCT:.+]], label %[[BBCF:.+]] +// CHECK-DAG: [[BBCT]]: +// CHECK-DAG: br label %[[BBCE:.+]] +// CHECK-DAG: [[BBCF]]: +// CHECK-DAG: [[UBV1:%.+]] = load i32, i32* [[UB]] +// CHECK-DAG: br label %[[BBCE]] +// CHECK: [[BBCE]]: +// CHECK: [[SELUB:%.+]] = phi i32 [ 16908288, %[[BBCT]] ], [ [[UBV1]], %[[BBCF]] ] +// CHECK: store i32 [[SELUB]], i32* [[UB]] +// CHECK: [[LBV0:%.+]] = load i32, i32* [[LB]] +// CHECK: store i32 [[LBV0]], i32* [[IV]] +// CHECK: br label %[[BBINNFOR:.+]] +// CHECK: [[BBINNFOR]]: +// CHECK: [[IVVAL0:%.+]] = load i32, i32* [[IV]] +// CHECK: [[UBV2:%.+]] = load i32, i32* [[UB]] +// CHECK: [[IVLEUB:%.+]] = icmp ule i32 [[IVVAL0]], [[UBV2]] +// CHECK: br i1 [[IVLEUB]], label %[[BBINNBODY:.+]], label %[[BBINNEND:.+]] +// CHECK: [[BBINNBODY]]: +// CHECK: {{.+}} = load i32, i32* [[IV]] +// ... loop body ... +// CHECK: br label %[[BBBODYCONT:.+]] +// CHECK: [[BBBODYCONT]]: +// CHECK: br label %[[BBINNINC:.+]] +// CHECK: [[BBINNINC]]: +// CHECK: [[IVVAL1:%.+]] = load i32, i32* [[IV]] +// CHECK: [[IVINC:%.+]] = add i32 [[IVVAL1]], 1 +// CHECK: store i32 [[IVINC]], i32* [[IV]] +// CHECK: br label %[[BBINNFOR]] +// CHECK: [[BBINNEND]]: +// CHECK: br label %[[LPEXIT:.+]] +// CHECK: [[LPEXIT]]: +// CHECK: call void @__kmpc_for_static_fini(%ident_t* [[DEF_LOC_DISTRIBUTE_0]], i32 [[GBL_TIDV]]) +// CHECK: ret void + +// CHECK-LABEL: test_precond +void test_precond() { + char a = 0; char i; + #pragma omp target + #pragma omp teams + #pragma omp distribute simd linear(i) + for(i = a; i < 10; ++i); +} + +// a is passed as a parameter to the outlined functions +// CHECK: define {{.*}}void @.omp_outlined.{{.*}}(i32* noalias [[GBL_TIDP:%.+]], i32* noalias [[BND_TID:%.+]], i8* dereferenceable({{[0-9]+}}) [[APARM:%.+]]) +// CHECK: store i8* [[APARM]], i8** [[APTRADDR:%.+]] +// ..many loads of %0.. +// CHECK: [[A2:%.+]] = load i8*, i8** [[APTRADDR]] +// CHECK: [[AVAL0:%.+]] = load i8, i8* [[A2]] +// CHECK: store i8 [[AVAL0]], i8* [[CAP_EXPR:%.+]], +// CHECK: [[AVAL1:%.+]] = load i8, i8* [[CAP_EXPR]] +// CHECK: load i8, i8* [[CAP_EXPR]] +// CHECK: [[AVAL2:%.+]] = load i8, i8* [[CAP_EXPR]] +// CHECK: [[ACONV:%.+]] = sext i8 [[AVAL2]] to i32 +// CHECK: [[ACMP:%.+]] = icmp slt i32 [[ACONV]], 10 +// CHECK: br i1 [[ACMP]], label %[[PRECOND_THEN:.+]], label %[[PRECOND_END:.+]] +// CHECK: [[PRECOND_THEN]] +// CHECK: call void @__kmpc_for_static_init_4 +// CHECK: call void @__kmpc_for_static_fini +// CHECK: [[PRECOND_END]] + +// no templates for now, as these require special handling in target regions and/or declare target + +// HCHECK-LABEL: fint +// HCHECK: call {{.*}}i32 {{.+}}ftemplate +// HCHECK: ret i32 + +// HCHECK: load i16, i16* +// HCHECK: store i16 % +// HCHECK: call i32 @__tgt_target_teams( +// HCHECK: call void @__kmpc_for_static_init_4( +template +T ftemplate() { + short aa = 0; + +#pragma omp target +#pragma omp teams +#pragma omp distribute simd dist_schedule(static, aa) + for (int i = 0; i < 100; i++) { + } + return T(); +} + +int fint(void) { return ftemplate(); } + +#endif + +// CHECK: !{!"llvm.loop.vectorize.width", i32 8} +// CHECK: !{!"llvm.loop.vectorize.enable", i1 true} +// CHECK: !{!"llvm.loop.vectorize.width", i32 32} diff --git a/clang/test/OpenMP/distribute_simd_firstprivate_codegen.cpp b/clang/test/OpenMP/distribute_simd_firstprivate_codegen.cpp new file mode 100644 index 0000000000000000000000000000000000000000..8db26a4c05d7901de133273e1bbf1a7eb604294e --- /dev/null +++ b/clang/test/OpenMP/distribute_simd_firstprivate_codegen.cpp @@ -0,0 +1,384 @@ +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 + +// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +template +struct S { + T f; + S(T a) : f(a) {} + S() : f() {} + operator T() { return T(); } + ~S() {} +}; + +// CHECK: [[S_FLOAT_TY:%.+]] = type { float } +// CHECK: [[S_INT_TY:%.+]] = type { i{{[0-9]+}} } +template +T tmain() { + S test; + T t_var = T(); + T vec[] = {1, 2}; + S s_arr[] = {1, 2}; + S &var = test; + #pragma omp target + #pragma omp teams +#pragma omp distribute simd firstprivate(t_var, vec, s_arr, s_arr, var, var) + for (int i = 0; i < 2; ++i) { + vec[i] = t_var; + s_arr[i] = var; + } + return T(); +} + +int main() { + static int svar; + volatile double g; + volatile double &g1 = g; + + #ifdef LAMBDA + // LAMBDA-LABEL: @main + // LAMBDA: call{{.*}} void [[OUTER_LAMBDA:@.+]]( + [&]() { + static float sfvar; + // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( + // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( + // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( + + // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]( + // LAMBDA: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED:@.+]] to {{.+}}) + #pragma omp target + #pragma omp teams +#pragma omp distribute simd firstprivate(g, g1, svar, sfvar) + for (int i = 0; i < 2; ++i) { + // LAMBDA-64: define{{.*}} internal{{.*}} void [[OMP_OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, i{{[0-9]+}} [[G_IN:%.+]], i{{[0-9]+}} [[G1_IN:%.+]], i{{[0-9]+}} [[SVAR_IN:%.+]], i{{[0-9]+}} [[SFVAR_IN:%.+]]) + // LAMBDA-32: define internal{{.*}} void [[OMP_OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, double*{{.*}} [[G_IN:%.+]], i{{[0-9]+}} [[G1_IN:%.+]], i{{[0-9]+}} [[SVAR_IN:%.+]], i{{[0-9]+}} [[SFVAR_IN:%.+]]) + // Private alloca's for conversion + // LAMBDA-64: [[G_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA-32: [[G_ADDR:%.+]] = alloca double*, + // LAMBDA: [[G1_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[SFVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[G1_REF:%.+]] = alloca double*, + // LAMBDA: [[TMP:%.+]] = alloca double*, + + // Actual private variables to be used in the body (tmp is used for the reference type) + // LAMBDA: [[G_PRIVATE:%.+]] = alloca double, + // LAMBDA: [[G1_PRIVATE:%.+]] = alloca double, + // LAMBDA: [[TMP_PRIVATE:%.+]] = alloca double*, + // LAMBDA: [[SVAR_PRIVATE:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[SFVAR_PRIVATE:%.+]] = alloca float, + + // Store input parameter addresses into private alloca's for conversion + // LAMBDA-64: store i{{[0-9]+}} [[G_IN]], i{{[0-9]+}}* [[G_ADDR]], + // LAMBDA-32: store double* [[G_IN]], double** [[G_ADDR]], + // LAMBDA: store i{{[0-9]+}} [[G1_IN]], i{{[0-9]+}}* [[G1_ADDR]], + // LAMBDA: store i{{[0-9]+}} [[SVAR_IN]], i{{[0-9]+}}* [[SVAR_ADDR]], + // LAMBDA: store i{{[0-9]+}} [[SFVAR_IN]], i{{[0-9]+}}* [[SFVAR_ADDR]], + + // LAMBDA-64-DAG: [[G_CONV:%.+]] = bitcast i{{[0-9]+}}* [[G_ADDR]] to double* + // LAMBDA-32-DAG: [[G_ADDR_VAL:%.+]] = load double*, double** [[G_ADDR]], + // LAMBDA-DAG: [[G1_CONV:%.+]] = bitcast i{{[0-9]+}}* [[G1_ADDR]] to double* + // LAMBDA-DAG: store double* [[G1_CONV]], double** [[G1_REF]], + // LAMBDA-64-DAG: [[SVAR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[SVAR_ADDR]] to i{{[0-9]+}}* + // LAMBDA-DAG: [[SFVAR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[SFVAR_ADDR]] to float* + // LAMBDA-DAG: [[G1_REF_VAL:%.+]] = load double*, double** [[G1_REF]], + // LAMBDA-DAG: store double* [[G1_REF_VAL]], double** [[TMP]], + // LAMBDA-64-DAG: [[G_CONV_VAL:%.+]] = load{{.*}} double, double* [[G_CONV]], + // LAMBDA-32-DAG: [[G_CONV_VAL:%.+]] = load{{.*}} double, double* [[G_ADDR_VAL]], + // LAMBDA-DAG: store double [[G_CONV_VAL]], double* [[G_PRIVATE]], + // LAMBDA-DAG: [[TMP_VAL:%.+]] = load double*, double** [[TMP]], + // LAMBDA-DAG: [[TMP_VAL_VAL:%.+]] = load{{.*}} double, double* [[TMP_VAL]], + // LAMBDA-DAG: store double [[TMP_VAL_VAL]], double* [[G1_PRIVATE]], + // LAMBDA-DAG: store double* [[G1_PRIVATE]], double** [[TMP_PRIVATE]], + // LAMBDA-64-DAG: [[SVAR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_CONV]], + // LAMBDA-32-DAG: [[SVAR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_ADDR]], + // LAMBDA-DAG: store i{{[0-9]+}} [[SVAR_CONV_VAL]], i{{[0-9]+}}* [[SVAR_PRIVATE]], + // LAMBDA-DAG: [[SFVAR_CONV_VAL:%.+]] = load float, float* [[SFVAR_CONV]], + // LAMBDA-DAG: store float [[SFVAR_CONV_VAL]], float* [[SFVAR_PRIVATE]], + // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( + g += 1; + g1 += 1; + svar += 3; + sfvar += 4.0; + // LAMBDA-DAG: [[G_VAL:%.+]] = load double, double* [[G_PRIVATE]], + // LAMBDA-DAG: [[G_NEXT:%.+]] = fadd double [[G_VAL]], 1.{{.+}} + // LAMBDA-DAG: store double [[G_NEXT]], double* [[G_PRIVATE]], + // LAMBDA-DAG: [[TMP_VAL1:%.+]] = load double*, double** [[TMP_PRIVATE]], + // LAMBDA-DAG: [[TMP_VAL_VAL1:%.+]] = load{{.*}} double, double* [[TMP_VAL1]], + // LAMBDA-DAG: [[TMP_ADD:%.+]] = fadd double [[TMP_VAL_VAL1]], 1.{{.+}} + // LAMBDA-DAG: store{{.*}} double [[TMP_ADD]], double* [[TMP_VAL1]], + // LAMBDA-DAG: [[SVAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_PRIVATE]], + // LAMBDA-DAG: [[SVAR_ADD:%.+]] = add{{.*}} i{{[0-9]+}} [[SVAR_VAL]], 3 + // LAMBDA-DAG: store i{{[0-9]+}} [[SVAR_ADD]], i{{[0-9]+}}* [[SVAR_PRIVATE]], + // LAMBDA-DAG: [[SFVAR_VAL:%.+]] = load float, float* [[SFVAR_PRIVATE]], + // LAMBDA-DAG: [[SFVAR_CONV_VAL1:%.+]] = fpext float [[SFVAR_VAL]] to double + // LAMBDA-DAG: [[SFVAR_ADD:%.+]] = fadd double [[SFVAR_CONV_VAL1]], 4.{{.+}} + // LAMBDA-DAG: [[SFVAR_CONV_VAL2:%.+]] = fptrunc double [[SFVAR_ADD]] to float + // LAMBDA-DAG: store float [[SFVAR_CONV_VAL2:%.+]], float* [[SFVAR_PRIVATE]], + + // call inner lambda (use refs to private alloca's) + // LAMBDA: [[GEP_0:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA: store double* [[G_PRIVATE]], double** [[GEP_0]], + // LAMBDA: [[GEP_1:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[TMP_PAR:%.+]] = load double*, double** [[TMP_PRIVATE]], + // LAMBDA: store double* [[TMP_PAR]], double** [[GEP_1]], + // LAMBDA: [[GEP_2:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: store i{{[0-9]+}}* [[SVAR_PRIVATE]], i{{[0-9]+}}** [[GEP_2]], + // LAMBDA: [[GEP_3:%.+]] = getelementptr{{.+}}, i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // LAMBDA: store float* [[SFVAR_PRIVATE]], float** [[GEP_3]], + // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]](%{{.+}}* {{.+}}) + // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( + [&]() { + // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) + // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], + g += 2; + g1 += 2; + svar += 4; + sfvar += 8.0; + // LAMBDA-DAG: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] + // LAMBDA-DAG: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA-DAG: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]], + // LAMBDA-DAG: [[G_REF_VAL:%.+]] = load double, double* [[G_REF]], + // LAMBDA-DAG: [[G_REF_ADD:%.+]] = fadd double [[G_REF_VAL]], 2.{{.+}} + // LAMBDA-DAG: store double [[G_REF_ADD]], double* [[G_REF]] + + // LAMBDA-DAG: [[TMP_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA-DAG: [[G1_REF:%.+]] = load double*, double** [[TMP_PTR_REF]] + // LAMBDA-DAG: [[G1_REF_VAL:%.+]] = load double, double* [[G1_REF]], + // LAMBDA-DAG: [[G1_ADD:%.+]] = fadd double [[G1_REF_VAL]], 2.{{.+}} + // LAMBDA-DAG: store double [[G1_ADD]], double* [[G1_REF]], + + // LAMBDA-DAG: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA-DAG: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] + // LAMBDA-DAG: [[SVAR_REF_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_REF]] + // LAMBDA-DAG: [[SVAR_ADD:%.+]] = add{{.*}} i{{[0-9]+}} [[SVAR_REF_VAL]], 4 + // LAMBDA-DAG: store i{{[0-9]+}} [[SVAR_ADD]], i{{[0-9]+}}* [[SVAR_REF]] + + // LAMBDA-DAG: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // LAMBDA-DAG: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] + // LAMBDA-DAG: [[SFVAR_REF_VAL:%.+]] = load float, float* [[SFVAR_REF]] + // LAMBDA-DAG: [[SFVAR_REF_CONV:%.+]] = fpext float [[SFVAR_REF_VAL]] to double + // LAMBDA-DAG: [[SFVAR_ADD:%.+]] = fadd double [[SFVAR_REF_CONV]], 8.{{.+}} + // LAMBDA-DAG: [[SFVAR_ADD_CONV:%.+]] = fptrunc double [[SFVAR_ADD]] to float + // LAMBDA-DAG: store float [[SFVAR_ADD_CONV]], float* [[SFVAR_REF]], + }(); + } + }(); + return 0; + #else + S test; + int t_var = 0; + int vec[] = {1, 2}; + S s_arr[] = {1, 2}; + S &var = test; + + #pragma omp target + #pragma omp teams + #pragma omp distribute simd firstprivate(t_var, vec, s_arr, s_arr, var, var, svar) + for (int i = 0; i < 2; ++i) { + vec[i] = t_var; + s_arr[i] = var; + } + return tmain(); + #endif +} + +// CHECK: define{{.*}} i{{[0-9]+}} @main() +// CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], +// CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( +// CHECK: call void [[OFFLOAD_FUN:@.+]]( +// CHECK: ret + +// CHECK: define{{.+}} [[OFFLOAD_FUN]]( +// CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 5, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, i{{[0-9]+}}, [2 x i{{[0-9]+}}]*, [2 x [[S_FLOAT_TY]]]*, [[S_FLOAT_TY]]*, i{{[0-9]+}})* [[OMP_OUTLINED:@.+]] to void +// CHECK: ret +// +// CHECK: define internal void [[OMP_OUTLINED]](i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, i{{[0-9]+}} [[T_VAR_IN:%.+]], [2 x i{{[0-9]+}}]*{{.*}} [[VEC_IN:%.+]], [2 x [[S_FLOAT_TY]]]*{{.*}} [[S_ARR_IN:%.+]], [[S_FLOAT_TY]]*{{.*}} [[VAR_IN:%.+]], i{{[0-9]+}} [[SVAR_IN:%.+]]) + +// CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, +// CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, +// CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, +// CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[TMP:%.+]] = alloca [[S_FLOAT_TY]]*, + +// discard omp loop variables +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, + +// CHECK-DAG: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, +// CHECK-DAG: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], +// CHECK-DAG: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_FLOAT_TY]]], +// CHECK-DAG: [[VAR_PRIV:%.+]] = alloca [[S_FLOAT_TY]], +// CHECK-DAG: [[TMP_PRIV:%.+]] = alloca [[S_FLOAT_TY]]*, +// CHECK: [[SVAR_PRIV:%.+]] = alloca i{{[0-9]+}}, + +// CHECK: store i{{[0-9]+}} [[T_VAR_IN]], i{{[0-9]+}}* [[T_VAR_ADDR]], +// CHECK: store [2 x i{{[0-9]+}}]* [[VEC_IN]], [2 x i{{[0-9]+}}]** [[VEC_ADDR]], +// CHECK: store [2 x [[S_FLOAT_TY]]]* [[S_ARR_IN]], [2 x [[S_FLOAT_TY]]]** [[S_ARR_ADDR]], +// CHECK: store [[S_FLOAT_TY]]* [[VAR_IN]], [[S_FLOAT_TY]]** [[VAR_ADDR]], +// CHECK: store i{{[0-9]+}} [[SVAR_IN]], i{{[0-9]+}}* [[SVAR_ADDR]], + +// init t_var +// CHECK-64-DAG: [[T_VAR_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[T_VAR_ADDR]] to i{{[0-9]+}}* +// CHECK-64-DAG: [[T_VAR_ADDR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_ADDR_CONV]], +// CHECK-32-DAG: [[T_VAR_ADDR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_ADDR]], +// CHECK-DAG: store i{{[0-9]+}} [[T_VAR_ADDR_CONV_VAL]], i{{[0-9]+}}* [[T_VAR_PRIV]], + +// init vec +// CHECK-DAG: [[VEC_ADDR_VAL:%.+]] = load [2 x i{{[0-9]+}}]*, [2 x i{{[0-9]+}}]** [[VEC_ADDR]], +// CHECK-DAG: [[VEC_ADDR_VAL_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_ADDR_VAL]] to i{{[0-9]+}}* +// CHECK-DAG: [[VEC_PRIV_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_PRIV]] to i{{[0-9]+}}* +// CHECK-DAG: call void @llvm.memcpy.{{.*}}(i{{[0-9]+}}* [[VEC_PRIV_BCAST]], i{{[0-9]+}}* [[VEC_ADDR_VAL_BCAST]],{{.+}}) + +// init s_arr +// CHECK-DAG: [[S_ARR_ADDR_VAL:%.+]] = load [2 x [[S_FLOAT_TY]]]*, [2 x [[S_FLOAT_TY]]]** [[S_ARR_ADDR]], +// CHECK-DAG: [[S_ARR_ADDR_BCAST:%.+]] = bitcast [2 x [[S_FLOAT_TY]]]* [[S_ARR_ADDR_VAL]] to [[S_FLOAT_TY]]* +// CHECK-DAG: [[S_ARR_PRIV_BGN:%.+]] = getelementptr{{.+}} [2 x [[S_FLOAT_TY]]], [2 x [[S_FLOAT_TY]]]* [[S_ARR_PRIV]]{{.+}} +// CHECK-DAG: [[S_ARR_PRIV_NEXT:%.+]] = getelementptr [[S_FLOAT_TY]], [[S_FLOAT_TY]]* [[S_ARR_PRIV_BGN]]{{.+}} +// CHECK-DAG: [[S_ARR_IS_EMPTY:%.+]] = icmp eq [[S_FLOAT_TY]]* [[S_ARR_PRIV_BGN]], [[S_ARR_PRIV_NEXT]] +// CHECK-DAG: br i1 [[S_ARR_IS_EMPTY]], label %[[S_ARR_CPY_DONE:.+]], label %[[S_ARR_CPY_BODY:.+]] + +// CHECK-DAG: [[S_ARR_CPY_BODY]]: +// CHECK-DAG: [[S_ARR_SRC_PAST:%.+]] = phi{{.+}} [ [[S_ARR_ADDR_BCAST]],{{.+}} ], [ [[S_ARR_SRC:%.+]],{{.+}} ] +// CHECK-DAG: [[S_ARR_DST_PAST:%.+]] = phi{{.+}} [ [[S_ARR_PRIV_BGN]],{{.+}} ], [ [[S_ARR_DST:%.+]],{{.+}} ] +// CHECK-DAG: [[S_ARR_SRC_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_SRC_PAST]] to{{.+}} +// CHECK-DAG: [[S_ARR_DST_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_DST_PAST]] to{{.+}} +// CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* [[S_ARR_DST_BCAST]], {{.+}}* [[S_ARR_SRC_BCAST]]{{.+}}) +// CHECK-DAG: [[S_ARR_SRC]] = getelementptr{{.+}} +// CHECK-DAG: [[S_ARR_DST]] = getelementptr{{.+}} +// CHECK-DAG: [[S_ARR_CPY_FIN:%.+]] = icmp{{.+}} [[S_ARR_DST]], [[S_ARR_PRIV_NEXT]] +// CHECK-DAG: br i1 [[S_ARR_CPY_FIN]], label %[[S_ARR_CPY_DONE]], label %[[S_ARR_CPY_BODY]] +// CHECK-DAG: [[S_ARR_CPY_DONE]]: + +// init var +// CHECK-DAG: [[VAR_ADDR_VAL:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[VAR_ADDR]], +// CHECK-DAG: store{{.+}} [[VAR_ADDR_VAL]],{{.+}} [[TMP]], +// CHECK-DAG: [[TMP_VAL:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[TMP]], +// CHECK-DAG: [[VAR_PRIV_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[VAR_PRIV]] to{{.+}} +// CHECK-DAG: [[TMP_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[TMP_VAL]] to{{.+}} +// CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* [[VAR_PRIV_BCAST]], {{.+}}* [[TMP_BCAST]],{{.+}}) +// CHECK-DAG: store [[S_FLOAT_TY]]* [[VAR_PRIV]], [[S_FLOAT_TY]]** [[TMP_PRIV]], + +// init svar +// CHECK-64-DAG: [[SVAR_ADDR_CONV:%.+]] = bitcast{{.+}} [[SVAR_ADDR]] to{{.+}} +// CHECK-64-DAG: [[SVAR_CONV_VAL:%.+]] = load{{.+}},{{.+}} [[SVAR_ADDR_CONV]], +// CHECK-32-DAG: [[SVAR_CONV_VAL:%.+]] = load{{.+}},{{.+}} [[SVAR_ADDR]], +// CHECK-DAG: store{{.+}} [[SVAR_CONV_VAL]],{{.+}} [[SVAR_PRIV]], + +// CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, +// CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, +// CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, +// CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, + +// CHECK: call void @__kmpc_for_static_init_4( +// CHECK: call void @__kmpc_for_static_fini( +// CHECK: ret void + +// Template +// CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() +// CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], +// CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( +// CHECK: call void [[OFFLOAD_FUN_1:@.+]]( +// CHECK: ret + +// CHECK: define{{.+}} [[OFFLOAD_FUN_1]]( +// CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 4, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*, i{{[0-9]+}}, [2 x i{{[0-9]+}}]*, [2 x [[S_INT_TY]]]*, [[S_INT_TY]]*)* [[OMP_OUTLINED_1:@.+]] to void +// CHECK: ret +// +// CHECK: define internal void [[OMP_OUTLINED_1]](i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, i{{[0-9]+}} [[T_VAR_IN:%.+]], [2 x i{{[0-9]+}}]*{{.*}} [[VEC_IN:%.+]], [2 x [[S_INT_TY]]]*{{.*}} [[S_ARR_IN:%.+]], [[S_INT_TY]]*{{.*}} [[VAR_IN:%.+]]) + +// CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, +// CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_INT_TY]]]*, +// CHECK: [[VAR_ADDR:%.+]] = alloca [[S_INT_TY]]*, +// CHECK: [[TMP:%.+]] = alloca [[S_INT_TY]]*, + +// discard omp loop variables +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, +// CHECK: {{.*}} = alloca i{{[0-9]+}}, + +// CHECK-DAG: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, +// CHECK-DAG: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], +// CHECK-DAG: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_INT_TY]]], +// CHECK-DAG: [[VAR_PRIV:%.+]] = alloca [[S_INT_TY]], +// CHECK-DAG: [[TMP_PRIV:%.+]] = alloca [[S_INT_TY]]*, + +// CHECK: store i{{[0-9]+}} [[T_VAR_IN]], i{{[0-9]+}}* [[T_VAR_ADDR]], +// CHECK: store [2 x i{{[0-9]+}}]* [[VEC_IN]], [2 x i{{[0-9]+}}]** [[VEC_ADDR]], +// CHECK: store [2 x [[S_INT_TY]]]* [[S_ARR_IN]], [2 x [[S_INT_TY]]]** [[S_ARR_ADDR]], +// CHECK: store [[S_INT_TY]]* [[VAR_IN]], [[S_INT_TY]]** [[VAR_ADDR]], + +// init t_var +// CHECK-64-DAG: [[T_VAR_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[T_VAR_ADDR]] to i{{[0-9]+}}* +// CHECK-64-DAG: [[T_VAR_ADDR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_ADDR_CONV]], +// CHECK-32-DAG: [[T_VAR_ADDR_CONV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_ADDR]], +// CHECK-DAG: store i{{[0-9]+}} [[T_VAR_ADDR_CONV_VAL]], i{{[0-9]+}}* [[T_VAR_PRIV]], + +// init vec +// CHECK-DAG: [[VEC_ADDR_VAL:%.+]] = load [2 x i{{[0-9]+}}]*, [2 x i{{[0-9]+}}]** [[VEC_ADDR]], +// CHECK-DAG: [[VEC_ADDR_VAL_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_ADDR_VAL]] to i{{[0-9]+}}* +// CHECK-DAG: [[VEC_PRIV_BCAST:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_PRIV]] to i{{[0-9]+}}* +// CHECK-DAG: call void @llvm.memcpy.{{.*}}(i{{[0-9]+}}* [[VEC_PRIV_BCAST]], i{{[0-9]+}}* [[VEC_ADDR_VAL_BCAST]],{{.+}}) + +// init s_arr +// CHECK-DAG: [[S_ARR_ADDR_VAL:%.+]] = load [2 x [[S_INT_TY]]]*, [2 x [[S_INT_TY]]]** [[S_ARR_ADDR]], +// CHECK-DAG: [[S_ARR_ADDR_BCAST:%.+]] = bitcast [2 x [[S_INT_TY]]]* [[S_ARR_ADDR_VAL]] to [[S_INT_TY]]* +// CHECK-DAG: [[S_ARR_PRIV_BGN:%.+]] = getelementptr{{.+}} [2 x [[S_INT_TY]]], [2 x [[S_INT_TY]]]* [[S_ARR_PRIV]]{{.+}} +// CHECK-DAG: [[S_ARR_PRIV_NEXT:%.+]] = getelementptr [[S_INT_TY]], [[S_INT_TY]]* [[S_ARR_PRIV_BGN]]{{.+}} +// CHECK-DAG: [[S_ARR_IS_EMPTY:%.+]] = icmp eq [[S_INT_TY]]* [[S_ARR_PRIV_BGN]], [[S_ARR_PRIV_NEXT]] +// CHECK-DAG: br i1 [[S_ARR_IS_EMPTY]], label %[[S_ARR_CPY_DONE:.+]], label %[[S_ARR_CPY_BODY:.+]] + +// CHECK-DAG: [[S_ARR_CPY_BODY]]: +// CHECK-DAG: [[S_ARR_SRC_PAST:%.+]] = phi{{.+}} [ [[S_ARR_ADDR_BCAST]],{{.+}} ], [ [[S_ARR_SRC:%.+]],{{.+}} ] +// CHECK-DAG: [[S_ARR_DST_PAST:%.+]] = phi{{.+}} [ [[S_ARR_PRIV_BGN]],{{.+}} ], [ [[S_ARR_DST:%.+]],{{.+}} ] +// CHECK-DAG: [[S_ARR_SRC_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_SRC_PAST]] to{{.+}} +// CHECK-DAG: [[S_ARR_DST_BCAST:%.+]] = bitcast{{.+}} [[S_ARR_DST_PAST]] to{{.+}} +// CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* [[S_ARR_DST_BCAST]], {{.+}}* [[S_ARR_SRC_BCAST]]{{.+}}) +// CHECK-DAG: [[S_ARR_SRC]] = getelementptr{{.+}} +// CHECK-DAG: [[S_ARR_DST]] = getelementptr{{.+}} +// CHECK-DAG: [[S_ARR_CPY_FIN:%.+]] = icmp{{.+}} [[S_ARR_DST]], [[S_ARR_PRIV_NEXT]] +// CHECK-DAG: br i1 [[S_ARR_CPY_FIN]], label %[[S_ARR_CPY_DONE]], label %[[S_ARR_CPY_BODY]] +// CHECK-DAG: [[S_ARR_CPY_DONE]]: + +// init var +// CHECK-DAG: [[VAR_ADDR_VAL:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[VAR_ADDR]], +// CHECK-DAG: store{{.+}} [[VAR_ADDR_VAL]],{{.+}} [[TMP]], +// CHECK-DAG: [[TMP_VAL:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[TMP]], +// CHECK-DAG: [[VAR_PRIV_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[VAR_PRIV]] to{{.+}} +// CHECK-DAG: [[TMP_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[TMP_VAL]] to{{.+}} +// CHECK-DAG: call{{.+}} @llvm.memcpy.{{.+}}({{.+}}* [[VAR_PRIV_BCAST]], {{.+}}* [[TMP_BCAST]],{{.+}}) +// CHECK-DAG: store [[S_INT_TY]]* [[VAR_PRIV]], [[S_INT_TY]]** [[TMP_PRIV]], + +// CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, +// CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, +// CHECK-DAG: store i{{[0-9]+}} 1, i{{[0-9]+}}* %.omp{{.+}}, +// CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* %.omp{{.+}}, + +// CHECK: call void @__kmpc_for_static_init_4( +// CHECK: call void @__kmpc_for_static_fini( +// CHECK: ret void + +// CHECK: !{!"llvm.loop.vectorize.enable", i1 true} +#endif diff --git a/clang/test/OpenMP/distribute_simd_lastprivate_codegen.cpp b/clang/test/OpenMP/distribute_simd_lastprivate_codegen.cpp new file mode 100644 index 0000000000000000000000000000000000000000..bc4604d1c259cbc667aa7dfcf834112cca98c0fc --- /dev/null +++ b/clang/test/OpenMP/distribute_simd_lastprivate_codegen.cpp @@ -0,0 +1,390 @@ +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 + +// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +template +struct S { + T f; + S(T a) : f(a) {} + S() : f() {} + operator T() { return T(); } + ~S() {} +}; + +// CHECK: [[S_FLOAT_TY:%.+]] = type { float } +// CHECK: [[S_INT_TY:%.+]] = type { i{{[0-9]+}} } +template +T tmain() { + S test; + T t_var = T(); + T vec[] = {1, 2}; + S s_arr[] = {1, 2}; + S &var = test; + #pragma omp target + #pragma omp teams +#pragma omp distribute simd lastprivate(t_var, vec, s_arr, s_arr, var, var) + for (int i = 0; i < 2; ++i) { + vec[i] = t_var; + s_arr[i] = var; + } + return T(); +} + +int main() { + static int svar; + volatile double g; + volatile double &g1 = g; + + #ifdef LAMBDA + // LAMBDA-LABEL: @main + // LAMBDA: call{{.*}} void [[OUTER_LAMBDA:@.+]]( + [&]() { + static float sfvar; + // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( + // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( + // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( + + // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]( + // LAMBDA: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, i32 4, {{.+}}* [[OMP_OUTLINED:@.+]] to {{.+}}) + #pragma omp target + #pragma omp teams +#pragma omp distribute simd lastprivate(g, g1, svar, sfvar) + for (int i = 0; i < 2; ++i) { + // LAMBDA: define{{.*}} internal{{.*}} void [[OMP_OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, double*{{.+}} [[G_IN:%.+]], double*{{.+}} [[G1_IN:%.+]], i{{[0-9]+}}*{{.+}} [[SVAR_IN:%.+]], float*{{.+}} [[SFVAR_IN:%.+]]) + // LAMBDA: [[G_PRIVATE_ADDR:%.+]] = alloca double*, + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = alloca double*, + // LAMBDA: [[SVAR_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}*, + // LAMBDA: [[SFVAR_PRIVATE_ADDR:%.+]] = alloca float*, + // LAMBDA: [[TMP_G1:%.+]] = alloca double*, + // loop variables + // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, + // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, + // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, + // LAMBDA: {{.+}} = alloca i{{[0-9]+}}, + // LAMBDA: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[G_PRIVATE:%.+]] = alloca double, + // LAMBDA: [[G1_PRIVATE:%.+]] = alloca double, + // LAMBDA: [[TMP_G1_PRIVATE:%.+]] = alloca double*, + // LAMBDA: [[SVAR_PRIVATE:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[SFVAR_PRIVATE:%.+]] = alloca float, + // LAMBDA: store double* [[G_IN]], double** [[G_PRIVATE_ADDR]], + // LAMBDA: store double* [[G1_IN]], double** [[G1_PRIVATE_ADDR]], + // LAMBDA: store i{{[0-9]+}}* [[SVAR_IN]], i{{[0-9]+}}** [[SVAR_PRIVATE_ADDR]], + // LAMBDA: store float* [[SFVAR_IN]], float** [[SFVAR_PRIVATE_ADDR]], + + // init private variables + // LAMBDA: [[G_IN_REF:%.+]] = load double*, double** [[G_PRIVATE_ADDR]], + // LAMBDA: [[SVAR_IN_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PRIVATE_ADDR]], + // LAMBDA: [[SFVAR_IN_REF:%.+]] = load float*, float** [[SFVAR_PRIVATE_ADDR]], + // LAMBDA: [[G1_IN_REF:%.+]] = load double*, double** [[G1_PRIVATE_ADDR]], + // LAMBDA: store double* [[G1_IN_REF]], double** [[TMP_G1]], + // LAMBDA: [[TMP_G1_VAL:%.+]] = load double*, double** [[TMP_G1]], + // LAMBDA: store double* [[G1_PRIVATE]], double** [[TMP_G1_PRIVATE]], + g = 1; + g1 = 1; + svar = 3; + sfvar = 4.0; + // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( + // LAMBDA: store double 1.0{{.+}}, double* [[G_PRIVATE]], + // LAMBDA: [[TMP_G1_REF:%.+]] = load double*, double** [[TMP_G1_PRIVATE]], + // LAMBDA: store{{.+}} double 1.0{{.+}}, double* [[TMP_G1_REF]], + // LAMBDA: store i{{[0-9]+}} 3, i{{[0-9]+}}* [[SVAR_PRIVATE]], + // LAMBDA: store float 4.0{{.+}}, float* [[SFVAR_PRIVATE]], + // LAMBDA: [[G_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA: store double* [[G_PRIVATE]], double** [[G_PRIVATE_ADDR_REF]], + // LAMBDA: [[TMP_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR_FROM_TMP:%.+]] = load double*, double** [[TMP_G1_PRIVATE]], + // LAMBDA: store double* [[G1_PRIVATE_ADDR_FROM_TMP]], double** [[TMP_PRIVATE_ADDR_REF]], + // LAMBDA: [[SVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: store i{{[0-9]+}}* [[SVAR_PRIVATE]], i{{[0-9]+}}** [[SVAR_PRIVATE_ADDR_REF]] + // LAMBDA: [[SFVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // LAMBDA: store float* [[SFVAR_PRIVATE]], float** [[SFVAR_PRIVATE_ADDR_REF]] + // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]](%{{.+}}* [[ARG]]) + // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( + + // linear counter + // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], + // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 + // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_COUNTER_BLOCK:.+]], label %[[OMP_COUNTER_DONE:.+]] + // LAMBDA: [[OMP_COUNTER_BLOCK]]: + // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* % + // LAMBDA: br label %[[OMP_COUNTER_DONE]] + // LAMBDA: [[OMP_COUNTER_DONE]]: + + + // LAMBDA: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], + // LAMBDA: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 + // LAMBDA: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] + + // LAMBDA: [[OMP_LASTPRIV_BLOCK]]: + // LAMBDA: [[G_PRIV_VAL:%.+]] = load double, double* [[G_PRIVATE]], + // LAMBDA: store{{.*}} double [[G_PRIV_VAL]], double* [[G_IN_REF]], + // LAMBDA: [[TMP_G1_PRIV_REF:%.+]] = load double*, double** [[TMP_G1_PRIVATE]], + // LAMBDA: [[TMP_G1_PRIV_VAL:%.+]] = load double, double* [[TMP_G1_PRIV_REF]], + // LAMBDA: store{{.*}} double [[TMP_G1_PRIV_VAL]], double* [[TMP_G1_VAL]], + + // LAMBDA: [[SVAR_PRIV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SVAR_PRIVATE]], + // LAMBDA: store i{{[0-9]+}} [[SVAR_PRIV_VAL]], i{{[0-9]+}}* [[SVAR_IN_REF]], + // LAMBDA: [[SFVAR_PRIV_VAL:%.+]] = load float, float* [[SFVAR_PRIVATE]], + // LAMBDA: store float [[SFVAR_PRIV_VAL]], float* [[SFVAR_IN_REF]], + // LAMBDA: br label %[[OMP_LASTPRIV_DONE]] + // LAMBDA: [[OMP_LASTPRIV_DONE]]: + // LAMBDA: ret + [&]() { + // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) + // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], + g = 2; + g1 = 2; + svar = 4; + sfvar = 8.0; + // LAMBDA: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] + // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]] + // LAMBDA: store double 2.0{{.+}}, double* [[G_REF]] + + // LAMBDA: [[TMP_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load double*, double** [[TMP_PTR_REF]] + // LAMBDA: store double 2.0{{.+}}, double* [[G1_REF]], + // LAMBDA: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] + // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SVAR_REF]] + // LAMBDA: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // LAMBDA: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] + // LAMBDA: store float 8.0{{.+}}, float* [[SFVAR_REF]] + }(); + } + }(); + return 0; + #else + S test; + int t_var = 0; + int vec[] = {1, 2}; + S s_arr[] = {1, 2}; + S &var = test; + + #pragma omp target + #pragma omp teams +#pragma omp distribute simd lastprivate(t_var, vec, s_arr, s_arr, var, var, svar) + for (int i = 0; i < 2; ++i) { + vec[i] = t_var; + s_arr[i] = var; + } + int i; + + return tmain(); + #endif +} + +// CHECK: define{{.*}} i{{[0-9]+}} @main() +// CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], +// CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( +// CHECK: call void [[OFFLOAD_FUN:@.+]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]* {{.+}}, [2 x [[S_FLOAT_TY]]]* {{.+}}, [[S_FLOAT_TY]]* {{.+}}, i{{[0-9]+}} {{.+}}) +// CHECK: ret + +// CHECK: define{{.+}} [[OFFLOAD_FUN]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]*{{.+}} {{.+}}, [2 x [[S_FLOAT_TY]]]*{{.+}} {{.+}}, [[S_FLOAT_TY]]*{{.+}} {{.+}}, i{{[0-9]+}} {{.+}}) +// CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams( +// CHECK: ret +// +// CHECK: define internal void [[OMP_OUTLINED:@.+]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}}, i{{[0-9]+}}*{{.+}} [[T_VAR_IN:%.+]], [2 x i{{[0-9]+}}]*{{.+}} [[VEC_IN:%.+]], [2 x [[S_FLOAT_TY]]]*{{.+}} [[S_ARR_IN:%.+]], [[S_FLOAT_TY]]*{{.+}} [[VAR_IN:%.+]], i{{[0-9]+}}*{{.*}} [[S_VAR_IN:%.+]]) +// CHECK: {{.+}} = alloca i{{[0-9]+}}*, +// CHECK: {{.+}} = alloca i{{[0-9]+}}*, +// CHECK: [[T_VAR_ADDR:%.+]] = alloca i{{[0-9]+}}*, +// CHECK: [[VEC_ADDR:%.+]] = alloca [2 x i{{[0-9]+}}]*, +// CHECK: [[S_ARR_ADDR:%.+]] = alloca [2 x [[S_FLOAT_TY]]]*, +// CHECK: [[VAR_ADDR:%.+]] = alloca [[S_FLOAT_TY]]*, +// CHECK: [[SVAR_ADDR:%.+]] = alloca i{{[0-9]+}}*, +// CHECK: [[TMP:%.*]] = alloca [[S_FLOAT_TY]]*, +// skip loop variables +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: [[OMP_IS_LAST:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], +// CHECK: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_FLOAT_TY]]], +// CHECK: [[VAR_PRIV:%.+]] = alloca [[S_FLOAT_TY]], +// CHECK: [[TMP_PRIV:%.+]] = alloca [[S_FLOAT_TY]]*, +// CHECK: [[S_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, + +// copy from parameters to local address variables +// CHECK: store i{{[0-9]+}}* [[T_VAR_IN]], i{{[0-9]+}}** [[T_VAR_ADDR]], +// CHECK: store [2 x i{{[0-9]+}}]* [[VEC_IN]], [2 x i{{[0-9]+}}]** [[VEC_ADDR]], +// CHECK: store [2 x [[S_FLOAT_TY]]]* [[S_ARR_IN]], [2 x [[S_FLOAT_TY]]]** [[S_ARR_ADDR]], +// CHECK: store [[S_FLOAT_TY]]* [[VAR_IN]], [[S_FLOAT_TY]]** [[VAR_ADDR]], +// CHECK: store i{{[0-9]+}}* [[S_VAR_IN]], i{{[0-9]+}}** [[SVAR_ADDR]], + +// load content of local address variables +// CHECK: [[T_VAR_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[T_VAR_ADDR]], +// CHECK: [[VEC_ADDR_REF:%.+]] = load [2 x i{{[0-9]+}}]*, [2 x i{{[0-9]+}}]** [[VEC_ADDR]], +// CHECK: [[S_ARR_ADDR_REF:%.+]] = load [2 x [[S_FLOAT_TY]]]*, [2 x [[S_FLOAT_TY]]]** [[S_ARR_ADDR]], +// CHECK: [[SVAR_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_ADDR]], +// CHECK: [[VAR_ADDR_REF:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[VAR_ADDR]], +// CHECK: store [[S_FLOAT_TY]]* [[VAR_ADDR_REF]], [[S_FLOAT_TY]]** [[TMP]], +// CHECK: store i{{[0-9]+}} 0, i{{[0-9]+}}* [[OMP_IS_LAST]], +// CHECK: [[TMP_REF:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[TMP]], +// the distribute loop +// CHECK: call void @__kmpc_for_static_init_4( +// assignment: vec[i] = t_var; +// CHECK: [[T_VAR_PRIV_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_PRIV]], +// CHECK: [[VEC_PTR:%.+]] = getelementptr inbounds [2 x i{{[0-9]+}}], [2 x i{{[0-9]+}}]* [[VEC_PRIV]], i{{[0-9]+}} 0, i{{[0-9]+}} {{.+}} +// CHECK: store i{{[0-9]+}} [[T_VAR_PRIV_VAL]], i{{[0-9]+}}* [[VEC_PTR]], + +// assignment: s_arr[i] = var; +// CHECK-DAG: [[S_ARR_PTR:%.+]] = getelementptr inbounds [2 x [[S_FLOAT_TY]]], [2 x [[S_FLOAT_TY]]]* [[S_ARR_PRIV]], +// CHECK-DAG: [[TMP_VAL:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[TMP_PRIV]], +// CHECK-DAG: [[S_ARR_PTR_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[S_ARR_PTR]] to i8* +// CHECK-DAG: [[TMP_VAL_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[TMP_VAL]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[S_ARR_PTR_BCAST]], i8* [[TMP_VAL_BCAST]], +// CHECK: call void @__kmpc_for_static_fini( + +// lastprivates +// CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST]], +// CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 +// CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] + +// CHECK: [[OMP_LASTPRIV_BLOCK]]: +// CHECK: [[T_VAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_PRIV]], +// CHECK: store i{{[0-9]+}} [[T_VAR_VAL]], i{{[0-9]+}}* [[T_VAR_ADDR_REF]], +// CHECK: [[BCAST_VEC_ADDR_REF:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_ADDR_REF]] to i8* +// CHECK: [[BCAST_VEC_PRIV:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_PRIV]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[BCAST_VEC_ADDR_REF]], i8* [[BCAST_VEC_PRIV]], +// CHECK: [[S_ARR_BEGIN:%.+]] = getelementptr inbounds [2 x [[S_FLOAT_TY]]], [2 x [[S_FLOAT_TY]]]* [[S_ARR_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 +// CHECK: [[S_ARR_PRIV_BCAST:%.+]] = bitcast [2 x [[S_FLOAT_TY]]]* [[S_ARR_PRIV]] to [[S_FLOAT_TY]]* +// CHECK: [[S_ARR_BEGIN_GEP:%.+]] = getelementptr [[S_FLOAT_TY]], [[S_FLOAT_TY]]* [[S_ARR_BEGIN]], i{{[0-9]+}} 2 +// CHECK: [[S_ARR_IS_EMPTY:%.+]] = icmp eq [[S_FLOAT_TY]]* [[S_ARR_BEGIN]], [[S_ARR_BEGIN_GEP]] +// CHECK: br i1 [[S_ARR_IS_EMPTY]], label %[[S_ARR_COPY_DONE:.+]], label %[[S_ARR_COPY_BLOCK:.+]] +// CHECK: [[S_ARR_COPY_BLOCK]]: +// CHECK: [[S_ARR_SRC_EL:%.+]] = phi [[S_FLOAT_TY]]*{{.+}} +// CHECK: [[S_ARR_DST_EL:%.+]] = phi [[S_FLOAT_TY]]*{{.+}} +// CHECK: [[S_ARR_DST_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[S_ARR_DST_EL]] to i8* +// CHECK: [[S_ARR_SRC_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[S_ARR_SRC_EL]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[S_ARR_DST_BCAST]], i8* [[S_ARR_SRC_BCAST]]{{.+}}) +// CHECK: [[S_ARR_DST_NEXT:%.+]] = getelementptr [[S_FLOAT_TY]], [[S_FLOAT_TY]]* [[S_ARR_DST_EL]], i{{[0-9]+}} 1 +// CHECK: [[S_ARR_SRC_NEXT:%.+]] = getelementptr{{.+}} +// CHECK: [[CPY_IS_FINISHED:%.+]] = icmp eq [[S_FLOAT_TY]]* [[S_ARR_DST_NEXT]], [[S_ARR_BEGIN_GEP]] +// CHECK: br i1 [[CPY_IS_FINISHED]], label %[[S_ARR_COPY_DONE]], label %[[S_ARR_COPY_BLOCK]] +// CHECK: [[S_ARR_COPY_DONE]]: +// CHECK: [[TMP_VAL1:%.+]] = load [[S_FLOAT_TY]]*, [[S_FLOAT_TY]]** [[TMP_PRIV]], +// CHECK: [[VAR_ADDR_REF_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[TMP_REF]] to i8* +// CHECK: [[TMP_VAL1_BCAST:%.+]] = bitcast [[S_FLOAT_TY]]* [[TMP_VAL1]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[VAR_ADDR_REF_BCAST]], i8* [[TMP_VAL1_BCAST]],{{.+}}) +// CHECK: [[SVAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[S_VAR_PRIV]], +// CHECK: store i{{[0-9]+}} [[SVAR_VAL]], i{{[0-9]+}}* [[SVAR_ADDR_REF]], +// CHECK: ret void + +// template tmain +// CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() +// CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], +// CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( +// CHECK: call void [[OFFLOAD_FUN_1:@.+]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]* {{.+}}, [2 x [[S_INT_TY]]]* {{.+}}, [[S_INT_TY]]* {{.+}}) +// CHECK: ret + + +// CHECK: define internal void [[OFFLOAD_FUN_1]]( +// CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 4, +// CHECK: ret + +// CHECK: define internal void [[OMP_OUTLINED_1:@.+]](i{{[0-9]+}}* noalias [[GTID_ADDR1:%.+]], i{{[0-9]+}}* noalias %{{.+}}, i{{[0-9]+}}*{{.+}} [[T_VAR_IN1:%.+]], [2 x i{{[0-9]+}}]*{{.+}} [[VEC_IN1:%.+]], [2 x [[S_INT_TY]]]*{{.+}} [[S_ARR_IN1:%.+]], [[S_INT_TY]]*{{.+}} [[VAR_IN1:%.+]]) +// skip alloca of global_tid and bound_tid +// CHECK: {{.+}} = alloca i{{[0-9]+}}*, +// CHECK: {{.+}} = alloca i{{[0-9]+}}*, +// CHECK: [[T_VAR_ADDR1:%.+]] = alloca i{{[0-9]+}}*, +// CHECK: [[VEC_ADDR1:%.+]] = alloca [2 x i{{[0-9]+}}]*, +// CHECK: [[S_ARR_ADDR1:%.+]] = alloca [2 x [[S_INT_TY]]]*, +// CHECK: [[VAR_ADDR1:%.+]] = alloca [[S_INT_TY]]*, +// CHECK: [[TMP:%.+]] = alloca [[S_INT_TY]]*, +// skip loop variables +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: {{.+}} = alloca i{{[0-9]+}}, +// CHECK: [[OMP_IS_LAST1:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[T_VAR_PRIV1:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[VEC_PRIV1:%.+]] = alloca [2 x i{{[0-9]+}}], +// CHECK: [[S_ARR_PRIV1:%.+]] = alloca [2 x [[S_INT_TY]]], +// CHECK: [[VAR_PRIV1:%.+]] = alloca [[S_INT_TY]], +// CHECK: [[TMP_PRIV1:%.+]] = alloca [[S_INT_TY]]*, + +// skip init of bound and global tid +// CHECK: store i{{[0-9]+}}* {{.*}}, +// CHECK: store i{{[0-9]+}}* {{.*}}, +// copy from parameters to local address variables +// CHECK: store i{{[0-9]+}}* [[T_VAR_IN1]], i{{[0-9]+}}** [[T_VAR_ADDR1]], +// CHECK: store [2 x i{{[0-9]+}}]* [[VEC_IN1]], [2 x i{{[0-9]+}}]** [[VEC_ADDR1]], +// CHECK: store [2 x [[S_INT_TY]]]* [[S_ARR_IN1]], [2 x [[S_INT_TY]]]** [[S_ARR_ADDR1]], +// CHECK: store [[S_INT_TY]]* [[VAR_IN1]], [[S_INT_TY]]** [[VAR_ADDR1]], + +// load content of local address variables +// CHECK: [[T_VAR_ADDR_REF1:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[T_VAR_ADDR1]], +// CHECK: [[VEC_ADDR_REF1:%.+]] = load [2 x i{{[0-9]+}}]*, [2 x i{{[0-9]+}}]** [[VEC_ADDR1]], +// CHECK: [[S_ARR_ADDR_REF1:%.+]] = load [2 x [[S_INT_TY]]]*, [2 x [[S_INT_TY]]]** [[S_ARR_ADDR1]], +// CHECK: [[VAR_ADDR1_REF:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[VAR_ADDR1]], +// CHECK-DAG: store [[S_INT_TY]]* [[VAR_ADDR1_REF]], [[S_INT_TY]]** [[TMP]], +// CHECK-DAG: store i{{[0-9]+}} 0, i{{[0-9]+}}* [[OMP_IS_LAST1]], +// CHECK-DAG: [[TMP_REF:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[TMP]], +// CHECK: call void @__kmpc_for_static_init_4( +// assignment: vec[i] = t_var; +// CHECK: [[IV_VAL1:%.+]] = +// CHECK: [[T_VAR_PRIV_VAL1:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_PRIV1]], +// CHECK: [[VEC_PTR1:%.+]] = getelementptr inbounds [2 x i{{[0-9]+}}], [2 x i{{[0-9]+}}]* [[VEC_PRIV1]], i{{[0-9]+}} 0, i{{[0-9]+}} {{.+}} +// CHECK: store i{{[0-9]+}} [[T_VAR_PRIV_VAL1]], i{{[0-9]+}}* [[VEC_PTR1]], + +// assignment: s_arr[i] = var; +// CHECK-DAG: [[S_ARR_PTR1:%.+]] = getelementptr inbounds [2 x [[S_INT_TY]]], [2 x [[S_INT_TY]]]* [[S_ARR_PRIV1]], +// CHECK-DAG: [[TMP_VAL1:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[TMP_PRIV1]], +// CHECK-DAG: [[S_ARR_PTR_BCAST1:%.+]] = bitcast [[S_INT_TY]]* [[S_ARR_PTR1]] to i8* +// CHECK-DAG: [[TMP_VAL_BCAST1:%.+]] = bitcast [[S_INT_TY]]* [[TMP_VAL1]] to i8* +// CHECK-DAG: call void @llvm.memcpy.{{.+}}(i8* [[S_ARR_PTR_BCAST1]], i8* [[TMP_VAL_BCAST1]], +// CHECK: call void @__kmpc_for_static_fini( + +// lastprivates +// CHECK: [[OMP_IS_LAST_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[OMP_IS_LAST1]], +// CHECK: [[IS_LAST_IT:%.+]] = icmp ne i{{[0-9]+}} [[OMP_IS_LAST_VAL]], 0 +// CHECK: br i1 [[IS_LAST_IT]], label %[[OMP_LASTPRIV_BLOCK:.+]], label %[[OMP_LASTPRIV_DONE:.+]] + +// CHECK: [[OMP_LASTPRIV_BLOCK]]: +// CHECK: [[T_VAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[T_VAR_PRIV1]], +// CHECK: store i{{[0-9]+}} [[T_VAR_VAL]], i{{[0-9]+}}* [[T_VAR_ADDR_REF1]], +// CHECK: [[BCAST_VEC_ADDR_REF:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_ADDR_REF1]] to i8* +// CHECK: [[BCAST_VEC_PRIV:%.+]] = bitcast [2 x i{{[0-9]+}}]* [[VEC_PRIV1]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[BCAST_VEC_ADDR_REF]], i8* [[BCAST_VEC_PRIV]], +// CHECK: [[S_ARR_BEGIN:%.+]] = getelementptr inbounds [2 x [[S_INT_TY]]], [2 x [[S_INT_TY]]]* [[S_ARR_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 +// CHECK: [[S_ARR_PRIV_BCAST:%.+]] = bitcast [2 x [[S_INT_TY]]]* [[S_ARR_PRIV1]] to [[S_INT_TY]]* +// CHECK: [[S_ARR_BEGIN_GEP:%.+]] = getelementptr [[S_INT_TY]], [[S_INT_TY]]* [[S_ARR_BEGIN]], i{{[0-9]+}} 2 +// CHECK: [[S_ARR_IS_EMPTY:%.+]] = icmp eq [[S_INT_TY]]* [[S_ARR_BEGIN]], [[S_ARR_BEGIN_GEP]] +// CHECK: br i1 [[S_ARR_IS_EMPTY]], label %[[S_ARR_COPY_DONE:.+]], label %[[S_ARR_COPY_BLOCK:.+]] +// CHECK: [[S_ARR_COPY_BLOCK]]: +// CHECK: [[S_ARR_SRC_EL:%.+]] = phi [[S_INT_TY]]*{{.+}} +// CHECK: [[S_ARR_DST_EL:%.+]] = phi [[S_INT_TY]]*{{.+}} +// CHECK: [[S_ARR_DST_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[S_ARR_DST_EL]] to i8* +// CHECK: [[S_ARR_SRC_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[S_ARR_SRC_EL]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[S_ARR_DST_BCAST]], i8* [[S_ARR_SRC_BCAST]]{{.+}}) +// CHECK: [[S_ARR_DST_NEXT:%.+]] = getelementptr [[S_INT_TY]], [[S_INT_TY]]* [[S_ARR_DST_EL]], i{{[0-9]+}} 1 +// CHECK: [[S_ARR_SRC_NEXT:%.+]] = getelementptr{{.+}} +// CHECK: [[CPY_IS_FINISHED:%.+]] = icmp eq [[S_INT_TY]]* [[S_ARR_DST_NEXT]], [[S_ARR_BEGIN_GEP]] +// CHECK: br i1 [[CPY_IS_FINISHED]], label %[[S_ARR_COPY_DONE]], label %[[S_ARR_COPY_BLOCK]] +// CHECK: [[S_ARR_COPY_DONE]]: +// CHECK: [[TMP_VAL1:%.+]] = load [[S_INT_TY]]*, [[S_INT_TY]]** [[TMP_PRIV1]], +// CHECK: [[VAR_ADDR_REF_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[TMP_REF]] to i8* +// CHECK: [[TMP_VAL1_BCAST:%.+]] = bitcast [[S_INT_TY]]* [[TMP_VAL1]] to i8* +// CHECK: call void @llvm.memcpy.{{.+}}(i8* [[VAR_ADDR_REF_BCAST]], i8* [[TMP_VAL1_BCAST]],{{.+}}) +// CHECK: ret void + +// CHECK: !{!"llvm.loop.vectorize.enable", i1 true} +#endif diff --git a/clang/test/OpenMP/distribute_simd_private_codegen.cpp b/clang/test/OpenMP/distribute_simd_private_codegen.cpp new file mode 100644 index 0000000000000000000000000000000000000000..ede807a0fb63b2062b7fbde99827b4d233964fc0 --- /dev/null +++ b/clang/test/OpenMP/distribute_simd_private_codegen.cpp @@ -0,0 +1,208 @@ +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-32 + +// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +template +struct S { + T f; + S(T a) : f(a) {} + S() : f() {} + operator T() { return T(); } + ~S() {} +}; + +// CHECK: [[S_FLOAT_TY:%.+]] = type { float } +// CHECK: [[S_INT_TY:%.+]] = type { i{{[0-9]+}} } +template +T tmain() { + S test; + T t_var = T(); + T vec[] = {1, 2}; + S s_arr[] = {1, 2}; + S &var = test; + #pragma omp target + #pragma omp teams +#pragma omp distribute simd private(t_var, vec, s_arr, s_arr, var, var) + for (int i = 0; i < 2; ++i) { + vec[i] = t_var; + s_arr[i] = var; + } + return T(); +} + +int main() { + static int svar; + volatile double g; + volatile double &g1 = g; + + #ifdef LAMBDA + // LAMBDA-LABEL: @main + // LAMBDA: call{{.*}} void [[OUTER_LAMBDA:@.+]]( + [&]() { + static float sfvar; + // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( + // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( + // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( + + // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]() + // LAMBDA: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, i32 0, {{.+}}* [[OMP_OUTLINED:@.+]] to {{.+}}) + #pragma omp target + #pragma omp teams +#pragma omp distribute simd private(g, g1, svar, sfvar) + for (int i = 0; i < 2; ++i) { + // LAMBDA: define{{.*}} internal{{.*}} void [[OMP_OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}}) + // LAMBDA: [[G_PRIVATE_ADDR:%.+]] = alloca double, + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = alloca double, + // LAMBDA: [[TMP_PRIVATE_ADDR:%.+]] = alloca double*, + // LAMBDA: [[SVAR_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[SFVAR_PRIVATE_ADDR:%.+]] = alloca float, + // LAMBDA: store double* [[G1_PRIVATE_ADDR]], double** [[TMP_PRIVATE_ADDR]], + g = 1; + g1 = 1; + svar = 3; + sfvar = 4.0; + // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( + // LAMBDA: store double 1.0{{.+}}, double* [[G_PRIVATE_ADDR]], + // LAMBDA: store i{{[0-9]+}} 3, i{{[0-9]+}}* [[SVAR_PRIVATE_ADDR]], + // LAMBDA: store float 4.0{{.+}}, float* [[SFVAR_PRIVATE_ADDR]], + // LAMBDA: [[G_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA: store double* [[G_PRIVATE_ADDR]], double** [[G_PRIVATE_ADDR_REF]], + // LAMBDA: [[TMP_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR_FROM_TMP:%.+]] = load double*, double** [[TMP_PRIVATE_ADDR]], + // LAMBDA: store double* [[G1_PRIVATE_ADDR_FROM_TMP]], double** [[TMP_PRIVATE_ADDR_REF]], + // LAMBDA: [[SVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: store i{{[0-9]+}}* [[SVAR_PRIVATE_ADDR]], i{{[0-9]+}}** [[SVAR_PRIVATE_ADDR_REF]] + // LAMBDA: [[SFVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // LAMBDA: store float* [[SFVAR_PRIVATE_ADDR]], float** [[SFVAR_PRIVATE_ADDR_REF]] + // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]](%{{.+}}* [[ARG]]) + // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( + [&]() { + // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) + // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], + g = 2; + g1 = 2; + svar = 4; + sfvar = 8.0; + // LAMBDA: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] + // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]] + // LAMBDA: store double 2.0{{.+}}, double* [[G_REF]] + + // LAMBDA: [[TMP_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load double*, double** [[TMP_PTR_REF]] + // LAMBDA: store double 2.0{{.+}}, double* [[G1_REF]], + // LAMBDA: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] + // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SVAR_REF]] + // LAMBDA: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // LAMBDA: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] + // LAMBDA: store float 8.0{{.+}}, float* [[SFVAR_REF]] + }(); + } + }(); + return 0; + #else + S test; + int t_var = 0; + int vec[] = {1, 2}; + S s_arr[] = {1, 2}; + S &var = test; + + #pragma omp target + #pragma omp teams +#pragma omp distribute simd private(t_var, vec, s_arr, s_arr, var, var, svar) + for (int i = 0; i < 2; ++i) { + vec[i] = t_var; + s_arr[i] = var; + } + int i; + + #pragma omp target + #pragma omp teams +#pragma omp distribute simd + for (i = 0; i < 2; ++i) { + ; + } + return tmain(); + #endif +} + +// CHECK: define{{.*}} i{{[0-9]+}} @main() +// CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], +// CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( +// CHECK: call void [[OFFLOAD_FUN:@.+]]( +// CHECK: ret + +// CHECK: define{{.+}} [[OFFLOAD_FUN]]() +// CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED:@.+]] to void +// CHECK: ret +// +// CHECK: define internal void [[OMP_OUTLINED]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}}) +// CHECK: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], +// CHECK: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_FLOAT_TY]]], +// CHECK-NOT: alloca [2 x [[S_FLOAT_TY]]], +// CHECK: [[VAR_PRIV:%.+]] = alloca [[S_FLOAT_TY]], +// CHECK-NOT: alloca [[S_FLOAT_TY]], +// CHECK: [[S_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, +// CHECK: store i{{[0-9]+}}* [[GTID_ADDR]], i{{[0-9]+}}** [[GTID_ADDR_REF:%.+]] +// CHECK-NOT: [[T_VAR_PRIV]] +// CHECK-NOT: [[VEC_PRIV]] +// CHECK: {{.+}}: +// CHECK: [[S_ARR_PRIV_ITEM:%.+]] = phi [[S_FLOAT_TY]]* +// CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR]]([[S_FLOAT_TY]]* [[S_ARR_PRIV_ITEM]]) +// CHECK-NOT: [[T_VAR_PRIV]] +// CHECK-NOT: [[VEC_PRIV]] +// CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR]]([[S_FLOAT_TY]]* [[VAR_PRIV]]) +// CHECK: call void @__kmpc_for_static_init_4( +// CHECK: call void @__kmpc_for_static_fini( +// CHECK: ret void + +// CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() +// CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], +// CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( +// CHECK: call void [[OFFLOAD_FUN_1:@.+]]( +// CHECK: ret + + +// CHECK: define internal void [[OFFLOAD_FUN_1]]() +// CHECK: call void (%{{.+}}*, i{{[0-9]+}}, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)*, ...) @__kmpc_fork_teams(%{{.+}}* @{{.+}}, i{{[0-9]+}} 0, void (i{{[0-9]+}}*, i{{[0-9]+}}*, ...)* bitcast (void (i{{[0-9]+}}*, i{{[0-9]+}}*)* [[OMP_OUTLINED_1:@.+]] to void +// CHECK: ret +// +// CHECK: define internal void [[OMP_OUTLINED_1]](i{{[0-9]+}}* noalias [[GTID_ADDR:%.+]], i{{[0-9]+}}* noalias %{{.+}}) +// CHECK: [[T_VAR_PRIV:%.+]] = alloca i{{[0-9]+}}, +// CHECK: [[VEC_PRIV:%.+]] = alloca [2 x i{{[0-9]+}}], +// CHECK: [[S_ARR_PRIV:%.+]] = alloca [2 x [[S_INT_TY]]], +// CHECK-NOT: alloca [2 x [[S_INT_TY]]], +// CHECK: [[VAR_PRIV:%.+]] = alloca [[S_INT_TY]], +// CHECK-NOT: alloca [[S_INT_TY]], +// CHECK: store i{{[0-9]+}}* [[GTID_ADDR]], i{{[0-9]+}}** [[GTID_ADDR_REF:%.+]] +// CHECK-NOT: [[T_VAR_PRIV]] +// CHECK-NOT: [[VEC_PRIV]] +// CHECK: {{.+}}: +// CHECK: [[S_ARR_PRIV_ITEM:%.+]] = phi [[S_INT_TY]]* +// CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR]]([[S_INT_TY]]* [[S_ARR_PRIV_ITEM]]) +// CHECK-NOT: [[T_VAR_PRIV]] +// CHECK-NOT: [[VEC_PRIV]] +// CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR]]([[S_INT_TY]]* [[VAR_PRIV]]) +// CHECK: call void @__kmpc_for_static_init_4( +// CHECK: call void @__kmpc_for_static_fini( +// CHECK: ret void + +// CHECK: !{!"llvm.loop.vectorize.enable", i1 true} +#endif diff --git a/clang/test/OpenMP/distribute_simd_reduction_codegen.cpp b/clang/test/OpenMP/distribute_simd_reduction_codegen.cpp new file mode 100644 index 0000000000000000000000000000000000000000..c061ba926b8aa4eb1ebd3c1fd92e19df26b75335 --- /dev/null +++ b/clang/test/OpenMP/distribute_simd_reduction_codegen.cpp @@ -0,0 +1,201 @@ +// RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 +// RUN: %clang_cc1 -DCHECK -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 +// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DCHECK -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 + +// RUN: %clang_cc1 -DLAMBDA -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -DLAMBDA -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix LAMBDA --check-prefix LAMBDA-64 + +// expected-no-diagnostics +#ifndef HEADER +#define HEADER + +template +T tmain() { + T t_var = T(); + T vec[] = {1, 2}; +#pragma omp target +#pragma omp teams +#pragma omp distribute simd reduction(+: t_var) + for (int i = 0; i < 2; ++i) { + t_var += (T) i; + } + return T(); +} + +int main() { + static int sivar; +#ifdef LAMBDA + // LAMBDA-LABEL: @main + // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( + [&]() { + // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) + // LAMBDA: call void @[[LOFFL1:.+]]( + // LAMBDA: ret +#pragma omp target +#pragma omp teams +#pragma omp distribute simd reduction(+: sivar) + for (int i = 0; i < 2; ++i) { + // LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} [[SIVAR_ARG:%.+]]) + // LAMBDA: [[SIVAR_ADDR:%.+]] = alloca i{{.+}}, + // LAMBDA: store{{.+}} [[SIVAR_ARG]], {{.+}} [[SIVAR_ADDR]], + // LAMBDA: [[SIVAR_CONV:%.+]] = bitcast{{.+}} [[SIVAR_ADDR]] to + // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 1, {{.+}} @[[LOUTL1:.+]] to {{.+}}, {{.+}} [[SIVAR_CONV]]) + // LAMBDA: ret void + + // LAMBDA: define internal void @[[LOUTL1]]({{.+}}, {{.+}}, {{.+}} [[SIVAR_ARG:%.+]]) + // Skip global and bound tid vars + // LAMBDA: {{.+}} = alloca i32*, + // LAMBDA: {{.+}} = alloca i32*, + // LAMBDA: [[SIVAR_ADDR:%.+]] = alloca i{{.+}}*, + // LAMBDA: alloca i{{.+}}, + // LAMBDA: alloca i{{.+}}, + // LAMBDA: alloca i{{.+}}, + // LAMBDA: alloca i{{.+}}, + // LAMBDA: alloca i{{.+}}, + // LAMBDA: [[SIVAR_PRIV:%.+]] = alloca i{{.+}}, + // LAMBDA: store{{.+}} [[SIVAR_ARG]], {{.+}} [[SIVAR_ADDR]], + // LAMBDA: [[SIVAR_REF:%.+]] = load{{.+}}, {{.+}} [[SIVAR_ADDR]] + // LAMBDA: store{{.+}} 0, {{.+}} [[SIVAR_PRIV]], + + // LAMBDA: call void @__kmpc_for_static_init_4( + // LAMBDA: store{{.+}}, {{.+}} [[SIVAR_PRIV]], + // LAMBDA: call void [[INNER_LAMBDA:@.+]]( + // LAMBDA: call void @__kmpc_for_static_fini( + // LAMBDA: [[LAST_ITER:%.+]] = load i32, i32* % + // LAMBDA: [[IS_LAST:%.+]] = icmp ne i32 [[LAST_ITER]], 0 + // LAMBDA: br i1 [[IS_LAST]], label %[[THEN:.+]], label %[[DONE:.+]] + // LAMBDA: [[THEN]] + // LAMBDA: store i32 2, i32* % + // LAMBDA: br label %[[DONE]] + // LAMBDA: [[DONE]] + // LAMBDA: [[SIVAR_ORIG_VAL:%.+]] = load i32, i32* [[SIVAR_REF]], + // LAMBDA: [[SIVAR_PRIV_VAL:%.+]] = load i32, i32* [[SIVAR_PRIV]], + // LAMBDA: [[ADD:%.+]] = add nsw i32 [[SIVAR_ORIG_VAL]], [[SIVAR_PRIV_VAL]] + // LAMBDA: store i32 [[ADD]], i32* [[SIVAR_REF]], + // LAMBDA: ret void + + sivar += i; + + [&]() { + // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) + // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], + + sivar += 4; + // LAMBDA: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] + + // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // LAMBDA: [[SIVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SIVAR_PTR_REF]] + // LAMBDA: [[SIVAR_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SIVAR_REF]] + // LAMBDA: [[SIVAR_INC:%.+]] = add{{.+}} [[SIVAR_VAL]], 4 + // LAMBDA: store i{{[0-9]+}} [[SIVAR_INC]], i{{[0-9]+}}* [[SIVAR_REF]] + }(); + } + }(); + return 0; +#else +#pragma omp target +#pragma omp teams +#pragma omp distribute simd reduction(+: sivar) + for (int i = 0; i < 2; ++i) { + sivar += i; + } + return tmain(); +#endif +} + +// CHECK: define {{.*}}i{{[0-9]+}} @main() +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) +// CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) +// CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() +// CHECK: ret + +// CHECK: define{{.*}} void @[[OFFL1]](i{{64|32}} [[SIVAR_ARG:%.+]]) +// CHECK: [[SIVAR_ADDR:%.+]] = alloca i{{.+}}, +// CHECK: store{{.+}} [[SIVAR_ARG]], {{.+}} [[SIVAR_ADDR]], +// CHECK-64: [[SIVAR_CONV:%.+]] = bitcast{{.+}} [[SIVAR_ADDR]] to +// CHECK-64: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 1, {{.+}} @[[OUTL1:.+]] to {{.+}}, {{.+}} [[SIVAR_CONV]]) +// CHECK-32: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 1, {{.+}} @[[OUTL1:.+]] to {{.+}}, {{.+}} [[SIVAR_ADDR]]) +// CHECK: ret void + +// CHECK: define internal void @[[OUTL1]]({{.+}}, {{.+}}, {{.+}} [[SIVAR_ARG:%.+]]) +// Skip global and bound tid vars +// CHECK: {{.+}} = alloca i32*, +// CHECK: {{.+}} = alloca i32*, +// CHECK: [[SIVAR_ADDR:%.+]] = alloca i{{.+}}*, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: [[SIVAR_PRIV:%.+]] = alloca i{{.+}}, +// CHECK: store{{.+}} [[SIVAR_ARG]], {{.+}} [[SIVAR_ADDR]], +// CHECK: [[SIVAR_REF:%.+]] = load{{.+}}, {{.+}} [[SIVAR_ADDR]] +// CHECK: store{{.+}} 0, {{.+}} [[SIVAR_PRIV]], + +// CHECK: call void @__kmpc_for_static_init_4( +// CHECK: store{{.+}}, {{.+}} [[SIVAR_PRIV]], +// CHECK: call void @__kmpc_for_static_fini( +// CHECK: [[LAST_ITER:%.+]] = load i32, i32* % +// CHECK: [[IS_LAST:%.+]] = icmp ne i32 [[LAST_ITER]], 0 +// CHECK: br i1 [[IS_LAST]], label %[[THEN:.+]], label %[[DONE:.+]] +// CHECK: [[THEN]] +// CHECK: store i32 2, i32* % +// CHECK: br label %[[DONE]] +// CHECK: [[DONE]] +// CHECK: [[SIVAR_ORIG_VAL:%.+]] = load i32, i32* [[SIVAR_REF]], +// CHECK: [[SIVAR_PRIV_VAL:%.+]] = load i32, i32* [[SIVAR_PRIV]], +// CHECK: [[ADD:%.+]] = add nsw i32 [[SIVAR_ORIG_VAL]], [[SIVAR_PRIV_VAL]] +// CHECK: store i32 [[ADD]], i32* [[SIVAR_REF]], +// CHECK: ret void + +// CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, +// CHECK: call void @[[TOFFL1:.+]]({{.+}}) +// CHECK: ret + +// CHECK: define{{.*}} void @[[TOFFL1]](i{{64|32}} [[TVAR_ARG:%.+]]) +// CHECK: [[TVAR_ADDR:%.+]] = alloca i{{.+}}, +// CHECK: store{{.+}} [[TVAR_ARG]], {{.+}} [[TVAR_ADDR]], +// CHECK-64: [[TVAR_CONV:%.+]] = bitcast{{.+}} [[TVAR_ADDR]] to +// CHECK-64: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 1, {{.+}} @[[TOUTL1:.+]] to {{.+}}, {{.+}} [[TVAR_CONV]]) +// CHECK-32: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 1, {{.+}} @[[TOUTL1:.+]] to {{.+}}, {{.+}} [[TVAR_ADDR]]) +// CHECK: ret void + +// CHECK: define internal void @[[TOUTL1]]({{.+}}, {{.+}}, {{.+}} [[TVAR_ARG:%.+]]) +// Skip global and bound tid vars +// CHECK: {{.+}} = alloca i32*, +// CHECK: {{.+}} = alloca i32*, +// CHECK: [[TVAR_ADDR:%.+]] = alloca i{{.+}}*, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: alloca i{{.+}}, +// CHECK: [[TVAR_PRIV:%.+]] = alloca i{{.+}}, +// CHECK: store{{.+}} [[TVAR_ARG]], {{.+}} [[TVAR_ADDR]], +// CHECK: [[TVAR_REF:%.+]] = load{{.+}}, {{.+}} [[TVAR_ADDR]] +// CHECK: store{{.+}} 0, {{.+}} [[TVAR_PRIV]], + +// CHECK: call void @__kmpc_for_static_init_4( +// CHECK: store{{.+}}, {{.+}} [[TVAR_PRIV]], +// CHECK: call void @__kmpc_for_static_fini( +// CHECK: [[LAST_ITER:%.+]] = load i32, i32* % +// CHECK: [[IS_LAST:%.+]] = icmp ne i32 [[LAST_ITER]], 0 +// CHECK: br i1 [[IS_LAST]], label %[[THEN:.+]], label %[[DONE:.+]] +// CHECK: [[THEN]] +// CHECK: store i32 2, i32* % +// CHECK: br label %[[DONE]] +// CHECK: [[DONE]] +// CHECK: [[TVAR_ORIG_VAL:%.+]] = load i32, i32* [[TVAR_REF]], +// CHECK: [[TVAR_PRIV_VAL:%.+]] = load i32, i32* [[TVAR_PRIV]], +// CHECK: [[ADD:%.+]] = add nsw i32 [[TVAR_ORIG_VAL]], [[TVAR_PRIV_VAL]] +// CHECK: store i32 [[ADD]], i32* [[TVAR_REF]], +// CHECK: ret void + +// CHECK: !{!"llvm.loop.vectorize.enable", i1 true} +#endif