Skip to content

[OpenACC][CIR] Implement 'self' lowering on compute constructs #135851

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Apr 15, 2025
Merged

[OpenACC][CIR] Implement 'self' lowering on compute constructs #135851

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
3ad1313
[OpenACC][CIR] Implement 'self' lowering on compute constructs
erichkeane Apr 15, 2025
3a7a457
Fix spelling of
erichkeane Apr 15, 2025
b890b7b
Nits from Andy
erichkeane Apr 15, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions clang/include/clang/AST/OpenACCClause.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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());
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These convenience functions made a lot of sense/got good use in codegen, so they seemed sensible to add.

}

bool hasConditionExpr() const {
assert(HasConditionExpr.has_value() &&
Expand Down
216 changes: 116 additions & 100 deletions clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,46 +32,51 @@ constexpr bool isOneOfTypes =
template <typename ToTest, typename T>
constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;

template <typename OpTy>
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We end up needing to template-ize this, since the visitor handler functions are not needing to if-constexpr instead of just the apply function.

class OpenACCClauseCIREmitter final
: public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
CIRGenModule &cgm;
: public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
OpTy &operation;
CIRGenFunction &cgf;
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Function instead of Module because we need to be able to emit expressions, which we do at the function.

CIRGenBuilderTy &builder;

// This is necessary since a few of the clauses emit differently based on the
// directive kind they are attached to.
OpenACCDirectiveKind dirKind;
// TODO(cir): 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.cgm.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);
}
}

Expand All @@ -89,64 +94,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.cgm.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,
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We need this because we have a few host-compilers that don't support CTAD fully yet.

OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc);
}

} // namespace

template <typename Op, typename TermOp>
Expand All @@ -158,24 +169,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);
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

See here for the magic around getting the 'insertion' point right. We now only view the clauses information 1x in the emitter (rather than collect data, then emit data for attributes), but all of the ops they 'insert' need to happen before the OpenACC operation.

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;
}

Expand All @@ -187,14 +201,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;
}

Expand Down Expand Up @@ -254,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();
}
30 changes: 28 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/kernels.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
{}

Expand Down Expand Up @@ -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
}
Loading
Loading