Skip to content

[OpenACC][CIR] Implement 'data' construct lowering #135038

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 9, 2025
Merged

[OpenACC][CIR] Implement 'data' construct lowering #135038

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
d175c7c
[OpenACC][CIR] Implement 'data' construct lowering, lower OACC->LLVMIR
erichkeane Apr 9, 2025
4f7d751
Clang-format
erichkeane Apr 9, 2025
97209e4
Remove lowering to LLVM-IR and tests
erichkeane Apr 9, 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
8 changes: 5 additions & 3 deletions clang/include/clang/AST/OpenACCClause.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1316,11 +1316,13 @@ template <class Impl> class OpenACCClauseVisitor {
switch (C->getClauseKind()) {
#define VISIT_CLAUSE(CLAUSE_NAME) \
case OpenACCClauseKind::CLAUSE_NAME: \
Visit##CLAUSE_NAME##Clause(*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
getDerived().Visit##CLAUSE_NAME##Clause( \
*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
return;
#define CLAUSE_ALIAS(ALIAS_NAME, CLAUSE_NAME, DEPRECATED) \
case OpenACCClauseKind::ALIAS_NAME: \
Visit##CLAUSE_NAME##Clause(*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
getDerived().Visit##CLAUSE_NAME##Clause( \
*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
return;
#include "clang/Basic/OpenACCClauses.def"

Expand All @@ -1333,7 +1335,7 @@ template <class Impl> class OpenACCClauseVisitor {
#define VISIT_CLAUSE(CLAUSE_NAME) \
void Visit##CLAUSE_NAME##Clause( \
const OpenACC##CLAUSE_NAME##Clause &Clause) { \
return getDerived().Visit##CLAUSE_NAME##Clause(Clause); \
return getDerived().VisitClause(Clause); \
}

#include "clang/Basic/OpenACCClauses.def"
Expand Down
11 changes: 5 additions & 6 deletions clang/lib/CIR/CodeGen/CIRGenFunction.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -571,14 +571,13 @@ class CIRGenFunction : public CIRGenTypeCache {
// OpenACC Emission
//===--------------------------------------------------------------------===//
private:
// Function to do the basic implementation of a 'compute' operation, including
// the clauses/etc. This might be generalizable in the future to work for
// other constructs, or at least be the base for construct emission.
// Function to do the basic implementation of an operation with an Associated
// Statement. Models AssociatedStmtConstruct.
template <typename Op, typename TermOp>
mlir::LogicalResult
emitOpenACCComputeOp(mlir::Location start, mlir::Location end,
llvm::ArrayRef<const OpenACCClause *> clauses,
const Stmt *structuredBlock);
emitOpenACCOpAssociatedStmt(mlir::Location start, mlir::Location end,
llvm::ArrayRef<const OpenACCClause *> clauses,
const Stmt *associatedStmt);

public:
mlir::LogicalResult
Expand Down
67 changes: 49 additions & 18 deletions clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,41 +27,68 @@ class OpenACCClauseCIREmitter final
: public OpenACCClauseVisitor<OpenACCClauseCIREmitter> {
CIRGenModule &cgm;

struct AttributeData {
// Value of the 'default' attribute, added on 'data' and 'compute'/etc
// constructs as a 'default-attr'.
std::optional<ClauseDefaultValue> defaultVal = std::nullopt;
} attrData;

void clauseNotImplemented(const OpenACCClause &c) {
cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
}

public:
OpenACCClauseCIREmitter(CIRGenModule &cgm) : cgm(cgm) {}

#define VISIT_CLAUSE(CN) \
void Visit##CN##Clause(const OpenACC##CN##Clause &clause) { \
clauseNotImplemented(clause); \
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;
}
}

// Apply any of the clauses that resulted in an 'attribute'.
template <typename Op> void applyAttributes(Op &op) {
if (attrData.defaultVal.has_value())
op.setDefaultAttr(*attrData.defaultVal);
}
#include "clang/Basic/OpenACCClauses.def"
};
} // namespace

