Skip to content

Commit a778554

Browse files
author
Johannes Lang
committed
Compatibility with CUDA 12.5
1 parent 263b222 commit a778554

24 files changed

Lines changed: 190 additions & 181 deletions

‎RG-Evo‎

0 Bytes
Binary file not shown.

‎RG-Evo-shared‎

0 Bytes
Binary file not shown.

‎include/EOMs/rk_data.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#pragma once
22
#include "core/config_build.hpp"
33
#if DMFE_WITH_CUDA
4-
#include "core/device_vector.hpp"
4+
#include "core/device_vector.cuh"
55
#endif
66
#include <cstddef>
77
#include <vector>

‎include/EOMs/time_steps.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77

88
#if DMFE_WITH_CUDA
99
#include <cuda_runtime.h>
10-
#include "core/device_vector.hpp"
10+
#include "core/device_vector.cuh"
1111
#include <thrust/device_ptr.h>
1212
#endif
1313

‎include/convolution/convolution.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
#include <vector>
44

55
#if DMFE_WITH_CUDA
6-
#include "core/device_vector.hpp"
6+
#include "core/device_vector.cuh"
77
#include <thrust/device_ptr.h>
88
#include <cuda_runtime.h>
99
#endif

‎include/core/device_utils.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77

88
#if DMFE_WITH_CUDA
99
#include <cuda_runtime.h>
10-
#include "core/device_vector.hpp"
10+
#include "core/device_vector.cuh"
1111
#include <thrust/device_ptr.h>
1212

1313
// Debug-only CUDA error check utilities. Use tiny inline helpers + macros to avoid overhead when not debugging.

‎include/core/device_vector.cuh‎

Lines changed: 160 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,160 @@
1+
#pragma once
2+
#include "core/config_build.hpp"
3+
#include <cuda_runtime.h>
4+
#include <thrust/device_vector.h>
5+
#include <thrust/device_ptr.h>
6+
#include <thrust/system_error.h>
7+
#include <thrust/system/cuda/error.h>
8+
#include <cstddef>
9+
#include <cstdint>
10+
#include <limits>
11+
#include <new>
12+
13+
namespace dmfe {
14+
namespace detail {
15+
inline bool async_pool_supported() {
16+
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
17+
static int supported = -1;
18+
if (supported >= 0) {
19+
return supported != 0;
20+
}
21+
int device = 0;
22+
cudaError_t dev_err = cudaGetDevice(&device);
23+
if (dev_err != cudaSuccess) {
24+
supported = 0;
25+
return false;
26+
}
27+
int attr = 0;
28+
cudaError_t err = cudaDeviceGetAttribute(&attr, cudaDevAttrMemoryPoolsSupported, device);
29+
if (err != cudaSuccess) {
30+
supported = 0;
31+
return false;
32+
}
33+
supported = attr;
34+
return supported != 0;
35+
#else
36+
return false;
37+
#endif
38+
}
39+
40+
inline void configure_async_pool_once() {
41+
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
42+
static bool configured = false;
43+
if (configured) {
44+
return;
45+
}
46+
configured = true;
47+
int device = 0;
48+
if (cudaGetDevice(&device) != cudaSuccess) {
49+
return;
50+
}
51+
cudaMemPool_t pool = nullptr;
52+
if (cudaDeviceGetDefaultMemPool(&pool, device) != cudaSuccess) {
53+
return;
54+
}
55+
std::uint64_t threshold = std::numeric_limits<std::uint64_t>::max();
56+
cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &threshold);
57+
#endif
58+
}
59+
60+
inline void* allocate_device_bytes(std::size_t bytes, cudaStream_t stream) {
61+
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
62+
if (stream != 0 && async_pool_supported()) {
63+
configure_async_pool_once();
64+
void* ptr = nullptr;
65+
cudaError_t err = cudaMallocAsync(&ptr, bytes, stream);
66+
if (err == cudaSuccess) {
67+
return ptr;
68+
}
69+
// Fall back if async allocation fails for any reason.
70+
}
71+
#endif
72+
void* ptr = nullptr;
73+
cudaError_t err = cudaMalloc(&ptr, bytes);
74+
if (err != cudaSuccess) {
75+
throw thrust::system_error(err, thrust::cuda_category(), "cudaMalloc failed");
76+
}
77+
return ptr;
78+
}
79+
80+
inline void deallocate_device_bytes(void* ptr, cudaStream_t stream) {
81+
if (!ptr) {
82+
return;
83+
}
84+
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
85+
if (stream != 0 && async_pool_supported()) {
86+
cudaError_t err = cudaFreeAsync(ptr, stream);
87+
if (err == cudaSuccess) {
88+
return;
89+
}
90+
// Fall back if async free fails for any reason.
91+
}
92+
#endif
93+
cudaError_t err = cudaFree(ptr);
94+
if (err != cudaSuccess) {
95+
throw thrust::system_error(err, thrust::cuda_category(), "cudaFree failed");
96+
}
97+
}
98+
} // namespace detail
99+
100+
template <typename T>
101+
class cuda_async_allocator {
102+
public:
103+
using value_type = T;
104+
using pointer = thrust::device_ptr<T>;
105+
using const_pointer = thrust::device_ptr<const T>;
106+
using reference = T&;
107+
using const_reference = const T&;
108+
using size_type = std::size_t;
109+
using difference_type = std::ptrdiff_t;
110+
111+
cuda_async_allocator() noexcept : stream_(0) {}
112+
explicit cuda_async_allocator(cudaStream_t stream) noexcept : stream_(stream) {}
113+
114+
template <class U>
115+
cuda_async_allocator(const cuda_async_allocator<U>& other) noexcept : stream_(other.stream()) {}
116+
117+
pointer allocate(std::size_t n) {
118+
if (n > max_size()) {
119+
throw std::bad_alloc();
120+
}
121+
std::size_t bytes = n * sizeof(T);
122+
return thrust::device_pointer_cast(static_cast<T*>(detail::allocate_device_bytes(bytes, stream_)));
123+
}
124+
125+
void deallocate(pointer ptr, std::size_t) noexcept {
126+
try {
127+
detail::deallocate_device_bytes(ptr.get(), stream_);
128+
} catch (...) {
129+
// Thrust allocators are required not to throw from deallocate.
130+
}
131+
}
132+
133+
std::size_t max_size() const noexcept {
134+
return std::numeric_limits<std::size_t>::max() / sizeof(T);
135+
}
136+
137+
cudaStream_t stream() const noexcept { return stream_; }
138+
139+
template <class U>
140+
struct rebind {
141+
using other = cuda_async_allocator<U>;
142+
};
143+
144+
bool operator==(const cuda_async_allocator& other) const noexcept {
145+
return stream_ == other.stream_;
146+
}
147+
148+
bool operator!=(const cuda_async_allocator& other) const noexcept {
149+
return !(*this == other);
150+
}
151+
152+
private:
153+
cudaStream_t stream_;
154+
};
155+
156+
template <typename T>
157+
using device_vector = thrust::device_vector<T, cuda_async_allocator<T>>;
158+
159+
} // namespace dmfe
160+

