Add CudaSpace

CMakeExt/CompilerFlags.cmake

@@ -147,7 +147,7 @@ endif()
 set (CXX_GDB_FLAG "-g"
     CACHE STRING "C++ compiler (clang++) debug symbols flag")
 if(OPENMP_FOUND)
-    set (CXX_OMP_FLAG ${OpenMP_CXX_FLAGS})
+        #set (CXX_OMP_FLAG ${OpenMP_CXX_FLAGS})
 endif()
 
 

CMakeExt/Cuda.cmake

@@ -0,0 +1 @@
+enable_language(CUDA)

CMakeLists.txt

@@ -83,6 +83,8 @@ option(ENABLE_HDF5
        "Specify whether HDF5 features are enabled" on)
 option(ENABLE_MEMKIND
        "Specify whether Memkind features are enabled" on)
+option(ENABLE_CUDA
+       "Specify whether CUDA library parts should be enabled" off)
 option(ENABLE_NASTYMPI
        "Specify whether the NastyMPI proxy should be enabled" off)
 
@@ -124,6 +126,9 @@ include(${CMAKE_SOURCE_DIR}/CMakeExt/IPM.cmake)
 include(${CMAKE_SOURCE_DIR}/CMakeExt/PLASMA.cmake)
 include(${CMAKE_SOURCE_DIR}/CMakeExt/HDF5.cmake)
 include(${CMAKE_SOURCE_DIR}/CMakeExt/Memkind.cmake)
+if (ENABLE_CUDA)
+  include(${CMAKE_SOURCE_DIR}/CMakeExt/Cuda.cmake)
+endif()
 
 if (ENABLE_MKL)
   include(${CMAKE_SOURCE_DIR}/CMakeExt/MKL.cmake)

dash/CMakeLists.txt

@@ -43,9 +43,14 @@ set(ENABLE_MEMKIND ${ENABLE_MEMKIND}
     PARENT_SCOPE)
 
 
+if(ENABLE_CUDA)
+  set(DASH_CUDA_LIBRARY_GLOB "src/*.cu")
+  set(DASH_CUDA_TEST_GLOB "test/*.cu")
+endif()
+
 # Source- and header files to be compiled (OBJ):
 file(GLOB_RECURSE DASH_LIBRARY_SOURCES
-     "src/*.c" "src/*.h" "src/*.cc" "src/*.cpp" )
+        "src/*.c" "src/*.h" "src/*.cc" "src/*.cpp" ${DASH_CUDA_LIBRARY_GLOB})
 file(GLOB_RECURSE DASH_LIBRARY_HEADERS
      "include/*.h")
 
@@ -54,7 +59,7 @@ file(GLOB_RECURSE DASH_TEST_SOURCES
 
 foreach(TESTCASE ${TESTCASES_LIST})
 file(GLOB_RECURSE DASH_TEST_SOURCES
-  ${DASH_TEST_SOURCES} "test/${TESTCASE}.h" "test/${TESTCASE}.cc")
+        ${DASH_TEST_SOURCES} "test/${TESTCASE}.h" "test/${TESTCASE}.cc" ${DASH_CUDA_TEST_GLOB})
 endforeach()
 
 # Directories containing the implementation of the library (-I):
