Skip to content

Commit 6601143

Browse files
jhuber6memfrob
jhuber6
authored and
memfrob
committed
[OpenMP] Rework OpenMP remarks
This patch rewrites and reworks a few of the existing remarks to make the mmore concise and consistent prior to writing the documentation for them. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D105898
1 parent 9beb3e4 commit 6601143

10 files changed

+89
-183
lines changed

‎clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c

Lines changed: 9 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -8,59 +8,44 @@ void baz(void) __attribute__((assume("omp_no_openmp")));
88

99
void bar1(void) {
1010
#pragma omp parallel // #0
11-
// all-remark@#0 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
12-
// safe-remark@#0 {{Parallel region is used in unknown ways; will not attempt to rewrite the state machine.}}
13-
// force-remark@#0 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__2_wrapper, kernel ID: <NONE>}}
11+
// safe-remark@#0 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine.}}
1412
{
1513
}
1614
}
1715
void bar2(void) {
1816
#pragma omp parallel // #1
19-
// all-remark@#1 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
20-
// safe-remark@#1 {{Parallel region is used in unknown ways; will not attempt to rewrite the state machine.}}
21-
// force-remark@#1 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__6_wrapper, kernel ID: <NONE>}}
17+
// safe-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine.}}
2218
{
2319
}
2420
}
2521

2622
void foo1(void) {
2723
#pragma omp target teams // #2
28-
// all-remark@#2 {{Generic-mode kernel is executed with a customized state machine [3 known parallel regions] (good).}}
29-
// all-remark@#2 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__1_wrapper, kernel ID: __omp_offloading}}
30-
// all-remark@#2 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__2_wrapper, kernel ID: __omp_offloading}}
24+
// all-remark@#2 {{Rewriting generic-mode kernel with a customized state machine.}}
25+
3126
{
32-
baz(); // all-remark {{Kernel will be executed in generic-mode due to this potential side-effect, consider to add `__attribute__((assume("ompx_spmd_amenable")))` to the called function '_Z3bazv'.}}
27+
baz(); // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
3328
#pragma omp parallel // #3
34-
// all-remark@#3 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
35-
// all-remark@#3 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__1_wrapper, kernel ID: __omp_offloading}}
3629
{
3730
}
3831
bar1();
3932
#pragma omp parallel // #4
40-
// all-remark@#4 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
41-
// all-remark@#4 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__2_wrapper, kernel ID: __omp_offloading}}
4233
{
4334
}
4435
}
4536
}
4637

