HIP: Heterogenous-computing Interface for Portability
hip_runtime_api.h
Go to the documentation of this file.
1 /*
2 Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
23 //#pragma once
24 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_API_H
25 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_API_H
26 
31 #include <stdint.h>
32 #include <stddef.h>
33 
34 #ifndef GENERIC_GRID_LAUNCH
35 #define GENERIC_GRID_LAUNCH 1
36 #endif
37 
38 #ifndef __HIP_VDI__
39 #define __HIP_VDI__ 0
40 #endif
41 
43 #include <hip/hip_runtime_api.h>
44 #include <hip/hcc_detail/driver_types.h>
47 
48 #if !__HIP_VDI__
49 #include <hsa/hsa.h>
50 #include <hip/hcc_detail/program_state.hpp>
51 #endif
52 
53 #if defined(_MSC_VER)
54 #define DEPRECATED(msg) __declspec(deprecated(msg))
55 #else // !defined(_MSC_VER)
56 #define DEPRECATED(msg) __attribute__ ((deprecated(msg)))
57 #endif // !defined(_MSC_VER)
58 
59 #define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases.For more details please refer https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list"
60 
61 #if defined(__HCC__) && (__hcc_workweek__ < 16155)
62 #error("This version of HIP requires a newer version of HCC.");
63 #endif
64 
65 #define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*)0x01)
66 #define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*)0x02)
67 #define HIP_LAUNCH_PARAM_END ((void*)0x03)
68 
69 #ifdef __cplusplus
70  #include <mutex>
71  #include <string>
72  #include <unordered_map>
73  #include <vector>
74 
75  #define __dparm(x) \
76  = x
77 #else
78  #define __dparm(x)
79 #endif
80 
81 namespace hip_impl {
82 hipError_t hip_init();
83 } // namespace hip_impl
84 
85 // Structure definitions:
86 #ifdef __cplusplus
87 extern "C" {
88 #endif
89 
90 //---
91 // API-visible structures
92 typedef struct ihipCtx_t* hipCtx_t;
93 
94 // Note many APIs also use integer deviceIds as an alternative to the device pointer:
95 typedef int hipDevice_t;
96 
97 typedef struct ihipStream_t* hipStream_t;
98 
99 // TODO: IPC implementation
100 
101 #define hipIpcMemLazyEnablePeerAccess 0
102 
103 #define HIP_IPC_HANDLE_SIZE 64
104 
105 typedef struct hipIpcMemHandle_st {
106  char reserved[HIP_IPC_HANDLE_SIZE];
108 
109 // TODO: IPC event handle currently unsupported
110 struct ihipIpcEventHandle_t;
111 typedef struct ihipIpcEventHandle_t* hipIpcEventHandle_t;
112 
113 
114 // END TODO
115 
116 typedef struct ihipModule_t* hipModule_t;
117 
118 typedef struct ihipModuleSymbol_t* hipFunction_t;
119 
120 typedef struct hipFuncAttributes {
121  int binaryVersion;
122  int cacheModeCA;
123  size_t constSizeBytes;
124  size_t localSizeBytes;
125  int maxDynamicSharedSizeBytes;
126  int maxThreadsPerBlock;
127  int numRegs;
128  int preferredShmemCarveout;
129  int ptxVersion;
130  size_t sharedSizeBytes;
132 
133 typedef struct ihipEvent_t* hipEvent_t;
134 
135 enum hipLimit_t {
136  hipLimitMallocHeapSize = 0x02,
137 };
138 
143 #define hipStreamDefault \
145  0x00
146 #define hipStreamNonBlocking 0x01
147 
148 
150 #define hipEventDefault 0x0
151 #define hipEventBlockingSync \
152  0x1
153 #define hipEventDisableTiming \
154  0x2
155 #define hipEventInterprocess 0x4
156 #define hipEventReleaseToDevice \
157  0x40000000
158 #define hipEventReleaseToSystem \
161  0x80000000
162 
165 
167 #define hipHostMallocDefault 0x0
168 #define hipHostMallocPortable 0x1
169 #define hipHostMallocMapped \
170  0x2
171 #define hipHostMallocWriteCombined 0x4
173 #define hipHostMallocCoherent \
174  0x40000000
175 #define hipHostMallocNonCoherent \
177  0x80000000
178 
180 #define hipDeviceMallocDefault 0x0
181 #define hipDeviceMallocFinegrained 0x1
182 
183 #define hipHostRegisterDefault 0x0
185 #define hipHostRegisterPortable 0x1
186 #define hipHostRegisterMapped \
187  0x2
188 #define hipHostRegisterIoMemory 0x4
190 
191 
192 #define hipDeviceScheduleAuto 0x0
193 #define hipDeviceScheduleSpin \
194  0x1
195 #define hipDeviceScheduleYield \
197  0x2
198 #define hipDeviceScheduleBlockingSync 0x4
200 #define hipDeviceScheduleMask 0x7
201 
202 #define hipDeviceMapHost 0x8
203 #define hipDeviceLmemResizeToMax 0x16
204 
205 #define hipArrayDefault 0x00
206 #define hipArrayLayered 0x01
207 #define hipArraySurfaceLoadStore 0x02
208 #define hipArrayCubemap 0x04
209 #define hipArrayTextureGather 0x08
210 
211 /*
212  * @brief hipJitOption
213  * @enum
214  * @ingroup Enumerations
215  */
216 typedef enum hipJitOption {
217  hipJitOptionMaxRegisters = 0,
218  hipJitOptionThreadsPerBlock,
219  hipJitOptionWallTime,
220  hipJitOptionInfoLogBuffer,
221  hipJitOptionInfoLogBufferSizeBytes,
222  hipJitOptionErrorLogBuffer,
223  hipJitOptionErrorLogBufferSizeBytes,
224  hipJitOptionOptimizationLevel,
225  hipJitOptionTargetFromContext,
226  hipJitOptionTarget,
227  hipJitOptionFallbackStrategy,
228  hipJitOptionGenerateDebugInfo,
229  hipJitOptionLogVerbose,
230  hipJitOptionGenerateLineInfo,
231  hipJitOptionCacheMode,
232  hipJitOptionSm3xOpt,
233  hipJitOptionFastCompile,
234  hipJitOptionNumOptions
235 } hipJitOption;
236 
237 
241 typedef enum hipFuncCache_t {
247 
248 
252 typedef enum hipSharedMemConfig {
259 
260 
265 typedef struct dim3 {
266  uint32_t x;
267  uint32_t y;
268  uint32_t z;
269 #ifdef __cplusplus
270  dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){};
271 #endif
272 } dim3;
273 
274 
275 // Doxygen end group GlobalDefs
279 //-------------------------------------------------------------------------------------------------
280 
281 
282 // The handle allows the async commands to use the stream even if the parent hipStream_t goes
283 // out-of-scope.
284 // typedef class ihipStream_t * hipStream_t;
285 
286 
287 /*
288  * Opaque structure allows the true event (pointed at by the handle) to remain "live" even if the
289  * surrounding hipEvent_t goes out-of-scope. This is handy for cases where the hipEvent_t goes
290  * out-of-scope but the true event is being written by some async queue or device */
291 // typedef struct hipEvent_t {
292 // struct ihipEvent_t *_handle;
293 //} hipEvent_t;
294 
295 
321 hipError_t hipDeviceSynchronize(void);
322 
323 
335 hipError_t hipDeviceReset(void);
336 
337 
369 hipError_t hipSetDevice(int deviceId);
370 
371 
385 hipError_t hipGetDevice(int* deviceId);
386 
387 
400 hipError_t hipGetDeviceCount(int* count);
401 
411 hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceId);
412 
426 hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId);
427 
428 
439 hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig);
440 
441 
452 hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig);
453 
464 hipError_t hipDeviceGetLimit(size_t* pValue, enum hipLimit_t limit);
465 
466 
477 hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t config);
478 
490 hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig);
491 
492 
504 hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config);
505 
530 hipError_t hipSetDeviceFlags(unsigned flags);
531 
540 hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop);
541 
542 // end doxygen Device
565 hipError_t hipGetLastError(void);
566 
567 
578 hipError_t hipPeekAtLastError(void);
579 
580 
589 const char* hipGetErrorName(hipError_t hip_error);
590 
591 
602 const char* hipGetErrorString(hipError_t hipError);
603 
604 // end doxygen Error
637 hipError_t hipStreamCreate(hipStream_t* stream);
638 
639 
657 hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags);
658 
659 
678 hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority);
679 
680 
695 hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority);
696 
697 
716 hipError_t hipStreamDestroy(hipStream_t stream);
717 
718 
734 hipError_t hipStreamQuery(hipStream_t stream);
735 
736 
756 hipError_t hipStreamSynchronize(hipStream_t stream);
757 
758 
778 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags);
779 
780 
794 hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags);
795 
796 
810 hipError_t hipStreamGetPriority(hipStream_t stream, int* priority);
811 
812 
816 typedef void (*hipStreamCallback_t)(hipStream_t stream, hipError_t status, void* userData);
817 
833 hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
834  unsigned int flags);
835 
836 
837 // end doxygen Stream
874 hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
875 
876 
888 hipError_t hipEventCreate(hipEvent_t* event);
889 
890 
918 #ifdef __cplusplus
919 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream = NULL);
920 #else
921 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream);
922 #endif
923 
940 hipError_t hipEventDestroy(hipEvent_t event);
941 
942 
960 hipError_t hipEventSynchronize(hipEvent_t event);
961 
962 
991 hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop);
992 
993 
1009 hipError_t hipEventQuery(hipEvent_t event);
1010 
1011 
1012 // end doxygen Events
1043 hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr);
1044 
1058 hipError_t hipMalloc(void** ptr, size_t size);
1059 
1074 hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags);
1075 
1088 DEPRECATED("use hipHostMalloc instead")
1089 hipError_t hipMallocHost(void** ptr, size_t size);
1090 
1104 hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
1105 
1119 DEPRECATED("use hipHostMalloc instead")
1120 hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags);
1121 
1133 hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags);
1134 
1144 hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr);
1145 
1182 hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags);
1183 
1192 hipError_t hipHostUnregister(void* hostPtr);
1193 
1213 hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height);
1214 
1228 hipError_t hipFree(void* ptr);
1229 
1240 DEPRECATED("use hipHostFree instead")
1241 hipError_t hipFreeHost(void* ptr);
1242 
1256 hipError_t hipHostFree(void* ptr);
1257 
1285 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
1286 
1304 hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes);
1305 
1323 hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes);
1324 
1342 hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes);
1343 
1361 hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream);
1362 
1380 hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream);
1381 
1399 hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
1400  hipStream_t stream);
1401 
1402 #if __HIP_VDI__
1403 hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
1404  hipModule_t hmod, const char* name);
1405 
1406 hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName);
1407 hipError_t hipGetSymbolSize(size_t* size, const void* symbolName);
1408 hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
1409  size_t sizeBytes, size_t offset __dparm(0),
1410  hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
1411 hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src,
1412  size_t sizeBytes, size_t offset,
1413  hipMemcpyKind kind, hipStream_t stream __dparm(0));
1414 hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName,
1415  size_t sizeBytes, size_t offset __dparm(0),
1416  hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost));
1417 hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
1418  size_t sizeBytes, size_t offset,
1419  hipMemcpyKind kind,
1420  hipStream_t stream __dparm(0));
1421 #else
1422 __attribute__((visibility("hidden")))
1423 hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
1424 
1425 
1436 inline
1437 __attribute__((visibility("hidden")))
1438 hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
1439  //HIP_INIT_API(hipGetSymbolAddress, devPtr, symbolName);
1440  hip_impl::hip_init();
1441  size_t size = 0;
1442  return hipModuleGetGlobal(devPtr, &size, 0, (const char*)symbolName);
1443 }
1444 
1445 
1456 inline
1457 __attribute__((visibility("hidden")))
1458 hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
1459  // HIP_INIT_API(hipGetSymbolSize, size, symbolName);
1460  hip_impl::hip_init();
1461  void* devPtr = nullptr;
1462  return hipModuleGetGlobal(&devPtr, size, 0, (const char*)symbolName);
1463 }
1464 
1465 #if defined(__cplusplus)
1466 } // extern "C"
1467 #endif
1468 
1469 namespace hip_impl {
1470 hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind,
1471  const char*);
1472 } // Namespace hip_impl.
1473 
1474 #if defined(__cplusplus)
1475 extern "C" {
1476 #endif
1477 
1501 inline
1502 __attribute__((visibility("hidden")))
1503 hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
1504  size_t sizeBytes, size_t offset __dparm(0),
1505  hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)) {
1506  if (!symbolName) return hipErrorInvalidSymbol;
1507 
1508  hipDeviceptr_t dst = NULL;
1509  hipGetSymbolAddress(&dst, (const char*)symbolName);
1510 
1511  return hip_impl::hipMemcpyToSymbol(dst, src, sizeBytes, offset, kind,
1512  (const char*)symbolName);
1513 }
1514 
1515 #if defined(__cplusplus)
1516 } // extern "C"
1517 #endif
1518 
1519 namespace hip_impl {
1520 hipError_t hipMemcpyToSymbolAsync(void*, const void*, size_t, size_t,
1521  hipMemcpyKind, hipStream_t, const char*);
1522 hipError_t hipMemcpyFromSymbol(void*, const void*, size_t, size_t,
1523  hipMemcpyKind, const char*);
1524 hipError_t hipMemcpyFromSymbolAsync(void*, const void*, size_t, size_t,
1525  hipMemcpyKind, hipStream_t, const char*);
1526 } // Namespace hip_impl.
1527 
1528 #if defined(__cplusplus)
1529 extern "C" {
1530 #endif
1531 
1557 inline
1558 __attribute__((visibility("hidden")))
1559 hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src,
1560  size_t sizeBytes, size_t offset,
1561  hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
1562  if (!symbolName) return hipErrorInvalidSymbol;
1563 
1564  hipDeviceptr_t dst = NULL;
1565  hipGetSymbolAddress(&dst, symbolName);
1566 
1567  return hip_impl::hipMemcpyToSymbolAsync(dst, src, sizeBytes, offset, kind,
1568  stream,
1569  (const char*)symbolName);
1570 }
1571 
1572 inline
1573 __attribute__((visibility("hidden")))
1574 hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName,
1575  size_t sizeBytes, size_t offset __dparm(0),
1576  hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)) {
1577  if (!symbolName) return hipErrorInvalidSymbol;
1578 
1579  hipDeviceptr_t src = NULL;
1580  hipGetSymbolAddress(&src, symbolName);
1581 
1582  return hip_impl::hipMemcpyFromSymbol(dst, src, sizeBytes, offset, kind,
1583  (const char*)symbolName);
1584 }
1585 
1586 inline
1587 __attribute__((visibility("hidden")))
1588 hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
1589  size_t sizeBytes, size_t offset,
1590  hipMemcpyKind kind,
1591  hipStream_t stream __dparm(0)) {
1592  if (!symbolName) return hipErrorInvalidSymbol;
1593 
1594  hipDeviceptr_t src = NULL;
1595  hipGetSymbolAddress(&src, symbolName);
1596 
1597  return hip_impl::hipMemcpyFromSymbolAsync(dst, src, sizeBytes, offset, kind,
1598  stream,
1599  (const char*)symbolName);
1600 }
1601 
1602 #endif // __HIP_VDI__
1603 
1631 hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
1632  hipStream_t stream __dparm(0));
1633 
1643 hipError_t hipMemset(void* dst, int value, size_t sizeBytes);
1644 
1654 hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes);
1655 
1665 hipError_t hipMemsetD32(hipDeviceptr_t dest, int value, size_t count);
1666 
1682 hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream __dparm(0));
1683 
1699 hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count,
1700  hipStream_t stream __dparm(0));
1701 
1713 hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height);
1714 
1727 hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height,hipStream_t stream __dparm(0));
1728 
1737 hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent );
1738 
1748 hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent ,hipStream_t stream __dparm(0));
1749 
1759 hipError_t hipMemGetInfo(size_t* free, size_t* total);
1760 
1761 
1762 hipError_t hipMemPtrGetInfo(void* ptr, size_t* size);
1763 
1764 
1777 hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width,
1778  size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault));
1779 hipError_t hipArrayCreate(hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray);
1780 
1781 hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray);
1782 
1783 hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent);
1784 
1793 hipError_t hipFreeArray(hipArray* array);
1794 
1807 hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc,
1808  struct hipExtent extent, unsigned int flags);
1825 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
1826  size_t height, hipMemcpyKind kind);
1827 hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy);
1828 
1846 hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
1847  size_t height, hipMemcpyKind kind, hipStream_t stream __dparm(0));
1848 
1865 hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
1866  size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
1867 
1884 hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
1885  size_t count, hipMemcpyKind kind);
1886 
1902 hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
1903  size_t count, hipMemcpyKind kind);
1904 
1918 hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count);
1919 
1933 hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count);
1934 
1945 hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p);
1946 
1947 // doxygen end Memory
1979 hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId);
1980 
1981 
1998 hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags);
1999 
2000 
2012 hipError_t hipDeviceDisablePeerAccess(int peerDeviceId);
2013 
2026 hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr);
2027 
2028 #ifndef USE_PEER_NON_UNIFIED
2029 #define USE_PEER_NON_UNIFIED 1
2030 #endif
2031 
2032 #if USE_PEER_NON_UNIFIED == 1
2033 
2044 hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId,
2045  size_t sizeBytes);
2046 
2059 hipError_t hipMemcpyPeerAsync(void* dst, int dstDeviceId, const void* src, int srcDevice,
2060  size_t sizeBytes, hipStream_t stream __dparm(0));
2061 #endif
2062 
2063 
2064 // doxygen end PeerToPeer
2083 // TODO-ctx - more description on error codes.
2084 hipError_t hipInit(unsigned int flags);
2085 
2086 
2106 DEPRECATED(DEPRECATED_MSG)
2107 hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device);
2108 
2119 DEPRECATED(DEPRECATED_MSG)
2120 hipError_t hipCtxDestroy(hipCtx_t ctx);
2121 
2132 DEPRECATED(DEPRECATED_MSG)
2133 hipError_t hipCtxPopCurrent(hipCtx_t* ctx);
2134 
2145 DEPRECATED(DEPRECATED_MSG)
2146 hipError_t hipCtxPushCurrent(hipCtx_t ctx);
2147 
2158 DEPRECATED(DEPRECATED_MSG)
2159 hipError_t hipCtxSetCurrent(hipCtx_t ctx);
2160 
2171 DEPRECATED(DEPRECATED_MSG)
2172 hipError_t hipCtxGetCurrent(hipCtx_t* ctx);
2173 
2185 DEPRECATED(DEPRECATED_MSG)
2186 hipError_t hipCtxGetDevice(hipDevice_t* device);
2187 
2205 DEPRECATED(DEPRECATED_MSG)
2206 hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion);
2207 
2221 DEPRECATED(DEPRECATED_MSG)
2222 hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig);
2223 
2237 DEPRECATED(DEPRECATED_MSG)
2238 hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig);
2239 
2253 DEPRECATED(DEPRECATED_MSG)
2254 hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config);
2255 
2269 DEPRECATED(DEPRECATED_MSG)
2270 hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig);
2271 
2283 DEPRECATED(DEPRECATED_MSG)
2284 hipError_t hipCtxSynchronize(void);
2285 
2296 DEPRECATED(DEPRECATED_MSG)
2297 hipError_t hipCtxGetFlags(unsigned int* flags);
2298 
2318 DEPRECATED(DEPRECATED_MSG)
2319 hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags);
2320 
2337 DEPRECATED(DEPRECATED_MSG)
2338 hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx);
2339 
2352 hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active);
2353 
2366 hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev);
2367 
2379 hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev);
2380 
2391 hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev);
2392 
2404 hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags);
2405 
2406 // doxygen end Context Management
2418 hipError_t hipDeviceGet(hipDevice_t* device, int ordinal);
2419 
2428 hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device);
2429 
2438 hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
2439 
2448 hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device);
2449 
2450 
2458 hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId);
2459 
2460 
2468 hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device);
2469 
2485 hipError_t hipDriverGetVersion(int* driverVersion);
2486 
2499 hipError_t hipRuntimeGetVersion(int* runtimeVersion);
2500 
2512 hipError_t hipModuleLoad(hipModule_t* module, const char* fname);
2513 
2524 hipError_t hipModuleUnload(hipModule_t module);
2525 
2536 hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname);
2537 
2547 hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func);
2548 
2550 
2551  Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
2552  Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt)
2553  : name(nullptr), address(address), byte_cnt(byte_cnt) {
2554  if (name)
2555  this->name = strdup(name);
2556  }
2557 
2558  Agent_global& operator=(Agent_global&& t) {
2559  if (this == &t) return *this;
2560 
2561  if (name) free(name);
2562  name = t.name;
2563  address = t.address;
2564  byte_cnt = t.byte_cnt;
2565 
2566  t.name = nullptr;
2567  t.address = nullptr;
2568  t.byte_cnt = 0;
2569 
2570  return *this;
2571  }
2572 
2573  Agent_global(Agent_global&& t)
2574  : name(nullptr), address(nullptr), byte_cnt(0) {
2575  *this = std::move(t);
2576  }
2577 
2578  // not needed, delete them to prevent bugs
2579  Agent_global(const Agent_global&) = delete;
2580  Agent_global& operator=(Agent_global& t) = delete;
2581 
2582  ~Agent_global() { if (name) free(name); }
2583 
2584  char* name;
2585  hipDeviceptr_t address;
2586  uint32_t byte_cnt;
2587 };
2588 
2589 #if !__HIP_VDI__
2590 #if defined(__cplusplus)
2591 } // extern "C"
2592 #endif
2593 
2594 namespace hip_impl {
2595 hsa_executable_t executable_for(hipModule_t);
2596 const char* hash_for(hipModule_t);
2597 
2598 template<typename ForwardIterator>
2599 std::pair<hipDeviceptr_t, std::size_t> read_global_description(
2600  ForwardIterator f, ForwardIterator l, const char* name) {
2601  const auto it = std::find_if(f, l, [=](const Agent_global& x) {
2602  return strcmp(x.name, name) == 0;
2603  });
2604 
2605  return it == l ?
2606  std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
2607 }
2608 
2609 std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
2610  hsa_executable_t executable);
2611 hsa_agent_t this_agent();
2612 
2613 inline
2614 __attribute__((visibility("hidden")))
2615 hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
2616  hipModule_t hmod, const char* name) {
2617  // the key of the map would the hash of code object associated with the
2618  // hipModule_t instance
2619  static std::unordered_map<
2620  std::string, std::vector<Agent_global>> agent_globals;
2621  std::string key(hash_for(hmod));
2622 
2623  if (agent_globals.count(key) == 0) {
2624  static std::mutex mtx;
2625  std::lock_guard<std::mutex> lck{mtx};
2626 
2627  if (agent_globals.count(key) == 0) {
2628  agent_globals.emplace(
2629  key, read_agent_globals(this_agent(), executable_for(hmod)));
2630  }
2631  }
2632 
2633  const auto it0 = agent_globals.find(key);
2634  if (it0 == agent_globals.cend()) {
2635  hip_throw(
2636  std::runtime_error{"agent_globals data structure corrupted."});
2637  }
2638 
2639  std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
2640  it0->second.cend(), name);
2641 
2642  // HACK for SWDEV-173477
2643  //
2644  // For code objects with global symbols of length 0, ROCR runtime would
2645  // ignore them even though they exist in the symbol table. Therefore the
2646  // result from read_agent_globals() can't be trusted entirely.
2647  //
2648  // As a workaround to tame applications which depend on the existence of
2649  // global symbols with length 0, always return hipSuccess here.
2650  //
2651  // This behavior shall be reverted once ROCR runtime has been fixed to
2652  // address SWDEV-173477
2653 
2654  //return *dptr ? hipSuccess : hipErrorNotFound;
2655  return hipSuccess;
2656 }
2657 
2658 inline
2659 __attribute__((visibility("hidden")))
2660 hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
2661  const char* name) {
2662  static std::unordered_map<
2663  hsa_agent_t, std::vector<Agent_global>> agent_globals;
2664  static std::once_flag f;
2665 
2666  std::call_once(f, []() {
2667  for (auto&& agent_executables : executables()) {
2668  std::vector<Agent_global> tmp0;
2669  for (auto&& executable : agent_executables.second) {
2670  auto tmp1 = read_agent_globals(agent_executables.first,
2671  executable);
2672 
2673  tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
2674  make_move_iterator(tmp1.end()));
2675  }
2676  agent_globals.emplace(agent_executables.first, move(tmp0));
2677  }
2678  });
2679 
2680  const auto it = agent_globals.find(this_agent());
2681 
2682  if (it == agent_globals.cend()) return hipErrorNotInitialized;
2683 
2684  std::tie(*dptr, *bytes) = read_global_description(it->second.cbegin(),
2685  it->second.cend(), name);
2686 
2687  return *dptr ? hipSuccess : hipErrorNotFound;
2688 }
2689 } // Namespace hip_impl.
2690 
2691 #if defined(__cplusplus)
2692 extern "C" {
2693 #endif
2694 
2705 inline
2706 __attribute__((visibility("hidden")))
2707 hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
2708  hipModule_t hmod, const char* name) {
2709  if (!dptr || !bytes) return hipErrorInvalidValue;
2710 
2711  if (!name) return hipErrorNotInitialized;
2712 
2713  const auto r = hmod ?
2714  hip_impl::read_agent_global_from_module(dptr, bytes, hmod, name) :
2715  hip_impl::read_agent_global_from_process(dptr, bytes, name);
2716 
2717  return r;
2718 }
2719 #endif // __HIP_VDI__
2720 
2721 hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
2731 hipError_t hipModuleLoadData(hipModule_t* module, const void* image);
2732 
2745 hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions,
2746  hipJitOption* options, void** optionValues);
2747 
2772 hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY,
2773  unsigned int gridDimZ, unsigned int blockDimX,
2774  unsigned int blockDimY, unsigned int blockDimZ,
2775  unsigned int sharedMemBytes, hipStream_t stream,
2776  void** kernelParams, void** extra);
2777 
2778 // doxygen end Version Management
2796 // TODO - expand descriptions:
2802 hipError_t hipProfilerStart();
2803 
2804 
2810 hipError_t hipProfilerStop();
2811 
2812 
2817 // TODO: implement IPC apis
2818 
2844 hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
2845 
2882 hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
2883 
2902 hipError_t hipIpcCloseMemHandle(void* devPtr);
2903 
2904 
2905 // hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
2906 // hipError_t hipIpcCloseMemHandle(void *devPtr);
2907 // // hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
2908 // hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
2909 
2910 
2931 hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem __dparm(0), hipStream_t stream __dparm(0));
2932 
2933 
2944 hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset);
2945 
2946 
2955 hipError_t hipLaunchByPtr(const void* func);
2956 
2957 
2958 
2964 #ifdef __cplusplus
2965 } /* extern "c" */
2966 #endif
2967 
2968 #include <hip/hcc_detail/hip_prof_api.h>
2969 
2970 #ifdef __cplusplus
2971 extern "C" {
2972 #endif
2973 
2976 hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg);
2977 hipError_t hipRemoveApiCallback(uint32_t id);
2978 hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg);
2979 hipError_t hipRemoveActivityCallback(uint32_t id);
2980 static inline const char* hipApiName(const uint32_t& id) { return hip_api_name(id); }
2981 const char* hipKernelNameRef(const hipFunction_t f);
2982 #ifdef __cplusplus
2983 } /* extern "C" */
2984 #endif
2985 
2986 #ifdef __cplusplus
2987 
2988 hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr,
2989  const hipChannelFormatDesc* desc, size_t size = UINT_MAX);
2990 
2991 hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t* offset,
2992  const void* devPtr, const struct hipChannelFormatDesc* desc,
2993  size_t size, textureReference* tex);
2994 
2995 /*
2996  * @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture
2997  *reference tex.
2998  *
2999  * @p desc describes how the memory is interpreted when fetching values from the texture. The @p
3000  *offset parameter is an optional byte offset as with the low-level hipBindTexture() function. Any
3001  *memory previously bound to tex is unbound.
3002  *
3003  * @param[in] offset - Offset in bytes
3004  * @param[out] tex - texture to bind
3005  * @param[in] devPtr - Memory area on device
3006  * @param[in] desc - Channel format
3007  * @param[in] size - Size of the memory area pointed to by devPtr
3008  * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
3009  **/
3010 template <class T, int dim, enum hipTextureReadMode readMode>
3011 hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex, const void* devPtr,
3012  const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) {
3013  return ihipBindTextureImpl(dim, readMode, offset, devPtr, &desc, size, &tex);
3014 }
3015 
3016 /*
3017  * @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture
3018  *reference tex.
3019  *
3020  * @p desc describes how the memory is interpreted when fetching values from the texture. The @p
3021  *offset parameter is an optional byte offset as with the low-level hipBindTexture() function. Any
3022  *memory previously bound to tex is unbound.
3023  *
3024  * @param[in] offset - Offset in bytes
3025  * @param[in] tex - texture to bind
3026  * @param[in] devPtr - Memory area on device
3027  * @param[in] size - Size of the memory area pointed to by devPtr
3028  * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
3029  **/
3030 template <class T, int dim, enum hipTextureReadMode readMode>
3031 hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& tex, const void* devPtr,
3032  size_t size = UINT_MAX) {
3033  return ihipBindTextureImpl(dim, readMode, offset, devPtr, &(tex.channelDesc), size, &tex);
3034 }
3035 
3036 // C API
3037 hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr,
3038  const hipChannelFormatDesc* desc, size_t width, size_t height,
3039  size_t pitch);
3040 
3041 hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size_t* offset,
3042  const void* devPtr, const struct hipChannelFormatDesc* desc,
3043  size_t width, size_t height, textureReference* tex);
3044 
3045 template <class T, int dim, enum hipTextureReadMode readMode>
3046 hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& tex,
3047  const void* devPtr, size_t width, size_t height, size_t pitch) {
3048  return ihipBindTexture2DImpl(dim, readMode, offset, devPtr, &(tex.channelDesc), width, height,
3049  &tex);
3050 }
3051 
3052 template <class T, int dim, enum hipTextureReadMode readMode>
3053 hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& tex,
3054  const void* devPtr, const struct hipChannelFormatDesc& desc,
3055  size_t width, size_t height, size_t pitch) {
3056  return ihipBindTexture2DImpl(dim, readMode, offset, devPtr, &desc, width, height, &tex);
3057 }
3058 
3059 // C API
3060 hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array,
3061  const hipChannelFormatDesc* desc);
3062 
3063 hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode,
3064  hipArray_const_t array,
3065  const struct hipChannelFormatDesc& desc,
3066  textureReference* tex);
3067 
3068 template <class T, int dim, enum hipTextureReadMode readMode>
3069 hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex, hipArray_const_t array) {
3070  return ihipBindTextureToArrayImpl(dim, readMode, array, tex.channelDesc, &tex);
3071 }
3072 
3073 template <class T, int dim, enum hipTextureReadMode readMode>
3074 hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex, hipArray_const_t array,
3075  const struct hipChannelFormatDesc& desc) {
3076  return ihipBindTextureToArrayImpl(dim, readMode, array, desc, &tex);
3077 }
3078 
3079 template <class T, int dim, enum hipTextureReadMode readMode>
3080 inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *tex,
3081  hipArray_const_t array,
3082  const struct hipChannelFormatDesc* desc) {
3083  return ihipBindTextureToArrayImpl(dim, readMode, array, *desc, tex);
3084 }
3085 
3086 // C API
3087 hipError_t hipBindTextureToMipmappedArray(const textureReference* tex,
3088  hipMipmappedArray_const_t mipmappedArray,
3089  const hipChannelFormatDesc* desc);
3090 
3091 template <class T, int dim, enum hipTextureReadMode readMode>
3092 hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
3093  hipMipmappedArray_const_t mipmappedArray) {
3094  return hipSuccess;
3095 }
3096 
3097 template <class T, int dim, enum hipTextureReadMode readMode>
3098 hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
3099  hipMipmappedArray_const_t mipmappedArray,
3100  const hipChannelFormatDesc& desc) {
3101  return hipSuccess;
3102 }
3103 
3104 /*
3105  * @brief Unbinds the textuer bound to @p tex
3106  *
3107  * @param[in] tex - texture to unbind
3108  *
3109  * @return #hipSuccess
3110  **/
3111 hipError_t hipUnbindTexture(const textureReference* tex);
3112 
3113 extern hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject);
3114 
3115 template <class T, int dim, enum hipTextureReadMode readMode>
3116 hipError_t hipUnbindTexture(struct texture<T, dim, readMode>& tex) {
3117  return ihipUnbindTextureImpl(tex.textureObject);
3118 }
3119 
3120 hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array);
3121 hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref);
3122 hipError_t hipGetTextureReference(const textureReference** texref, const void* symbol);
3123 
3124 hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc,
3125  const hipTextureDesc* pTexDesc,
3126  const hipResourceViewDesc* pResViewDesc);
3127 
3128 hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
3129 
3130 hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc,
3131  hipTextureObject_t textureObject);
3132 hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc,
3133  hipTextureObject_t textureObject);
3134 hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc,
3135  hipTextureObject_t textureObject);
3136 hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsigned int flags);
3137 
3138 hipError_t hipTexRefSetAddressMode(textureReference* tex, int dim, hipTextureAddressMode am);
3139 
3140 hipError_t hipTexRefSetFilterMode(textureReference* tex, hipTextureFilterMode fm);
3141 
3142 hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags);
3143 
3144 hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int NumPackedComponents);
3145 
3146 hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDeviceptr_t devPtr,
3147  size_t size);
3148 
3149 hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc,
3150  hipDeviceptr_t devPtr, size_t pitch);
3151 
3152 hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
3153 
3154 hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
3155 
3156 // doxygen end Texture
3162 #endif
3163 
3164 
3182 // end-group HCC_Specific
3188 // doxygen end HIP API
3193 #endif
hipError_t hipHostFree(void *ptr)
Free memory allocated by the hcc hip host memory allocation API This API performs an implicit hipDevi...
Definition: hip_memory.cpp:1926
prefer larger L1 cache and smaller shared memory
Definition: hip_runtime_api.h:244
hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, const char *kname)
Function with kname will be extracted if present in module.
Definition: hip_module.cpp:455
hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx)
Disable direct access from current context&#39;s virtual address space to memory allocations physically l...
Definition: hip_peer.cpp:218
hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent)
Fills synchronously the memory area pointed to by pitchedDevPtr with the constant value...
Definition: hip_memory.cpp:1799
hipError_t hipStreamGetPriority(hipStream_t stream, int *priority)
Query the priority of a stream.
Definition: hip_stream.cpp:234
hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig)
Set Cache configuration for a specific function.
Definition: hip_device.cpp:84
hipError_t hipDriverGetVersion(int *driverVersion)
Returns the approximate HIP driver version.
Definition: hip_context.cpp:88
hipError_t hipPeekAtLastError(void)
Return last error returned by any HIP runtime API call.
Definition: hip_error.cpp:41
hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags)
Set flags for the primary context.
Definition: hip_context.cpp:324
hipError_t hipDevicePrimaryCtxRetain(hipCtx_t *pctx, hipDevice_t dev)
Retain the primary context on the GPU.
Definition: hip_context.cpp:299
struct dim3 dim3
hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device)
Returns the compute capability of the device.
Definition: hip_device.cpp:369
hipError_t hipDeviceGetByPCIBusId(int *device, const char *pciBusId)
Returns a handle to a compute device.
Definition: hip_device.cpp:427
hipError_t hipExtMallocWithFlags(void **ptr, size_t sizeBytes, unsigned int flags)
Allocate memory on the default accelerator.
Definition: hip_memory.cpp:267
hipError_t hipMemGetAddressRange(hipDeviceptr_t *pbase, size_t *psize, hipDeviceptr_t dptr)
Get information on memory allocations.
Definition: hip_memory.cpp:1989
TODO-doc.
hipError_t hipFreeHost(void *ptr)
Free memory allocated by the hcc hip host memory allocation API. [Deprecated].
Definition: hip_memory.cpp:1959
hipError_t hipMemcpyToArray(hipArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, hipMemcpyKind kind)
Copies data between host and device.
Definition: hip_memory.cpp:1285
unsigned long long hipSurfaceObject_t
Definition: hip_surface_types.h:36
hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues)
builds module from code object which resides in host memory. Image is pointer to that location...
Definition: hip_module.cpp:567
Definition: driver_types.h:232
hipError_t hipMallocPitch(void **ptr, size_t *pitch, size_t width, size_t height)
Definition: hip_memory.cpp:440
hipError_t hipMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream __dparm(0))
Copies data between host and device.
hipError_t hipMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind)
Copies data between host and device.
Definition: hip_memory.cpp:1629
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void *userData, unsigned int flags)
Adds a callback to be called on the host after all currently enqueued items in the stream have comple...
Definition: hip_stream.cpp:254
hipError_t hipMemcpyFromArray(void *dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind)
Copies data between host and device.
Definition: hip_memory.cpp:1304
uint32_t x
x
Definition: hip_runtime_api.h:266
hipError_t hipMemcpyAtoH(void *dst, hipArray *srcArray, size_t srcOffset, size_t count)
Copies data between host and device.
Definition: hip_memory.cpp:1341
hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags)
Enable direct access from current device&#39;s virtual address space to memory allocations physically loc...
Definition: hip_peer.cpp:191
hipError_t hipCtxPopCurrent(hipCtx_t *ctx)
Pop the current/default context and return the popped context.
Definition: hip_context.cpp:136
const char * hipGetErrorString(hipError_t hipError)
Return handy text string message to explain the error which occurred.
Definition: hip_error.cpp:54
hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig *pConfig)
Returns bank width of shared memory for current device.
Definition: hip_device.cpp:125
Definition: hip_runtime_api.h:137
hipError_t hipDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority)
Returns numerical values that correspond to the least and greatest stream priority.
Definition: hip_stream.cpp:118
prefer equal size L1 cache and shared memory
Definition: hip_runtime_api.h:245
hipError_t hipHostGetDevicePointer(void **devPtr, void *hstPtr, unsigned int flags)
Get Device pointer from Host Pointer allocated through hipHostMalloc.
hipError_t hipFreeArray(hipArray *array)
Frees an array on the device.
Definition: hip_memory.cpp:1961
hipError_t hipStreamCreateWithPriority(hipStream_t *stream, unsigned int flags, int priority)
Create an asynchronous stream with the specified priority.
Definition: hip_stream.cpp:109
hipError_t hipMemsetAsync(void *dst, int value, size_t sizeBytes, hipStream_t stream __dparm(0))
Fills the first sizeBytes bytes of the memory area pointed to by dev with the constant byte value val...
hipError_t hipEventSynchronize(hipEvent_t event)
Wait for an event to complete.
Definition: hip_event.cpp:167
Definition: driver_types.h:173
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void *src, size_t sizeBytes)
Copy data from Host to Device.
Definition: hip_memory.cpp:1130
hipError_t hipModuleUnload(hipModule_t module)
Frees the module.
Definition: hip_module.cpp:114
Definition: hip_module.cpp:89
hipError_t hipSetDeviceFlags(unsigned flags)
The current device behavior is changed according the flags passed.
hipError_t hipEventQuery(hipEvent_t event)
Query event status.
Definition: hip_event.cpp:257
hipError_t hipDeviceDisablePeerAccess(int peerDeviceId)
Disable direct access from current device&#39;s virtual address space to memory allocations physically lo...
Definition: hip_peer.cpp:184
hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig *pConfig)
Get Shared memory bank configuration.
Definition: hip_context.cpp:244
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device)
Create a context and set it as current/ default context.
Definition: hip_context.cpp:55
#define hipArrayDefault
Default HIP array allocation flag.
Definition: hip_runtime_api.h:205
hipError_t hipMallocArray(hipArray **array, const hipChannelFormatDesc *desc, size_t width, size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault))
Allocate an array on the device.
hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config)
Set Shared memory bank configuration.
Definition: hip_context.cpp:236
hipError_t hipCtxSetCurrent(hipCtx_t ctx)
Set the passed context as current/default.
Definition: hip_context.cpp:181
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *devPtr)
Gets an interprocess memory handle for an existing device memory allocation.
Definition: hip_memory.cpp:2010
Definition: hip_runtime_api.h:120
hipError_t hipMemset2DAsync(void *dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream __dparm(0))
Fills asynchronously the memory area pointed to by dst with the constant value.
Definition: driver_types.h:107
hipError_t hipHostMalloc(void **ptr, size_t size, unsigned int flags)
Allocate device accessible page locked host memory.
Definition: hip_memory.cpp:304
Definition: hip_hcc_internal.h:880
Definition: texture_types.h:73
hipError_t hipDeviceGetLimit(size_t *pValue, enum hipLimit_t limit)
Get Resource limits of current device.
Definition: hip_device.cpp:96
hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
builds module from code object which resides in host memory. Image is pointer to that location...
Definition: hip_module.cpp:548
Definition: driver_types.h:70
hipDeviceAttribute_t
Definition: hip_runtime_api.h:259
hipError_t hipEventDestroy(hipEvent_t event)
Destroy the specified event.
Definition: hip_event.cpp:155
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
Create an asynchronous stream.
Definition: hip_stream.cpp:95
hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem __dparm(0), hipStream_t stream __dparm(0))
Configure a kernel launch.
hipError_t hipChooseDevice(int *device, const hipDeviceProp_t *prop)
Device which matches hipDeviceProp_t is returned.
Definition: hip_device.cpp:453
hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig)
Set L1/Shared cache partition.
Definition: hip_context.cpp:228
hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t stream, void **kernelParams, void **extra)
launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelp...
Definition: hip_runtime_api.h:265
uint32_t y
y
Definition: hip_runtime_api.h:267
void(* hipStreamCallback_t)(hipStream_t stream, hipError_t status, void *userData)
Definition: hip_runtime_api.h:816
hipError_t hipModuleLoad(hipModule_t *module, const char *fname)
Loads code object from file into a hipModule_t.
Definition: hip_module.cpp:553
hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev)
Resets the primary context on the GPU.
Definition: hip_context.cpp:311
hipError_t hipEventCreateWithFlags(hipEvent_t *event, unsigned flags)
Create an event with the specified flags.
Definition: hip_event.cpp:97
hipError_t hipHostAlloc(void **ptr, size_t size, unsigned int flags)
Allocate device accessible page locked host memory [Deprecated].
Definition: hip_memory.cpp:379
hipError_t hipMallocHost(void **ptr, size_t size)
Allocate pinned host memory [Deprecated].
Definition: hip_memory.cpp:375
hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
Return the elapsed time between two events.
Definition: hip_event.cpp:200
hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig)
Set L1/Shared cache partition.
Definition: hip_device.cpp:76
hipError_t hipDeviceCanAccessPeer(int *canAccessPeer, int deviceId, int peerDeviceId)
Determine if a device can access a peer&#39;s memory.
Definition: hip_peer.cpp:177
hipError_t hipGetDeviceCount(int *count)
Return number of compute-capable devices.
Definition: hip_device.cpp:71
hipError_t hipMemset(void *dst, int value, size_t sizeBytes)
Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant byte value va...
Definition: hip_memory.cpp:1715
Definition: driver_types.h:245
hipError_t hipStreamDestroy(hipStream_t stream)
Destroys the specified stream.
Definition: hip_stream.cpp:191
hipError_t hipHostGetFlags(unsigned int *flagsPtr, void *hostPtr)
Return flags associated with host pointer.
Definition: hip_memory.cpp:888
hipError_t hipStreamSynchronize(hipStream_t stream)
Wait for all commands in stream to complete.
Definition: hip_stream.cpp:180
hipError_t hipIpcOpenMemHandle(void **devPtr, hipIpcMemHandle_t handle, unsigned int flags)
Opens an interprocess memory handle exported from another process and returns a device pointer usable...
Definition: hip_memory.cpp:2046
hipError_t hipCtxGetCacheConfig(hipFuncCache_t *cacheConfig)
Set Cache configuration for a specific function.
Definition: hip_context.cpp:220
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes)
Copy data from Device to Device.
Definition: hip_memory.cpp:1168
hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream __dparm(0))
Fills asynchronously the memory area pointed to by pitchedDevPtr with the constant value...
Definition: hip_runtime_api.h:254
no preference for shared memory or L1 (default)
Definition: hip_runtime_api.h:242
hipError_t hipCtxSynchronize(void)
Blocks until the default context has completed all preceding requested tasks.
Definition: hip_context.cpp:252
hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream __dparm(0))
Fills the memory area pointed to by dev with the constant integer value for specified number of times...
hipError_t hipCtxGetCurrent(hipCtx_t *ctx)
Get the handle of the current/ default context.
Definition: hip_context.cpp:170
hipError_t hipMalloc3DArray(hipArray **array, const struct hipChannelFormatDesc *desc, struct hipExtent extent, unsigned int flags)
Allocate an array on the device.
Definition: hip_memory.cpp:800
hipError_t hipMemcpyDtoHAsync(void *dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
Copy data from Device to Host asynchronously.
Definition: hip_memory.cpp:1226
hipError_t hipDeviceSynchronize(void)
Waits on all active streams on current device.
Definition: hip_device.cpp:144
hipError_t hipCtxPushCurrent(hipCtx_t ctx)
Push the context to be set as current/ default context.
Definition: hip_context.cpp:157
hipError_t hipMemcpyDtoH(void *dst, hipDeviceptr_t src, size_t sizeBytes)
Copy data from Device to Host.
Definition: hip_memory.cpp:1149
Defines the different newt vector types for HIP runtime.
hipError_t hipMemcpyAsync(void *dst, const void *src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream __dparm(0))
Copy data from src to dst asynchronously.
hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device)
Returns an identifer string for the device.
Definition: hip_device.cpp:381
hipError_t hipGetDeviceProperties(hipDeviceProp_t *prop, int deviceId)
Returns device properties.
Definition: hip_device.cpp:316
hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
Register host memory so it can be accessed from the current device.
Definition: hip_memory.cpp:916
Definition: driver_types.h:61
hipError_t hipMalloc(void **ptr, size_t size)
Allocate memory on the default accelerator.
Definition: hip_memory.cpp:239
const char * hipGetErrorName(hipError_t hip_error)
Return name of the specified error code in text form.
Definition: hip_error.cpp:48
hipError_t hipMemset2D(void *dst, size_t pitch, int value, size_t width, size_t height)
Fills the memory area pointed to by dst with the constant value.
Definition: hip_memory.cpp:1731
Definition: driver_types.h:225
hipFuncCache_t
Definition: hip_runtime_api.h:241
hipError_t hipGetLastError(void)
Return last error returned by any HIP runtime API call and resets the stored error code to #hipSucces...
Definition: hip_error.cpp:32
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
Make the specified compute stream wait for an event.
Definition: hip_stream.cpp:126
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags)
Return flags associated with this stream.
Definition: hip_stream.cpp:219
Defines surface types for HIP runtime.
hipError_t hipMemGetInfo(size_t *free, size_t *total)
Query memory info. Return snapshot of free memory, and total allocatable memory on the device...
Definition: hip_memory.cpp:1835
hipError_t hipCtxGetDevice(hipDevice_t *device)
Get the handle of the device associated with current/default context.
Definition: hip_context.cpp:194
hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev)
Release the primary context on the GPU.
Definition: hip_context.cpp:288
hipError_t hipFree(void *ptr)
Free memory allocated by the hcc hip memory allocation API. This API performs an implicit hipDeviceSy...
Definition: hip_memory.cpp:1894
uint32_t z
z
Definition: hip_runtime_api.h:268
hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int *apiVersion)
Returns the approximate HIP api version.
Definition: hip_context.cpp:210
hipError_t hipDeviceReset(void)
The state of current device is discarded and updated to a fresh state.
Definition: hip_device.cpp:149
hipError_t hipInit(unsigned int flags)
Explicitly initializes the HIP runtime.
Definition: hip_context.cpp:42
hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
Copies data between host and device.
Definition: hip_memory.cpp:1360
hipError_t hipRuntimeGetVersion(int *runtimeVersion)
Returns the approximate HIP Runtime version.
Definition: hip_context.cpp:100
hipError_t hipFuncGetAttributes(hipFuncAttributes *attr, const void *func)
Definition: hip_module.cpp:486
hipError_t hipMemcpyHtoA(hipArray *dstArray, size_t dstOffset, const void *srcHost, size_t count)
Copies data between host and device.
Definition: hip_memory.cpp:1323
hipError_t hipSetupArgument(const void *arg, size_t size, size_t offset)
Set a kernel argument.
Definition: hip_clang.cpp:183
hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
Copy data from Device to Device asynchronously.
Definition: hip_memory.cpp:1218
Definition: hip_runtime_api.h:83
hipError_t hipSetDevice(int deviceId)
Set default device to be used for subsequent hip API calls from this thread.
Definition: hip_device.cpp:133
hipError_t hipDeviceGet(hipDevice_t *device, int ordinal)
Returns a handle to a compute device.
Definition: hip_context.cpp:73
hipError_t hipDeviceTotalMem(size_t *bytes, hipDevice_t device)
Returns the total amount of memory on the device.
Definition: hip_device.cpp:415
Definition: hip_runtime_api.h:2549
hipError_t hipFuncSetCacheConfig(const void *func, hipFuncCache_t config)
Set Cache configuration for a specific function.
Definition: hip_device.cpp:109
The compiler selects a device-specific value for the banking.
Definition: hip_runtime_api.h:253
Definition: hip_runtime_api.h:81
hipError_t hipMemcpyPeerAsync(void *dst, int dstDeviceId, const void *src, int srcDevice, size_t sizeBytes, hipStream_t stream __dparm(0))
Copies memory from one device to memory on another device.
hipError_t hipCtxGetFlags(unsigned int *flags)
Return flags used for creating default context.
Definition: hip_context.cpp:257
__attribute__((visibility("hidden"))) hipError_t hipModuleGetGlobal(void **
Copies the memory address of symbol symbolName to devPtr.
hipError_t hipCtxDestroy(hipCtx_t ctx)
Destroy a HIP context.
Definition: hip_context.cpp:112
hipError_t hipRegisterApiCallback(uint32_t id, void *fun, void *arg)
Definition: hip_intercept.cpp:33
hipSharedMemConfig
Definition: hip_runtime_api.h:252
Definition: driver_types.h:38
Definition: hip_hcc_internal.h:700
hipError_t hipDeviceGetAttribute(int *pi, hipDeviceAttribute_t attr, int deviceId)
Query for a specific device attribute.
Definition: hip_device.cpp:289
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void *src, size_t sizeBytes, hipStream_t stream)
Copy data from Host to Device asynchronously.
Definition: hip_memory.cpp:1211
hipError_t hipHostUnregister(void *hostPtr)
Un-register host pointer.
Definition: hip_memory.cpp:981
Definition: hip_hcc_internal.h:519
hipError_t hipMemcpyPeer(void *dst, int dstDeviceId, const void *src, int srcDeviceId, size_t sizeBytes)
Copies memory from one device to memory on another device.
Definition: hip_peer.cpp:198
hipError_t hipStreamCreate(hipStream_t *stream)
Create an asynchronous stream.
Definition: hip_stream.cpp:102
hipError_t hipMemcpy(void *dst, const void *src, size_t sizeBytes, hipMemcpyKind kind)
Copy data from src to dst.
Definition: hip_memory.cpp:1103
hipError_t hipEventCreate(hipEvent_t *event)
Definition: hip_event.cpp:103
Definition: driver_types.h:82
Definition: hip_runtime_api.h:105
hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int *flags, int *active)
Get the state of the primary context.
Definition: hip_context.cpp:266
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags)
Enables direct access to memory allocations in a peer context.
Definition: hip_peer.cpp:212
Definition: driver_types.h:201
hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config)
The bank width of shared memory on current device is set.
Definition: hip_device.cpp:117
hipError_t hipGetDevice(int *deviceId)
Return the default device id for the calling host thread.
Definition: hip_device.cpp:32
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
Record an event in the specified stream.
Definition: hip_event.cpp:110
Definition: hip_hcc_internal.h:367
hipError_t hipDeviceGetPCIBusId(char *pciBusId, int len, int device)
Returns a PCI Bus Id string for the device, overloaded to take int device ID.
Definition: hip_device.cpp:395
prefer larger shared memory and smaller L1 cache
Definition: hip_runtime_api.h:243
hipError_t hipMemcpy2DToArray(hipArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind)
Copies data between host and device.
Definition: hip_memory.cpp:1233
hipError_t hipStreamQuery(hipStream_t stream)
Return #hipSuccess if all of the operations in the specified stream have completed, or #hipErrorNotReady if not.
Definition: hip_stream.cpp:157
Definition: hip_runtime_api.h:256
hipError_t hipIpcCloseMemHandle(void *devPtr)
Close memory mapped with hipIpcOpenMemHandle.
Definition: hip_memory.cpp:2076
hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeBytes)
Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant byte value va...
Definition: hip_memory.cpp:1767
hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void *ptr)
Return attributes for the specified pointer.
Definition: hip_memory.cpp:161
hipError_t hipProfilerStop()
Stop recording of profiling information. When using this API, start the profiler with profiling disab...
Definition: hip_hcc.cpp:2453
hipError_t hipMemsetD32(hipDeviceptr_t dest, int value, size_t count)
Fills the memory area pointed to by dest with the constant integer value for specified number of time...
Definition: hip_memory.cpp:1783
hipError_t hipProfilerStart()
Start recording of profiling information When using this API, start the profiler with profiling disab...
Definition: hip_hcc.cpp:2443
hipError_t hipLaunchByPtr(const void *func)
Launch a kernel.
Definition: hip_clang.cpp:201
Definition: texture_types.h:93