‎include/core/device_vector.hpp‎

Lines changed: 1 addition & 164 deletions
Original file line numberDiff line numberDiff line change
@@ -1,171 +1,8 @@
11
#pragma once
22
#include "core/config_build.hpp"
3-
4-
#if DMFE_WITH_CUDA
5-
#include <cuda_runtime.h>
6-
#include <thrust/device_vector.h>
7-
#include <thrust/device_ptr.h>
8-
#include <thrust/system_error.h>
9-
#include <thrust/system/cuda/error.h>
10-
#include <cstddef>
11-
#include <cstdint>
12-
#include <limits>
13-
#include <new>
14-
15-
namespace dmfe {
16-
namespace detail {
17-
inline bool async_pool_supported() {
18-
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
19-
static int supported = -1;
20-
if (supported >= 0) {
21-
return supported != 0;
22-
}
23-
int device = 0;
24-
cudaError_t dev_err = cudaGetDevice(&device);
25-
if (dev_err != cudaSuccess) {
26-
supported = 0;
27-
return false;
28-
}
29-
int attr = 0;
30-
cudaError_t err = cudaDeviceGetAttribute(&attr, cudaDevAttrMemoryPoolsSupported, device);
31-
if (err != cudaSuccess) {
32-
supported = 0;
33-
return false;
34-
}
35-
supported = attr;
36-
return supported != 0;
37-
#else
38-
return false;
39-
#endif
40-
}
41-
42-
inline void configure_async_pool_once() {
43-
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
44-
static bool configured = false;
45-
if (configured) {
46-
return;
47-
}
48-
configured = true;
49-
int device = 0;
50-
if (cudaGetDevice(&device) != cudaSuccess) {
51-
return;
52-
}
53-
cudaMemPool_t pool = nullptr;
54-
if (cudaDeviceGetDefaultMemPool(&pool, device) != cudaSuccess) {
55-
return;
56-
}
57-
std::uint64_t threshold = std::numeric_limits<std::uint64_t>::max();
58-
cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &threshold);
59-
#endif
60-
}
61-
62-
inline void* allocate_device_bytes(std::size_t bytes, cudaStream_t stream) {
63-
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
64-
if (stream != 0 && async_pool_supported()) {
65-
configure_async_pool_once();
66-
void* ptr = nullptr;
67-
cudaError_t err = cudaMallocAsync(&ptr, bytes, stream);
68-
if (err == cudaSuccess) {
69-
return ptr;
70-
}
71-
// Fall back if async allocation fails for any reason.
72-
}
73-
#endif
74-
void* ptr = nullptr;
75-
cudaError_t err = cudaMalloc(&ptr, bytes);
76-
if (err != cudaSuccess) {
77-
throw thrust::system_error(err, thrust::cuda_category(), "cudaMalloc failed");
78-
}
79-
return ptr;
80-
}
81-
82-
inline void deallocate_device_bytes(void* ptr, cudaStream_t stream) {
83-
if (!ptr) {
84-
return;
85-
}
86-
#if defined(CUDART_VERSION) && (CUDART_VERSION >= 11020)
87-
if (stream != 0 && async_pool_supported()) {
88-
cudaError_t err = cudaFreeAsync(ptr, stream);
89-
if (err == cudaSuccess) {
90-
return;
91-
}
92-
// Fall back if async free fails for any reason.
93-
}
94-
#endif
95-
cudaError_t err = cudaFree(ptr);
96-
if (err != cudaSuccess) {
97-
throw thrust::system_error(err, thrust::cuda_category(), "cudaFree failed");
98-
}
99-
}
100-
} // namespace detail
101-
102-
template <typename T>
103-
class cuda_async_allocator {
104-
public:
105-
using value_type = T;
106-
using pointer = thrust::device_ptr<T>;
107-
using const_pointer = thrust::device_ptr<const T>;
108-
using reference = T&;
109-
using const_reference = const T&;
110-
using size_type = std::size_t;
111-
using difference_type = std::ptrdiff_t;
112-
113-
cuda_async_allocator() noexcept : stream_(0) {}
114-
explicit cuda_async_allocator(cudaStream_t stream) noexcept : stream_(stream) {}
115-
116-
template <class U>
117-
cuda_async_allocator(const cuda_async_allocator<U>& other) noexcept : stream_(other.stream()) {}
118-
119-
pointer allocate(std::size_t n) {
120-
if (n > max_size()) {
121-
throw std::bad_alloc();
122-
}
123-
std::size_t bytes = n * sizeof(T);
124-
return thrust::device_pointer_cast(static_cast<T*>(detail::allocate_device_bytes(bytes, stream_)));
125-
}
126-
127-
void deallocate(pointer ptr, std::size_t) noexcept {
128-
try {
129-
detail::deallocate_device_bytes(ptr.get(), stream_);
130-
} catch (...) {
131-
// Thrust allocators are required not to throw from deallocate.
132-
}
133-
}
134-
135-
std::size_t max_size() const noexcept {
136-
return std::numeric_limits<std::size_t>::max() / sizeof(T);
137-
}
138-
139-
cudaStream_t stream() const noexcept { return stream_; }
140-
141-
template <class U>
142-
struct rebind {
143-
using other = cuda_async_allocator<U>;
144-
};
145-
146-
bool operator==(const cuda_async_allocator& other) const noexcept {
147-
return stream_ == other.stream_;
148-
}
149-
150-
bool operator!=(const cuda_async_allocator& other) const noexcept {
151-
return !(*this == other);
152-
}
153-
154-
private:
155-
cudaStream_t stream_;
156-
};
157-
158-
template <typename T>
159-
using device_vector = thrust::device_vector<T, cuda_async_allocator<T>>;
160-
161-
} // namespace dmfe
162-
163-
#else
1643
#include <vector>
1654

1665
namespace dmfe {
1676
template <typename T>
1687
using device_vector = std::vector<T>;
169-
}
170-
171-
#endif
8+
}

‎include/core/host_device_utils.hpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,10 @@
11
#pragma once
22
#include <vector>
3+
#if DMFE_WITH_CUDA
4+
#include "core/device_vector.cuh"
5+
#else
36
#include "core/device_vector.hpp"
7+
#endif
48
#include <thrust/device_ptr.h>
59
#include <cassert>
610

‎include/core/vector_utils.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,7 @@
55
#include <algorithm>
66

77
#if DMFE_WITH_CUDA
8-
#include "core/device_vector.hpp"
8+
#include "core/device_vector.cuh"
99
#include <thrust/device_ptr.h>
1010

1111
// Slice helpers

0 commit comments

Comments
 (0)