jdoerfert updated this revision to Diff 277951.
jdoerfert added a comment.
Improve wording (again)
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D83707/new/
https://reviews.llvm.org/D83707
Files:
clang/test/OpenMP/remarks_parallel_in_target_state_machine.c
llvm/lib/Transforms/IPO/OpenMPOpt.cpp
Index: llvm/lib/Transforms/IPO/OpenMPOpt.cpp
===================================================================
--- llvm/lib/Transforms/IPO/OpenMPOpt.cpp
+++ llvm/lib/Transforms/IPO/OpenMPOpt.cpp
@@ -1033,6 +1033,7 @@
// Check if the function is uses in a __kmpc_kernel_prepare_parallel call at
// all.
bool UnknownUse = false;
+ bool KernelPrepareUse = false;
unsigned NumDirectCalls = 0;
SmallVector<Use *, 2> ToBeReplacedStateMachineUses;
@@ -1049,31 +1050,90 @@
}
if (OpenMPOpt::getCallIfRegularCall(*U.getUser(),
&KernelPrepareParallelRFI)) {
+ KernelPrepareUse = true;
ToBeReplacedStateMachineUses.push_back(&U);
return;
}
UnknownUse = true;
});
- // If this ever hits, we should investigate.
- if (UnknownUse || NumDirectCalls != 1)
+ // Do not emit a remark if we haven't seen a __kmpc_kernel_prepare_parallel
+ // use.
+ if (!KernelPrepareUse)
continue;
- // TODO: This is not a necessary restriction and should be lifted.
- if (ToBeReplacedStateMachineUses.size() != 2)
+ {
+ auto Remark = [&](OptimizationRemark OR) {
+ return OR << "Found a parallel region that is called in a target "
+ "region but not part of a combined target construct nor "
+ "nesed 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.";
+ };
+ emitRemarkOnFunction(F, "OpenMPParallelRegionInNonSPMD", Remark);
+ }
+
+ // If this ever hits, we should investigate.
+ // TODO: Checking the number of uses is not a necessary restriction and
+ // should be lifted.
+ if (UnknownUse || NumDirectCalls != 1 ||
+ ToBeReplacedStateMachineUses.size() != 2) {
+ {
+ auto Remark = [&](OptimizationRemark OR) {
+ return OR << "Parallel region is used in "
+ << (UnknownUse ? "unknown" : "unexpected")
+ << " ways; will not attempt to rewrite the state machine.";
+ };
+ emitRemarkOnFunction(F, "OpenMPParallelRegionInNonSPMD", Remark);
+ }
continue;
+ }
// Even if we have __kmpc_kernel_prepare_parallel calls, we (for now) give
// up if the function is not called from a unique kernel.
Kernel K = getUniqueKernelFor(*F);
- if (!K)
+ if (!K) {
+ {
+ auto Remark = [&](OptimizationRemark OR) {
+ return OR << "Parallel region is not known to be called from a "
+ "unique single target region, maybe the surrounding "
+ "function has external linkage?; will not attempt to "
+ "rewrite the state machine use.";
+ };
+ emitRemarkOnFunction(F, "OpenMPParallelRegionInMultipleKernesl",
+ Remark);
+ }
continue;
+ }
// We now know F is a parallel body function called only from the kernel K.
// We also identified the state machine uses in which we replace the
// function pointer by a new global symbol for identification purposes. This
// ensures only direct calls to the function are left.
+ {
+ auto RemarkParalleRegion = [&](OptimizationRemark OR) {
+ return OR << "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: "
+ << ore::NV("OpenMPParallelRegion", F->getName())
+ << ", kernel ID: "
+ << ore::NV("OpenMPTargetRegion", K->getName()) << ")";
+ };
+ emitRemarkOnFunction(F, "OpenMPParallelRegionInNonSPMD",
+ RemarkParalleRegion);
+ auto RemarkKernel = [&](OptimizationRemark OR) {
+ return OR << "Target region containing the parallel region that is "
+ "specialized. (parallel region ID: "
+ << ore::NV("OpenMPParallelRegion", F->getName())
+ << ", kernel ID: "
+ << ore::NV("OpenMPTargetRegion", K->getName()) << ")";
+ };
+ emitRemarkOnFunction(K, "OpenMPParallelRegionInNonSPMD", RemarkKernel);
+ }
+
Module &M = *F->getParent();
Type *Int8Ty = Type::getInt8Ty(M.getContext());
Index: clang/test/OpenMP/remarks_parallel_in_target_state_machine.c
===================================================================
--- /dev/null
+++ clang/test/OpenMP/remarks_parallel_in_target_state_machine.c
@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -verify=host -Rpass=openmp -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -Rpass=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
+// RUN: %clang_cc1 -fexperimental-new-pass-manager -verify -Rpass=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
+
+// host-no-diagnostics
+
+void bar(void) {
+#pragma omp parallel // #1
+ // expected-remark@#1 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nesed 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.}}
+ // expected-remark@#1 {{Parallel region is not known to be called from a unique single target region, maybe the surrounding function has external linkage?; will not attempt to rewrite the state machine use.}}
+ {
+ }
+}
+
+void foo(void) {
+#pragma omp target teams // #2
+ // expected-remark@#2 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__1_wrapper, kernel ID: __omp_offloading_22}}
+ // expected-remark@#2 {{Target region containing the parallel region that is specialized. (parallel region ID: __omp_outlined__3_wrapper, kernel ID: __omp_offloading_22}}
+ {
+#pragma omp parallel // #3
+ // expected-remark@#3 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nesed 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.}}
+ // 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_22}}
+ {
+ }
+ bar();
+#pragma omp parallel // #4
+ // expected-remark@#4 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nesed 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.}}
+ // 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__3_wrapper, kernel ID: __omp_offloading_22}}
+ {
+ }
+ }
+}
+
+// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num moved to}}
+// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
