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

[tianshilei1992]

Summary

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.

[jdoerfert]

LG, though we might want to make sure we mark it as readonly and inaccessible_mem_only in the openmp-opt pass (see OMPKinds.td)

[tianshilei1992]

update OMPKinds.def

[tianshilei1992]

fix lit.cfg

[barannikov88]

@tianshilei1992
This broke the build:

FAILED: lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o 
CCACHE_CPP2=yes CCACHE_HASHDIR=yes /usr/bin/ccache /usr/bin/clang++ -DGTEST_HAS_RTTI=0 -D_DEBUG -D_GNU_SOURCE -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -I/home/sergei/llvm-project/cmake-build-release/lib/Frontend/OpenMP -I/home/sergei/llvm-project/llvm/lib/Frontend/OpenMP -I/home/sergei/llvm-project/cmake-build-release/include -I/home/sergei/llvm-project/llvm/include -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wc++98-compat-extra-semi -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wsuggest-override -Wstring-conversion -Wmisleading-indentation -fdiagnostics-color -ffunction-sections -fdata-sections -O3 -DNDEBUG -fPIC  -fno-exceptions -fno-rtti -UNDEBUG -std=c++14 -MD -MT lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o -MF lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o.d -o lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o -c /home/sergei/llvm-project/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
In file included from /home/sergei/llvm-project/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp:354:
/home/sergei/llvm-project/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def:684:1: error: use of undeclared identifier 'OMPRTL_omp_get_wtime'
__OMP_RTL_ATTRS(omp_get_wtime, GetterAttrs, AttributeSet(), ParamAttrs())
^
/home/sergei/llvm-project/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def:610:17: note: expanded from macro '__OMP_RTL_ATTRS'
  OMP_RTL_ATTRS(OMPRTL_##Name, FnAttrSet, RetAttrSet, ArgAttrSets)
                ^
<scratch space>:78:1: note: expanded from here
OMPRTL_omp_get_wtime
^
1 error generated.

[tianshilei1992]

In D130179#3672041, @barannikov88 wrote:

@tianshilei1992
This broke the build:

FAILED: lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o 
CCACHE_CPP2=yes CCACHE_HASHDIR=yes /usr/bin/ccache /usr/bin/clang++ -DGTEST_HAS_RTTI=0 -D_DEBUG -D_GNU_SOURCE -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -I/home/sergei/llvm-project/cmake-build-release/lib/Frontend/OpenMP -I/home/sergei/llvm-project/llvm/lib/Frontend/OpenMP -I/home/sergei/llvm-project/cmake-build-release/include -I/home/sergei/llvm-project/llvm/include -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wc++98-compat-extra-semi -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wsuggest-override -Wstring-conversion -Wmisleading-indentation -fdiagnostics-color -ffunction-sections -fdata-sections -O3 -DNDEBUG -fPIC  -fno-exceptions -fno-rtti -UNDEBUG -std=c++14 -MD -MT lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o -MF lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o.d -o lib/Frontend/OpenMP/CMakeFiles/LLVMFrontendOpenMP.dir/OMPIRBuilder.cpp.o -c /home/sergei/llvm-project/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
In file included from /home/sergei/llvm-project/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp:354:
/home/sergei/llvm-project/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def:684:1: error: use of undeclared identifier 'OMPRTL_omp_get_wtime'
__OMP_RTL_ATTRS(omp_get_wtime, GetterAttrs, AttributeSet(), ParamAttrs())
^
/home/sergei/llvm-project/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def:610:17: note: expanded from macro '__OMP_RTL_ATTRS'
  OMP_RTL_ATTRS(OMPRTL_##Name, FnAttrSet, RetAttrSet, ArgAttrSets)
                ^
<scratch space>:78:1: note: expanded from here
OMPRTL_omp_get_wtime
^
1 error generated.

Thanks for the info. I reverted it.

[tianshilei1992]

fix build issue

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 __OMP_TYPE(Int8)
 __OMP_TYPE(Int16)
 __OMP_TYPE(Int32)
 __OMP_TYPE(Int64)
 __OMP_TYPE(Int8Ptr)
 __OMP_TYPE(Int16Ptr)
 __OMP_TYPE(Int32Ptr)
 __OMP_TYPE(Int64Ptr)
+__OMP_TYPE(Double)
 
 OMP_TYPE(SizeTy, M.getDataLayout().getIntPtrType(Ctx))
 
 #define __OMP_PTR_TYPE(NAME, BASE) OMP_TYPE(NAME, BASE->getPointerTo())
 
 __OMP_PTR_TYPE(VoidPtr, Int8)
 __OMP_PTR_TYPE(VoidPtrPtr, VoidPtr)
 __OMP_PTR_TYPE(VoidPtrPtrPtr, VoidPtrPtr)

 __OMP_RTL(omp_in_final, false, Int32, )
 __OMP_RTL(omp_get_proc_bind, false, Int32, )
 __OMP_RTL(omp_get_num_places, false, Int32, )
 __OMP_RTL(omp_get_num_procs, false, Int32, )
 __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)
 __OMP_RTL(omp_set_nested, false, Void, Int32)
 __OMP_RTL(omp_set_schedule, false, Void, Int32, Int32)
 __OMP_RTL(omp_set_max_active_levels, false, Void, Int32)
 
 __OMP_RTL(__kmpc_master, false, Int32, IdentPtr, Int32)

 __OMP_RTL_ATTRS(omp_get_place_proc_ids, GetterArgWriteAttrs, AttributeSet(),
                 ParamAttrs(AttributeSet(), AttributeSet(EnumAttr(NoCapture),
                                                         EnumAttr(WriteOnly))))
 __OMP_RTL_ATTRS(omp_get_place_num, GetterAttrs, AttributeSet(), ParamAttrs())
 __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())
 __OMP_RTL_ATTRS(omp_set_nested, SetterAttrs, AttributeSet(), ParamAttrs())
 __OMP_RTL_ATTRS(omp_set_schedule, SetterAttrs, AttributeSet(), ParamAttrs())
 __OMP_RTL_ATTRS(omp_set_max_active_levels, SetterAttrs, AttributeSet(),
                 ParamAttrs())
 
