[Openmp-commits] [openmp] 602e0eb - [OpenMP][DeviceRTL] Fix the issue that multiple calls to `omp_get_wtime` is optimized out by mistake

Shilei Tian via Openmp-commits openmp-commits at lists.llvm.org
Fri Jul 22 10:46:49 PDT 2022


Author: Shilei Tian
Date: 2022-07-22T13:46:45-04:00
New Revision: 602e0eb9f0c69db455185b5fbe6766ad2573f6e6

URL: https://github.com/llvm/llvm-project/commit/602e0eb9f0c69db455185b5fbe6766ad2573f6e6
DIFF: https://github.com/llvm/llvm-project/commit/602e0eb9f0c69db455185b5fbe6766ad2573f6e6.diff

LOG: [OpenMP][DeviceRTL] Fix the issue that multiple calls to `omp_get_wtime` is optimized out by mistake

Multiple calls to `omp_get_wtime` could be optimized out due to the function
is mistakenly marked as `readnone`. This patch fixes the issue, and also add the
support to run optimization on `libomptarget` tests.

Reviewed By: jdoerfert

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

Added: 
    openmp/libomptarget/test/offloading/wtime.c

Modified: 
    llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
    openmp/libomptarget/DeviceRTL/src/Misc.cpp
    openmp/libomptarget/test/lit.cfg

Removed: 
    


################################################################################
diff  --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
index 9d1ab57729b74..7a70af9ce792f 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -37,6 +37,7 @@ __OMP_TYPE(Int8Ptr)
 __OMP_TYPE(Int16Ptr)
 __OMP_TYPE(Int32Ptr)
 __OMP_TYPE(Int64Ptr)
+__OMP_TYPE(Double)
 
 OMP_TYPE(SizeTy, M.getDataLayout().getIntPtrType(Ctx))
 
@@ -236,6 +237,7 @@ __OMP_RTL(omp_get_place_proc_ids, false, Void, Int32, Int32Ptr)
 __OMP_RTL(omp_get_place_num, false, Int32, )
 __OMP_RTL(omp_get_partition_num_places, false, Int32, )
 __OMP_RTL(omp_get_partition_place_nums, false, Void, Int32Ptr)
+__OMP_RTL(omp_get_wtime, false, Double,)
 
 __OMP_RTL(omp_set_num_threads, false, Void, Int32)
 __OMP_RTL(omp_set_dynamic, false, Void, Int32)
@@ -681,6 +683,7 @@ __OMP_RTL_ATTRS(omp_get_partition_num_places, GetterAttrs, AttributeSet(),
                 ParamAttrs())
 __OMP_RTL_ATTRS(omp_get_partition_place_nums, GetterAttrs, AttributeSet(),
                 ParamAttrs())
+__OMP_RTL_ATTRS(omp_get_wtime, GetterAttrs, AttributeSet(), ParamAttrs())
 
 __OMP_RTL_ATTRS(omp_set_num_threads, SetterAttrs, AttributeSet(), ParamAttrs())
 __OMP_RTL_ATTRS(omp_set_dynamic, SetterAttrs, AttributeSet(), ParamAttrs())
