Skip to content

Commit 92d0d39

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] deviceptr clause lowering for local 'declare'
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 5dbe83c commit 92d0d39

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)