[clang] [OpenACC][CIR] Implement 'data' construct lowering, lower OACC->LLVMIR (PR #135038)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 9 08:12:15 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Erich Keane (erichkeane)
<details>
<summary>Changes</summary>
This patch does two things primarily:
1- It does the lowering of the OpenACC 'data' construct, which requires getting the `default` clause (as `data` requires at least 1 of a list of clauses, and this is the easiest one). The lowering of the clauses appears to happen in 1 of 2 ways: a- as an operand. or b- as an attribute.
This patch adds infrastructure to lower as an attribute, as that is how 'data' works.
2- This patch adds the infrastructure/calls to do the OpenACCDialect->LLVM-IR lowering. Unfortunately only a handful of constructs are actually functional in the OpenACC dialect, of which `data` is one (hence the choice to do it here, and why I chose to do it as one patch). SO, like the Flang OpenACC implementation, attempts to lower below CIR/OpenACC Dialect will likely fail.
In addition to those, it changes the OpenACCClauseVisitor a bit, which previously just required that each of the derived classes have all of the clauses covered. This patch modifies it so that the visitor directly calls the derived class from its visitor function, which leaves the base-class ones the ability to defer to a generic function. This was previously like this because I had some use cases that I didn't end up using, and the 'generic' function here seems much more useful.
---
Full diff: https://github.com/llvm/llvm-project/pull/135038.diff
9 Files Affected:
- (modified) clang/include/clang/AST/OpenACCClause.h (+5-3)
- (modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+5-6)
- (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+49-18)
- (modified) clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt (+1)
- (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+2)
- (added) clang/test/CIR/CodeGenOpenACC/data.c (+64)
- (modified) clang/test/CIR/CodeGenOpenACC/kernels.c (+14-2)
- (modified) clang/test/CIR/CodeGenOpenACC/parallel.c (+14-2)
- (modified) clang/test/CIR/CodeGenOpenACC/serial.c (+14-2)
``````````diff
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index fda1837594c99..3687af76a559f 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -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"
@@ -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"
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index fb5ec6a868a1b..c30fcc2a05f87 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -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
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 7a8879add784a..e7e56d3602e3a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -27,6 +27,12 @@ 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());
}
@@ -34,34 +40,55 @@ class OpenACCClauseCIREmitter final
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;
@@ -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");
@@ -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");
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
index 7baff3412a84e..634b4042c9cb3 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
@@ -20,5 +20,6 @@ add_clang_library(clangCIRLoweringDirectToLLVM
MLIRCIR
MLIRBuiltinToLLVMIRTranslation
MLIRLLVMToLLVMIRTranslation
+ MLIROpenACCToLLVMIRTranslation
MLIRIR
)
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 7ca36409c9cac..14cb63e7c58a4 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -26,6 +26,7 @@
#include "mlir/Pass/PassManager.h"
#include "mlir/Target/LLVMIR/Dialect/Builtin/BuiltinToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/OpenACC/OpenACCToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
#include "mlir/Transforms/DialectConversion.h"
#include "clang/CIR/Dialect/IR/CIRDialect.h"
@@ -1492,6 +1493,7 @@ lowerDirectlyFromCIRToLLVMIR(mlir::ModuleOp mlirModule, LLVMContext &llvmCtx) {
mlir::registerBuiltinDialectTranslation(*mlirCtx);
mlir::registerLLVMDialectTranslation(*mlirCtx);
mlir::registerCIRDialectTranslation(*mlirCtx);
+ mlir::registerOpenACCDialectTranslation(*mlirCtx);
llvm::TimeTraceScope translateScope("translateModuleToLLVMIR");
diff --git a/clang/test/CIR/CodeGenOpenACC/data.c b/clang/test/CIR/CodeGenOpenACC/data.c
new file mode 100644
index 0000000000000..025b7747539f3
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/data.c
@@ -0,0 +1,64 @@
+// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s --check-prefix=CIR
+// RUN: %clang_cc1 -fopenacc -emit-llvm -fclangir %s -o - | FileCheck %s -check-prefix=LLVM
+
+void acc_data(void) {
+ // CIR: cir.func @acc_data() {
+ // LLVM: define void @acc_data() {
+
+#pragma acc data default(none)
+ {
+ int i = 0;
+ ++i;
+ }
+ // CIR-NEXT: acc.data {
+ // CIR-NEXT: cir.alloca
+ // CIR-NEXT: cir.const
+ // CIR-NEXT: cir.store
+ // CIR-NEXT: cir.load
+ // CIR-NEXT: cir.unary
+ // CIR-NEXT: cir.store
+ // CIR-NEXT: acc.terminator
+ // CIR-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}
+ //
+ // LLVM: call void @__tgt_target_data_begin_mapper
+ // LLVM-NEXT: br label %[[ACC_DATA:.+]]
+ // LLVM: [[ACC_DATA]]:
+ // LLVM-NEXT: store i32 0
+ // LLVM-NEXT: load i32
+ // LLVM-NEXT: add nsw i32 %{{.*}}, 1
+ // LLVM-NEXT: store i32
+ // LLVM-NEXT: br label %[[ACC_DATA_END:.+]]
+ //
+ // LLVM: [[ACC_DATA_END]]:
+ // LLVM: call void @__tgt_target_data_end_mapper
+
+#pragma acc data default(present)
+ {
+ int i = 0;
+ ++i;
+ }
+ // CIR-NEXT: acc.data {
+ // CIR-NEXT: cir.alloca
+ // CIR-NEXT: cir.const
+ // CIR-NEXT: cir.store
+ // CIR-NEXT: cir.load
+ // CIR-NEXT: cir.unary
+ // CIR-NEXT: cir.store
+ // CIR-NEXT: acc.terminator
+ // CIR-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}
+
+ // LLVM: call void @__tgt_target_data_begin_mapper
+ // LLVM-NEXT: br label %[[ACC_DATA:.+]]
+ // LLVM: [[ACC_DATA]]:
+ // LLVM-NEXT: store i32 0
+ // LLVM-NEXT: load i32
+ // LLVM-NEXT: add nsw i32 %{{.*}}, 1
+ // LLVM-NEXT: store i32
+ // LLVM-NEXT: br label %[[ACC_DATA_END:.+]]
+ //
+ // LLVM: [[ACC_DATA_END]]:
+ // LLVM: call void @__tgt_target_data_end_mapper
+
+ // CIR-NEXT: cir.return
+ // LLVM: ret void
+}
diff --git a/clang/test/CIR/CodeGenOpenACC/kernels.c b/clang/test/CIR/CodeGenOpenACC/kernels.c
index 91684859f7115..0c950fe3d0f9c 100644
--- a/clang/test/CIR/CodeGenOpenACC/kernels.c
+++ b/clang/test/CIR/CodeGenOpenACC/kernels.c
@@ -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 {
@@ -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
diff --git a/clang/test/CIR/CodeGenOpenACC/parallel.c b/clang/test/CIR/CodeGenOpenACC/parallel.c
index 7c1509a129980..e18270435460c 100644
--- a/clang/test/CIR/CodeGenOpenACC/parallel.c
+++ b/clang/test/CIR/CodeGenOpenACC/parallel.c
@@ -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 {
@@ -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
diff --git a/clang/test/CIR/CodeGenOpenACC/serial.c b/clang/test/CIR/CodeGenOpenACC/serial.c
index 9897cd3d4e8d9..72a0995549da3 100644
--- a/clang/test/CIR/CodeGenOpenACC/serial.c
+++ b/clang/test/CIR/CodeGenOpenACC/serial.c
@@ -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 {
@@ -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
``````````
</details>
https://github.com/llvm/llvm-project/pull/135038
More information about the cfe-commits
mailing list