[OpenACC][CIR] deviceptr clause lowering for local 'declare' #169085
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 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.
erichkeane
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesThis 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. Full diff: https://github.com/llvm/llvm-project/pull/169085.diff 2 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index c7e6a256c3868..c5c6bcd0153a4 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -933,7 +933,8 @@ class OpenACCClauseCIREmitter final
void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
- mlir::acc::KernelsOp, mlir::acc::DataOp>) {
+ mlir::acc::KernelsOp, mlir::acc::DataOp,
+ mlir::acc::DeclareEnterOp>) {
for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::DevicePtrOp>(
var, mlir::acc::DataClause::acc_deviceptr, {},
@@ -942,9 +943,7 @@ class OpenACCClauseCIREmitter final
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
- // TODO: When we've implemented this for everything, switch this to an
- // unreachable. declare remains.
- return clauseNotImplemented(clause);
+ llvm_unreachable("Unknown construct kind in VisitDevicePtrClause");
}
}
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp b/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp
new file mode 100644
index 0000000000000..d8021ef9a9dc5
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/declare-deviceptr.cpp
@@ -0,0 +1,108 @@
+// 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(HasSideEffects *ArgHSE, int *ArgInt) {
+ // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE
+ // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt
+ // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE
+ // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ HasSideEffects *LocalHSE;
+ int *LocalInt;
+#pragma acc declare deviceptr(ArgHSE, ArgInt, LocalHSE, LocalInt)
+ // 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"}
+ // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
+ // 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"}
+ // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
+ // 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>>)
+
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+ void MemFunc2(HasSideEffects *ArgHSE, int *ArgInt);
+};
+
+void use() {
+ Struct s;
+ s.MemFunc1(nullptr, nullptr);
+}
+
+void Struct::MemFunc2(HasSideEffects *ArgHSE, int *ArgInt) {
+ // CHECK: cir.func {{.*}}MemFunc2{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE
+ // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt
+ // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE
+ // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ HasSideEffects *LocalHSE;
+ int *LocalInt;
+#pragma acc declare deviceptr(ArgHSE, ArgInt)
+ // 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"}
+ // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
+ // 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>>)
+
+#pragma acc declare deviceptr(LocalHSE, LocalInt)
+ // 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"}
+ // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
+ // 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>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])
+}
+
+extern "C" void do_thing();
+
+void NormalFunc(HasSideEffects *ArgHSE, int *ArgInt) {
+ // CHECK: cir.func {{.*}}NormalFunc{{.*}}(%[[ARG_HSE:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}}, %[[ARG_INT:.*]]: !cir.ptr<!s32i> {{.*}})
+ // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["ArgHSE
+ // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["ArgInt
+ // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !cir.ptr<!rec_HasSideEffects>{{.*}}["LocalHSE
+ // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !cir.ptr<!s32i>{{.*}}["LocalInt
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.store
+ HasSideEffects *LocalHSE;
+ int *LocalInt;
+#pragma acc declare deviceptr(ArgHSE, ArgInt)
+ // 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"}
+ // CHECK-NEXT: %[[DEV_PTR_ARG_INT:.*]] = acc.deviceptr varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ArgInt"}
+ // 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>>)
+ {
+ // CHECK-NEXT: cir.scope {
+#pragma acc declare deviceptr(LocalHSE, LocalInt)
+ // 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"}
+ // CHECK-NEXT: %[[DEV_PTR_LOC_INT:.*]] = acc.deviceptr varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "LocalInt"}
+ // 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>>)
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]])
+
+ }
+ // CHECK-NEXT: }
+
+ // Make sure that cleanup gets put in the right scope.
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]])
+}
+
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 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.
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.