__OMP_RTL_ATTRS(__kmpc_doacross_post, BarrierAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), ReadOnlyPtrAttrs))
 __OMP_RTL_ATTRS(__kmpc_doacross_wait, BarrierAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), ReadOnlyPtrAttrs))
 __OMP_RTL_ATTRS(__kmpc_doacross_fini, BarrierAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs))
 
 __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))
 
 __OMP_RTL_ATTRS(__kmpc_alloc, DefaultAttrs, ReturnPtrAttrs, ParamAttrs())
 __OMP_RTL_ATTRS(__kmpc_aligned_alloc, DefaultAttrs, ReturnPtrAttrs,
                 ParamAttrs())
 __OMP_RTL_ATTRS(__kmpc_free, AllocAttrs, AttributeSet(), ParamAttrs())

openmp/libomptarget/DeviceRTL/src/Misc.cpp

 
 double getWTick() {
   // Timer precision is 1ns
   return ((double)1E-9);
 }
 
 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();
 }
 
 #pragma omp end declare variant
 
 } // namespace impl
 } // namespace _OMP
 

openmp/libomptarget/test/lit.cfg

 # - for valid targets that do not exist in the system, return false, so that the
 #   same test can be used for different targets.
 
 # Scan all the valid targets.
 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))
         config.substitutions.append(("%libomptarget-compilexx-and-run-generic",
             "%libomptarget-compilexx-and-run-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-compile-and-run-generic",
             "%libomptarget-compile-and-run-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-compilexx-generic",
             "%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",
             "%libomptarget-run-fail-" + libomptarget_target))
         config.substitutions.append(("%clangxx-generic",
             "%clangxx-" + libomptarget_target))
         config.substitutions.append(("%clang-generic",
             "%clang-" + libomptarget_target))
if libomptarget_target == config.libomptarget_current_target:
             "%libomptarget-compile-" + libomptarget_target + " && " + \
             "%libomptarget-run-" + libomptarget_target))
         config.substitutions.append(("%libomptarget-compilexx-" + \
             libomptarget_target, \
             "%clangxx-" + libomptarget_target + " %s -o %t"))
         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"))
         config.substitutions.append(("%libomptarget-run-fail-" + \
             libomptarget_target, \
             "%not --crash %t"))
         config.substitutions.append(("%clangxx-" + libomptarget_target, \
                                      "%clangxx %openmp_flags %cuda_flags %flags -fopenmp-targets=" +\
else:
             libomptarget_target, \
             "echo ignored-command"))
         config.substitutions.append(("%libomptarget-compilexx-" + \
             libomptarget_target, \
             "echo ignored-command"))
         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"))
         config.substitutions.append(("%libomptarget-run-fail-" + \
             libomptarget_target, \
             "echo ignored-command"))
         config.substitutions.append(("%clang-" + libomptarget_target, \
             "echo ignored-command"))

openmp/libomptarget/test/offloading/wtime.c

+// 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]+