4738
void foo2(void) {
4839
#pragma omp target teams // #5
49-
// all-remark@#5 {{Generic-mode kernel is executed with a customized state machine [4 known parallel regions] (good).}}
50-
// all-remark@#5 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__5_wrapper, kernel ID: __omp_offloading}}
51-
// all-remark@#5 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__4_wrapper, kernel ID: __omp_offloading}}
40+
// all-remark@#5 {{Rewriting generic-mode kernel with a customized state machine.}}
5241
{
53-
baz(); // all-remark {{Kernel will be executed in generic-mode due to this potential side-effect, consider to add `__attribute__((assume("ompx_spmd_amenable")))` to the called function '_Z3bazv'.}}
42+
baz(); // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
5443
#pragma omp parallel // #6
55-
// all-remark@#6 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
56-
// all-remark@#6 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__4_wrapper, kernel ID: __omp_offloading}}
5744
{
5845
}
5946
bar1();
6047
bar2();
6148
#pragma omp parallel // #7
62-
// all-remark@#7 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
63-
// all-remark@#7 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__5_wrapper, kernel ID: __omp_offloading}}
6449
{
6550
}
6651
bar1();
@@ -70,21 +55,15 @@ void foo2(void) {
7055

7156
void foo3(void) {
7257
#pragma omp target teams // #8
73-
// all-remark@#8 {{Generic-mode kernel is executed with a customized state machine [4 known parallel regions] (good).}}
74-
// all-remark@#8 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__7_wrapper, kernel ID: __omp_offloading}}
75-
// all-remark@#8 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__8_wrapper, kernel ID: __omp_offloading}}
58+
// all-remark@#8 {{Rewriting generic-mode kernel with a customized state machine.}}
7659
{
77-
baz(); // all-remark {{Kernel will be executed in generic-mode due to this potential side-effect, consider to add `__attribute__((assume("ompx_spmd_amenable")))` to the called function '_Z3bazv'.}}
60+
baz(); // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
7861
#pragma omp parallel // #9
79-
// all-remark@#9 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
80-
// all-remark@#9 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__7_wrapper, kernel ID: __omp_offloading}}
8162
{
8263
}
8364
bar1();
8465
bar2();
8566
#pragma omp parallel // #10
86-
// all-remark@#10 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
87-
// all-remark@#10 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__8_wrapper, kernel ID: __omp_offloading}}
8867
{
8968
}
9069
bar1();
@@ -104,5 +83,4 @@ void spmd(void) {
10483
}
10584
}
10685

107-
// all-remark@* 5 {{OpenMP runtime call __kmpc_global_thread_num moved to beginning of OpenMP region}}
10886
// all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}

‎clang/test/OpenMP/remarks_parallel_in_target_state_machine.c

Lines changed: 6 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -8,28 +8,21 @@ void baz(void) __attribute__((assume("omp_no_openmp")));
88

99
void bar(void) {
1010
#pragma omp parallel // #1 \
11-
// expected-remark@#1 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}} \
12-
// expected-remark@#1 {{Parallel region is used in unknown ways; will not attempt to rewrite the state machine.}}
11+
// expected-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine.}}
1312
{
1413
}
1514
}
1615

1716
void foo(void) {
18-
#pragma omp target teams // #2 \
19-
// expected-remark@#2 {{Generic-mode kernel is executed with a customized state machine [3 known parallel regions] (good).}}
20-
// expected-remark@#2 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__1_wrapper, kernel ID: __omp_offloading}} \
21-
// expected-remark@#2 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__2_wrapper, kernel ID: __omp_offloading}}
17+
#pragma omp target teams // #2
18+
// expected-remark@#2 {{Rewriting generic-mode kernel with a customized state machine.}}
2219
{
23-
baz(); // expected-remark {{Kernel will be executed in generic-mode due to this potential side-effect, consider to add `__attribute__((assume("ompx_spmd_amenable")))` to the called function '_Z3bazv'.}}
24-
#pragma omp parallel // #3 \
25-
// expected-remark@#3 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}} \
26-
// expected-remark@#3 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__1_wrapper, kernel ID: __omp_offloading}}
20+
baz(); // expected-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
21+
#pragma omp parallel
2722
{
2823
}
2924
bar();
30-
#pragma omp parallel // #4 \
31-
// expected-remark@#4 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}} \
32-
// expected-remark@#4 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__2_wrapper, kernel ID: __omp_offloading}}
25+
#pragma omp parallel
3326
{
3427
}
3528
}
@@ -47,5 +40,4 @@ void spmd(void) {
4740
}
4841
}
4942

50-
// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num moved to beginning of OpenMP region}}
5143
// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}

‎llvm/lib/Transforms/IPO/AttributorAttributes.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5328,10 +5328,10 @@ ChangeStatus AAHeapToStackFunction::updateImpl(Attributor &A) {
53285328

53295329
// Emit a missed remark if this is missed OpenMP globalization.
53305330
auto Remark = [&](OptimizationRemarkMissed ORM) {
5331-
return ORM << "Could not move globalized variable to the stack as "
5332-
"variable is potentially captured in call; mark "
5333-
"parameter as "
5334-
"`__attribute__((noescape))` to override.";
5331+
return ORM
5332+
<< "Could not move globalized variable to the stack. "
5333+
"Variable is potentially captured in call. Mark "
5334+
"parameter as `__attribute__((noescape))` to override.";
53355335
};
53365336

53375337
if (AI.LibraryFunctionId == LibFunc___kmpc_alloc_shared)

0 commit comments

Comments
 (0)