Index: cfe/trunk/include/clang/AST/OpenMPClause.h =================================================================== --- cfe/trunk/include/clang/AST/OpenMPClause.h +++ cfe/trunk/include/clang/AST/OpenMPClause.h @@ -3479,7 +3479,7 @@ /// In this example directive '#pragma omp teams' has clause 'num_teams' /// with single expression 'n'. /// -class OMPNumTeamsClause : public OMPClause { +class OMPNumTeamsClause : public OMPClause, public OMPClauseWithPreInit { friend class OMPClauseReader; /// \brief Location of '('. SourceLocation LParenLoc; @@ -3495,20 +3495,27 @@ /// \brief Build 'num_teams' clause. /// /// \param E Expression associated with this clause. + /// \param HelperE Helper Expression associated with this clause. + /// \param CaptureRegion Innermost OpenMP region where expressions in this + /// clause must be captured. /// \param StartLoc Starting location of the clause. /// \param LParenLoc Location of '('. /// \param EndLoc Ending location of the clause. /// - OMPNumTeamsClause(Expr *E, SourceLocation StartLoc, SourceLocation LParenLoc, + OMPNumTeamsClause(Expr *E, Stmt *HelperE, OpenMPDirectiveKind CaptureRegion, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) - : OMPClause(OMPC_num_teams, StartLoc, EndLoc), LParenLoc(LParenLoc), - NumTeams(E) {} + : OMPClause(OMPC_num_teams, StartLoc, EndLoc), OMPClauseWithPreInit(this), + LParenLoc(LParenLoc), NumTeams(E) { + setPreInitStmt(HelperE, CaptureRegion); + } /// \brief Build an empty clause. /// OMPNumTeamsClause() - : OMPClause(OMPC_num_teams, SourceLocation(), SourceLocation()), - LParenLoc(SourceLocation()), NumTeams(nullptr) {} + : OMPClause(OMPC_num_teams, SourceLocation(), SourceLocation()), + OMPClauseWithPreInit(this), LParenLoc(SourceLocation()), + NumTeams(nullptr) {} /// \brief Sets the location of '('. void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; } /// \brief Returns the location of '('. Index: cfe/trunk/include/clang/AST/RecursiveASTVisitor.h =================================================================== --- cfe/trunk/include/clang/AST/RecursiveASTVisitor.h +++ cfe/trunk/include/clang/AST/RecursiveASTVisitor.h @@ -2995,6 +2995,7 @@ template bool RecursiveASTVisitor::VisitOMPNumTeamsClause( OMPNumTeamsClause *C) { + TRY_TO(VisitOMPClauseWithPreInit(C)); TRY_TO(TraverseStmt(C->getNumTeams())); return true; } Index: cfe/trunk/lib/AST/OpenMPClause.cpp =================================================================== --- cfe/trunk/lib/AST/OpenMPClause.cpp +++ cfe/trunk/lib/AST/OpenMPClause.cpp @@ -52,6 +52,8 @@ return static_cast(C); case OMPC_num_threads: return static_cast(C); + case OMPC_num_teams: + return static_cast(C); case OMPC_default: case OMPC_proc_bind: case OMPC_final: @@ -79,7 +81,6 @@ case OMPC_threads: case OMPC_simd: case OMPC_map: - case OMPC_num_teams: case OMPC_thread_limit: case OMPC_priority: case OMPC_grainsize: Index: cfe/trunk/lib/AST/StmtProfile.cpp =================================================================== --- cfe/trunk/lib/AST/StmtProfile.cpp +++ cfe/trunk/lib/AST/StmtProfile.cpp @@ -497,6 +497,7 @@ VisitOMPClauseList(C); } void OMPClauseProfiler::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) { + VistOMPClauseWithPreInit(C); if (C->getNumTeams()) Profiler->VisitStmt(C->getNumTeams()); } Index: cfe/trunk/lib/Sema/SemaOpenMP.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaOpenMP.cpp +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp @@ -6771,6 +6771,69 @@ llvm_unreachable("Unknown OpenMP directive"); } break; + case OMPC_num_teams: + switch (DKind) { + case OMPD_target_teams: + CaptureRegion = OMPD_target; + break; + case OMPD_cancel: + case OMPD_parallel: + case OMPD_parallel_sections: + case OMPD_parallel_for: + case OMPD_parallel_for_simd: + case OMPD_target: + case OMPD_target_simd: + case OMPD_target_parallel: + case OMPD_target_parallel_for: + case OMPD_target_parallel_for_simd: + case OMPD_target_teams_distribute: + case OMPD_target_teams_distribute_simd: + case OMPD_target_teams_distribute_parallel_for: + case OMPD_target_teams_distribute_parallel_for_simd: + case OMPD_teams_distribute_parallel_for: + case OMPD_teams_distribute_parallel_for_simd: + case OMPD_distribute_parallel_for: + case OMPD_distribute_parallel_for_simd: + case OMPD_task: + case OMPD_taskloop: + case OMPD_taskloop_simd: + case OMPD_target_data: + case OMPD_target_enter_data: + case OMPD_target_exit_data: + case OMPD_target_update: + case OMPD_teams: + case OMPD_teams_distribute: + case OMPD_teams_distribute_simd: + // Do not capture num_teams-clause expressions. + break; + case OMPD_threadprivate: + case OMPD_taskyield: + case OMPD_barrier: + case OMPD_taskwait: + case OMPD_cancellation_point: + case OMPD_flush: + case OMPD_declare_reduction: + case OMPD_declare_simd: + case OMPD_declare_target: + case OMPD_end_declare_target: + case OMPD_simd: + case OMPD_for: + case OMPD_for_simd: + case OMPD_sections: + case OMPD_section: + case OMPD_single: + case OMPD_master: + case OMPD_critical: + case OMPD_taskgroup: + case OMPD_distribute: + case OMPD_ordered: + case OMPD_atomic: + case OMPD_distribute_simd: + llvm_unreachable("Unexpected OpenMP directive with num_teams-clause"); + case OMPD_unknown: + llvm_unreachable("Unknown OpenMP directive"); + } + break; case OMPC_schedule: case OMPC_dist_schedule: case OMPC_firstprivate: @@ -6804,7 +6867,6 @@ case OMPC_threads: case OMPC_simd: case OMPC_map: - case OMPC_num_teams: case OMPC_thread_limit: case OMPC_priority: case OMPC_grainsize: @@ -10860,6 +10922,8 @@ SourceLocation LParenLoc, SourceLocation EndLoc) { Expr *ValExpr = NumTeams; + Stmt *HelperValStmt = nullptr; + OpenMPDirectiveKind CaptureRegion = OMPD_unknown; // OpenMP [teams Constrcut, Restrictions] // The num_teams expression must evaluate to a positive integer value. @@ -10867,7 +10931,16 @@ /*StrictlyPositive=*/true)) return nullptr; - return new (Context) OMPNumTeamsClause(ValExpr, StartLoc, LParenLoc, EndLoc); + OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective(); + CaptureRegion = getOpenMPCaptureRegionForClause(DKind, OMPC_num_teams); + if (CaptureRegion != OMPD_unknown) { + llvm::MapVector Captures; + ValExpr = tryBuildCapture(*this, ValExpr, Captures).get(); + HelperValStmt = buildPreInits(Context, Captures); + } + + return new (Context) OMPNumTeamsClause(ValExpr, HelperValStmt, CaptureRegion, + StartLoc, LParenLoc, EndLoc); } OMPClause *Sema::ActOnOpenMPThreadLimitClause(Expr *ThreadLimit, Index: cfe/trunk/lib/Serialization/ASTReaderStmt.cpp =================================================================== --- cfe/trunk/lib/Serialization/ASTReaderStmt.cpp +++ cfe/trunk/lib/Serialization/ASTReaderStmt.cpp @@ -2300,6 +2300,7 @@ } void OMPClauseReader::VisitOMPNumTeamsClause(OMPNumTeamsClause *C) { + VisitOMPClauseWithPreInit(C); C->setNumTeams(Reader->Record.readSubExpr()); C->setLParenLoc(Reader->ReadSourceLocation()); } Index: cfe/trunk/lib/Serialization/ASTWriterStmt.cpp =================================================================== --- cfe/trunk/lib/Serialization/ASTWriterStmt.cpp +++ cfe/trunk/lib/Serialization/ASTWriterStmt.cpp @@ -2067,6 +2067,7 @@ } void OMPClauseWriter::VisitOMPNumTeamsClause(OMPNumTeamsClause *C) { + VisitOMPClauseWithPreInit(C); Record.AddStmt(C->getNumTeams()); Record.AddSourceLocation(C->getLParenLoc()); } Index: cfe/trunk/test/OpenMP/target_teams_num_teams_codegen.cpp =================================================================== --- cfe/trunk/test/OpenMP/target_teams_num_teams_codegen.cpp +++ cfe/trunk/test/OpenMP/target_teams_num_teams_codegen.cpp @@ -0,0 +1,344 @@ +// Test host codegen. +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 + +// Test target codegen - host bc file has to be created first. +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 --check-prefix TCHECK --check-prefix TCHECK-64 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 --check-prefix TCHECK --check-prefix TCHECK-64 +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 --check-prefix TCHECK --check-prefix TCHECK-32 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -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 --check-prefix TCHECK --check-prefix TCHECK-32 + +// 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:@.+]] = 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: [[S1:%.+]] = type { double } +// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } +// CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } +// CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } + +// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 } + +// We have 6 target regions + +// CHECK-DAG: @{{.*}} = private constant i8 0 +// CHECK-DAG: @{{.*}} = private constant i8 0 +// CHECK-DAG: @{{.*}} = private constant i8 0 +// CHECK-DAG: @{{.*}} = private constant i8 0 +// CHECK-DAG: @{{.*}} = private constant i8 0 +// CHECK-DAG: @{{.*}} = private constant i8 0 + +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] + +// Check if offloading descriptor is created. +// CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] +// CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] +// CHECK: [[DEVBEGIN:@.+]] = external constant i8 +// CHECK: [[DEVEND:@.+]] = external constant i8 +// CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }] +// CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] } + +// Check target registration is registered as a Ctor. +// CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* [[REGFN:@.+]] to void ()*), i8* null }] + + +template +tx ftemplate(int n) { + tx a = 0; + + #pragma omp target teams num_teams(tx(20)) + { + } + + short b = 1; + #pragma omp target teams num_teams(b) + { + a += b; + } + + return a; +} + +static +int fstatic(int n) { + + #pragma omp target teams num_teams(n) + { + } + + #pragma omp target teams num_teams(32+n) + { + } + + return n+1; +} + +struct S1 { + double a; + + int r1(int n){ + int b = 1; + + #pragma omp target teams num_teams(n-b) + { + this->a = (double)b + 1.5; + } + + #pragma omp target teams num_teams(1024) + { + this->a = 2.5; + } + + return (int)a; + } +}; + +// CHECK: define {{.*}}@{{.*}}bar{{.*}} +int bar(int n){ + int a = 0; + + S1 S; + // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}}) + a += S.r1(n); + + // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}}) + a += fstatic(n); + + // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}}) + a += ftemplate(n); + + return a; +} + + + +// +// CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]]) +// +// CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align +// CHECK: store i32 1, i32* [[B:%.+]], align +// CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align +// CHECK: [[BV:%.+]] = load i32, i32* [[B]], align +// CHECK: [[SUB:%.+]] = sub nsw i32 [[NV]], [[BV]] +// CHECK: store i32 [[SUB]], i32* [[CAPE_ADDR:%.+]], align +// CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* +// CHECK-64: store i32 [[CEV]], i32* [[CONV]], align +// CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align +// CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align +// CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 0) +// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align +// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] +// +// CHECK: [[FAIL]] +// CHECK: call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) +// CHECK: br label {{%?}}[[END]] +// CHECK: [[END]] +// +// +// +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 1024, i32 0) +// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align +// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] +// +// CHECK: [[FAIL]] +// CHECK: call void [[HVT2:@.+]]([[S1]]* {{[^,]+}}) +// CHECK: br label {{%?}}[[END]] +// CHECK: [[END]] +// + + + + + + +// +// CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]]) +// +// CHECK-DAG: store i32 [[PARM]], i32* [[N_ADDR:%.+]], align +// CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align +// CHECK: store i32 [[NV]], i32* [[CAPE_ADDR:%.+]], align +// CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* +// CHECK-64: store i32 [[CEV]], i32* [[CONV]], align +// CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align +// CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align +// CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 [[TEAMS]], i32 0) +// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align +// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] +// +// CHECK: [[FAIL]] +// CHECK: call void [[HVT3:@.+]](i[[SZ]] [[ARG]]) +// CHECK: br label {{%?}}[[END]] +// CHECK: [[END]] +// +// +// +// CHECK: [[NV:%.+]] = load i32, i32* [[N_ADDR]], align +// CHECK: [[ADD:%.+]] = add nsw i32 32, [[NV]] +// CHECK: store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align +// CHECK: [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32* +// CHECK-64: store i32 [[CEV]], i32* [[CONV]], align +// CHECK-32: store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align +// CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align +// CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 [[TEAMS]], i32 0) +// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align +// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] +// +// CHECK: [[FAIL]] +// CHECK: call void [[HVT4:@.+]](i[[SZ]] [[ARG]]) +// CHECK: br label {{%?}}[[END]] +// CHECK: [[END]] +// + + + + + + +// +// CHECK: define {{.*}}[[FTEMPLATE]] +// +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 20, i32 0) +// CHECK-NEXT: store i32 [[RET]], i32* [[RHV:%.+]], align +// CHECK-NEXT: [[RET2:%.+]] = load i32, i32* [[RHV]], align +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] +// +// CHECK: [[FAIL]] +// CHECK: call void [[HVT5:@.+]]() +// CHECK: br label {{%?}}[[END]] +// +// CHECK: [[END]] +// +// +// +// CHECK: store i16 1, i16* [[B:%.+]], align +// CHECK: [[BV:%.+]] = load i16, i16* [[B]], align +// CHECK: store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align +// CHECK: [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align +// CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16* +// CHECK: store i16 [[CEV]], i16* [[CONV]], align +// CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align +// CHECK: [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align +// CHECK: [[TEAMS:%.+]] = sext i16 [[T]] to i32 +// +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 0) +// CHECK: store i32 [[RET]], i32* [[RHV:%.+]], align +// CHECK: [[RET2:%.+]] = load i32, i32* [[RHV]], align +// CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0 +// CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] +// +// CHECK: [[FAIL]] +// CHECK: call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]]) +// CHECK: br label {{%?}}[[END]] +// CHECK: [[END]] +// + + + + + + +// Check that the offloading functions are emitted and that the parallel function +// is appropriately guarded. + +// CHECK: define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]]) +// CHECK-DAG: store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align +// CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* +// CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV]], align +// CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) +// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2, +// +// + + +// CHECK: define internal void [[HVT2]]([[S1]]* {{%.+}}) +// CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 1024, i32 0) +// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 1, +// +// + + + + + + + + +// CHECK: define internal void [[HVT3]](i[[SZ]] [[PARM:%.+]]) +// CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align +// CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* +// CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV]], align +// CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) +// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, +// +// +// CHECK: define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]]) +// CHECK-DAG: store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align +// CHECK-64: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32* +// CHECK-64: [[NT:%.+]] = load i32, i32* [[CONV]], align +// CHECK-32: [[NT:%.+]] = load i32, i32* [[CAPE_ADDR]], align +// CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) +// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, +// +// + + + + + +// CHECK: define internal void [[HVT5]]( +// CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 20, i32 0) +// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0, +// +// + + +// CHECK: define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]]) +// CHECK-DAG: store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align +// CHECK: [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16* +// CHECK: [[T:%.+]] = load i16, i16* [[CONV]], align +// CHECK: [[NT:%.+]] = sext i16 [[T]] to i32 +// CHECK: call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 0) +// CHECK: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2, +// +// + + + +#endif Index: cfe/trunk/tools/libclang/CIndex.cpp =================================================================== --- cfe/trunk/tools/libclang/CIndex.cpp +++ cfe/trunk/tools/libclang/CIndex.cpp @@ -2169,6 +2169,7 @@ } void OMPClauseEnqueue::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) { + VisitOMPClauseWithPreInit(C); Visitor->AddStmt(C->getNumTeams()); }