[llvm] 7eb899c - [OpenMP] Add more verbose remarks for runtime folding

[Joseph Huber via llvm-commits]

Author: Joseph Huber
Date: 2021-09-10T17:36:06-04:00
New Revision: 7eb899cbcdd15bd57e5630bb0c43163d0c992a82

URL: https://github.com/llvm/llvm-project/commit/7eb899cbcdd15bd57e5630bb0c43163d0c992a82
DIFF: https://github.com/llvm/llvm-project/commit/7eb899cbcdd15bd57e5630bb0c43163d0c992a82.diff

LOG: [OpenMP] Add more verbose remarks for runtime folding

We peform runtime folding, but do not currently emit remarks when it is
performed. This is because it comes from the runtime library and is
beyond the users control. However, people may still wish to view  this
and similar information easily, so we can enable this behaviour using a
special flag to enable verbose remarks.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D109627

Added: 
    openmp/docs/remarks/OMP180.rst

Modified: 
    llvm/lib/Transforms/IPO/OpenMPOpt.cpp
    openmp/docs/remarks/OptimizationRemarks.rst

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Transforms/IPO/OpenMPOpt.cpp b/llvm/lib/Transforms/IPO/OpenMPOpt.cpp
index 18bf1fc5e5af8..456c94090a149 100644
--- a/llvm/lib/Transforms/IPO/OpenMPOpt.cpp
+++ b/llvm/lib/Transforms/IPO/OpenMPOpt.cpp
@@ -102,6 +102,11 @@ static cl::opt<bool> AlwaysInlineDeviceFunctions(
     cl::desc("Inline all applicible functions on the device."), cl::Hidden,
     cl::init(false));
 
+static cl::opt<bool>
+    EnableVerboseRemarks("openmp-opt-verbose-remarks", cl::ZeroOrMore,
+                         cl::desc("Enables more verbose remarks."), cl::Hidden,
+                         cl::init(false));
+
 STATISTIC(NumOpenMPRuntimeCallsDeduplicated,
           "Number of OpenMP runtime calls deduplicated");
 STATISTIC(NumOpenMPParallelRegionsDeleted,
@@ -4034,11 +4039,25 @@ struct AAFoldRuntimeCallCallSiteReturned : AAFoldRuntimeCall {
     ChangeStatus Changed = ChangeStatus::UNCHANGED;
 
     if (SimplifiedValue.hasValue() && SimplifiedValue.getValue()) {
-      Instruction &CB = *getCtxI();
-      A.changeValueAfterManifest(CB, **SimplifiedValue);
-      A.deleteAfterManifest(CB);
+      Instruction &I = *getCtxI();
+      A.changeValueAfterManifest(I, **SimplifiedValue);
+      A.deleteAfterManifest(I);
+
+      CallBase *CB = dyn_cast<CallBase>(&I);
+      auto Remark = [&](OptimizationRemark OR) {
+        if (auto *C = dyn_cast<ConstantInt>(*SimplifiedValue))
+          return OR << "Replacing OpenMP runtime call "
+                    << CB->getCalledFunction()->getName() << " with "
+                    << ore::NV("FoldedValue", C->getZExtValue()) << ".";
+        else
+          return OR << "Replacing OpenMP runtime call "
+                    << CB->getCalledFunction()->getName() << ".";
+      };
+
+      if (CB && EnableVerboseRemarks)
+        A.emitRemark<OptimizationRemark>(CB, "OMP180", Remark);
 
-      LLVM_DEBUG(dbgs() << TAG << "Folding runtime call: " << CB << " with "
+      LLVM_DEBUG(dbgs() << TAG << "Replacing runtime call: " << I << " with "
                         << **SimplifiedValue << "\n");
 
       Changed = ChangeStatus::CHANGED;

diff  --git a/openmp/docs/remarks/OMP180.rst b/openmp/docs/remarks/OMP180.rst
new file mode 100644
index 0000000000000..756cdcb558fd0
--- /dev/null
+++ b/openmp/docs/remarks/OMP180.rst
@@ -0,0 +1,38 @@
+.. _omp180:
+
+Replacing OpenMP runtime call <call> with <value>.
+====================================================================
+
+This optimization remark indicates that analysis determined an OpenMP runtime 
+calls can be replaced with a constant value. This can occur when an OpenMP 
+runtime call that queried some internal state was found to always return a 
+single value after analysis.
+
+Example
+-------
+
+This optimization will trigger for most target regions to simplify the runtime 
+once certain constants are known. This will trigger for internal runtime 
+functions so it requires enabling verbose remarks with 
+`-openmp-opt-verbose-remarks`.
+
+.. code-block:: c++
+
+  void foo() {
+  #pragma omp target parallel
+    { }
+  }
+
+.. code-block:: console
+
+  $ clang test.c -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass=openmp-opt \
+    -mllvm -openmp-opt-verbose-remarks
+  remark: Replacing runtime call __kmpc_is_spmd_exec_mode with 1. [OMP180] [-Rpass=openmp-opt]
+  remark: Replacing runtime call __kmpc_is_spmd_exec_mode with 1. [OMP180] [-Rpass=openmp-opt]
+  remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=openmp-opt]
+  remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=openmp-opt]
+
+Diagnostic Scope
+----------------
+
+OpenMP optimization remark.

diff  --git a/openmp/docs/remarks/OptimizationRemarks.rst b/openmp/docs/remarks/OptimizationRemarks.rst
index 1463a61be40e2..bf964e2ffc74a 100644
--- a/openmp/docs/remarks/OptimizationRemarks.rst
+++ b/openmp/docs/remarks/OptimizationRemarks.rst
@@ -39,6 +39,7 @@ OpenMP Remarks
    OMP150
    OMP160
    OMP170
+   OMP180
 
 .. list-table::
    :widths: 15 15 70
@@ -107,3 +108,6 @@ OpenMP Remarks
    * - :ref:`OMP170 <omp170>`
      - Optimization
      - OpenMP runtime call <call> deduplicated.
+   * - :ref:`OMP180 <omp180>`
+     - Optimization
+     - Replacing OpenMP runtime call <call> with <value>.