Skip to content

Commit 4006df9

Browse files
authored
[OpenACC][CIR] Implement 'nohost' lowering. (#170369)
This clause is pretty small/trivial and is a simple 'set a bool' value on the IR node, so its implementation is quite simple. We create the Operation with this as 'false', so the 'nohost' marks it as true always.
1 parent f0e1254 commit 4006df9

File tree

2 files changed

+11
-7
lines changed

2 files changed

+11
-7
lines changed

clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -329,6 +329,10 @@ class OpenACCRoutineClauseEmitter final
329329
void VisitVectorClause(const OpenACCVectorClause &clause) {
330330
routineOp.addVector(builder.getContext(), lastDeviceTypeValues);
331331
}
332+
333+
void VisitNoHostClause(const OpenACCNoHostClause &clause) {
334+
routineOp.setNohost(/*attrValue=*/true);
335+
}
332336
};
333337
} // namespace
334338

clang/test/CIR/CodeGenOpenACC/routine-clauses.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
22

3-
#pragma acc routine seq
3+
#pragma acc routine seq nohost
44
void Func1() {}
55

66
void Func2() {}
@@ -10,16 +10,16 @@ void Func2() {}
1010
void Func3() {}
1111

1212
void Func4() {}
13-
#pragma acc routine(Func4) worker
13+
#pragma acc routine(Func4) worker nohost
1414

15-
#pragma acc routine vector
15+
#pragma acc routine nohost vector
1616
void Func5() {}
1717

1818
void Func6() {}
19-
#pragma acc routine(Func6) vector
19+
#pragma acc routine(Func6) nohost vector
2020

2121
// CHECK: cir.func{{.*}} @[[F1_NAME:.*Func1[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F1_R_NAME:.*]]]>}
22-
// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq
22+
// CHECK: acc.routine @[[F1_R_NAME]] func(@[[F1_NAME]]) seq nohost
2323

2424
// CHECK: cir.func{{.*}} @[[F2_NAME:.*Func2[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F2_R_NAME:.*]]]>}
2525

@@ -34,5 +34,5 @@ void Func6() {}
3434
// CHECK: cir.func{{.*}} @[[F6_NAME:.*Func6[^\(]*]]({{.*}}){{.*}} attributes {acc.routine_info = #acc.routine_info<[@[[F6_R_NAME:.*]]]>}
3535

3636
// CHECK: acc.routine @[[F2_R_NAME]] func(@[[F2_NAME]]) seq
37-
// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) worker
38-
// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) vector
37+
// CHECK: acc.routine @[[F4_R_NAME]] func(@[[F4_NAME]]) worker nohost
38+
// CHECK: acc.routine @[[F6_R_NAME]] func(@[[F6_NAME]]) vector nohost

0 commit comments

Comments
 (0)