Skip to content

Commit b73767d

Browse files
erichkeaneaadeshps-mcw
erichkeane
authored and
aadeshps-mcw
committed
[OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) (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.
1 parent 48c3902 commit b73767d

File tree

4 files changed

+175
-3
lines changed

4 files changed

+175
-3
lines changed

‎clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp‎

Lines changed: 29 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,13 +11,41 @@
1111
//===----------------------------------------------------------------------===//
1212

1313
#include "CIRGenFunction.h"
14+
#include "mlir/Dialect/OpenACC/OpenACC.h"
1415
#include "clang/AST/DeclOpenACC.h"
1516

1617
using namespace clang;
1718
using namespace clang::CIRGen;
1819

20+
namespace {
21+
struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
22+
mlir::acc::DeclareEnterOp enterOp;
23+
24+
OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(enterOp) {}
25+
26+
void emit(CIRGenFunction &cgf) override {
27+
mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(),
28+
enterOp, {});
29+
30+
// TODO(OpenACC): Some clauses require that we add info about them to the
31+
// DeclareExitOp. However, we don't have any of those implemented yet, so
32+
// we should add infrastructure here to do that once we have one
33+
// implemented.
34+
}
35+
};
36+
} // namespace
37+
1938
void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
20-
getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Declare Construct");
39+
mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc());
40+
auto enterOp = mlir::acc::DeclareEnterOp::create(
41+
builder, exprLoc, mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()),
42+
{});
43+
44+
emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(),
45+
d.clauses());
46+
47+
ehStack.pushCleanup<OpenACCDeclareCleanup>(CleanupKind::NormalCleanup,
48+
enterOp);
2149
}
2250

2351
void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {

‎clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -876,6 +876,18 @@ class OpenACCClauseCIREmitter final
876876
}
877877
}
878878

879+
void VisitLinkClause(const OpenACCLinkClause &clause) {
880+
if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
881+
for (const Expr *var : clause.getVarList())
882+
addDataOperand<mlir::acc::DeclareLinkOp>(
883+
var, mlir::acc::DataClause::acc_declare_link, {},
884+
/*structured=*/true,
885+
/*implicit=*/false);
886+
} else {
887+
llvm_unreachable("Unknown construct kind in VisitLinkClause");
888+
}
889+
}
890+
879891
void VisitDeleteClause(const OpenACCDeleteClause &clause) {
880892
if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
881893
for (const Expr *var : clause.getVarList())
@@ -1151,6 +1163,7 @@ EXPL_SPEC(mlir::acc::AtomicReadOp)
11511163
EXPL_SPEC(mlir::acc::AtomicWriteOp)
11521164
EXPL_SPEC(mlir::acc::AtomicCaptureOp)
11531165
EXPL_SPEC(mlir::acc::AtomicUpdateOp)
1166+
EXPL_SPEC(mlir::acc::DeclareEnterOp)
11541167
#undef EXPL_SPEC
11551168

11561169
template <typename ComputeOp, typename LoopOp>

‎clang/test/CIR/CodeGenOpenACC/declare-link.cpp‎

Lines changed: 130 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,130 @@
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
2+
3+
struct HasSideEffects {
4+
HasSideEffects();
5+
~HasSideEffects();
6+
};
7+
8+
// TODO: OpenACC: Implement 'global', NS lowering.
9+
10+
struct Struct {
11+
static const HasSideEffects StaticMemHSE;
12+
static const HasSideEffects StaticMemHSEArr[5];
13+
static const int StaticMemInt;
14+
15+
// TODO: OpenACC: Implement static-local lowering.
16+
17+
void MemFunc1() {
18+
// CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) {
19+
// CHECK-NEXT: cir.alloca{{.*}}["this"
20+
// CHECK-NEXT: cir.store
21+
// CHECK-NEXT: cir.load
22+
extern HasSideEffects LocalHSE;
23+
extern HasSideEffects LocalHSEArr[5];
24+
extern int LocalInt;
25+
#pragma acc declare link(LocalHSE, LocalInt, LocalHSEArr[1:1])
26+
27+
// CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE : !cir.ptr<!rec_HasSideEffects>
28+
// CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"}
29+
//
30+
// CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt : !cir.ptr<!s32i>
31+
// CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
32+
//
33+
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
34+
// CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
35+
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
36+
// CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
37+
// CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
38+
// CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
39+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
40+
// CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
41+
// 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]"}
42+
//
43+
// 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>>)
44+
//
45+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
46+
}
47+
48+
void MemFunc2();
49+
};
50+
void use() {
51+
Struct s;
52+
s.MemFunc1();
53+
}
54+
55+
void Struct::MemFunc2() {
56+
// CHECK: cir.func {{.*}}MemFunc2{{.*}}({{.*}}) {
57+
// CHECK-NEXT: cir.alloca{{.*}}["this"
58+
// CHECK-NEXT: cir.store
59+
// CHECK-NEXT: cir.load
60+
extern HasSideEffects LocalHSE2;
61+
extern HasSideEffects LocalHSEArr2[5];
62+
extern int LocalInt2;
63+
64+
#pragma acc declare link(LocalHSE2, LocalInt2, LocalHSEArr2[1:1])
65+
// CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE2 : !cir.ptr<!rec_HasSideEffects>
66+
// CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE2"}
67+
//
68+
// CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt2 : !cir.ptr<!s32i>
69+
// CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt2"}
70+
//
71+
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
72+
// CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
73+
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
74+
// CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
75+
// CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
76+
// CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
77+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
78+
// CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr2 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
79+
// 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]"}
80+
//
81+
// 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>>)
82+
//
83+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
84+
}
85+
86+
extern "C" void do_thing();
87+
88+
void NormalFunc() {
89+
// CHECK: cir.func {{.*}}NormalFunc{{.*}}()
90+
extern HasSideEffects LocalHSE3;
91+
extern HasSideEffects LocalHSEArr3[5];
92+
extern int LocalInt3;
93+
// CHECK-NEXT: cir.scope
94+
{
95+
extern HasSideEffects InnerHSE;
96+
#pragma acc declare link(LocalHSE3, LocalInt3, LocalHSEArr3[1:1], InnerHSE)
97+
// CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE3 : !cir.ptr<!rec_HasSideEffects>
98+
// CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE3"}
99+
//
100+
// CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt3 : !cir.ptr<!s32i>
101+
// CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt3"}
102+
//
103+
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
104+
// CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
105+
// CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
106+
// CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
107+
// CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
108+
// CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
109+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
110+
// CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr3 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
111+
// 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]"}
112+
//
113+
// CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @InnerHSE : !cir.ptr<!rec_HasSideEffects>
114+
// CHECK-NEXT: %[[INNERHSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "InnerHSE"}
115+
//
116+
// 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>)
117+
//
118+
// CHECK
119+
120+
do_thing();
121+
// CHECK-NEXT: cir.call @do_thing
122+
123+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
124+
}
125+
// CHECK-NEXT: }
126+
127+
do_thing();
128+
// CHECK-NEXT: cir.call @do_thing
129+
}
130+

‎clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,8 @@
11
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
22

33
void HelloWorld(int *A) {
4+
extern int *E;
45

5-
// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
6-
#pragma acc declare create(A)
6+
// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: create}}
7+
#pragma acc declare link(E) create(A)
78
}

0 commit comments

Comments
 (0)