[OpenACC][CIR] Implement 'data' construct lowering #135038
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThis patch does two things primarily: 1- It does the lowering of the OpenACC 'data' construct, which requires getting the 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 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:
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
|
|
@llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesThis patch does two things primarily: 1- It does the lowering of the OpenACC 'data' construct, which requires getting the 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 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:
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
|
This translation was done way back when we started with the OpenACC dialect and it is likely not what will be supported in the long term for OpenACC codegen. The idea behind OpenACC code generation was to be done in MLIR until the LLVM IR dialect and not directly emit LLVM IR as it is done for OpenMP. SO I'm not sure it is needed to test it with your current lowering. |
erichkeane
changed the title
That is good feedback, thanks! I've modified the patch to remove the lowering, plus modified the commit title and message. |
clementval
approved these changes
Apr 9, 2025
Good. I should probably clean that up someday so it doesn't lead to confusion. |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/27/builds/8369 Here is the relevant piece of the build log for the reference |
AllinLeeYL
pushed a commit
to AllinLeeYL/llvm-project
that referenced
this pull request
Apr 10, 2025
This patch 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. In addition to that, 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.
var-const
pushed a commit
to ldionne/llvm-project
that referenced
this pull request
Apr 17, 2025
This patch 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. In addition to that, 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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This patch does the lowering of the OpenACC 'data' construct, which requires getting the
defaultclause (asdatarequires 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.
In addition to that, 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.