33 #include <hip/hip_runtime.h> 36 #include <cuda_runtime.h> 103 template <
typename T>
116 throw std::invalid_argument(
"attempting to construct a TargetAllocator with an invalid target");
126 return this->_target;
136 throw std::invalid_argument(
"attempting to set TargetAllocator to an invalid target");
139 this->_target = new_target;
155 using type =
typename std::conditional<std::is_void<T>::value, char, T>::type;
157 std::size_t n_bytes = n *
sizeof(type);
158 switch (this->_target) {
161 std::allocator<type> std_alloc;
164 ptr = std_alloc.allocate(n);
170 std::memset((
void *) ptr, 0, n_bytes);
177 #ifdef HAS_ROCM_MODEL 185 if (this->_always_managed) {
186 error = hipMallocManaged(&hip_ptr, n_bytes);
190 int managed_memory = 0;
191 hipGetDevice(&device);
192 hipDeviceGetAttribute(&managed_memory, hipDeviceAttributeManagedMemory, device);
195 if (managed_memory) {
196 error = hipMallocManaged(&hip_ptr, n_bytes);
199 error = hipMalloc(&hip_ptr, n_bytes);
202 if (error != hipSuccess) {
203 throw std::runtime_error(hipGetErrorString(error));
205 error = hipMemset(hip_ptr, 0, n_bytes);
206 if (error != hipSuccess) {
207 throw std::runtime_error(hipGetErrorString(error));
212 throw std::invalid_argument(
"target MLIR_ROCM is unavailable");
216 #ifdef HAS_CUDA_MODEL 221 error = cudaMallocManaged(&cuda_ptr, n_bytes);
222 if (error != cudaSuccess) {
223 throw std::runtime_error(cudaGetErrorString(error));
225 error = cudaMemset(cuda_ptr, 0, n_bytes);
226 if (error != cudaSuccess) {
227 throw std::runtime_error(cudaGetErrorString(error));
232 throw std::invalid_argument(
"target MLIR_CUDA is unavailable");
236 throw std::invalid_argument(
"unknown allocation target");
251 using type =
typename std::conditional<std::is_void<T>::value, char, T>::type;
252 switch (this->_target) {
255 std::allocator<type> std_alloc;
256 std_alloc.deallocate((type*) ptr, n);
260 #ifdef HAS_ROCM_MODEL 263 error = hipFree(ptr);
264 if (error != hipSuccess) {
265 throw std::runtime_error(hipGetErrorString(error));
269 throw std::invalid_argument(
"target MLIR_ROCM is unavailable");
273 #ifdef HAS_CUDA_MODEL 276 error = cudaFree(ptr);
277 if (error != cudaSuccess) {
278 throw std::runtime_error(cudaGetErrorString(error));
282 throw std::invalid_argument(
"target MLIR_CUDA is unavailable");
286 throw std::invalid_argument(
"unknown allocation target");
292 bool _always_managed;
std::string get_string_from_target(Target const target)
Get a string representation of a given target.
TargetAllocator(Target target, bool always_managed=false)
Construct a TargetAllocator.
vectorized CPU code generated with MLIR
Allocator structure for dynamically allocating memory on multiple targets.
bool is_gpu(Target const target)
Checks if this is a GPU target.
special value to handle unknown targets
T value_type
type to allocate
std::string get_target_list_string()
Returns a string containing the list of available targets.
a token to indicate the maximum number of targets
bool is_concrete(Target const target)
Checks if target is a real, concrete target.
void set_target(Target new_target)
Set a new target for this allocator.
Target get_target() const
Get the target for this allocator.
baseline CPU model generated with the original opencarp code generator
CUDA code for NVIDIA GPUs generated with MLIR.
T * allocate(std::size_t n, bool do_zero=true)
Allocate memory for type T.
Target get_target_from_string(std::string const str)
Returns a value from the Target enum from a given string.
void deallocate(T *ptr, std::size_t n=0)
Deallocate memory pointed by ptr.
Target
enum that represents different targets to run ionic models on.
T * allocate_on_target(Target target, std::size_t n, bool always_managed=false, bool do_zero=true)
Utility function for allocating memory on a target. See TargetAllocator.
void deallocate_on_target(Target target, T *ptr)
Utility function for deallocating memory on a target. See TargetAllocator.
ROCM code for AMD GPUs generated with MLIR.