template <typename Op, typename TermOp>
mlir::LogicalResult CIRGenFunction::emitOpenACCComputeOp(
mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
mlir::Location start, mlir::Location end,
llvm::ArrayRef<const OpenACCClause *> clauses,
const Stmt *structuredBlock) {
llvm::ArrayRef<const OpenACCClause *> clauses, const Stmt *associatedStmt) {
mlir::LogicalResult res = mlir::success();

llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;

// Clause-emitter must be here because it might modify operands.
OpenACCClauseCIREmitter clauseEmitter(getCIRGenModule());
clauseEmitter.VisitClauseList(clauses);

llvm::SmallVector<mlir::Type> retTy;
llvm::SmallVector<mlir::Value> operands;
auto op = builder.create<Op>(start, retTy, operands);

// Apply the attributes derived from the clauses.
clauseEmitter.applyAttributes(op);

mlir::Block &block = op.getRegion().emplaceBlock();
mlir::OpBuilder::InsertionGuard guardCase(builder);
builder.setInsertionPointToEnd(&block);

LexicalScope ls{*this, start, builder.getInsertionBlock()};
res = emitStmt(structuredBlock, /*useCurrentScope=*/true);
res = emitStmt(associatedStmt, /*useCurrentScope=*/true);

builder.create<TermOp>(end);
return res;
Expand All @@ -74,19 +101,28 @@ CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) {

switch (s.getDirectiveKind()) {
case OpenACCDirectiveKind::Parallel:
return emitOpenACCComputeOp<ParallelOp, mlir::acc::YieldOp>(
return emitOpenACCOpAssociatedStmt<ParallelOp, mlir::acc::YieldOp>(
start, end, s.clauses(), s.getStructuredBlock());
case OpenACCDirectiveKind::Serial:
return emitOpenACCComputeOp<SerialOp, mlir::acc::YieldOp>(
return emitOpenACCOpAssociatedStmt<SerialOp, mlir::acc::YieldOp>(
start, end, s.clauses(), s.getStructuredBlock());
case OpenACCDirectiveKind::Kernels:
return emitOpenACCComputeOp<KernelsOp, mlir::acc::TerminatorOp>(
return emitOpenACCOpAssociatedStmt<KernelsOp, mlir::acc::TerminatorOp>(
start, end, s.clauses(), s.getStructuredBlock());
default:
llvm_unreachable("invalid compute construct kind");
}
}

mlir::LogicalResult
CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
mlir::Location start = getLoc(s.getSourceRange().getEnd());
mlir::Location end = getLoc(s.getSourceRange().getEnd());

return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
start, end, s.clauses(), s.getStructuredBlock());
}

mlir::LogicalResult
CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct");
Expand All @@ -97,11 +133,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Combined Construct");
return mlir::failure();
}
mlir::LogicalResult
CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Data Construct");
return mlir::failure();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
const OpenACCEnterDataConstruct &s) {
getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
Expand Down
37 changes: 37 additions & 0 deletions clang/test/CIR/CodeGenOpenACC/data.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s

void acc_data(void) {
// CHECK: cir.func @acc_data() {

#pragma acc data default(none)
{
int i = 0;
++i;
}
// CHECK-NEXT: acc.data {
// CHECK-NEXT: cir.alloca
// CHECK-NEXT: cir.const
// CHECK-NEXT: cir.store
// CHECK-NEXT: cir.load
// CHECK-NEXT: cir.unary
// CHECK-NEXT: cir.store
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}

#pragma acc data default(present)
{
int i = 0;
++i;
}
// CHECK-NEXT: acc.data {
// CHECK-NEXT: cir.alloca
// CHECK-NEXT: cir.const
// CHECK-NEXT: cir.store
// CHECK-NEXT: cir.load
// CHECK-NEXT: cir.unary
// CHECK-NEXT: cir.store
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}

// CHECK-NEXT: cir.return
}
16 changes: 14 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/kernels.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,9 +6,21 @@ void acc_kernels(void) {
{}

// CHECK-NEXT: acc.kernels {
// CHECK-NEXT:acc.terminator
// CHECK-NEXT: acc.terminator
// CHECK-NEXT:}

#pragma acc kernels default(none)
{}
// CHECK-NEXT: acc.kernels {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}

#pragma acc kernels default(present)
{}
// CHECK-NEXT: acc.kernels {
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}

#pragma acc kernels
while(1){}
// CHECK-NEXT: acc.kernels {
Expand All @@ -23,7 +35,7 @@ void acc_kernels(void) {
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
// CHECK-NEXT:acc.terminator
// CHECK-NEXT: acc.terminator
// CHECK-NEXT:}

// CHECK-NEXT: cir.return
Expand Down
16 changes: 14 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/parallel.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,9 +5,21 @@ void acc_parallel(void) {
#pragma acc parallel
{}
// CHECK-NEXT: acc.parallel {
// CHECK-NEXT:acc.yield
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}

#pragma acc parallel default(none)
{}
// CHECK-NEXT: acc.parallel {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}

#pragma acc parallel default(present)
{}
// CHECK-NEXT: acc.parallel {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}

#pragma acc parallel
while(1){}
// CHECK-NEXT: acc.parallel {
Expand All @@ -22,7 +34,7 @@ void acc_parallel(void) {
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
// CHECK-NEXT:acc.yield
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}

// CHECK-NEXT: cir.return
Expand Down
16 changes: 14 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/serial.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,9 +6,21 @@ void acc_serial(void) {
{}

// CHECK-NEXT: acc.serial {
// CHECK-NEXT:acc.yield
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}

#pragma acc serial default(none)
{}
// CHECK-NEXT: acc.serial {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}

#pragma acc serial default(present)
{}
// CHECK-NEXT: acc.serial {
// CHECK-NEXT: acc.yield
// CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}

#pragma acc serial
while(1){}
// CHECK-NEXT: acc.serial {
Expand All @@ -23,7 +35,7 @@ void acc_serial(void) {
// CHECK-NEXT: }
// cir.scope end:
// CHECK-NEXT: }
// CHECK-NEXT:acc.yield
// CHECK-NEXT: acc.yield
// CHECK-NEXT:}

// CHECK-NEXT: cir.return
Expand Down
Loading