Repository navigation
Expand file tree
/
Copy pathmemory.cpp
More file actions
196 lines (139 loc) · 5.47 KB
/
Copy pathmemory.cpp
File metadata and controls
196 lines (139 loc) · 5.47 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
//
// SPDX-License-Identifier: BSD-3-Clause
#include <ginkgo/core/base/memory.hpp>
#include <cuda.h>
#include <cuda_runtime.h>
#include <ginkgo/core/base/exception_helpers.hpp>
#include "cuda/base/scoped_device_id.hpp"
namespace gko {
#define GKO_ASSERT_NO_CUDA_ALLOCATION_ERRORS(_operation, _size) \
{ \
auto error_code = _operation; \
if (error_code == cudaErrorMemoryAllocation) { \
throw AllocationError(__FILE__, __LINE__, "cuda", _size); \
} else { \
GKO_ASSERT_NO_CUDA_ERRORS(error_code); \
} \
}
#if GKO_VERBOSE_LEVEL >= 1
#define GKO_EXIT_ON_CUDA_ERROR(_operation) \
{ \
const auto error_code = _operation; \
if (error_code != cudaSuccess) { \
int device_id{-1}; \
cudaGetDevice(&device_id); \
std::cerr << "Unrecoverable CUDA error on device " << device_id \
<< " in " << __func__ << ":" << __LINE__ << ": " \
<< cudaGetErrorName(error_code) << ": " \
<< cudaGetErrorString(error_code) << std::endl \
<< "Exiting program" << std::endl; \
std::exit(error_code); \
} \
}
#else
#define GKO_EXIT_ON_CUDA_ERROR(_operation) \
{ \
const auto error_code = _operation; \
if (error_code != cudaSuccess) { \
std::exit(error_code); \
} \
}
#endif
void* CudaAllocator::allocate(size_type num_bytes)
{
void* ptr{};
GKO_ASSERT_NO_CUDA_ALLOCATION_ERRORS(cudaMalloc(&ptr, num_bytes),
num_bytes);
return ptr;
}
void CudaAllocator::deallocate(void* ptr)
{
GKO_EXIT_ON_CUDA_ERROR(cudaFree(ptr));
}
#if CUDA_VERSION >= 11020
CudaAsyncAllocator::CudaAsyncAllocator(cudaStream_t stream) : stream_{stream} {}
void* CudaAsyncAllocator::allocate(size_type num_bytes)
{
void* ptr{};
GKO_ASSERT_NO_CUDA_ALLOCATION_ERRORS(
cudaMallocAsync(&ptr, num_bytes, stream_), num_bytes);
return ptr;
}
void CudaAsyncAllocator::deallocate(void* ptr)
{
GKO_EXIT_ON_CUDA_ERROR(cudaFreeAsync(ptr, stream_));
}
#else // Fall back to regular allocation
CudaAsyncAllocator::CudaAsyncAllocator(cudaStream_t stream) : stream_{stream}
{
#if GKO_VERBOSE_LEVEL >= 1
std::cerr << "This version of CUDA does not support cudaMallocAsync, "
"please use CudaAllocator instead of CudaAsyncAllocator.\n";
#endif
}
void* CudaAsyncAllocator::allocate(size_type num_bytes)
{
void* ptr{};
GKO_ASSERT_NO_CUDA_ALLOCATION_ERRORS(cudaMalloc(&ptr, num_bytes),
num_bytes);
return ptr;
}
void CudaAsyncAllocator::deallocate(void* ptr)
{
GKO_EXIT_ON_CUDA_ERROR(cudaFree(ptr));
}
#endif
bool CudaAsyncAllocator::check_environment(int device_id,
CUstream_st* stream) const
{
return stream == stream_;
}
CudaUnifiedAllocator::CudaUnifiedAllocator(int device_id)
: CudaUnifiedAllocator{device_id, cudaMemAttachGlobal}
{}
CudaUnifiedAllocator::CudaUnifiedAllocator(int device_id, unsigned int flags)
: device_id_{device_id}, flags_{flags}
{}
void* CudaUnifiedAllocator::allocate(size_type num_bytes)
{
// we need to set the device ID in case this gets used in a host executor
detail::cuda_scoped_device_id_guard g(device_id_);
void* ptr{};
GKO_ASSERT_NO_CUDA_ALLOCATION_ERRORS(
cudaMallocManaged(&ptr, num_bytes, flags_), num_bytes);
return ptr;
}
void CudaUnifiedAllocator::deallocate(void* ptr)
{
// we need to set the device ID in case this gets used in a host executor
detail::cuda_scoped_device_id_guard g(device_id_);
GKO_EXIT_ON_CUDA_ERROR(cudaFree(ptr));
}
bool CudaUnifiedAllocator::check_environment(int device_id,
CUstream_st* stream) const
{
return device_id == device_id_;
}
CudaHostAllocator::CudaHostAllocator(int device_id) : device_id_{device_id} {}
void* CudaHostAllocator::allocate(size_type num_bytes)
{
// we need to set the device ID in case this gets used in a host executor
detail::cuda_scoped_device_id_guard g(device_id_);
void* ptr{};
GKO_ASSERT_NO_CUDA_ALLOCATION_ERRORS(cudaMallocHost(&ptr, num_bytes),
num_bytes);
return ptr;
}
void CudaHostAllocator::deallocate(void* ptr)
{
// we need to set the device ID in case this gets used in a host executor
detail::cuda_scoped_device_id_guard g(device_id_);
GKO_EXIT_ON_CUDA_ERROR(cudaFreeHost(ptr));
}
bool CudaHostAllocator::check_environment(int device_id,
CUstream_st* stream) const
{
return device_id == device_id_;
}
} // namespace gko