@@ -919,7 +922,7 @@ __OMP_RTL_ATTRS(__kmpc_doacross_fini, BarrierAttrs, AttributeSet(),
 
 __OMP_RTL_ATTRS(__kmpc_alloc_shared, AttributeSet(
   EnumAttr(NoUnwind),
-  EnumAttr(NoSync), 
+  EnumAttr(NoSync),
   AllocSizeAttr(0, None)), ReturnPtrAttrs, ParamAttrs())
 __OMP_RTL_ATTRS(__kmpc_free_shared, DeviceAllocAttrs, AttributeSet(),
                 ParamAttrs(NoCaptureAttrs))

diff  --git a/openmp/libomptarget/DeviceRTL/src/Misc.cpp b/openmp/libomptarget/DeviceRTL/src/Misc.cpp
index 554a13ae4794e..7166925db4362 100644
--- a/openmp/libomptarget/DeviceRTL/src/Misc.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Misc.cpp
@@ -52,7 +52,7 @@ double getWTick() {
 
 double getWTime() {
   unsigned long long nsecs;
-  asm("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
+  asm volatile("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
   return (double)nsecs * getWTick();
 }
 

diff  --git a/openmp/libomptarget/test/lit.cfg b/openmp/libomptarget/test/lit.cfg
index 5436d324e06d6..f0eadea1f67d6 100644
--- a/openmp/libomptarget/test/lit.cfg
+++ b/openmp/libomptarget/test/lit.cfg
@@ -128,7 +128,7 @@ for libomptarget_target in config.libomptarget_all_targets:
     # Is this target in the current system? If so create a compile, run and test
     # command. Otherwise create command that return false.
     if libomptarget_target == config.libomptarget_current_target:
-        config.substitutions.append(("%libomptarget-compilexx-run-and-check-generic", 
+        config.substitutions.append(("%libomptarget-compilexx-run-and-check-generic",
             "%libomptarget-compilexx-run-and-check-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-compile-run-and-check-generic",
             "%libomptarget-compile-run-and-check-" + libomptarget_target))
@@ -140,6 +140,18 @@ for libomptarget_target in config.libomptarget_all_targets:
             "%libomptarget-compilexx-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-compile-generic",
             "%libomptarget-compile-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileoptxx-run-and-check-generic",
+            "%libomptarget-compileoptxx-run-and-check-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileopt-run-and-check-generic",
+            "%libomptarget-compileopt-run-and-check-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileoptxx-and-run-generic",
+            "%libomptarget-compileoptxx-and-run-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileopt-and-run-generic",
+            "%libomptarget-compileopt-and-run-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileoptxx-generic",
+            "%libomptarget-compileoptxx-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileopt-generic",
+            "%libomptarget-compileopt-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-run-generic",
             "%libomptarget-run-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-run-fail-generic",
@@ -174,6 +186,28 @@ for libomptarget_target in config.libomptarget_all_targets:
         config.substitutions.append(("%libomptarget-compile-" + \
             libomptarget_target, \
             "%clang-" + libomptarget_target + " %s -o %t"))
+        config.substitutions.append(("%libomptarget-compileoptxx-run-and-check-" + \
+            libomptarget_target, \
+            "%libomptarget-compileoptxx-and-run-" + libomptarget_target + \
+            " | " + config.libomptarget_filecheck + " %s"))
+        config.substitutions.append(("%libomptarget-compileopt-run-and-check-" + \
+            libomptarget_target, \
+            "%libomptarget-compileopt-and-run-" + libomptarget_target + \
+            " | " + config.libomptarget_filecheck + " %s"))
+        config.substitutions.append(("%libomptarget-compileoptxx-and-run-" + \
+            libomptarget_target, \
+            "%libomptarget-compileoptxx-" + libomptarget_target + " && " + \
+            "%libomptarget-run-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileopt-and-run-" + \
+            libomptarget_target, \
+            "%libomptarget-compileopt-" + libomptarget_target + " && " + \
+            "%libomptarget-run-" + libomptarget_target))
+        config.substitutions.append(("%libomptarget-compileoptxx-" + \
+            libomptarget_target, \
+            "%clangxx-" + libomptarget_target + " -O3 %s -o %t"))
+        config.substitutions.append(("%libomptarget-compileopt-" + \
+            libomptarget_target, \
+            "%clang-" + libomptarget_target + " -O3 %s -o %t"))
         config.substitutions.append(("%libomptarget-run-" + \
             libomptarget_target, \
             "%t"))
@@ -207,6 +241,24 @@ for libomptarget_target in config.libomptarget_all_targets:
         config.substitutions.append(("%libomptarget-compile-" + \
             libomptarget_target, \
             "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compileopt-run-and-check-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compileoptxx-run-and-check-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compileopt-and-run-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compileoptxx-and-run-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compileoptxx-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compileopt-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
         config.substitutions.append(("%libomptarget-run-" + \
             libomptarget_target, \
             "echo ignored-command"))

diff  --git a/openmp/libomptarget/test/offloading/wtime.c b/openmp/libomptarget/test/offloading/wtime.c
new file mode 100644
index 0000000000000..bc2e1f79907eb
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/wtime.c
@@ -0,0 +1,24 @@
+// RUN: %libomptarget-compileopt-run-and-check-generic
+
+// UNSUPPORTED: amdgcn-amd-amdhsa
+// UNSUPPORTED: amdgcn-amd-amdhsa-oldDriver
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+
+#include <omp.h>
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+  int data[1024];
+#pragma omp target
+  {
+    double start = omp_get_wtime();
+    for (int i = 0; i < 1024; ++i)
+      data[i] = i;
+    double end = omp_get_wtime();
+    double duration = end - start;
+    printf("duration: %lfs\n", duration);
+  }
+  return 0;
+}
+
+// CHECK: duration: [1-9]+


        


More information about the Openmp-commits mailing list