Omptarget Updates #3169

crtrott · 2020-07-08T16:08:35Z

Significant Update to OpenMPTarget

requires C++17
works on Intel GPUs
Compilation and tests fixed/worked around so that full build/test passes.
Fixed some stuff for reference counting.

dalg24 · 2020-07-08T16:42:02Z

cmake/kokkos_compiler_id.cmake

@@ -63,6 +63,16 @@ IF(KOKKOS_CXX_COMPILER_ID STREQUAL Clang)
  IF (INTERNAL_HAVE_CRAY_COMPILER) #not actually Clang
    SET(KOKKOS_CLANG_IS_CRAY TRUE)
  ENDIF()
+  # The clang based Intel compiler reports as Clang to most versions of CMake


We are going towards all vendor compiler being identified by CMake as Clang...

dalg24 · 2020-07-08T16:44:54Z

cmake/kokkos_enable_devices.cmake

@@ -75,6 +78,7 @@ SET(ClangOpenMPFlag -fopenmp=libomp)

  COMPILER_SPECIFIC_FLAGS(
    Clang      ${ClangOpenMPFlag} -Wno-openmp-mapping
+    IntelClang -fiopenmp -Wno-openmp-mapping


Why are you not using ClangOpenMPFlag if you bother setting it earlier?

its not fopenmp=libomp, its fiopenmp with an "i", because its not really clang ….

Right but you define a variable line 48 and never use it

probably because it wasn't used there before some recent update, which came into this branch via a rebase

core/perf_test/CMakeLists.txt

core/src/OpenMPTarget/Kokkos_OpenMPTarget_Parallel.hpp

core/src/OpenMPTarget/Kokkos_OpenMPTarget_Parallel_MDRange.hpp

core/src/impl/Kokkos_ClockTic.hpp

core/src/impl/Kokkos_SharedAlloc.hpp

dhollman

mostly looks fine

dhollman · 2020-07-08T23:34:03Z

containers/src/Kokkos_ScatterView.hpp

-                    Kokkos::Experimental::ScatterNonAtomic>
-    : Sum<ValueType, DeviceType> {
+                    Kokkos::Experimental::ScatterNonAtomic> {
+  ValueType& value;


Did you mean for this to be public? If so, putting public: two lines below makes it look like you didn't intend to do that.

dhollman · 2020-07-08T23:34:28Z

containers/src/Kokkos_ScatterView.hpp

-                    Kokkos::Experimental::ScatterAtomic>
-    : Sum<ValueType, DeviceType> {
+                    Kokkos::Experimental::ScatterAtomic> {
+  ValueType& value;


dhollman · 2020-07-08T23:38:23Z

containers/src/Kokkos_ScatterView.hpp

  }
+  KOKKOS_FORCEINLINE_FUNCTION void operator--() { update(ValueType(-1)); }


You probably should have a reason for not using {} initialization with numeric literals to avoid accidental narrowing.

Wait, what if ValueType is unsigned? It should still be OK, right?

Wait, what if ValueType is unsigned? It should still be OK, right?

It will do the right thing

dhollman · 2020-07-08T23:39:29Z

containers/src/Kokkos_ScatterView.hpp

  KOKKOS_FORCEINLINE_FUNCTION void operator+=(ValueType const& rhs) {
-    this->join(this->reference(), rhs);
-  }
-  KOKKOS_FORCEINLINE_FUNCTION void operator++() {
-    this->join(this->reference(), 1);
-  }
-  KOKKOS_FORCEINLINE_FUNCTION void operator++(int) {
-    this->join(this->reference(), 1);
+    update(rhs);
  }
+  KOKKOS_FORCEINLINE_FUNCTION void operator++() { update(1); }
+  KOKKOS_FORCEINLINE_FUNCTION void operator++(int) { update(1); }
  KOKKOS_FORCEINLINE_FUNCTION void operator-=(ValueType const& rhs) {
-    this->join(this->reference(), ValueType(-rhs));
-  }
-  KOKKOS_FORCEINLINE_FUNCTION void operator--() {
-    this->join(this->reference(), ValueType(-1));
-  }
-  KOKKOS_FORCEINLINE_FUNCTION void operator--(int) {
-    this->join(this->reference(), ValueType(-1));
+    update(ValueType(-rhs));
  }
+  KOKKOS_FORCEINLINE_FUNCTION void operator--() { update(ValueType(-1)); }
+  KOKKOS_FORCEINLINE_FUNCTION void operator--(int) { update(ValueType(-1)); }


Why do these return void?

they did before. And yeah we probably should fix that as general ScatterView overhaul. Also #3162

dhollman · 2020-07-08T23:48:13Z

core/src/impl/Kokkos_Atomic_Generic_Secondary.hpp

+template <typename T>
+KOKKOS_INLINE_FUNCTION void atomic_mul(volatile T* const dest, const T val) {
+  (void)atomic_fetch_mul(dest, val);
+}
+
+template <typename T>
+KOKKOS_INLINE_FUNCTION void atomic_div(volatile T* const dest, const T val) {
+  (void)atomic_fetch_div(dest, val);
+}
+


What does this have to do with OMPTarget?

nothing which is why its also in this PR: #3162

dhollman · 2020-07-08T23:53:26Z

core/src/impl/Kokkos_SharedAlloc.hpp

+#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ENABLE_CXX17)
+#define KOKKOS_IMPL_IF_ON_HOST if constexpr (omp_is_initial_device() == true)
+#else
+#define KOKKOS_IMPL_IF_ON_HOST if (true)


Comment explaining why this is sufficient on pre-C++17, maybe?

it is not, but OpenMPTarget now requires C++17, I can remove the KOKKOS_ENABLE_CXX17 here.

dhollman · 2020-07-08T23:55:59Z

core/unit_test/TestReduce.hpp

-  Kokkos::parallel_reduce(nw, functor_type(nw), result1_v, result2,
+  Kokkos::parallel_reduce("int_combined-reduce_mixed",
+                          Kokkos::RangePolicy<TEST_EXECSPACE>(0, nw),
+                          functor_type(nw), result1_v, result2,
                          Kokkos::Sum<int64_t, Kokkos::HostSpace>{result3_v});


Is this a driveby change on combined reducers? If so, should it be in a separate pull request?

its not necessarily a drive-by change. It is a change specifically necessary for OpenMPTarget since OpenMPTarget doesn't have something like UVM or HostPinned memory which was used previously for these tests.

masterleinad · 2020-07-10T18:03:00Z

The formatting is off and the CUDA-11.0-NVCC-C++17-RDC CI complains:

/var/jenkins/workspace/Kokkos/install/include/Cuda/Kokkos_Cuda_Parallel.hpp(471): error: calling a constexpr __host__ function("operator()") from a __device__ function("exec_range") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.
          detected during:
            instantiation of "std::enable_if<std::is_same<TagType, void>::value, void>::type Kokkos::Impl::ParallelFor<FunctorType, Kokkos::RangePolicy<Traits...>, Kokkos::Cuda>::exec_range<TagType>(Kokkos::Impl::ParallelFor<FunctorType, Kokkos::RangePolicy<Traits...>, Kokkos::Cuda>::Member) const [with FunctorType=lambda [](int)->void, Traits=<Kokkos::CudaUVMSpace::execution_space>, TagType=void]" 
(493): here
            instantiation of "void Kokkos::Impl::ParallelFor<FunctorType, Kokkos::RangePolicy<Traits...>, Kokkos::Cuda>::operator()() const [with FunctorType=lambda [](int)->void, Traits=<Kokkos::CudaUVMSpace::execution_space>]" 
/var/jenkins/workspace/Kokkos/install/include/Cuda/Kokkos_Cuda_KernelLaunch.hpp(124): here
            instantiation of "void Kokkos::Impl::cuda_parallel_launch_local_memory(DriverType) [with DriverType=Kokkos::Impl::ParallelFor<lambda [](int)->void, Kokkos::RangePolicy<Kokkos::CudaUVMSpace::execution_space>, Kokkos::CudaUVMSpace::execution_space>]" 
/var/jenkins/workspace/Kokkos/install/include/Cuda/Kokkos_Cuda_KernelLaunch.hpp(478): here
            instantiation of "cudaFuncAttributes Kokkos::Impl::CudaParallelLaunch<DriverType, Kokkos::LaunchBounds<0U, 0U>, Kokkos::Impl::Experimental::CudaLaunchMechanism::LocalMemory>::get_cuda_func_attributes() [with DriverType=Kokkos::Impl::ParallelFor<lambda [](int)->void, Kokkos::RangePolicy<Kokkos::CudaUVMSpace::execution_space>, Kokkos::CudaUVMSpace::execution_space>]" 
(501): here
            instantiation of "void Kokkos::Impl::ParallelFor<FunctorType, Kokkos::RangePolicy<Traits...>, Kokkos::Cuda>::execute() const [with FunctorType=lambda [](int)->void, Traits=<Kokkos::CudaUVMSpace::execution_space>]" 
/var/jenkins/workspace/Kokkos/install/include/Kokkos_Parallel.hpp(176): here
            instantiation of "void Kokkos::parallel_for(const ExecPolicy &, const FunctorType &, const std::__cxx11::string &, std::enable_if<Kokkos::is_execution_policy<ExecPolicy>::value, void>::type *) [with ExecPolicy=Kokkos::RangePolicy<Kokkos::CudaUVMSpace::execution_space>, FunctorType=lambda [](int)->void]" 
/var/jenkins/workspace/Kokkos/install/include/Kokkos_Parallel.hpp(218): here
            instantiation of "void Kokkos::parallel_for(const std::__cxx11::string &, const ExecPolicy &, const FunctorType &) [with ExecPolicy=Kokkos::RangePolicy<Kokkos::CudaUVMSpace::execution_space>, FunctorType=lambda [](int)->void]" 
/var/jenkins/workspace/Kokkos/core/unit_test/TestViewSubview.hpp(973): here
            instantiation of "void TestViewSubview::Impl::test_1d_assign_impl<Space,LayoutSub,Layout,LayoutOrg,MemTraits>() [with Space=Kokkos::CudaUVMSpace, LayoutSub=Kokkos::LayoutLeft, Layout=Kokkos::LayoutLeft, LayoutOrg=Kokkos::LayoutLeft, MemTraits=void]" 
/var/jenkins/workspace/Kokkos/core/unit_test/TestViewSubview.hpp(1734): here
            instantiation of "void TestViewSubview::test_1d_assign<Space,MemTraits>() [with Space=Kokkos::CudaUVMSpace, MemTraits=void]" 
/var/jenkins/workspace/Kokkos/core/unit_test/cuda/TestCuda_SubView_c01.cpp(51): here

This checks via a constexpr if condition whether you are on the device and if so doesn't use the allocation tracking.

Needed declare target pragmas on some static functions

Same as HIP some stuff is not implemented. Now make in the root build directory works.

…ompile Also applied clang-format

Also fix cmake formatting??

crtrott · 2020-07-11T00:35:09Z

I ignored the ScatterView comments since that came in via a different PR, and we agreed to do another overhaul on that.

dalg24 · 2020-07-11T04:06:05Z

core/perf_test/CMakeLists.txt

@@ -49,11 +49,19 @@ SET(SOURCES
  )

 IF(Kokkos_ENABLE_HIP)
-# FIXME requires TeamPolicy
+# FIXME HIP requires TeamPolicy


We use the FIXME_HIP token

Apparently you don't because it only said FIXME before ;-)

dalg24 · 2020-07-11T04:07:13Z

core/perf_test/CMakeLists.txt

+ENDIF()
+
+IF(Kokkos_ENABLE_OPENMPTARGET)
+# FIXME OPENMPTARGET requires TeamPolicy Reductions and Custom Reduction


You are mixing FIXME OPENMPTARGET with WORKAROUND OPENMPTARGET. The point was to have unique token that can be searched for.

I never used WORKAROUND in a cmake file, so I followed what was in the file.

dalg24 reviewed Jul 8, 2020

View reviewed changes

crtrott added the Blocks Promotion Overview issue for release-blocking bugs label Jul 8, 2020

dhollman approved these changes Jul 8, 2020

View reviewed changes

crtrott added 12 commits July 10, 2020 15:57

Make CMake recognize Intel Clang compiler

8bd8669

SharedAllocRecord: put in constexpr OpenMPTarget detection

6b3f408

This checks via a constexpr if condition whether you are on the device and if so doesn't use the allocation tracking.

OpenMPTarget: workaround for intel compiler

a6fb4c7

Needed declare target pragmas on some static functions

OpenMPTarget: adds workaround for MDRange on intel compiler

4d5849e

IntelClang: register as a compiler thingy

099b800

OpenMPTarget: fix memory_fence and clock_tick

8829bb3

OpenMPTarget: ALways use MDRange fallback: fixes clang/nvidia release

328e7b0

OpenMPTarget: remove some things from being compiled for perf tests

e4d2d18

Same as HIP some stuff is not implemented. Now make in the root build directory works.

Fix some tests to work in absence of a shared memory space

60a2077

OpenMPTarget: Require C++17

1143ca5

OpenMPTarget: make tests which don't work yet with OpenMPTarget not c…

973338b

…ompile Also applied clang-format

OpenMPTarget: Update CI

c8a00a5

crtrott force-pushed the omptarget-intel branch from 89a4c28 to c8a00a5 Compare July 10, 2020 23:05

Address review comments, about adding more comments

d3cb283

Also fix cmake formatting??

crtrott force-pushed the omptarget-intel branch from be6ed54 to d3cb283 Compare July 11, 2020 00:16

crtrott merged commit eb93176 into kokkos:develop Jul 11, 2020

crtrott deleted the omptarget-intel branch July 11, 2020 03:21

dalg24 reviewed Jul 11, 2020

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Omptarget Updates #3169

Omptarget Updates #3169

crtrott commented Jul 8, 2020

dalg24 Jul 8, 2020

dalg24 Jul 8, 2020

crtrott Jul 9, 2020

dalg24 Jul 11, 2020

crtrott Jul 13, 2020

dhollman left a comment

dhollman Jul 8, 2020

dhollman Jul 8, 2020

dhollman Jul 8, 2020

mhoemmen Jul 9, 2020

dalg24 Jul 9, 2020

dhollman Jul 8, 2020

crtrott Jul 9, 2020

dhollman Jul 8, 2020

crtrott Jul 9, 2020

dhollman Jul 8, 2020

crtrott Jul 9, 2020

dhollman Jul 8, 2020

crtrott Jul 9, 2020

masterleinad commented Jul 10, 2020

crtrott commented Jul 11, 2020

dalg24 Jul 11, 2020

crtrott Jul 11, 2020

dalg24 Jul 11, 2020

crtrott Jul 11, 2020

		}
		KOKKOS_FORCEINLINE_FUNCTION void operator--() { update(ValueType(-1)); }

Omptarget Updates #3169

Omptarget Updates #3169

Conversation

crtrott commented Jul 8, 2020

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dhollman left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

masterleinad commented Jul 10, 2020

crtrott commented Jul 11, 2020

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment