10#ifdef ALPAKA_ACC_GPU_HIP_ENABLED
12# include <hip/hip_runtime_api.h>
13# include <hip/hip_version.h>
20 static constexpr char name[] =
"Hip";
24 using DeviceAttr_t = ::hipDeviceAttribute_t;
25 using DeviceProp_t = ::hipDeviceProp_t;
26 using Error_t = ::hipError_t;
27 using Event_t = ::hipEvent_t;
28 using Extent_t = ::hipExtent;
29 using Flag_t =
unsigned int;
30 using FuncAttributes_t = ::hipFuncAttributes;
31 using HostFn_t = void (*)(
void* data);
32 using Limit_t = ::hipLimit_t;
33 using Memcpy3DParms_t = ::hipMemcpy3DParms;
34 using MemcpyKind_t = ::hipMemcpyKind;
35 using PitchedPtr_t = ::hipPitchedPtr;
36 using Pos_t = ::hipPos;
37 using Stream_t = ::hipStream_t;
40 static constexpr Error_t success = ::hipSuccess;
41 static constexpr Error_t errorNotReady = ::hipErrorNotReady;
42 static constexpr Error_t errorHostMemoryAlreadyRegistered = ::hipErrorHostMemoryAlreadyRegistered;
43 static constexpr Error_t errorHostMemoryNotRegistered = ::hipErrorHostMemoryNotRegistered;
44 static constexpr Error_t errorUnsupportedLimit = ::hipErrorUnsupportedLimit;
45 static constexpr Error_t errorUnknown = ::hipErrorUnknown;
47 static constexpr Flag_t eventDefault = hipEventDefault;
48 static constexpr Flag_t eventBlockingSync = hipEventBlockingSync;
49 static constexpr Flag_t eventDisableTiming = hipEventDisableTiming;
50 static constexpr Flag_t eventInterprocess = hipEventInterprocess;
52 static constexpr Flag_t hostMallocDefault = hipHostMallocDefault;
53 static constexpr Flag_t hostMallocMapped = hipHostMallocMapped;
54 static constexpr Flag_t hostMallocPortable = hipHostMallocPortable;
55 static constexpr Flag_t hostMallocWriteCombined = hipHostMallocWriteCombined;
56 static constexpr Flag_t hostMallocCoherent = hipHostMallocCoherent;
57 static constexpr Flag_t hostMallocNonCoherent = hipHostMallocNonCoherent;
59 static constexpr Flag_t hostRegisterDefault = hipHostRegisterDefault;
60 static constexpr Flag_t hostRegisterPortable = hipHostRegisterPortable;
61 static constexpr Flag_t hostRegisterMapped = hipHostRegisterMapped;
62 static constexpr Flag_t hostRegisterIoMemory = hipHostRegisterIoMemory;
64 static constexpr MemcpyKind_t memcpyDefault = ::hipMemcpyDefault;
65 static constexpr MemcpyKind_t memcpyDeviceToDevice = ::hipMemcpyDeviceToDevice;
66 static constexpr MemcpyKind_t memcpyDeviceToHost = ::hipMemcpyDeviceToHost;
67 static constexpr MemcpyKind_t memcpyHostToDevice = ::hipMemcpyHostToDevice;
69 static constexpr Flag_t streamDefault = hipStreamDefault;
70 static constexpr Flag_t streamNonBlocking = hipStreamNonBlocking;
72 static constexpr DeviceAttr_t deviceAttributeMaxBlockDimX = ::hipDeviceAttributeMaxBlockDimX;
73 static constexpr DeviceAttr_t deviceAttributeMaxBlockDimY = ::hipDeviceAttributeMaxBlockDimY;
74 static constexpr DeviceAttr_t deviceAttributeMaxBlockDimZ = ::hipDeviceAttributeMaxBlockDimZ;
75 static constexpr DeviceAttr_t deviceAttributeMaxGridDimX = ::hipDeviceAttributeMaxGridDimX;
76 static constexpr DeviceAttr_t deviceAttributeMaxGridDimY = ::hipDeviceAttributeMaxGridDimY;
77 static constexpr DeviceAttr_t deviceAttributeMaxGridDimZ = ::hipDeviceAttributeMaxGridDimZ;
78 static constexpr DeviceAttr_t deviceAttributeMaxSharedMemoryPerBlock
79 = ::hipDeviceAttributeMaxSharedMemoryPerBlock;
80 static constexpr DeviceAttr_t deviceAttributeMaxThreadsPerBlock = ::hipDeviceAttributeMaxThreadsPerBlock;
81 static constexpr DeviceAttr_t deviceAttributeMultiprocessorCount = ::hipDeviceAttributeMultiprocessorCount;
82 static constexpr DeviceAttr_t deviceAttributeWarpSize = ::hipDeviceAttributeWarpSize;
84# if HIP_VERSION >= 40'500'000
85 static constexpr Limit_t limitPrintfFifoSize = ::hipLimitPrintfFifoSize;
87 static constexpr Limit_t limitPrintfFifoSize
88 =
static_cast<Limit_t
>(0x01);
90 static constexpr Limit_t limitMallocHeapSize = ::hipLimitMallocHeapSize;
100 static void hostFunction(
void* data)
102 auto ptr =
reinterpret_cast<HostFnAdaptor*
>(data);
103 ptr->func_(ptr->data_);
107 static void streamCallback(Stream_t, Error_t,
void* data)
109 auto ptr =
reinterpret_cast<HostFnAdaptor*
>(data);
110 ptr->func_(ptr->data_);
116 static inline Error_t deviceGetAttribute(
int* value, DeviceAttr_t attr,
int device)
118 return ::hipDeviceGetAttribute(value, attr, device);
121 static inline Error_t deviceGetLimit(
size_t* pValue, Limit_t limit)
123# if HIP_VERSION < 40'500'000
124 if(limit == limitPrintfFifoSize)
127 return errorUnsupportedLimit;
130 return ::hipDeviceGetLimit(pValue, limit);
133 static inline Error_t deviceReset()
135 return ::hipDeviceReset();
138 static inline Error_t deviceSetLimit(Limit_t ,
size_t )
141 return errorUnsupportedLimit;
144 static inline Error_t deviceSynchronize()
146 return ::hipDeviceSynchronize();
149 static inline Error_t eventCreate(Event_t* event)
151 return ::hipEventCreate(event);
154 static inline Error_t eventCreateWithFlags(Event_t* event, Flag_t flags)
156 return ::hipEventCreateWithFlags(event, flags);
159 static inline Error_t eventDestroy(Event_t event)
161 return ::hipEventDestroy(event);
164 static inline Error_t eventQuery(Event_t event)
166 return ::hipEventQuery(event);
169 static inline Error_t eventRecord(Event_t event, Stream_t stream)
171 return ::hipEventRecord(event, stream);
174 static inline Error_t eventSynchronize(Event_t event)
176 return ::hipEventSynchronize(event);
179 static inline Error_t
free(
void* devPtr)
181 return ::hipFree(devPtr);
184 static inline Error_t freeAsync([[maybe_unused]]
void* devPtr, [[maybe_unused]] Stream_t stream)
187# if HIP_VERSION >= 50'300'000
191 return ::hipFreeAsync(devPtr, stream);
203 static inline Error_t funcGetAttributes(FuncAttributes_t* attr,
void const* func)
205 return ::hipFuncGetAttributes(attr, func);
209 static inline Error_t funcGetAttributes(FuncAttributes_t* attr, T* func)
212# pragma GCC diagnostic push
213# pragma GCC diagnostic ignored "-Wconditionally-supported"
215 return ::hipFuncGetAttributes(attr,
reinterpret_cast<void const*
>(func));
217# pragma GCC diagnostic pop
221 static inline int getCurrentDevice()
224 using TApi = alpaka::ApiHipRt;
230 static inline Error_t getDeviceCount(
int* count)
232 return ::hipGetDeviceCount(count);
235 static inline Error_t getDeviceProperties(DeviceProp_t* prop,
int device)
237 return ::hipGetDeviceProperties(prop, device);
240 static inline char const* getErrorName(Error_t error)
242 return ::hipGetErrorName(error);
245 static inline char const* getErrorString(Error_t error)
247 return ::hipGetErrorString(error);
250 static inline Error_t getLastError()
252 return ::hipGetLastError();
255 static inline Error_t getSymbolAddress(
void** devPtr,
void const* symbol)
257 return ::hipGetSymbolAddress(devPtr, symbol);
261 static inline Error_t getSymbolAddress(
void** devPtr, T
const& symbol)
263 return ::hipGetSymbolAddress(devPtr, symbol);
266 static inline Error_t hostGetDevicePointer(
void** pDevice,
void* pHost, Flag_t flags)
268 return ::hipHostGetDevicePointer(pDevice, pHost, flags);
271 static inline Error_t hostFree(
void* ptr)
273 return ::hipHostFree(ptr);
276 static inline Error_t hostMalloc(
void** ptr,
size_t size, Flag_t flags)
278 return ::hipHostMalloc(ptr, size, flags);
281 static inline Error_t hostRegister(
void* ptr,
size_t size, Flag_t flags)
283 return ::hipHostRegister(ptr, size, flags);
286 static inline Error_t hostUnregister(
void* ptr)
288 return ::hipHostUnregister(ptr);
291 static inline Error_t launchHostFunc(Stream_t stream, HostFn_t fn,
void* userData)
294# if HIP_VERSION >= 50'400'000
296 return ::hipLaunchHostFunc(stream, HostFnAdaptor::hostFunction,
new HostFnAdaptor{fn, userData});
299 return ::hipStreamAddCallback(stream, HostFnAdaptor::streamCallback,
new HostFnAdaptor{fn, userData}, 0);
303 static inline Error_t
malloc(
void** devPtr,
size_t size)
305 return ::hipMalloc(devPtr, size);
308 static inline Error_t malloc3D(PitchedPtr_t* pitchedDevPtr, Extent_t extent)
310 return ::hipMalloc3D(pitchedDevPtr, extent);
313 static inline Error_t mallocAsync(
314 [[maybe_unused]]
void** devPtr,
315 [[maybe_unused]]
size_t size,
316 [[maybe_unused]] Stream_t stream)
319# if HIP_VERSION >= 50'600'000
320 return ::hipMallocAsync(devPtr, size, stream);
321# elif HIP_VERSION >= 50'300'000
325 return ::hipMallocAsync(devPtr, size, stream);
339 static inline Error_t mallocPitch(
void** devPtr,
size_t* pitch,
size_t width,
size_t height)
341 return ::hipMallocPitch(devPtr, pitch, width, height);
344 static inline Error_t memGetInfo(
size_t* free,
size_t* total)
346 return ::hipMemGetInfo(free, total);
349 static inline Error_t
memcpy(
void* dst,
void const* src,
size_t count, MemcpyKind_t kind)
351 return ::hipMemcpy(dst, src, count, kind);
354 static inline Error_t memcpy2DAsync(
364 return ::hipMemcpy2DAsync(dst, dpitch, src, spitch, width, height, kind, stream);
367 static inline Error_t memcpy3DAsync(Memcpy3DParms_t
const* p, Stream_t stream)
369 return ::hipMemcpy3DAsync(p, stream);
372 static inline Error_t memcpyAsync(
void* dst,
void const* src,
size_t count, MemcpyKind_t kind, Stream_t stream)
374 return ::hipMemcpyAsync(dst, src, count, kind, stream);
377 static inline Error_t memset2DAsync(
385 return ::hipMemset2DAsync(devPtr, pitch, value, width, height, stream);
388 static inline Error_t memset3DAsync(PitchedPtr_t pitchedDevPtr,
int value, Extent_t extent, Stream_t stream)
390 return ::hipMemset3DAsync(pitchedDevPtr, value, extent, stream);
393 static inline Error_t memsetAsync(
void* devPtr,
int value,
size_t count, Stream_t stream)
395 return ::hipMemsetAsync(devPtr, value, count, stream);
398 static inline Error_t setDevice(
int device)
400 return ::hipSetDevice(device);
403 static inline Error_t streamCreate(Stream_t* pStream)
405 return ::hipStreamCreate(pStream);
408 static inline Error_t streamCreateWithFlags(Stream_t* pStream, Flag_t flags)
410 return ::hipStreamCreateWithFlags(pStream, flags);
413 static inline Error_t streamDestroy(Stream_t stream)
415 return ::hipStreamDestroy(stream);
418 static inline Error_t streamQuery(Stream_t stream)
420 return ::hipStreamQuery(stream);
423 static inline Error_t streamSynchronize(Stream_t stream)
425 return ::hipStreamSynchronize(stream);
428 static inline Error_t streamWaitEvent(Stream_t stream, Event_t event, Flag_t flags)
430 return ::hipStreamWaitEvent(stream, event, flags);
433 static inline PitchedPtr_t makePitchedPtr(
void* d,
size_t p,
size_t xsz,
size_t ysz)
435 return ::make_hipPitchedPtr(d, p, xsz, ysz);
438 static inline Pos_t makePos(
size_t x,
size_t y,
size_t z)
440 return ::make_hipPos(x, y, z);
443 static inline Extent_t makeExtent(
size_t w,
size_t h,
size_t d)
445 return ::make_hipExtent(w, h, d);
#define ALPAKA_VERSION_NUMBER(major, minor, patch)
The alpaka accelerator library.
ALPAKA_FN_HOST auto free(TAlloc const &alloc, T const *const ptr) -> void
Frees the memory identified by the given pointer.
ALPAKA_FN_HOST auto memcpy(TQueue &queue, alpaka::detail::DevGlobalImplGeneric< TTag, TTypeDst > &viewDst, TViewSrc const &viewSrc) -> void
ALPAKA_FN_HOST auto malloc(TAlloc const &alloc, std::size_t const &sizeElems) -> T *