https://github.com/erichkeane updated
https://github.com/llvm/llvm-project/pull/135038
>From d175c7cb2de79731f5b5009bb08cc76f971b3e0a Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Tue, 8 Apr 2025 17:54:00 -0700
Subject: [PATCH 1/2] [OpenACC][CIR] Implement 'data' construct lowering, lower
OACC->LLVMIR
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.
---
clang/include/clang/AST/OpenACCClause.h | 8 ++-
clang/lib/CIR/CodeGen/CIRGenFunction.h | 13 ++--
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 66 ++++++++++++++-----
.../CIR/Lowering/DirectToLLVM/CMakeLists.txt | 1 +
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 2 +
clang/test/CIR/CodeGenOpenACC/data.c | 64 ++++++++++++++++++
clang/test/CIR/CodeGenOpenACC/kernels.c | 16 ++++-
clang/test/CIR/CodeGenOpenACC/parallel.c | 16 ++++-
clang/test/CIR/CodeGenOpenACC/serial.c | 16 ++++-
9 files changed, 169 insertions(+), 33 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenACC/data.c
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..abb91052c779e 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);
+ mlir::LogicalResult 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..66a45d4bf2660 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,56 @@ 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) {
+ 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 +102,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 +134,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
>From 4f7d751f9df717b9424be236686e76667cd22f85 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Wed, 9 Apr 2025 08:11:49 -0700
Subject: [PATCH 2/2] Clang-format
---
clang/lib/CIR/CodeGen/CIRGenFunction.h | 8 ++++----
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 3 +--
2 files changed, 5 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h
b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index abb91052c779e..c30fcc2a05f87 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -574,10 +574,10 @@ class CIRGenFunction : public CIRGenTypeCache {
// Function to do the basic implementation of an operation with an Associated
// Statement. Models AssociatedStmtConstruct.
template <typename Op, typename TermOp>
- mlir::LogicalResult emitOpenACCOpAssociatedStmt(
- mlir::Location start, mlir::Location end,
- llvm::ArrayRef<const OpenACCClause *> clauses,
- const Stmt *associatedStmt);
+ mlir::LogicalResult
+ 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 66a45d4bf2660..e7e56d3602e3a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -68,8 +68,7 @@ class OpenACCClauseCIREmitter final
template <typename Op, typename TermOp>
mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
mlir::Location start, mlir::Location end,
- llvm::ArrayRef<const OpenACCClause *> clauses,
- const Stmt *associatedStmt) {
+ llvm::ArrayRef<const OpenACCClause *> clauses, const Stmt *associatedStmt)
{
mlir::LogicalResult res = mlir::success();
llvm::SmallVector<mlir::Type> retTy;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
