forked from OSchip/llvm-project
[OpenMP] Support for num_teams-clause on the 'target teams' directive.
The num_teams-clause on the combined directive applies to the 'teams' region of this construct. We modify the NumTeamsClause class to capture the clause expression within the 'target' region. Reviewers: ABataev Differential Revision: https://reviews.llvm.org/D29085 llvm-svn: 293048
This commit is contained in:
Arpith Chacko Jacob
parent
46897a46ee
commit
bc126344e1
|
|
@ -3479,7 +3479,7 @@ public:
|
|||
/// 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 @@ public:
|
|||
/// \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 '('.
|
||||
|
|
|
|||
|
|
@ -2995,6 +2995,7 @@ bool RecursiveASTVisitor<Derived>::VisitOMPMapClause(OMPMapClause *C) {
|
|||
template <typename Derived>
|
||||
bool RecursiveASTVisitor<Derived>::VisitOMPNumTeamsClause(
|
||||
OMPNumTeamsClause *C) {
|
||||
TRY_TO(VisitOMPClauseWithPreInit(C));
|
||||
TRY_TO(TraverseStmt(C->getNumTeams()));
|
||||
return true;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -52,6 +52,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
|
|||
return static_cast<const OMPIfClause *>(C);
|
||||
case OMPC_num_threads:
|
||||
return static_cast<const OMPNumThreadsClause *>(C);
|
||||
case OMPC_num_teams:
|
||||
return static_cast<const OMPNumTeamsClause *>(C);
|
||||
case OMPC_default:
|
||||
case OMPC_proc_bind:
|
||||
case OMPC_final:
|
||||
|
|
@ -79,7 +81,6 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
|
|||
case OMPC_threads:
|
||||
case OMPC_simd:
|
||||
case OMPC_map:
|
||||
case OMPC_num_teams:
|
||||
case OMPC_thread_limit:
|
||||
case OMPC_priority:
|
||||
case OMPC_grainsize:
|
||||
|
|
|
|||
|
|
@ -497,6 +497,7 @@ void OMPClauseProfiler::VisitOMPMapClause(const OMPMapClause *C) {
|
|||
VisitOMPClauseList(C);
|
||||
}
|
||||
void OMPClauseProfiler::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) {
|
||||
VistOMPClauseWithPreInit(C);
|
||||
if (C->getNumTeams())
|
||||
Profiler->VisitStmt(C->getNumTeams());
|
||||
}
|
||||
|
|
|
|||
|
|
@ -6771,6 +6771,69 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
|||
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 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
|
|||
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 @@ OMPClause *Sema::ActOnOpenMPNumTeamsClause(Expr *NumTeams,
|
|||
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 @@ OMPClause *Sema::ActOnOpenMPNumTeamsClause(Expr *NumTeams,
|
|||
/*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<Expr *, DeclRefExpr *> 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,
|
||||
|
|
|
|||
|
|
@ -2300,6 +2300,7 @@ void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
|
|||
}
|
||||
|
||||
void OMPClauseReader::VisitOMPNumTeamsClause(OMPNumTeamsClause *C) {
|
||||
VisitOMPClauseWithPreInit(C);
|
||||
C->setNumTeams(Reader->Record.readSubExpr());
|
||||
C->setLParenLoc(Reader->ReadSourceLocation());
|
||||
}
|
||||
|
|
|
|||
|
|
@ -2067,6 +2067,7 @@ void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
|
|||
}
|
||||
|
||||
void OMPClauseWriter::VisitOMPNumTeamsClause(OMPNumTeamsClause *C) {
|
||||
VisitOMPClauseWithPreInit(C);
|
||||
Record.AddStmt(C->getNumTeams());
|
||||
Record.AddSourceLocation(C->getLParenLoc());
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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<typename tx>
|
||||
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<int>(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
|
||||
|
|
@ -2169,6 +2169,7 @@ void OMPClauseEnqueue::VisitOMPDeviceClause(const OMPDeviceClause *C) {
|
|||
}
|
||||
|
||||
void OMPClauseEnqueue::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) {
|
||||
VisitOMPClauseWithPreInit(C);
|
||||
Visitor->AddStmt(C->getNumTeams());
|
||||
}
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue