Skip to content

Commit 67bd966

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Implement 'device_type' lowering for Routine (#170893)
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.
1 parent 190b8d0 commit 67bd966

File tree

4 files changed

+99
-13
lines changed

4 files changed

+99
-13
lines changed

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

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -355,6 +355,13 @@ class OpenACCRoutineClauseEmitter final
355355
curValue.getZExtValue());
356356
}
357357
}
358+
359+
void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
360+
lastDeviceTypeValues.clear();
361+
362+
for (const DeviceTypeArgument &arg : clause.getArchitectures())
363+
lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo()));
364+
}
358365
};
359366
} // namespace
360367

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

Lines changed: 0 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -112,19 +112,6 @@ class OpenACCClauseCIREmitter final
112112
return createConstantInt(cgf.cgm.getLoc(loc), width, value);
113113
}
114114

115-
mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
116-
// '*' case leaves no identifier-info, just a nullptr.
117-
if (!ii)
118-
return mlir::acc::DeviceType::Star;
119-
return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
120-
.CaseLower("default", mlir::acc::DeviceType::Default)
121-
.CaseLower("host", mlir::acc::DeviceType::Host)
122-
.CaseLower("multicore", mlir::acc::DeviceType::Multicore)
123-
.CasesLower({"nvidia", "acc_device_nvidia"},
124-
mlir::acc::DeviceType::Nvidia)
125-
.CaseLower("radeon", mlir::acc::DeviceType::Radeon);
126-
}
127-
128115
mlir::acc::GangArgType decodeGangType(OpenACCGangKind gk) {
129116
switch (gk) {
130117
case OpenACCGangKind::Num:

‎clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -40,4 +40,17 @@ convertOpenACCModifiers(OpenACCModifierKind modifiers) {
4040
mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
4141
return mlirModifiers;
4242
}
43+
44+
inline mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
45+
// '*' case leaves no identifier-info, just a nullptr.
46+
if (!ii)
47+
return mlir::acc::DeviceType::Star;
48+
return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
49+
.CaseLower("default", mlir::acc::DeviceType::Default)
50+
.CaseLower("host", mlir::acc::DeviceType::Host)
51+
.CaseLower("multicore", mlir::acc::DeviceType::Multicore)
52+
.CasesLower({"nvidia", "acc_device_nvidia"},
53+
mlir::acc::DeviceType::Nvidia)
54+
.CaseLower("radeon", mlir::acc::DeviceType::Radeon);
55+
}
4356
} // namespace clang::CIRGen

‎clang/test/CIR/CodeGenOpenACC/routine-device_type.cpp‎

Lines changed: 79 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,79 @@
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
2+
3+
#pragma acc routine nohost device_type(nvidia, radeon) seq
4+
void Func1() {}
5+
void Func2() {}
6+
#pragma acc routine(Func2) device_type(radeon) seq
7+
8+
#pragma acc routine device_type(multicore) worker device_type(nvidia, radeon) seq
9+
void Func3() {}
10+
void Func4() {}
11+
#pragma acc routine(Func4) device_type(nvidia) seq device_type(radeon) vector
12+
13+
#pragma acc routine device_type(multicore) gang device_type(nvidia, radeon) gang
14+
void Func5() {}
15+
void Func6() {}
16+
#pragma acc routine(Func6) device_type(multicore) gang(dim:1) device_type(radeon) gang
17+
18+
#pragma acc routine device_type(host) gang device_type(nvidia, radeon) gang(dim:1)
19+
void Func7() {}
20+
void Func8() {}
21+
#pragma acc routine(Func8) device_type(radeon) gang(dim:2)
22+
23+
#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3)
24+
void Func9() {}
25+
void Func10() {}
26+
#pragma acc routine(Func10) device_type(nvidia) gang device_type(radeon) gang(dim:3)
27+
28+
#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3) device_type(multicore) gang
29+
void Func11() {}
30+
void Func12() {}
31+
#pragma acc routine(Func12) device_type(nvidia) gang(dim:2) device_type(radeon) gang(dim:3)
32+
33+
#pragma acc routine device_type(nvidia) gang(dim:2) device_type(radeon) gang
34+
void Func13() {}
35+
void Func14() {}
36+
#pragma acc routine(Func14) device_type(nvidia) gang(dim:2) device_type(radeon) gang
37+
38+
// CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
39+
// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq ([#acc.device_type<nvidia>, #acc.device_type<radeon>]) nohost
40+
41+
// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>}
42+
43+
// CHECK: cir.func{{.*}} @[[F3_NAME:.*Func3[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F3_R_NAME:.*]]]>}
44+
// CHECK: acc.routine @[[F3_R_NAME]] func(@[[F3_NAME]]) worker ([#acc.device_type<multicore>]) seq ([#acc.device_type<nvidia>, #acc.device_type<radeon>])
45+
46+
// CHECK: cir.func{{.*}} @[[F4_NAME:.*Func4[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F4_R_NAME:.*]]]>}
47+
48+
// CHECK: cir.func{{.*}} @[[F5_NAME:.*Func5[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F5_R_NAME:.*]]]>}
49+
// CHECK: acc.routine @[[F5_R_NAME]] func(@[[F5_NAME]]) gang([#acc.device_type<multicore>, #acc.device_type<nvidia>, #acc.device_type<radeon>])
50+
51+
// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
52+
//
53+
// CHECK: cir.func{{.*}} @[[F7_NAME:.*Func7[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F7_R_NAME:.*]]]>}
54+
// 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>])
55+
56+
// CHECK: cir.func{{.*}} @[[F8_NAME:.*Func8[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F8_R_NAME:.*]]]>}
57+
//
58+
// CHECK: cir.func{{.*}} @[[F9_NAME:.*Func9[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F9_R_NAME:.*]]]>}
59+
// CHECK: acc.routine @[[F9_R_NAME]] func(@[[F9_NAME]]) gang(dim: 2 : i64 [#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
60+
//
61+
// CHECK: cir.func{{.*}} @[[F10_NAME:.*Func10[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F10_R_NAME:.*]]]>}
62+
63+
// CHECK: cir.func{{.*}} @[[F11_NAME:.*Func11[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F11_R_NAME:.*]]]>}
64+
// 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>])
65+
//
66+
// CHECK: cir.func{{.*}} @[[F12_NAME:.*Func12[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F12_R_NAME:.*]]]>}
67+
//
68+
// CHECK: cir.func{{.*}} @[[F13_NAME:.*Func13[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F13_R_NAME:.*]]]>}
69+
// CHECK: acc.routine @[[F13_R_NAME]] func(@[[F13_NAME]]) gang([#acc.device_type<radeon>], dim: 2 : i64 [#acc.device_type<nvidia>])
70+
//
71+
// CHECK: cir.func{{.*}} @[[F14_NAME:.*Func14[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F14_R_NAME:.*]]]>}
72+
73+
// CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq ([#acc.device_type<radeon>])
74+
// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) vector ([#acc.device_type<radeon>]) seq ([#acc.device_type<nvidia>])
75+
// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) gang([#acc.device_type<radeon>], dim: 1 : i64 [#acc.device_type<multicore>])
76+
// CHECK: acc.routine @[[F8_R_NAME]] func(@[[F8_NAME]]) gang(dim: 2 : i64 [#acc.device_type<radeon>])
77+
// CHECK: acc.routine @[[F10_R_NAME]] func(@[[F10_NAME]]) gang([#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
78+
// CHECK: acc.routine @[[F12_R_NAME]] func(@[[F12_NAME]]) gang(dim: 2 : i64 [#acc.device_type<nvidia>], dim: 3 : i64 [#acc.device_type<radeon>])
79+
// CHECK: acc.routine @[[F14_R_NAME]] func(@[[F14_NAME]]) gang([#acc.device_type<radeon>], dim: 2 : i64 [#acc.device_type<nvidia>])

0 commit comments

Comments
 (0)