From 7ed7f34402d9cda94cf07e5c79a1514a570f92ee Mon Sep 17 00:00:00 2001 From: Pascal Jungblut Date: Wed, 15 May 2019 16:24:55 +0200 Subject: [PATCH] Add CudaSpace --- CMakeExt/CompilerFlags.cmake | 2 +- CMakeExt/Cuda.cmake | 1 + CMakeLists.txt | 5 +++ dash/CMakeLists.txt | 11 ++++-- dash/include/dash/memory/CudaSpace.h | 39 ++++++++++++++++++ dash/include/dash/memory/MemorySpace.h | 5 +++ dash/src/memory/CudaSpace.cu | 55 ++++++++++++++++++++++++++ dash/test/memory/CudaSpaceTest.cu | 29 ++++++++++++++ 8 files changed, 143 insertions(+), 4 deletions(-) create mode 100644 CMakeExt/Cuda.cmake create mode 100644 dash/include/dash/memory/CudaSpace.h create mode 100644 dash/src/memory/CudaSpace.cu create mode 100644 dash/test/memory/CudaSpaceTest.cu diff --git a/CMakeExt/CompilerFlags.cmake b/CMakeExt/CompilerFlags.cmake index 2ef087b00..13cf9ddc0 100644 --- a/CMakeExt/CompilerFlags.cmake +++ b/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() diff --git a/CMakeExt/Cuda.cmake b/CMakeExt/Cuda.cmake new file mode 100644 index 000000000..c5744a1ca --- /dev/null +++ b/CMakeExt/Cuda.cmake @@ -0,0 +1 @@ +enable_language(CUDA) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7434b4f29..d2916fdd9 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/dash/CMakeLists.txt b/dash/CMakeLists.txt index e86df9908..881e5319b 100644 --- a/dash/CMakeLists.txt +++ b/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 diff --git a/dash/include/dash/memory/CudaSpace.h b/dash/include/dash/memory/CudaSpace.h new file mode 100644 index 000000000..d686d2c15 --- /dev/null +++ b/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 + +// This memory space can only be used with nvcc +#ifdef __CUDACC__ +namespace dash { + +class CudaSpace + : public dash::MemorySpace { + +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 diff --git a/dash/include/dash/memory/MemorySpace.h b/dash/include/dash/memory/MemorySpace.h index f94c879ec..a88ac894e 100644 --- a/dash/include/dash/memory/MemorySpace.h +++ b/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 #include #include @@ -44,6 +45,10 @@ template <> MemorySpace* get_default_memory_space(); +template <> +MemorySpace* +get_default_memory_space(); + template <> MemorySpace* get_default_memory_space(); diff --git a/dash/src/memory/CudaSpace.cu b/dash/src/memory/CudaSpace.cu new file mode 100644 index 000000000..25c4f63bd --- /dev/null +++ b/dash/src/memory/CudaSpace.cu @@ -0,0 +1,55 @@ +#include +#include +#include +#include +#include +#include +#include + +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(&other); + + return nullptr != other_p; +} + +template <> +dash::MemorySpace* +dash::get_default_memory_space() +{ + static dash::CudaSpace cuda_space_singleton; + return &cuda_space_singleton; +} + diff --git a/dash/test/memory/CudaSpaceTest.cu b/dash/test/memory/CudaSpaceTest.cu new file mode 100644 index 000000000..4fd4afa5e --- /dev/null +++ b/dash/test/memory/CudaSpaceTest.cu @@ -0,0 +1,29 @@ +#include "GlobStaticMemTest.h" + +#include +#include +#include +#include + +TEST_F(GlobStaticMemTest, CudaSpace) +{ + using value_t = int; + using memory_t = dash::GlobStaticMem; + using allocator_t = dash::GlobalAllocator; + + 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); +}