[OpenACC][CIR] Implement 'device_type' lowering for Routine #170893
Merged
+99
−13
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
The 'device_type' clause modifies how the clauses that are legal after it (seq, worker, vector, gang, bind) work. Previous patches were aware of how that was going to happen, thanks to experience with doing the same work on other constructs/clauses, so this is mostly just a repeat of those. Tests for the first 4 and interactions with them are included, but 'bind' is not yet implemented, so its device_type tests will be added when it is lowered.
erichkeane
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThe 'device_type' clause modifies how the clauses that are legal after it (seq, worker, vector, gang, bind) work. Previous patches were aware of how that was going to happen, thanks to experience with doing the same work on other constructs/clauses, so this is mostly just a repeat of those. Tests for the first 4 and interactions with them are included, but 'bind' is not yet implemented, so its device_type tests will be added when it is lowered. Full diff: https://github.com/llvm/llvm-project/pull/170893.diff 4 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 050f110c6e365..56d4631f7845e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -355,6 +355,13 @@ class OpenACCRoutineClauseEmitter final
curValue.getZExtValue());
}
}
+
+ void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
+ lastDeviceTypeValues.clear();
+
+ for (const DeviceTypeArgument &arg : clause.getArchitectures())
+ lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo()));
+ }
};
} // namespace
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 2d4ed23a46d1c..8e7384ae66d8e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -112,19 +112,6 @@ class OpenACCClauseCIREmitter final
return createConstantInt(cgf.cgm.getLoc(loc), width, value);
}
- mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
- // '*' case leaves no identifier-info, just a nullptr.
- if (!ii)
- return mlir::acc::DeviceType::Star;
- return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
- .CaseLower("default", mlir::acc::DeviceType::Default)
- .CaseLower("host", mlir::acc::DeviceType::Host)
- .CaseLower("multicore", mlir::acc::DeviceType::Multicore)
- .CasesLower({"nvidia", "acc_device_nvidia"},
- mlir::acc::DeviceType::Nvidia)
- .CaseLower("radeon", mlir::acc::DeviceType::Radeon);
- }
-
mlir::acc::GangArgType decodeGangType(OpenACCGangKind gk) {
switch (gk) {
case OpenACCGangKind::Num:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h
index 5bcc9f57d67b1..639d14804087e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h
@@ -40,4 +40,17 @@ convertOpenACCModifiers(OpenACCModifierKind modifiers) {
mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
return mlirModifiers;
}
+
+inline mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
+ // '*' case leaves no identifier-info, just a nullptr.
+ if (!ii)
+ return mlir::acc::DeviceType::Star;
+ return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
+ .CaseLower("default", mlir::acc::DeviceType::Default)
+ .CaseLower("host", mlir::acc::DeviceType::Host)
+ .CaseLower("multicore", mlir::acc::DeviceType::Multicore)
+ .CasesLower({"nvidia", "acc_device_nvidia"},
+ mlir::acc::DeviceType::Nvidia)
+ .CaseLower("radeon", mlir::acc::DeviceType::Radeon);
+}
} // namespace clang::CIRGen
diff --git a/clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp b/clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp
new file mode 100644
index 0000000000000..61c985bd81f56
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp
@@ -0,0 +1,79 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+#pragma acc routine nohost device_type(nvidia, radeon) seq
+void Func1() {}
+void Func2() {}
+#pragma acc routine(Func2) device_type(radeon) seq
+
+#pragma acc routine device_type(multicore) worker device_type(nvidia, radeon) seq
+void Func3() {}
+void Func4() {}
+#pragma acc routine(Func4) device_type(nvidia) seq device_type(radeon) vector
+
+#pragma acc routine device_type(multicore) gang device_type(nvidia, radeon) gang
+void Func5() {}
+void Func6() {}
+#pragma acc routine(Func6) device_type(multicore) gang(dim:1) device_type(radeon) gang
+
+#pragma acc routine device_type(host) gang device_type(nvidia, radeon) gang(dim:1)
+void Func7() {}
+void Func8() {}
+#pragma acc routine(Func8) device_type(radeon) gang(dim:2)
+
+#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3)
+void Func9() {}
+void Func10() {}
+#pragma acc routine(Func10) device_type(nvidia) gang device_type(radeon) gang(dim:3)
+
+#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3) device_type(multicore) gang
+void Func11() {}
+void Func12() {}
+#pragma acc routine(Func12) device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3)
+
+#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang
+void Func13() {}
+void Func14() {}
+#pragma acc routine(Func14) device_type(nvidia) gang(dim:2) device_type(radeon) gang
+
+// CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq ([#acc.device_type<nvidia>, #acc.device_type<radeon>]) nohost
+
+// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>}
+
+// CHECK: cir.func{{.*}} @[[F3_NAME:.*Func3[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F3_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F3_R_NAME]] func(@[[F3_NAME]]) worker ([#acc.device_type<multicore>]) seq ([#acc.device_type<nvidia>, #acc.device_type<radeon>])
+
+// CHECK: cir.func{{.*}} @[[F4_NAME:.*Func4[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F4_R_NAME:.*]]]>}
+
+// CHECK: cir.func{{.*}} @[[F5_NAME:.*Func5[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F5_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F5_R_NAME]] func(@[[F5_NAME]]) gang([#acc.device_type<multicore>, #acc.device_type<nvidia>, #acc.device_type<radeon>])
+
+// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F7_NAME:.*Func7[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F7_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F7_R_NAME]] func(@[[F7_NAME]]) gang([#acc.device_type<host>], dim: 1 : i64 [#acc.device_type<nvidia>], dim: 1 : i64 [#acc.device_type<radeon>])
+
+// CHECK: cir.func{{.*}} @[[F8_NAME:.*Func8[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F8_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F9_NAME:.*Func9[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F9_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F9_R_NAME]] func(@[[F9_NAME]]) gang(dim: 2 : i64 [#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
+//
+// CHECK: cir.func{{.*}} @[[F10_NAME:.*Func10[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F10_R_NAME:.*]]]>}
+
+// CHECK: cir.func{{.*}} @[[F11_NAME:.*Func11[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F11_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F11_R_NAME]] func(@[[F11_NAME]]) gang([#acc.device_type<multicore>], dim: 2 : i64 [#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
+//
+// CHECK: cir.func{{.*}} @[[F12_NAME:.*Func12[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F12_R_NAME:.*]]]>}
+//
+// CHECK: cir.func{{.*}} @[[F13_NAME:.*Func13[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F13_R_NAME:.*]]]>}
+// CHECK: acc.routine @[[F13_R_NAME]] func(@[[F13_NAME]]) gang([#acc.device_type<radeon>], dim: 2 : i64 [#acc.device_type<nvidia>])
+//
+// CHECK: cir.func{{.*}} @[[F14_NAME:.*Func14[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F14_R_NAME:.*]]]>}
+
+// CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq ([#acc.device_type<radeon>])
+// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) vector ([#acc.device_type<radeon>]) seq ([#acc.device_type<nvidia>])
+// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) gang([#acc.device_type<radeon>], dim: 1 : i64 [#acc.device_type<multicore>])
+// CHECK: acc.routine @[[F8_R_NAME]] func(@[[F8_NAME]]) gang(dim: 2 : i64 [#acc.device_type<radeon>])
+// CHECK: acc.routine @[[F10_R_NAME]] func(@[[F10_NAME]]) gang([#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
+// CHECK: acc.routine @[[F12_R_NAME]] func(@[[F12_NAME]]) gang(dim: 2 : i64 [#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
+// CHECK: acc.routine @[[F14_R_NAME]] func(@[[F14_NAME]]) gang([#acc.device_type<radeon>], dim: 2 : i64 [#acc.device_type<nvidia>])
|
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.
The 'device_type' clause modifies how the clauses that are legal after it (seq, worker, vector, gang, bind) work. Previous patches were aware of how that was going to happen, thanks to experience with doing the same work on other constructs/clauses, so this is mostly just a repeat of those.
Tests for the first 4 and interactions with them are included, but 'bind' is not yet implemented, so its device_type tests will be added when it is lowered.