[OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) #168793
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
'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses.
erichkeane
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) Changes'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses. Full diff: https://github.com/llvm/llvm-project/pull/168793.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index b588a50aa0404..f6680cbaa8c78 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -12,12 +12,42 @@
#include "CIRGenFunction.h"
#include "clang/AST/DeclOpenACC.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
using namespace clang;
using namespace clang::CIRGen;
+namespace {
+ struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
+ mlir::acc::DeclareEnterOp enterOp;
+
+ OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp)
+ : enterOp(enterOp) {}
+
+ void emit(CIRGenFunction &cgf) override {
+ mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(),
+ enterOp, {});
+
+ // TODO(OpenACC): Some clauses require that we add info about them to the
+ // DeclareExitOp. However, we don't have any of those implemented yet, so
+ // we should add infrastructure here to do that once we have one
+ // implemented.
+ }
+
+ };
+} // namespace
+
void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
- getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Declare Construct");
+ mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc());
+ auto enterOp = mlir::acc::DeclareEnterOp::create(
+ builder, exprLoc,
+ mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()), {});
+
+ emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(),
+ d.clauses());
+
+ ehStack.pushCleanup<OpenACCDeclareCleanup>(CleanupKind::NormalCleanup,
+ enterOp);
}
void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 527dfd21db8a5..c7e6a256c3868 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -876,6 +876,18 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitLinkClause(const OpenACCLinkClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::DeclareLinkOp>(
+ var, mlir::acc::DataClause::acc_declare_link, {},
+ /*structured=*/true,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitLinkClause");
+ }
+ }
+
void VisitDeleteClause(const OpenACCDeleteClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
for (const Expr *var : clause.getVarList())
@@ -1151,6 +1163,7 @@ EXPL_SPEC(mlir::acc::AtomicReadOp)
EXPL_SPEC(mlir::acc::AtomicWriteOp)
EXPL_SPEC(mlir::acc::AtomicCaptureOp)
EXPL_SPEC(mlir::acc::AtomicUpdateOp)
+EXPL_SPEC(mlir::acc::DeclareEnterOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
new file mode 100644
index 0000000000000..8494a2354c7db
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
@@ -0,0 +1,130 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+struct HasSideEffects {
+ HasSideEffects();
+ ~HasSideEffects();
+};
+
+// TODO: OpenACC: Implement 'global', NS lowering.
+
+struct Struct {
+ static const HasSideEffects StaticMemHSE;
+ static const HasSideEffects StaticMemHSEArr[5];
+ static const int StaticMemInt;
+
+ // TODO: OpenACC: Implement static-local lowering.
+
+ void MemFunc1() {
+ // CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) {
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ extern HasSideEffects LocalHSE;
+ extern HasSideEffects LocalHSEArr[5];
+ extern int LocalInt;
+#pragma acc declare link(LocalHSE, LocalInt, LocalHSEArr[1:1])
+
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+
+ void MemFunc2();
+};
+void use() {
+ Struct s;
+ s.MemFunc1();
+}
+
+void Struct::MemFunc2() {
+ // CHECK: cir.func {{.*}}MemFunc2{{.*}}({{.*}}) {
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ extern HasSideEffects LocalHSE2;
+ extern HasSideEffects LocalHSEArr2[5];
+ extern int LocalInt2;
+
+#pragma acc declare link(LocalHSE2, LocalInt2, LocalHSEArr2[1:1])
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE2 : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE2"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt2 : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt2"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr2 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr2[1:1]"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+}
+
+extern "C" void do_thing();
+
+void NormalFunc() {
+ // CHECK: cir.func {{.*}}NormalFunc{{.*}}()
+ extern HasSideEffects LocalHSE3;
+ extern HasSideEffects LocalHSEArr3[5];
+ extern int LocalInt3;
+ // CHECK-NEXT: cir.scope
+ {
+ extern HasSideEffects InnerHSE;
+#pragma acc declare link(LocalHSE3, LocalInt3, LocalHSEArr3[1:1], InnerHSE)
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE3 : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE3"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt3 : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt3"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr3 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr3[1:1]"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @InnerHSE : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[INNERHSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "InnerHSE"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]], %[[INNERHSE_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>, !cir.ptr<!rec_HasSideEffects>)
+ //
+ // CHECK
+
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+ // CHECK-NEXT: }
+
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+}
+
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
index e85c26718acb8..c8b85a12f84e7 100644
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
@@ -1,7 +1,8 @@
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
void HelloWorld(int *A) {
+ extern int *E;
-// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
-#pragma acc declare create(A)
+// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: create}}
+#pragma acc declare link(E) create(A)
}
|
Member
|
@llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) Changes'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses. Full diff: https://github.com/llvm/llvm-project/pull/168793.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index b588a50aa0404..f6680cbaa8c78 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -12,12 +12,42 @@
#include "CIRGenFunction.h"
#include "clang/AST/DeclOpenACC.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
using namespace clang;
using namespace clang::CIRGen;
+namespace {
+ struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
+ mlir::acc::DeclareEnterOp enterOp;
+
+ OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp)
+ : enterOp(enterOp) {}
+
+ void emit(CIRGenFunction &cgf) override {
+ mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(),
+ enterOp, {});
+
+ // TODO(OpenACC): Some clauses require that we add info about them to the
+ // DeclareExitOp. However, we don't have any of those implemented yet, so
+ // we should add infrastructure here to do that once we have one
+ // implemented.
+ }
+
+ };
+} // namespace
+
void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
- getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Declare Construct");
+ mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc());
+ auto enterOp = mlir::acc::DeclareEnterOp::create(
+ builder, exprLoc,
+ mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()), {});
+
+ emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(),
+ d.clauses());
+
+ ehStack.pushCleanup<OpenACCDeclareCleanup>(CleanupKind::NormalCleanup,
+ enterOp);
}
void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 527dfd21db8a5..c7e6a256c3868 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -876,6 +876,18 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitLinkClause(const OpenACCLinkClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::DeclareLinkOp>(
+ var, mlir::acc::DataClause::acc_declare_link, {},
+ /*structured=*/true,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitLinkClause");
+ }
+ }
+
void VisitDeleteClause(const OpenACCDeleteClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
for (const Expr *var : clause.getVarList())
@@ -1151,6 +1163,7 @@ EXPL_SPEC(mlir::acc::AtomicReadOp)
EXPL_SPEC(mlir::acc::AtomicWriteOp)
EXPL_SPEC(mlir::acc::AtomicCaptureOp)
EXPL_SPEC(mlir::acc::AtomicUpdateOp)
+EXPL_SPEC(mlir::acc::DeclareEnterOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
new file mode 100644
index 0000000000000..8494a2354c7db
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
@@ -0,0 +1,130 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+struct HasSideEffects {
+ HasSideEffects();
+ ~HasSideEffects();
+};
+
+// TODO: OpenACC: Implement 'global', NS lowering.
+
+struct Struct {
+ static const HasSideEffects StaticMemHSE;
+ static const HasSideEffects StaticMemHSEArr[5];
+ static const int StaticMemInt;
+
+ // TODO: OpenACC: Implement static-local lowering.
+
+ void MemFunc1() {
+ // CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) {
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ extern HasSideEffects LocalHSE;
+ extern HasSideEffects LocalHSEArr[5];
+ extern int LocalInt;
+#pragma acc declare link(LocalHSE, LocalInt, LocalHSEArr[1:1])
+
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+
+ void MemFunc2();
+};
+void use() {
+ Struct s;
+ s.MemFunc1();
+}
+
+void Struct::MemFunc2() {
+ // CHECK: cir.func {{.*}}MemFunc2{{.*}}({{.*}}) {
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ extern HasSideEffects LocalHSE2;
+ extern HasSideEffects LocalHSEArr2[5];
+ extern int LocalInt2;
+
+#pragma acc declare link(LocalHSE2, LocalInt2, LocalHSEArr2[1:1])
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE2 : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE2"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt2 : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt2"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr2 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr2[1:1]"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+}
+
+extern "C" void do_thing();
+
+void NormalFunc() {
+ // CHECK: cir.func {{.*}}NormalFunc{{.*}}()
+ extern HasSideEffects LocalHSE3;
+ extern HasSideEffects LocalHSEArr3[5];
+ extern int LocalInt3;
+ // CHECK-NEXT: cir.scope
+ {
+ extern HasSideEffects InnerHSE;
+#pragma acc declare link(LocalHSE3, LocalInt3, LocalHSEArr3[1:1], InnerHSE)
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE3 : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE3"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt3 : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt3"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr3 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr3[1:1]"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @InnerHSE : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[INNERHSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "InnerHSE"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]], %[[INNERHSE_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>, !cir.ptr<!rec_HasSideEffects>)
+ //
+ // CHECK
+
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+ // CHECK-NEXT: }
+
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+}
+
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
index e85c26718acb8..c8b85a12f84e7 100644
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
@@ -1,7 +1,8 @@
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
void HelloWorld(int *A) {
+ extern int *E;
-// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
-#pragma acc declare create(A)
+// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: create}}
+#pragma acc declare link(E) create(A)
}
|
erichkeane
commented
Nov 19, 2025
| // CHECK-NEXT: cir.alloca{{.*}}["this" | ||
| // CHECK-NEXT: cir.store | ||
| // CHECK-NEXT: cir.load | ||
| extern HasSideEffects LocalHSE2; |
Collaborator
Author
There was a problem hiding this comment.
I'll note all the get_global stuff is because link requires extern variables in the same scope. It perhaps wasn't the best idea to start with link, but the rest are more complicated source changes, so this was a bit of a necessary evil. I promise the next ones will be better :)
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
🐧 Linux x64 Test Results
|
andykaylor
approved these changes
Nov 21, 2025
| using namespace clang::CIRGen; | ||
|
|
||
| namespace { | ||
| struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { |
Contributor
There was a problem hiding this comment.
What are the expected interactions between OpenACC and exceptions? Are exceptions even allowed in an OpenACC region?
I realize this is needed for normal cleanup, but it made me think about broader implications.
Collaborator
Author
There was a problem hiding this comment.
I believe the standard is pretty silent on exceptions, but most devices would fail exceptions. This SHOULD interact fine with exceptions though, at least based on my understanding of cleanups. It at least does get in sync with all the other cleanups as far as I can tell.
Collaborator
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/169/builds/17324 Here is the relevant piece of the build log for the reference |
aadeshps-mcw
pushed a commit
to aadeshps-mcw/llvm-project
that referenced
this pull request
Nov 26, 2025
…e) (llvm#168793) 'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses.
Priyanshu3820
pushed a commit
to Priyanshu3820/llvm-project
that referenced
this pull request
Nov 26, 2025
…e) (llvm#168793) 'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses.
augusto2112
pushed a commit
to augusto2112/llvm-project
that referenced
this pull request
Dec 3, 2025
…e) (llvm#168793) 'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'. A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical. 'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
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.
'declare' is a declaration directive, so it can appear at 3 places: Global/NS scope, class scope, or local scope. This patch implements ONLY the 'local' scope lowering for 'declare'.
A 'declare' is lowered as a 'declare_enter' and 'declare_exit' operation, plus data operands like all others. Sema restricts the form of some of these, but they are otherwise identical.
'declare' DOES require at least 1 clause for the examples to make sense, so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY attached to the 'declare_enter', and doesn't require any additional work besides a very small addition to how we handle clauses.