@@ -485,7 +490,7 @@ if (BUILD_TESTS)
       set_target_properties(
         ${DASH_TEST} PROPERTIES
         COMPILE_FLAGS
-        "${VARIANT_ADDITIONAL_COMPILE_FLAGS} -Wno-unused -Wno-sign-compare"
+        "${VARIANT_ADDITIONAL_COMPILE_FLAGS}"
       )
       set_target_properties(
         ${DASH_TEST} PROPERTIES

dash/include/dash/memory/CudaSpace.h

@@ -0,0 +1,39 @@
+#ifndef DASH__MEMORY__CUDA_SPACE_H__INCLUDED
+#define DASH__MEMORY__CUDA_SPACE_H__INCLUDED
+
+
+#include <dash/memory/MemorySpaceBase.h>
+
+// This memory space can only be used with nvcc
+#ifdef __CUDACC__
+namespace dash {
+
+class CudaSpace
+  : public dash::MemorySpace<memory_domain_local, memory_space_cuda_tag> {
+
+public:
+  using void_pointer       = void*;
+  using const_void_pointer = const void*;
+
+public:
+  CudaSpace()                      = default;
+  CudaSpace(CudaSpace const& other) = default;
+  CudaSpace(CudaSpace&& other)      = default;
+  CudaSpace& operator=(CudaSpace const& other) = default;
+  CudaSpace& operator=(CudaSpace&& other) = default;
+  ~CudaSpace()
+  {
+  }
+
+protected:
+  void* do_allocate(size_t bytes, size_t alignment) override;
+  void  do_deallocate(void* p, size_t bytes, size_t alignment) override;
+  bool  do_is_equal(std::pmr::memory_resource const& other) const
+      noexcept override;
+};
+
+}  // namespace dash
+
+#endif
+
+#endif

dash/include/dash/memory/MemorySpace.h

@@ -1,6 +1,7 @@
 #ifndef DASH__MEMORY__MEMORY_SPACE_H__INCLUDED
 #define DASH__MEMORY__MEMORY_SPACE_H__INCLUDED
 
+#include <dash/memory/CudaSpace.h>
 #include <dash/memory/HBWSpace.h>
 #include <dash/memory/HostSpace.h>
 
@@ -44,6 +45,10 @@ template <>
 MemorySpace<memory_domain_local, memory_space_hbw_tag>*
 get_default_memory_space<memory_domain_local, memory_space_hbw_tag>();
 
+template <>
+MemorySpace<memory_domain_local, memory_space_cuda_tag>*
+get_default_memory_space<memory_domain_local, memory_space_cuda_tag>();
+
 template <>
 MemorySpace<memory_domain_global, memory_space_host_tag>*
 get_default_memory_space<memory_domain_global, memory_space_host_tag>();

dash/src/memory/CudaSpace.cu

@@ -0,0 +1,55 @@
+#include <dash/Exception.h>
+#include <dash/internal/Logging.h>
+#include <dash/memory/MemorySpaceBase.h>
+#include <dash/memory/MemorySpace.h>
+#include <dash/memory/CudaSpace.h>
+#include <new>
+#include <assert.h>
+
+void* dash::CudaSpace::do_allocate(size_t bytes, size_t alignment)
+{
+  // Cuda guarantees alignment at 256 bytes but not more.
+  assert(alignment <= 256);
+  void_pointer ptr;
+  auto ret = cudaMallocManaged(&ptr, bytes) ;
+  if (ret != cudaSuccess) {
+    DASH_LOG_ERROR(
+        "CudaPace.do_allocate",
+        "Cannot allocate managed memory",
+        bytes,
+        alignment);
+    DASH_LOG_ERROR("CudaPace.do_allocate", cudaGetErrorString(ret));
+
+    std::bad_alloc();
+  }
+  return ptr;
+}
+
+void dash::CudaSpace::do_deallocate(void* p, size_t bytes, size_t alignment)
+{
+  if (cudaFree(p) != cudaSuccess) {
+    DASH_LOG_ERROR(
+        "CudaPace.do_deallocate",
+        "Cannot deallocate managed memory",
+        p,
+        bytes,
+        alignment);
+  }
+}
+
+bool dash::CudaSpace::do_is_equal(
+    std::pmr::memory_resource const& other) const noexcept
+{
+  const CudaSpace* other_p = dynamic_cast<const CudaSpace*>(&other);
+
+  return nullptr != other_p;
+}
+
+template <>
+dash::MemorySpace<dash::memory_domain_local, dash::memory_space_cuda_tag>*
+dash::get_default_memory_space<dash::memory_domain_local, dash::memory_space_cuda_tag>()
+{
+  static dash::CudaSpace cuda_space_singleton;
+  return &cuda_space_singleton;
+}
+

dash/test/memory/CudaSpaceTest.cu

@@ -0,0 +1,29 @@
+#include "GlobStaticMemTest.h"
+
+#include <dash/GlobRef.h>
+#include <dash/allocator/GlobalAllocator.h>
+#include <dash/memory/UniquePtr.h>
+#include <dash/memory/CudaSpace.h>
+
+TEST_F(GlobStaticMemTest, CudaSpace)
+{
+  using value_t     = int;
+  using memory_t    = dash::GlobStaticMem<dash::CudaSpace>;
+  using allocator_t = dash::GlobalAllocator<value_t, memory_t>;
+
+  memory_t    memory{dash::Team::All()};
+  allocator_t alloc{&memory};
+
+  auto const gptr = alloc.allocate(10);
+
+  EXPECT_TRUE_U(gptr);
+
+  auto *lbegin = dash::local_begin(gptr, dash::Team::All().myid());
+  auto *lend   = std::next(lbegin, 10);
+
+  std::uninitialized_fill(lbegin, lend, dash::myid());
+
+  ASSERT_EQ_U(*lbegin, dash::myid());
+
+  alloc.deallocate(gptr, 10);
+}