Skip to content
This repository was archived by the owner on Mar 20, 2023. It is now read-only.

Commit 43b595a

Browse files
committed
Do not use __managed__ in GPU builds. Add utils.
This fixes CPU execution of GPU builds on machines that do not have GPUs, which were previously segfaulting due to the use of the __managed__ keyword. Now the handling of Random123 state is more explicit for host/device. Also add a pair of helper functions coreneuron::[de]allocate_unified() that wrap cudaMallocManaged in GPU builds if --gpu was passed at runtime and fall back to new/delete otherwise, and a method coreneuron::unified_memory_enabled() that queries whether this condition is met. Additionally add a C++ allocator template coreneuron::unified_allocator<T> that wraps these functions, a templated coreneuron::alloc_deleter<T> for use with std::unique_ptr<T, D>, and a helper coreneuron::allocate_unique(...). Cleanup Random123 code by dropping an unused nrnran123_mutconstruct method. Tweak compilation/CMake scripts to remove libcudacoreneuron.a and instead build CUDA sources inside libcoreneuron.a. This sidesteps circular dependency issues that would otherwise be introduced by this commit. Modify CMake so `clang-format` target formats CUDA (.cu) files too.
1 parent 170a0bb commit 43b595a

8 files changed

Lines changed: 274 additions & 71 deletions

File tree

CMake/OpenAccHelper.cmake

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
# =============================================================================
1010
if(CORENRN_ENABLE_GPU)
1111
# Enable cudaProfiler{Start,Stop}() behind the Instrumentor::phase... APIs
12-
add_compile_definitions(CORENEURON_CUDA_PROFILING)
12+
add_compile_definitions(CORENEURON_CUDA_PROFILING CORENEURON_ENABLE_GPU)
1313
# cuda unified memory support
1414
if(CORENRN_ENABLE_CUDA_UNIFIED_MEMORY)
1515
add_compile_definitions(CORENEURON_UNIFIED_MEMORY)
@@ -49,8 +49,8 @@ if(CORENRN_ENABLE_GPU)
4949
endif()
5050
# -acc enables OpenACC support, -cuda links CUDA libraries and (very importantly!) seems to be
5151
# required to make the NVHPC compiler do the device code linking. Otherwise the explicit CUDA
52-
# device code (.cu files in libcudacoreneuron) has to be linked in a separate, earlier, step,
53-
# which apparently causes problems with interoperability with OpenACC. Passing -cuda to nvc++ when
52+
# device code (.cu files in libcoreneuron) has to be linked in a separate, earlier, step, which
53+
# apparently causes problems with interoperability with OpenACC. Passing -cuda to nvc++ when
5454
# compiling (as opposed to linking) seems to enable CUDA C++ support, which has other consequences
5555
# due to e.g. __CUDACC__ being defined. See https://github.com/BlueBrain/CoreNeuron/issues/607 for
5656
# more information about this. -gpu=cudaX.Y ensures that OpenACC code is compiled with the same
@@ -80,7 +80,7 @@ if(CORENRN_ENABLE_GPU)
8080
GLOBAL
8181
PROPERTY
8282
CORENEURON_LIB_LINK_FLAGS
83-
"${NVHPC_ACC_COMP_FLAGS} ${NVHPC_ACC_LINK_FLAGS} -rdynamic -lrt -Wl,--whole-archive -L${CMAKE_HOST_SYSTEM_PROCESSOR} -lcorenrnmech -L${CMAKE_INSTALL_PREFIX}/lib -lcoreneuron -lcudacoreneuron -Wl,--no-whole-archive"
83+
"${NVHPC_ACC_COMP_FLAGS} ${NVHPC_ACC_LINK_FLAGS} -rdynamic -lrt -Wl,--whole-archive -L${CMAKE_HOST_SYSTEM_PROCESSOR} -lcorenrnmech -L${CMAKE_INSTALL_PREFIX}/lib -lcoreneuron -Wl,--no-whole-archive"
8484
)
8585
else()
8686
set_property(GLOBAL PROPERTY CORENEURON_LIB_LINK_FLAGS

CMakeLists.txt

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,11 @@ list(APPEND CMAKE_MODULE_PATH ${CORENEURON_PROJECT_SOURCE_DIR}/CMake
5353
# =============================================================================
5454
set(CODING_CONV_PREFIX "CORENRN")
5555
set(CORENRN_3RDPARTY_DIR "external")
56+
# Adds .cu with respect to the current default in hpc-coding-conventions, and drops various patterns
57+
# that don't match anything in CoreNEURON. Note there's only one .c file left!
58+
set(CORENRN_ClangFormat_FILES_RE
59+
"^.*\\\\.cu?$$" "^.*\\\\.[chi]pp$$"
60+
CACHE STRING "List of regular expressions matching C/C++ filenames" FORCE)
5661
set(CORENRN_ClangFormat_EXCLUDES_RE
5762
"${CORENRN_PROJECT_SOURCE_DIR}/external/.*$$"
5863
CACHE STRING "list of regular expressions to exclude C/C++ files from formatting" FORCE)

coreneuron/CMakeLists.txt

Lines changed: 8 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,7 @@ if(CORENRN_ENABLE_GPU)
108108
# ~~~
109109
# artificial cells and some other cpp files (using Random123) should be compiled
110110
# without OpenACC to avoid use of GPU Random123 streams
111+
# OL210813: this shouldn't be needed anymore, but it may have a small performance benefit
111112
# ~~~
112113
set(OPENACC_EXCLUDED_FILES
113114
${CMAKE_CURRENT_BINARY_DIR}/netstim.cpp
@@ -118,24 +119,19 @@ if(CORENRN_ENABLE_GPU)
118119
${CMAKE_CURRENT_SOURCE_DIR}/io/setup_fornetcon.cpp
119120
${CMAKE_CURRENT_SOURCE_DIR}/io/corenrn_data_return.cpp
120121
${CMAKE_CURRENT_SOURCE_DIR}/io/global_vars.cpp)
121-
122+
# OL210831: why?
122123
set_source_files_properties(${DIMPLIC_CODE_FILE} ${NMODL_INBUILT_MOD_OUTPUTS}
123124
PROPERTIES COMPILE_FLAGS "")
124125

125126
set_source_files_properties(${OPENACC_EXCLUDED_FILES} PROPERTIES COMPILE_FLAGS
126127
"-DDISABLE_OPENACC")
127-
128-
add_library(cudacoreneuron STATIC ${CORENEURON_CUDA_FILES})
129-
set_target_properties(cudacoreneuron PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
130-
set(link_cudacoreneuron cudacoreneuron)
131128
# nrnran123.cpp is a symlink to nrnran123.cu, in GPU builds we compile this as CUDA code (so we
132129
# want to remove the .cpp here), while in non-GPU builds we compile it as plain C++. Unfortunately
133130
# CMake <v3.20 does not pass explicit -x <lang> options based on the LANGUAGE property
134131
# (https://cmake.org/cmake/help/latest/policy/CMP0119.html), so using a single .cu file and
135132
# setting LANGUAGE=CXX in non-GPU builds does not work.
136133
list(REMOVE_ITEM CORENEURON_CODE_FILES "${CMAKE_CURRENT_SOURCE_DIR}/utils/randoms/nrnran123.cpp")
137-
else()
138-
set(link_cudacoreneuron "")
134+
list(APPEND CORENEURON_CODE_FILES ${CORENEURON_CUDA_FILES})
139135
endif()
140136

141137
# =============================================================================
@@ -171,31 +167,23 @@ add_library(
171167
${CORENEURON_CODE_FILES} ${cudacorenrn_objs} ${NMODL_INBUILT_MOD_OUTPUTS})
172168
# Prevent CMake from running a device code link step when assembling libcoreneuron.a in GPU builds.
173169
# The device code linking needs to be deferred to the final step, where it is done by `nvc++ -cuda`.
174-
# OL210811: it's not clear why this is needed, given that coreneuron is a static library in GPU
175-
# builds.
176-
set_target_properties(coreneuron PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS OFF)
170+
set_target_properties(coreneuron PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
177171

178172
# need to have _kinderiv.h for mod2c generated files and nrnivmodl-core and nmodl building
179173
add_dependencies(coreneuron kin_deriv_header nrnivmodl-core)
180174

181175
# scopmath is created separately for nrnivmodl-core workflow
182176
add_library(scopmath STATIC ${CORENEURON_HEADER_FILES} ${SCOPMATH_CODE_FILES})
183177

184-
target_link_libraries(
185-
coreneuron
186-
${reportinglib_LIBRARY}
187-
${sonatareport_LIBRARY}
188-
${link_cudacoreneuron}
189-
${CALIPER_LIB}
190-
${likwid_LIBRARIES}
191-
${MPI_C_LIBRARIES})
178+
target_link_libraries(coreneuron ${reportinglib_LIBRARY} ${sonatareport_LIBRARY} ${CALIPER_LIB}
179+
${likwid_LIBRARIES} ${MPI_C_LIBRARIES})
192180
target_include_directories(coreneuron SYSTEM
193181
PRIVATE ${CORENEURON_PROJECT_SOURCE_DIR}/external/Random123/include)
194182
target_include_directories(coreneuron SYSTEM
195183
PRIVATE ${CORENEURON_PROJECT_SOURCE_DIR}/external/CLI11/include)
196184

197185
set_target_properties(
198-
coreneuron scopmath ${link_cudacoreneuron}
186+
coreneuron scopmath
199187
PROPERTIES ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib
200188
LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib
201189
POSITION_INDEPENDENT_CODE ON)
@@ -253,7 +241,7 @@ file(COPY apps/coreneuron.cpp DESTINATION ${CMAKE_BINARY_DIR}/share/coreneuron)
253241

254242
# coreneuron main libraries
255243
install(
256-
TARGETS coreneuron ${link_cudacoreneuron}
244+
TARGETS coreneuron
257245
EXPORT coreneuron
258246
LIBRARY DESTINATION lib
259247
ARCHIVE DESTINATION lib

coreneuron/apps/main1.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -140,9 +140,6 @@ void nrn_init_and_load_data(int argc,
140140
// initialise default coreneuron parameters
141141
initnrn();
142142

143-
// create mutex for nrn123, protect instance_count_
144-
nrnran123_mutconstruct();
145-
146143
// set global variables
147144
// precedence is: set by user, globals.dat, 34.0
148145
celsius = corenrn_param.celsius;

coreneuron/utils/memory.cpp

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
/*
2+
# =============================================================================
3+
# Copyright (c) 2016 - 2021 Blue Brain Project/EPFL
4+
#
5+
# See top-level LICENSE file for details.
6+
# =============================================================================.
7+
*/
8+
#include "coreneuron/apps/corenrn_parameters.hpp"
9+
#include "coreneuron/utils/memory.h"
10+
11+
#ifdef CORENEURON_ENABLE_GPU
12+
#include <cuda_runtime_api.h>
13+
#endif
14+
15+
#include <cassert>
16+
17+
namespace coreneuron {
18+
bool unified_memory_enabled() {
19+
#ifdef CORENEURON_ENABLE_GPU
20+
return corenrn_param.gpu;
21+
#endif
22+
return false;
23+
}
24+
25+
void* allocate_unified(std::size_t num_bytes) {
26+
#ifdef CORENEURON_ENABLE_GPU
27+
// The build supports GPU execution, check if --gpu was passed to actually
28+
// enable it. We should not call CUDA APIs in GPU builds if --gpu was not passed.
29+
if (corenrn_param.gpu) {
30+
// Allocate managed/unified memory.
31+
void* ptr{nullptr};
32+
auto const code = cudaMallocManaged(&ptr, num_bytes);
33+
assert(code == cudaSuccess);
34+
return ptr;
35+
}
36+
#endif
37+
// Either the build does not have GPU support or --gpu was not passed.
38+
// Allocate using standard operator new.
39+
// When we have C++17 support then propagate `alignment` here.
40+
return ::operator new(num_bytes);
41+
}
42+
43+
void deallocate_unified(void* ptr, std::size_t num_bytes) {
44+
// See comments in allocate_unified to understand the different branches.
45+
#ifdef CORENEURON_ENABLE_GPU
46+
if (corenrn_param.gpu) {
47+
// Don't assert success because it will fail if called at application
48+
// teardown, e.g. by a global std::unique_ptr destructor...this is not very nice.
49+
cudaFree(ptr);
50+
return;
51+
}
52+
#endif
53+
::operator delete(ptr, num_bytes);
54+
}
55+
} // namespace coreneuron

coreneuron/utils/memory.h

Lines changed: 93 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111

1212
#include <cstdint>
1313
#include <cstring>
14+
#include <memory>
1415

1516
#include "coreneuron/utils/nrn_assert.h"
1617
#include "coreneuron/nrniv/nrniv_decl.h"
@@ -21,6 +22,98 @@
2122
#define NRN_SOA_BYTE_ALIGN (8 * sizeof(double))
2223
#endif
2324

25+
namespace coreneuron {
26+
/** @brief Check if allocate_unified will return a unified memory address.
27+
*
28+
* If false, [de]allocate_unified simply forward to new/delete. It is
29+
* convenient to include this method here to avoid having to access
30+
* corenrn_param directly.
31+
*/
32+
bool unified_memory_enabled();
33+
34+
/** @brief Allocate unified memory in GPU builds iff GPU enabled, otherwise new
35+
*/
36+
void* allocate_unified(std::size_t num_bytes);
37+
38+
/** @brief Deallocate memory allocated by `allocate_unified`.
39+
*/
40+
void deallocate_unified(void* ptr, std::size_t num_bytes);
41+
42+
/** @brief C++ allocator that uses [de]allocate_unified.
43+
*/
44+
template <typename T>
45+
struct unified_allocator {
46+
using value_type = T;
47+
48+
unified_allocator() = default;
49+
50+
template <typename U>
51+
unified_allocator(unified_allocator<U> const&) noexcept {}
52+
53+
value_type* allocate(std::size_t n) {
54+
return static_cast<value_type*>(allocate_unified(n * sizeof(value_type)));
55+
}
56+
57+
void deallocate(value_type* p, std::size_t n) noexcept {
58+
deallocate_unified(p, n * sizeof(value_type));
59+
}
60+
};
61+
62+
template <typename T, typename U>
63+
bool operator==(unified_allocator<T> const&, unified_allocator<U> const&) noexcept {
64+
return true;
65+
}
66+
67+
template <typename T, typename U>
68+
bool operator!=(unified_allocator<T> const& x, unified_allocator<U> const& y) noexcept {
69+
return !(x == y);
70+
}
71+
72+
/** @brief Allocator-aware deleter for use with std::unique_ptr.
73+
*
74+
* This is copied from https://stackoverflow.com/a/23132307. See also
75+
* http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2017/p0316r0.html,
76+
* http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2020/p0211r3.html, and
77+
* boost::allocate_unique<...>.
78+
* Hopefully std::allocate_unique will be included in C++23.
79+
*/
80+
template <typename Alloc>
81+
struct alloc_deleter {
82+
alloc_deleter() = default; // OL210813 addition
83+
alloc_deleter(const Alloc& a)
84+
: a(a) {}
85+
86+
typedef typename std::allocator_traits<Alloc>::pointer pointer;
87+
88+
void operator()(pointer p) const {
89+
Alloc aa(a);
90+
std::allocator_traits<Alloc>::destroy(aa, std::addressof(*p));
91+
std::allocator_traits<Alloc>::deallocate(aa, p, 1);
92+
}
93+
94+
private:
95+
Alloc a;
96+
};
97+
98+
template <typename T, typename Alloc, typename... Args>
99+
auto allocate_unique(const Alloc& alloc, Args&&... args) {
100+
using AT = std::allocator_traits<Alloc>;
101+
static_assert(std::is_same<typename AT::value_type, std::remove_cv_t<T>>{}(),
102+
"Allocator has the wrong value_type");
103+
104+
Alloc a(alloc);
105+
auto p = AT::allocate(a, 1);
106+
try {
107+
AT::construct(a, std::addressof(*p), std::forward<Args>(args)...);
108+
using D = alloc_deleter<Alloc>;
109+
return std::unique_ptr<T, D>(p, D(a));
110+
} catch (...) {
111+
AT::deallocate(a, p, 1);
112+
throw;
113+
}
114+
}
115+
} // namespace coreneuron
116+
24117
/// for gpu builds with unified memory support
25118
/// OL210812: why do we include __CUDACC__ here?
26119
#if (defined(__CUDACC__) || defined(CORENEURON_UNIFIED_MEMORY))

0 commit comments

Comments
 (0)