https://github.com/erichkeane updated
https://github.com/llvm/llvm-project/pull/135851
>From 3ad13136ba9357873792392b61d14f5b131a472a Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Tue, 15 Apr 2025 10:59:38 -0700
Subject: [PATCH 1/3] [OpenACC][CIR] Implement 'self' lowering on compute
constructs
This is our first attempt at lowering a clause that is an 'operand' in
the OpenACC operand, so it does quite a bit of refactoring. My previous
plans on how to emit the clauses was not viable, so we instead do
'create the op, then use the visitor to fill in the operands'. This
resulted in the 'applyAttributes' function getting removed and a few
other functions simplified.
Additionally, it requires setting the insertion point a little to make
sure we're inserting 'around' the operation correctly.
Finally, since the OpenACC dialect only understands the MLIR types, we
had to introduce a use of the unrealized-conversion-cast, which we'll
probably getting good use out of in the future.
---
clang/include/clang/AST/OpenACCClause.h | 5 +
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 199 +++++++++++---------
clang/test/CIR/CodeGenOpenACC/kernels.c | 30 ++-
clang/test/CIR/CodeGenOpenACC/parallel.c | 30 ++-
clang/test/CIR/CodeGenOpenACC/serial.c | 30 ++-
5 files changed, 197 insertions(+), 97 deletions(-)
diff --git a/clang/include/clang/AST/OpenACCClause.h
b/clang/include/clang/AST/OpenACCClause.h
index 3687af76a559f..681567228cbb0 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -430,6 +430,11 @@ class OpenACCSelfClause final
}
bool isConditionExprClause() const { return HasConditionExpr.has_value(); }
+ bool isVarListClause() const { return !isConditionExprClause(); }
+ bool isEmptySelfClause() const {
+ return (isConditionExprClause() && !hasConditionExpr()) ||
+ (!isConditionExprClause() && getVarList().empty());
+ }
bool hasConditionExpr() const {
assert(HasConditionExpr.has_value() &&
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 152f996ed0fed..92f3ad2a68eb8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -32,46 +32,52 @@ constexpr bool isOneOfTypes =
template <typename ToTest, typename T>
constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
+template <typename OpTy>
class OpenACCClauseCIREmitter final
- : public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
- CIRGenModule &cgm;
+ : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
+ OpTy &operation;
+ CIRGenFunction &cgf;
+ CIRGenBuilderTy &builder;
+
// This is necessary since a few of the clauses emit differently based on the
// directive kind they are attached to.
OpenACCDirectiveKind dirKind;
+ // This source location should be able to go away once the NYI diagnostics
are
+ // gone.
SourceLocation dirLoc;
- struct AttributeData {
- // Value of the 'default' attribute, added on 'data' and 'compute'/etc
- // constructs as a 'default-attr'.
- std::optional<ClauseDefaultValue> defaultVal = std::nullopt;
- // For directives that have their device type architectures listed in
- // attributes (init/shutdown/etc), the list of architectures to be emitted.
- llvm::SmallVector<mlir::acc::DeviceType> deviceTypeArchs{};
- } attrData;
-
void clauseNotImplemented(const OpenACCClause &c) {
- cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
+ cgf.getCIRGenModule().errorNYI(c.getSourceRange(), "OpenACC Clause",
+ c.getClauseKind());
}
public:
- OpenACCClauseCIREmitter(CIRGenModule &cgm, OpenACCDirectiveKind dirKind,
- SourceLocation dirLoc)
- : cgm(cgm), dirKind(dirKind), dirLoc(dirLoc) {}
+ OpenACCClauseCIREmitter(OpTy &operation, CIRGenFunction &cgf,
+ CIRGenBuilderTy &builder,
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc)
+ : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind),
+ dirLoc(dirLoc) {}
void VisitClause(const OpenACCClause &clause) {
clauseNotImplemented(clause);
}
void VisitDefaultClause(const OpenACCDefaultClause &clause) {
- switch (clause.getDefaultClauseKind()) {
- case OpenACCDefaultClauseKind::None:
- attrData.defaultVal = ClauseDefaultValue::None;
- break;
- case OpenACCDefaultClauseKind::Present:
- attrData.defaultVal = ClauseDefaultValue::Present;
- break;
- case OpenACCDefaultClauseKind::Invalid:
- break;
+ // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
+ // operations listed in the rest of the arguments.
+ if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, DataOp>)
{
+ switch (clause.getDefaultClauseKind()) {
+ case OpenACCDefaultClauseKind::None:
+ operation.setDefaultAttr(ClauseDefaultValue::None);
+ break;
+ case OpenACCDefaultClauseKind::Present:
+ operation.setDefaultAttr(ClauseDefaultValue::Present);
+ break;
+ case OpenACCDefaultClauseKind::Invalid:
+ break;
+ }
+ } else {
+ return clauseNotImplemented(clause);
}
}
@@ -89,64 +95,70 @@ class OpenACCClauseCIREmitter final
}
void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) {
+ llvm::SmallVector<mlir::Attribute> deviceTypes;
+ std::optional<mlir::ArrayAttr> existingDeviceTypes =
+ operation.getDeviceTypes();
+
+ // Ensure we keep the existing ones, and in the correct 'new' order.
+ if (existingDeviceTypes) {
+ for (const mlir::Attribute &Attr : *existingDeviceTypes)
+ deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(),
+ cast<mlir::acc::DeviceTypeAttr>(Attr).getValue()));
+ }
- switch (dirKind) {
- case OpenACCDirectiveKind::Init:
- case OpenACCDirectiveKind::Set:
- case OpenACCDirectiveKind::Shutdown: {
- // Device type has a list that is either a 'star' (emitted as 'star'),
- // or an identifer list, all of which get added for attributes.
-
- for (const DeviceTypeArgument &arg : clause.getArchitectures())
- attrData.deviceTypeArchs.push_back(decodeDeviceType(arg.first));
- break;
- }
- default:
+ for (const DeviceTypeArgument &arg : clause.getArchitectures()) {
+ deviceTypes.push_back(mlir::acc::DeviceTypeAttr::get(
+ builder.getContext(), decodeDeviceType(arg.first)));
+ }
+ operation.removeDeviceTypesAttr();
+ operation.setDeviceTypesAttr(
+ mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
+ } else if constexpr (isOneOfTypes<OpTy, SetOp>) {
+ assert(!operation.getDeviceTypeAttr() && "already have device-type?");
+ assert(clause.getArchitectures().size() <= 1);
+
+ if (!clause.getArchitectures().empty())
+ operation.setDeviceType(
+ decodeDeviceType(clause.getArchitectures()[0].first));
+ } else {
return clauseNotImplemented(clause);
}
}
- // Apply any of the clauses that resulted in an 'attribute'.
- template <typename Op>
- void applyAttributes(CIRGenBuilderTy &builder, Op &op) {
-
- if (attrData.defaultVal.has_value()) {
- // FIXME: OpenACC: as we implement this for other directive kinds, we
have
- // to expand this list.
- // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
- // operations listed in the rest of the arguments.
- if constexpr (isOneOfTypes<Op, ParallelOp, SerialOp, KernelsOp, DataOp>)
- op.setDefaultAttr(*attrData.defaultVal);
- else
- cgm.errorNYI(dirLoc, "OpenACC 'default' clause lowering for ",
dirKind);
- }
-
- if (!attrData.deviceTypeArchs.empty()) {
- // FIXME: OpenACC: as we implement this for other directive kinds, we
have
- // to expand this list, or more likely, have a 'noop' branch as most
other
- // uses of this apply to the operands instead.
- // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
- if constexpr (isOneOfTypes<Op, InitOp, ShutdownOp>) {
- llvm::SmallVector<mlir::Attribute> deviceTypes;
- for (mlir::acc::DeviceType DT : attrData.deviceTypeArchs)
- deviceTypes.push_back(
- mlir::acc::DeviceTypeAttr::get(builder.getContext(), DT));
-
- op.setDeviceTypesAttr(
- mlir::ArrayAttr::get(builder.getContext(), deviceTypes));
- } else if constexpr (isOneOfTypes<Op, SetOp>) {
- assert(attrData.deviceTypeArchs.size() <= 1 &&
- "Set can only have a single architecture");
- if (!attrData.deviceTypeArchs.empty())
- op.setDeviceType(attrData.deviceTypeArchs[0]);
+ void VisitSelfClause(const OpenACCSelfClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
+ if (clause.isEmptySelfClause()) {
+ operation.setSelfAttr(true);
+ } else if (clause.isConditionExprClause()) {
+ assert(clause.hasConditionExpr());
+ mlir::Value condition =
+ cgf.evaluateExprAsBool(clause.getConditionExpr());
+
+ mlir::Location exprLoc = cgf.getCIRGenModule().getLoc(
+ clause.getConditionExpr()->getBeginLoc());
+ mlir::IntegerType targetType = mlir::IntegerType::get(
+ &cgf.getMLIRContext(), /*width=*/1,
+ mlir::IntegerType::SignednessSemantics::Signless);
+ auto ConversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
+ exprLoc, targetType, condition);
+ operation.getSelfCondMutable().append(ConversionOp.getResult(0));
} else {
- cgm.errorNYI(dirLoc, "OpenACC 'device_type' clause lowering for ",
- dirKind);
+ llvm_unreachable("var-list version of self shouldn't get here");
}
+ } else {
+ return clauseNotImplemented(clause);
}
}
};
+template <typename OpTy>
+auto makeClauseEmitter(OpTy &op, CIRGenFunction &cgf, CIRGenBuilderTy &builder,
+ OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
+ return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc);
+}
+
} // namespace
template <typename Op, typename TermOp>
@@ -158,24 +170,27 @@ mlir::LogicalResult
CIRGenFunction::emitOpenACCOpAssociatedStmt(
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
-
- // Clause-emitter must be here because it might modify operands.
- OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc);
- clauseEmitter.VisitClauseList(clauses);
-
auto op = builder.create<Op>(start, retTy, operands);
- // Apply the attributes derived from the clauses.
- clauseEmitter.applyAttributes(builder, op);
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ // Sets insertion point before the 'op', since every new expression needs
to
+ // be before the operation.
+ builder.setInsertionPoint(op);
+ makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
+ .VisitClauseList(clauses);
+ }
- mlir::Block &block = op.getRegion().emplaceBlock();
- mlir::OpBuilder::InsertionGuard guardCase(builder);
- builder.setInsertionPointToEnd(&block);
+ {
+ mlir::Block &block = op.getRegion().emplaceBlock();
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ builder.setInsertionPointToEnd(&block);
- LexicalScope ls{*this, start, builder.getInsertionBlock()};
- res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
+ LexicalScope ls{*this, start, builder.getInsertionBlock()};
+ res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
- builder.create<TermOp>(end);
+ builder.create<TermOp>(end);
+ }
return res;
}
@@ -187,14 +202,16 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCOp(
llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
-
- // Clause-emitter must be here because it might modify operands.
- OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule(), dirKind, dirLoc);
- clauseEmitter.VisitClauseList(clauses);
-
auto op = builder.create<Op>(start, retTy, operands);
- // Apply the attributes derived from the clauses.
- clauseEmitter.applyAttributes(builder, op);
+
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ // Sets insertion point before the 'op', since every new expression needs
to
+ // be before the operation.
+ builder.setInsertionPoint(op);
+ makeClauseEmitter(op, *this, builder, dirKind, dirLoc)
+ .VisitClauseList(clauses);
+ }
return res;
}
diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c
b/clang/test/CIR/CodeGenOpenACC/kernels.c
index 0c950fe3d0f9c..934daf9e8ecc0 100644
--- a/clang/test/CIR/CodeGenOpenACC/kernels.c
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_kernels(void) {
- // CHECK: cir.func @acc_kernels() {
+void acc_kernels(int cond) {
+ // CHECK: cir.func @acc_kernels(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond",
init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc kernels
{}
@@ -38,5 +40,29 @@ void acc_kernels(void) {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT:}
+#pragma acc kernels self
+ {}
+ // CHECK-NEXT: acc.kernels {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc kernels self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>,
!s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] :
!s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast
%[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
+#pragma acc kernels self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] :
!s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast
%[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c
b/clang/test/CIR/CodeGenOpenACC/parallel.c
index e18270435460c..c7a4bda6faa74 100644
--- a/clang/test/CIR/CodeGenOpenACC/parallel.c
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_parallel(void) {
- // CHECK: cir.func @acc_parallel() {
+void acc_parallel(int cond) {
+ // CHECK: cir.func @acc_parallel(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond",
init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc parallel
{}
// CHECK-NEXT: acc.parallel {
@@ -37,5 +39,29 @@ void acc_parallel(void) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc parallel self
+ {}
+ // CHECK-NEXT: acc.parallel {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc parallel self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>,
!s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] :
!s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast
%[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
+#pragma acc parallel self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] :
!s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast
%[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.parallel self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c
b/clang/test/CIR/CodeGenOpenACC/serial.c
index 72a0995549da3..38a38ad6c9514 100644
--- a/clang/test/CIR/CodeGenOpenACC/serial.c
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -1,7 +1,9 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
-void acc_serial(void) {
- // CHECK: cir.func @acc_serial() {
+void acc_serial(int cond) {
+ // CHECK: cir.func @acc_serial(%[[ARG:.*]]: !s32i{{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond",
init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
#pragma acc serial
{}
@@ -38,5 +40,29 @@ void acc_serial(void) {
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}
+#pragma acc serial self
+ {}
+ // CHECK-NEXT: acc.serial {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } attributes {selfAttr}
+
+#pragma acc serial self(cond)
+ {}
+ // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>,
!s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] :
!s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast
%[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
+#pragma acc serial self(0)
+ {}
+ // CHECK-NEXT: %[[ZERO_LITERAL:.*]] = cir.const #cir.int<0> : !s32i
+ // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ZERO_LITERAL]] :
!s32i), !cir.bool
+ // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast
%[[BOOL_CAST]] : !cir.bool to i1
+ // CHECK-NEXT: acc.serial self(%[[CONV_CAST]]) {
+ // CHECK-NEXT: acc.yield
+ // CHECK-NEXT: } loc
+
// CHECK-NEXT: cir.return
}
>From 3a7a45720ff9b0dde3bff8abf0fd9a3527ffbb33 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Tue, 15 Apr 2025 14:15:37 -0700
Subject: [PATCH 2/3] Fix spelling of
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 92f3ad2a68eb8..484369c2d6e4e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -141,9 +141,9 @@ class OpenACCClauseCIREmitter final
mlir::IntegerType targetType = mlir::IntegerType::get(
&cgf.getMLIRContext(), /*width=*/1,
mlir::IntegerType::SignednessSemantics::Signless);
- auto ConversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
+ auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
exprLoc, targetType, condition);
- operation.getSelfCondMutable().append(ConversionOp.getResult(0));
+ operation.getSelfCondMutable().append(conversionOp.getResult(0));
} else {
llvm_unreachable("var-list version of self shouldn't get here");
}
>From b890b7b8330e04b3fb0b164a0c26b8cda2f0921a Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Tue, 15 Apr 2025 15:22:30 -0700
Subject: [PATCH 3/3] Nits from Andy
---
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 29 ++++++++++-----------
1 file changed, 14 insertions(+), 15 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 484369c2d6e4e..3bcc6f908a841 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -42,13 +42,12 @@ class OpenACCClauseCIREmitter final
// This is necessary since a few of the clauses emit differently based on the
// directive kind they are attached to.
OpenACCDirectiveKind dirKind;
- // This source location should be able to go away once the NYI diagnostics
are
- // gone.
+ // TODO(cir): This source location should be able to go away once the NYI
+ // diagnostics are gone.
SourceLocation dirLoc;
void clauseNotImplemented(const OpenACCClause &c) {
- cgf.getCIRGenModule().errorNYI(c.getSourceRange(), "OpenACC Clause",
- c.getClauseKind());
+ cgf.cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
}
public:
@@ -136,8 +135,8 @@ class OpenACCClauseCIREmitter final
mlir::Value condition =
cgf.evaluateExprAsBool(clause.getConditionExpr());
- mlir::Location exprLoc = cgf.getCIRGenModule().getLoc(
- clause.getConditionExpr()->getBeginLoc());
+ mlir::Location exprLoc =
+ cgf.cgm.getLoc(clause.getConditionExpr()->getBeginLoc());
mlir::IntegerType targetType = mlir::IntegerType::get(
&cgf.getMLIRContext(), /*width=*/1,
mlir::IntegerType::SignednessSemantics::Signless);
@@ -271,46 +270,46 @@ mlir::LogicalResult
CIRGenFunction::emitOpenACCShutdownConstruct(
mlir::LogicalResult
CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Loop Construct");
return mlir::failure();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
const OpenACCCombinedConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Combined Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Combined Construct");
return mlir::failure();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
const OpenACCEnterDataConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC EnterData
Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
return mlir::failure();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
const OpenACCExitDataConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
return mlir::failure();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
const OpenACCHostDataConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
return mlir::failure();
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Wait Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Wait Construct");
return mlir::failure();
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Update Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
return mlir::failure();
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
return mlir::failure();
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) {
- getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Cache Construct");
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Cache Construct");
return mlir::failure();
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
