Skip to content

Commit 9a56e55

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] deviceptr clause lowering for local 'declare' (#169085)
This is very similar to the 'link' that was done in the last patch, except this works on all storage, but only on pointers. This also shows a bit more of how the enter/exit pairs work in the test. Implementation itself is very simple, as it is just properly handling it in the clause handler.
1 parent ad9bc6a commit 9a56e55

File tree

2 files changed

+111
-4
lines changed

2 files changed

+111
-4
lines changed

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

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -933,7 +933,8 @@ class OpenACCClauseCIREmitter final
933933

934934
void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
935935
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
936-
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
936+
mlir::acc::KernelsOp, mlir::acc::DataOp,
937+
mlir::acc::DeclareEnterOp>) {
937938
for (const Expr *var : clause.getVarList())
938939
addDataOperand<mlir::acc::DevicePtrOp>(
939940
var, mlir::acc::DataClause::acc_deviceptr, {},
@@ -942,9 +943,7 @@ class OpenACCClauseCIREmitter final
942943
} else if constexpr (isCombinedType<OpTy>) {
943944
applyToComputeOp(clause);
944945
} else {
945-
// TODO: When we've implemented this for everything, switch this to an
946-
// unreachable. declare remains.
947-
return clauseNotImplemented(clause);
946+
llvm_unreachable("Unknown construct kind in VisitDevicePtrClause");
948947
}
949948
}
950949

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

Lines changed: 108 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,108 @@
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(HasSideEffects *ArgHSE, int *ArgInt) {
18+
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
19+
// CHECK-NEXT: cir.alloca{{.*}}["this"
20+
// CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE
21+
// CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt
22+
// CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE
23+
// CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt
24+
// CHECK-NEXT: cir.store
25+
// CHECK-NEXT: cir.store
26+
// CHECK-NEXT: cir.store
27+
// CHECK-NEXT: cir.load
28+
HasSideEffects *LocalHSE;
29+
int *LocalInt;
30+
#pragma acc declare deviceptr(ArgHSE, ArgInt, LocalHSE, LocalInt)
31+
// CHECK-NEXT: %[[DEV_PTR_ARG_HSE:.*]] = acc.deviceptr varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSE"}
32+
// CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
33+
// CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
34+
// CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
35+
// CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]], %[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
36+
37+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
38+
}
39+
void MemFunc2(HasSideEffects *ArgHSE, int *ArgInt);
40+
};
41+
42+
void use() {
43+
Struct s;
44+
s.MemFunc1(nullptr, nullptr);
45+
}
46+
47+
void Struct::MemFunc2(HasSideEffects *ArgHSE, int *ArgInt) {
48+
// CHECK: cir.func {{.*}}MemFunc2{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
49+
// CHECK-NEXT: cir.alloca{{.*}}["this"
50+
// CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE
51+
// CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt
52+
// CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE
53+
// CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt
54+
// CHECK-NEXT: cir.store
55+
// CHECK-NEXT: cir.store
56+
// CHECK-NEXT: cir.store
57+
// CHECK-NEXT: cir.load
58+
HasSideEffects *LocalHSE;
59+
int *LocalInt;
60+
#pragma acc declare deviceptr(ArgHSE, ArgInt)
61+
// CHECK-NEXT: %[[DEV_PTR_ARG_HSE:.*]] = acc.deviceptr varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSE"}
62+
// CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
63+
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
64+
65+
#pragma acc declare deviceptr(LocalHSE, LocalInt)
66+
// CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
67+
// CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
68+
// CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
69+
//
70+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])
71+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])
72+
}
73+
74+
extern "C" void do_thing();
75+
76+
void NormalFunc(HasSideEffects *ArgHSE, int *ArgInt) {
77+
// CHECK: cir.func {{.*}}NormalFunc{{.*}}(%[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
78+
// CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE
79+
// CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt
80+
// CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE
81+
// CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt
82+
// CHECK-NEXT: cir.store
83+
// CHECK-NEXT: cir.store
84+
HasSideEffects *LocalHSE;
85+
int *LocalInt;
86+
#pragma acc declare deviceptr(ArgHSE, ArgInt)
87+
// CHECK-NEXT: %[[DEV_PTR_ARG_HSE:.*]] = acc.deviceptr varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "ArgHSE"}
88+
// CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
89+
// CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_ARG_HSE]], %[[DEV_PTR_ARG_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
90+
{
91+
// CHECK-NEXT: cir.scope {
92+
#pragma acc declare deviceptr(LocalHSE, LocalInt)
93+
// CHECK-NEXT: %[[DEV_PTR_LOC_HSE:.*]] = acc.deviceptr varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>) -> !cir.ptr<!cir.ptr<!rec_HasSideEffects>> {name = "LocalHSE"}
94+
// CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
95+
// CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[DEV_PTR_LOC_HSE]], %[[DEV_PTR_LOC_INT]] : !cir.ptr<!cir.ptr<!rec_HasSideEffects>>, !cir.ptr<!cir.ptr<!s32i>>)
96+
do_thing();
97+
// CHECK-NEXT: cir.call @do_thing
98+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])
99+
100+
}
101+
// CHECK-NEXT: }
102+
103+
// Make sure that cleanup gets put in the right scope.
104+
do_thing();
105+
// CHECK-NEXT: cir.call @do_thing
106+
// CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])
107+
}
108+

0 commit comments

Comments
 (0)