diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index c502cf8be5..52d01fce1f 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -313,7 +313,24 @@ enum hip_api_id_t { HIP_API_ID_hipGraphKernelNodeGetAttribute = 300, HIP_API_ID_hipGraphKernelNodeSetAttribute = 301, HIP_API_ID_hipLaunchHostFunc = 302, - HIP_API_ID_LAST = 302, + HIP_API_ID_hipDeviceGetDefaultMemPool = 303, + HIP_API_ID_hipDeviceGetMemPool = 304, + HIP_API_ID_hipDeviceSetMemPool = 305, + HIP_API_ID_hipFreeAsync = 306, + HIP_API_ID_hipMallocAsync = 307, + HIP_API_ID_hipMallocFromPoolAsync = 308, + HIP_API_ID_hipMemPoolCreate = 309, + HIP_API_ID_hipMemPoolDestroy = 310, + HIP_API_ID_hipMemPoolExportPointer = 311, + HIP_API_ID_hipMemPoolExportToShareableHandle = 312, + HIP_API_ID_hipMemPoolGetAccess = 313, + HIP_API_ID_hipMemPoolGetAttribute = 314, + HIP_API_ID_hipMemPoolImportFromShareableHandle = 315, + HIP_API_ID_hipMemPoolImportPointer = 316, + HIP_API_ID_hipMemPoolSetAccess = 317, + HIP_API_ID_hipMemPoolSetAttribute = 318, + HIP_API_ID_hipMemPoolTrimTo = 319, + HIP_API_ID_LAST = 319, HIP_API_ID_hipArray3DGetDescriptor = HIP_API_ID_NONE, HIP_API_ID_hipArrayGetDescriptor = HIP_API_ID_NONE, @@ -399,7 +416,9 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipDeviceGetAttribute: return "hipDeviceGetAttribute"; case HIP_API_ID_hipDeviceGetByPCIBusId: return "hipDeviceGetByPCIBusId"; case HIP_API_ID_hipDeviceGetCacheConfig: return "hipDeviceGetCacheConfig"; + case HIP_API_ID_hipDeviceGetDefaultMemPool: return "hipDeviceGetDefaultMemPool"; case HIP_API_ID_hipDeviceGetLimit: return "hipDeviceGetLimit"; + case HIP_API_ID_hipDeviceGetMemPool: return "hipDeviceGetMemPool"; case HIP_API_ID_hipDeviceGetName: return "hipDeviceGetName"; case HIP_API_ID_hipDeviceGetP2PAttribute: return "hipDeviceGetP2PAttribute"; case HIP_API_ID_hipDeviceGetPCIBusId: return "hipDeviceGetPCIBusId"; @@ -413,6 +432,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipDevicePrimaryCtxSetFlags: return "hipDevicePrimaryCtxSetFlags"; case HIP_API_ID_hipDeviceReset: return "hipDeviceReset"; case HIP_API_ID_hipDeviceSetCacheConfig: return "hipDeviceSetCacheConfig"; + case HIP_API_ID_hipDeviceSetMemPool: return "hipDeviceSetMemPool"; case HIP_API_ID_hipDeviceSetSharedMemConfig: return "hipDeviceSetSharedMemConfig"; case HIP_API_ID_hipDeviceSynchronize: return "hipDeviceSynchronize"; case HIP_API_ID_hipDeviceTotalMem: return "hipDeviceTotalMem"; @@ -437,6 +457,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipExternalMemoryGetMappedBuffer: return "hipExternalMemoryGetMappedBuffer"; case HIP_API_ID_hipFree: return "hipFree"; case HIP_API_ID_hipFreeArray: return "hipFreeArray"; + case HIP_API_ID_hipFreeAsync: return "hipFreeAsync"; case HIP_API_ID_hipFreeHost: return "hipFreeHost"; case HIP_API_ID_hipFreeMipmappedArray: return "hipFreeMipmappedArray"; case HIP_API_ID_hipFuncGetAttribute: return "hipFuncGetAttribute"; @@ -544,6 +565,8 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMalloc3D: return "hipMalloc3D"; case HIP_API_ID_hipMalloc3DArray: return "hipMalloc3DArray"; case HIP_API_ID_hipMallocArray: return "hipMallocArray"; + case HIP_API_ID_hipMallocAsync: return "hipMallocAsync"; + case HIP_API_ID_hipMallocFromPoolAsync: return "hipMallocFromPoolAsync"; case HIP_API_ID_hipMallocHost: return "hipMallocHost"; case HIP_API_ID_hipMallocManaged: return "hipMallocManaged"; case HIP_API_ID_hipMallocMipmappedArray: return "hipMallocMipmappedArray"; @@ -553,6 +576,17 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMemAllocPitch: return "hipMemAllocPitch"; case HIP_API_ID_hipMemGetAddressRange: return "hipMemGetAddressRange"; case HIP_API_ID_hipMemGetInfo: return "hipMemGetInfo"; + case HIP_API_ID_hipMemPoolCreate: return "hipMemPoolCreate"; + case HIP_API_ID_hipMemPoolDestroy: return "hipMemPoolDestroy"; + case HIP_API_ID_hipMemPoolExportPointer: return "hipMemPoolExportPointer"; + case HIP_API_ID_hipMemPoolExportToShareableHandle: return "hipMemPoolExportToShareableHandle"; + case HIP_API_ID_hipMemPoolGetAccess: return "hipMemPoolGetAccess"; + case HIP_API_ID_hipMemPoolGetAttribute: return "hipMemPoolGetAttribute"; + case HIP_API_ID_hipMemPoolImportFromShareableHandle: return "hipMemPoolImportFromShareableHandle"; + case HIP_API_ID_hipMemPoolImportPointer: return "hipMemPoolImportPointer"; + case HIP_API_ID_hipMemPoolSetAccess: return "hipMemPoolSetAccess"; + case HIP_API_ID_hipMemPoolSetAttribute: return "hipMemPoolSetAttribute"; + case HIP_API_ID_hipMemPoolTrimTo: return "hipMemPoolTrimTo"; case HIP_API_ID_hipMemPrefetchAsync: return "hipMemPrefetchAsync"; case HIP_API_ID_hipMemPtrGetInfo: return "hipMemPtrGetInfo"; case HIP_API_ID_hipMemRangeGetAttribute: return "hipMemRangeGetAttribute"; @@ -704,7 +738,9 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipDeviceGetAttribute", name) == 0) return HIP_API_ID_hipDeviceGetAttribute; if (strcmp("hipDeviceGetByPCIBusId", name) == 0) return HIP_API_ID_hipDeviceGetByPCIBusId; if (strcmp("hipDeviceGetCacheConfig", name) == 0) return HIP_API_ID_hipDeviceGetCacheConfig; + if (strcmp("hipDeviceGetDefaultMemPool", name) == 0) return HIP_API_ID_hipDeviceGetDefaultMemPool; if (strcmp("hipDeviceGetLimit", name) == 0) return HIP_API_ID_hipDeviceGetLimit; + if (strcmp("hipDeviceGetMemPool", name) == 0) return HIP_API_ID_hipDeviceGetMemPool; if (strcmp("hipDeviceGetName", name) == 0) return HIP_API_ID_hipDeviceGetName; if (strcmp("hipDeviceGetP2PAttribute", name) == 0) return HIP_API_ID_hipDeviceGetP2PAttribute; if (strcmp("hipDeviceGetPCIBusId", name) == 0) return HIP_API_ID_hipDeviceGetPCIBusId; @@ -718,6 +754,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipDevicePrimaryCtxSetFlags", name) == 0) return HIP_API_ID_hipDevicePrimaryCtxSetFlags; if (strcmp("hipDeviceReset", name) == 0) return HIP_API_ID_hipDeviceReset; if (strcmp("hipDeviceSetCacheConfig", name) == 0) return HIP_API_ID_hipDeviceSetCacheConfig; + if (strcmp("hipDeviceSetMemPool", name) == 0) return HIP_API_ID_hipDeviceSetMemPool; if (strcmp("hipDeviceSetSharedMemConfig", name) == 0) return HIP_API_ID_hipDeviceSetSharedMemConfig; if (strcmp("hipDeviceSynchronize", name) == 0) return HIP_API_ID_hipDeviceSynchronize; if (strcmp("hipDeviceTotalMem", name) == 0) return HIP_API_ID_hipDeviceTotalMem; @@ -742,6 +779,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipExternalMemoryGetMappedBuffer", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedBuffer; if (strcmp("hipFree", name) == 0) return HIP_API_ID_hipFree; if (strcmp("hipFreeArray", name) == 0) return HIP_API_ID_hipFreeArray; + if (strcmp("hipFreeAsync", name) == 0) return HIP_API_ID_hipFreeAsync; if (strcmp("hipFreeHost", name) == 0) return HIP_API_ID_hipFreeHost; if (strcmp("hipFreeMipmappedArray", name) == 0) return HIP_API_ID_hipFreeMipmappedArray; if (strcmp("hipFuncGetAttribute", name) == 0) return HIP_API_ID_hipFuncGetAttribute; @@ -849,6 +887,8 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMalloc3D", name) == 0) return HIP_API_ID_hipMalloc3D; if (strcmp("hipMalloc3DArray", name) == 0) return HIP_API_ID_hipMalloc3DArray; if (strcmp("hipMallocArray", name) == 0) return HIP_API_ID_hipMallocArray; + if (strcmp("hipMallocAsync", name) == 0) return HIP_API_ID_hipMallocAsync; + if (strcmp("hipMallocFromPoolAsync", name) == 0) return HIP_API_ID_hipMallocFromPoolAsync; if (strcmp("hipMallocHost", name) == 0) return HIP_API_ID_hipMallocHost; if (strcmp("hipMallocManaged", name) == 0) return HIP_API_ID_hipMallocManaged; if (strcmp("hipMallocMipmappedArray", name) == 0) return HIP_API_ID_hipMallocMipmappedArray; @@ -858,6 +898,17 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMemAllocPitch", name) == 0) return HIP_API_ID_hipMemAllocPitch; if (strcmp("hipMemGetAddressRange", name) == 0) return HIP_API_ID_hipMemGetAddressRange; if (strcmp("hipMemGetInfo", name) == 0) return HIP_API_ID_hipMemGetInfo; + if (strcmp("hipMemPoolCreate", name) == 0) return HIP_API_ID_hipMemPoolCreate; + if (strcmp("hipMemPoolDestroy", name) == 0) return HIP_API_ID_hipMemPoolDestroy; + if (strcmp("hipMemPoolExportPointer", name) == 0) return HIP_API_ID_hipMemPoolExportPointer; + if (strcmp("hipMemPoolExportToShareableHandle", name) == 0) return HIP_API_ID_hipMemPoolExportToShareableHandle; + if (strcmp("hipMemPoolGetAccess", name) == 0) return HIP_API_ID_hipMemPoolGetAccess; + if (strcmp("hipMemPoolGetAttribute", name) == 0) return HIP_API_ID_hipMemPoolGetAttribute; + if (strcmp("hipMemPoolImportFromShareableHandle", name) == 0) return HIP_API_ID_hipMemPoolImportFromShareableHandle; + if (strcmp("hipMemPoolImportPointer", name) == 0) return HIP_API_ID_hipMemPoolImportPointer; + if (strcmp("hipMemPoolSetAccess", name) == 0) return HIP_API_ID_hipMemPoolSetAccess; + if (strcmp("hipMemPoolSetAttribute", name) == 0) return HIP_API_ID_hipMemPoolSetAttribute; + if (strcmp("hipMemPoolTrimTo", name) == 0) return HIP_API_ID_hipMemPoolTrimTo; if (strcmp("hipMemPrefetchAsync", name) == 0) return HIP_API_ID_hipMemPrefetchAsync; if (strcmp("hipMemPtrGetInfo", name) == 0) return HIP_API_ID_hipMemPtrGetInfo; if (strcmp("hipMemRangeGetAttribute", name) == 0) return HIP_API_ID_hipMemRangeGetAttribute; @@ -1132,11 +1183,21 @@ typedef struct hip_api_data_s { hipFuncCache_t* cacheConfig; hipFuncCache_t cacheConfig__val; } hipDeviceGetCacheConfig; + struct { + hipMemPool_t* mem_pool; + hipMemPool_t mem_pool__val; + int device; + } hipDeviceGetDefaultMemPool; struct { size_t* pValue; size_t pValue__val; enum hipLimit_t limit; } hipDeviceGetLimit; + struct { + hipMemPool_t* mem_pool; + hipMemPool_t mem_pool__val; + int device; + } hipDeviceGetMemPool; struct { char* name; char name__val; @@ -1196,6 +1257,10 @@ typedef struct hip_api_data_s { struct { hipFuncCache_t cacheConfig; } hipDeviceSetCacheConfig; + struct { + int device; + hipMemPool_t mem_pool; + } hipDeviceSetMemPool; struct { hipSharedMemConfig config; } hipDeviceSetSharedMemConfig; @@ -1330,6 +1395,10 @@ typedef struct hip_api_data_s { hipArray* array; hipArray array__val; } hipFreeArray; + struct { + void* dev_ptr; + hipStream_t stream; + } hipFreeAsync; struct { void* ptr; } hipFreeHost; @@ -1984,6 +2053,19 @@ typedef struct hip_api_data_s { size_t height; unsigned int flags; } hipMallocArray; + struct { + void** dev_ptr; + void* dev_ptr__val; + size_t size; + hipStream_t stream; + } hipMallocAsync; + struct { + void** dev_ptr; + void* dev_ptr__val; + size_t size; + hipMemPool_t mem_pool; + hipStream_t stream; + } hipMallocFromPoolAsync; struct { void** ptr; void* ptr__val; @@ -2045,6 +2127,67 @@ typedef struct hip_api_data_s { size_t* total; size_t total__val; } hipMemGetInfo; + struct { + hipMemPool_t* mem_pool; + hipMemPool_t mem_pool__val; + const hipMemPoolProps* pool_props; + hipMemPoolProps pool_props__val; + } hipMemPoolCreate; + struct { + hipMemPool_t mem_pool; + } hipMemPoolDestroy; + struct { + hipMemPoolPtrExportData* export_data; + hipMemPoolPtrExportData export_data__val; + void* dev_ptr; + } hipMemPoolExportPointer; + struct { + void* shared_handle; + hipMemPool_t mem_pool; + hipMemAllocationHandleType handle_type; + unsigned int flags; + } hipMemPoolExportToShareableHandle; + struct { + hipMemAccessFlags* flags; + hipMemAccessFlags flags__val; + hipMemPool_t mem_pool; + hipMemLocation* location; + hipMemLocation location__val; + } hipMemPoolGetAccess; + struct { + hipMemPool_t mem_pool; + hipMemPoolAttr attr; + void* value; + } hipMemPoolGetAttribute; + struct { + hipMemPool_t* mem_pool; + hipMemPool_t mem_pool__val; + void* shared_handle; + hipMemAllocationHandleType handle_type; + unsigned int flags; + } hipMemPoolImportFromShareableHandle; + struct { + void** dev_ptr; + void* dev_ptr__val; + hipMemPool_t mem_pool; + hipMemPoolPtrExportData* export_data; + hipMemPoolPtrExportData export_data__val; + } hipMemPoolImportPointer; + struct { + hipMemPool_t mem_pool; + const hipMemAccessDesc* desc_list; + hipMemAccessDesc desc_list__val; + size_t count; + } hipMemPoolSetAccess; + struct { + hipMemPool_t mem_pool; + hipMemPoolAttr attr; + void* value; + } hipMemPoolSetAttribute; + struct { + hipMemPool_t mem_pool; + size_t min_bytes_to_hold; + } hipMemPoolTrimTo; struct { const void* dev_ptr; size_t count; @@ -2910,11 +3053,21 @@ typedef struct hip_api_data_s { #define INIT_hipDeviceGetCacheConfig_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceGetCacheConfig.cacheConfig = (hipFuncCache_t*)cacheConfig; \ }; +// hipDeviceGetDefaultMemPool[('hipMemPool_t*', 'mem_pool'), ('int', 'device')] +#define INIT_hipDeviceGetDefaultMemPool_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDeviceGetDefaultMemPool.mem_pool = (hipMemPool_t*)mem_pool; \ + cb_data.args.hipDeviceGetDefaultMemPool.device = (int)device; \ +}; // hipDeviceGetLimit[('size_t*', 'pValue'), ('hipLimit_t', 'limit')] #define INIT_hipDeviceGetLimit_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceGetLimit.pValue = (size_t*)pValue; \ cb_data.args.hipDeviceGetLimit.limit = (hipLimit_t)limit; \ }; +// hipDeviceGetMemPool[('hipMemPool_t*', 'mem_pool'), ('int', 'device')] +#define INIT_hipDeviceGetMemPool_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDeviceGetMemPool.mem_pool = (hipMemPool_t*)mem_pool; \ + cb_data.args.hipDeviceGetMemPool.device = (int)device; \ +}; // hipDeviceGetName[('char*', 'name'), ('int', 'len'), ('hipDevice_t', 'device')] #define INIT_hipDeviceGetName_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceGetName.name = (char*)name; \ @@ -2979,6 +3132,11 @@ typedef struct hip_api_data_s { #define INIT_hipDeviceSetCacheConfig_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceSetCacheConfig.cacheConfig = (hipFuncCache_t)cacheConfig; \ }; +// hipDeviceSetMemPool[('int', 'device'), ('hipMemPool_t', 'mem_pool')] +#define INIT_hipDeviceSetMemPool_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDeviceSetMemPool.device = (int)device; \ + cb_data.args.hipDeviceSetMemPool.mem_pool = (hipMemPool_t)mem_pool; \ +}; // hipDeviceSetSharedMemConfig[('hipSharedMemConfig', 'config')] #define INIT_hipDeviceSetSharedMemConfig_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipDeviceSetSharedMemConfig.config = (hipSharedMemConfig)config; \ @@ -3117,6 +3275,11 @@ typedef struct hip_api_data_s { #define INIT_hipFreeArray_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipFreeArray.array = (hipArray*)array; \ }; +// hipFreeAsync[('void*', 'dev_ptr'), ('hipStream_t', 'stream')] +#define INIT_hipFreeAsync_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipFreeAsync.dev_ptr = (void*)dev_ptr; \ + cb_data.args.hipFreeAsync.stream = (hipStream_t)stream; \ +}; // hipFreeHost[('void*', 'ptr')] #define INIT_hipFreeHost_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipFreeHost.ptr = (void*)ptr; \ @@ -3763,6 +3926,19 @@ typedef struct hip_api_data_s { cb_data.args.hipMallocArray.height = (size_t)height; \ cb_data.args.hipMallocArray.flags = (unsigned int)flags; \ }; +// hipMallocAsync[('void**', 'dev_ptr'), ('size_t', 'size'), ('hipStream_t', 'stream')] +#define INIT_hipMallocAsync_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMallocAsync.dev_ptr = (void**)dev_ptr; \ + cb_data.args.hipMallocAsync.size = (size_t)size; \ + cb_data.args.hipMallocAsync.stream = (hipStream_t)stream; \ +}; +// hipMallocFromPoolAsync[('void**', 'dev_ptr'), ('size_t', 'size'), ('hipMemPool_t', 'mem_pool'), ('hipStream_t', 'stream')] +#define INIT_hipMallocFromPoolAsync_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMallocFromPoolAsync.dev_ptr = (void**)dev_ptr; \ + cb_data.args.hipMallocFromPoolAsync.size = (size_t)size; \ + cb_data.args.hipMallocFromPoolAsync.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMallocFromPoolAsync.stream = (hipStream_t)stream; \ +}; // hipMallocHost[('void**', 'ptr'), ('size_t', 'size')] #define INIT_hipMallocHost_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMallocHost.ptr = (void**)ptr; \ @@ -3820,6 +3996,69 @@ typedef struct hip_api_data_s { cb_data.args.hipMemGetInfo.free = (size_t*)free; \ cb_data.args.hipMemGetInfo.total = (size_t*)total; \ }; +// hipMemPoolCreate[('hipMemPool_t*', 'mem_pool'), ('const hipMemPoolProps*', 'pool_props')] +#define INIT_hipMemPoolCreate_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolCreate.mem_pool = (hipMemPool_t*)mem_pool; \ + cb_data.args.hipMemPoolCreate.pool_props = (const hipMemPoolProps*)pool_props; \ +}; +// hipMemPoolDestroy[('hipMemPool_t', 'mem_pool')] +#define INIT_hipMemPoolDestroy_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolDestroy.mem_pool = (hipMemPool_t)mem_pool; \ +}; +// hipMemPoolExportPointer[('hipMemPoolPtrExportData*', 'export_data'), ('void*', 'dev_ptr')] +#define INIT_hipMemPoolExportPointer_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolExportPointer.export_data = (hipMemPoolPtrExportData*)export_data; \ + cb_data.args.hipMemPoolExportPointer.dev_ptr = (void*)ptr; \ +}; +// hipMemPoolExportToShareableHandle[('void*', 'shared_handle'), ('hipMemPool_t', 'mem_pool'), ('hipMemAllocationHandleType', 'handle_type'), ('unsigned int', 'flags')] +#define INIT_hipMemPoolExportToShareableHandle_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolExportToShareableHandle.shared_handle = (void*)shared_handle; \ + cb_data.args.hipMemPoolExportToShareableHandle.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolExportToShareableHandle.handle_type = (hipMemAllocationHandleType)handle_type; \ + cb_data.args.hipMemPoolExportToShareableHandle.flags = (unsigned int)flags; \ +}; +// hipMemPoolGetAccess[('hipMemAccessFlags*', 'flags'), ('hipMemPool_t', 'mem_pool'), ('hipMemLocation*', 'location')] +#define INIT_hipMemPoolGetAccess_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolGetAccess.flags = (hipMemAccessFlags*)flags; \ + cb_data.args.hipMemPoolGetAccess.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolGetAccess.location = (hipMemLocation*)location; \ +}; +// hipMemPoolGetAttribute[('hipMemPool_t', 'mem_pool'), ('hipMemPoolAttr', 'attr'), ('void*', 'value')] +#define INIT_hipMemPoolGetAttribute_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolGetAttribute.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolGetAttribute.attr = (hipMemPoolAttr)attr; \ + cb_data.args.hipMemPoolGetAttribute.value = (void*)value; \ +}; +// hipMemPoolImportFromShareableHandle[('hipMemPool_t*', 'mem_pool'), ('void*', 'shared_handle'), ('hipMemAllocationHandleType', 'handle_type'), ('unsigned int', 'flags')] +#define INIT_hipMemPoolImportFromShareableHandle_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolImportFromShareableHandle.mem_pool = (hipMemPool_t*)mem_pool; \ + cb_data.args.hipMemPoolImportFromShareableHandle.shared_handle = (void*)shared_handle; \ + cb_data.args.hipMemPoolImportFromShareableHandle.handle_type = (hipMemAllocationHandleType)handle_type; \ + cb_data.args.hipMemPoolImportFromShareableHandle.flags = (unsigned int)flags; \ +}; +// hipMemPoolImportPointer[('void**', 'dev_ptr'), ('hipMemPool_t', 'mem_pool'), ('hipMemPoolPtrExportData*', 'export_data')] +#define INIT_hipMemPoolImportPointer_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolImportPointer.dev_ptr = (void**)ptr; \ + cb_data.args.hipMemPoolImportPointer.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolImportPointer.export_data = (hipMemPoolPtrExportData*)export_data; \ +}; +// hipMemPoolSetAccess[('hipMemPool_t', 'mem_pool'), ('const hipMemAccessDesc*', 'desc_list'), ('size_t', 'count')] +#define INIT_hipMemPoolSetAccess_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolSetAccess.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolSetAccess.desc_list = (const hipMemAccessDesc*)desc_list; \ + cb_data.args.hipMemPoolSetAccess.count = (size_t)count; \ +}; +// hipMemPoolSetAttribute[('hipMemPool_t', 'mem_pool'), ('hipMemPoolAttr', 'attr'), ('void*', 'value')] +#define INIT_hipMemPoolSetAttribute_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolSetAttribute.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolSetAttribute.attr = (hipMemPoolAttr)attr; \ + cb_data.args.hipMemPoolSetAttribute.value = (void*)value; \ +}; +// hipMemPoolTrimTo[('hipMemPool_t', 'mem_pool'), ('size_t', 'min_bytes_to_hold')] +#define INIT_hipMemPoolTrimTo_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemPoolTrimTo.mem_pool = (hipMemPool_t)mem_pool; \ + cb_data.args.hipMemPoolTrimTo.min_bytes_to_hold = (size_t)min_bytes_to_hold; \ +}; // hipMemPrefetchAsync[('const void*', 'dev_ptr'), ('size_t', 'count'), ('int', 'device'), ('hipStream_t', 'stream')] #define INIT_hipMemPrefetchAsync_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemPrefetchAsync.dev_ptr = (const void*)dev_ptr; \ @@ -4772,10 +5011,18 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipDeviceGetCacheConfig: if (data->args.hipDeviceGetCacheConfig.cacheConfig) data->args.hipDeviceGetCacheConfig.cacheConfig__val = *(data->args.hipDeviceGetCacheConfig.cacheConfig); break; +// hipDeviceGetDefaultMemPool[('hipMemPool_t*', 'mem_pool'), ('int', 'device')] + case HIP_API_ID_hipDeviceGetDefaultMemPool: + if (data->args.hipDeviceGetDefaultMemPool.mem_pool) data->args.hipDeviceGetDefaultMemPool.mem_pool__val = *(data->args.hipDeviceGetDefaultMemPool.mem_pool); + break; // hipDeviceGetLimit[('size_t*', 'pValue'), ('hipLimit_t', 'limit')] case HIP_API_ID_hipDeviceGetLimit: if (data->args.hipDeviceGetLimit.pValue) data->args.hipDeviceGetLimit.pValue__val = *(data->args.hipDeviceGetLimit.pValue); break; +// hipDeviceGetMemPool[('hipMemPool_t*', 'mem_pool'), ('int', 'device')] + case HIP_API_ID_hipDeviceGetMemPool: + if (data->args.hipDeviceGetMemPool.mem_pool) data->args.hipDeviceGetMemPool.mem_pool__val = *(data->args.hipDeviceGetMemPool.mem_pool); + break; // hipDeviceGetName[('char*', 'name'), ('int', 'len'), ('hipDevice_t', 'device')] case HIP_API_ID_hipDeviceGetName: data->args.hipDeviceGetName.name = (data->args.hipDeviceGetName.name) ? strdup(data->args.hipDeviceGetName.name) : NULL; @@ -4825,6 +5072,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipDeviceSetCacheConfig[('hipFuncCache_t', 'cacheConfig')] case HIP_API_ID_hipDeviceSetCacheConfig: break; +// hipDeviceSetMemPool[('int', 'device'), ('hipMemPool_t', 'mem_pool')] + case HIP_API_ID_hipDeviceSetMemPool: + break; // hipDeviceSetSharedMemConfig[('hipSharedMemConfig', 'config')] case HIP_API_ID_hipDeviceSetSharedMemConfig: break; @@ -4919,6 +5169,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipFreeArray: if (data->args.hipFreeArray.array) data->args.hipFreeArray.array__val = *(data->args.hipFreeArray.array); break; +// hipFreeAsync[('void*', 'dev_ptr'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipFreeAsync: + break; // hipFreeHost[('void*', 'ptr')] case HIP_API_ID_hipFreeHost: break; @@ -5350,6 +5603,14 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipMallocArray.array) data->args.hipMallocArray.array__val = *(data->args.hipMallocArray.array); if (data->args.hipMallocArray.desc) data->args.hipMallocArray.desc__val = *(data->args.hipMallocArray.desc); break; +// hipMallocAsync[('void**', 'dev_ptr'), ('size_t', 'size'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMallocAsync: + if (data->args.hipMallocAsync.dev_ptr) data->args.hipMallocAsync.dev_ptr__val = *(data->args.hipMallocAsync.dev_ptr); + break; +// hipMallocFromPoolAsync[('void**', 'dev_ptr'), ('size_t', 'size'), ('hipMemPool_t', 'mem_pool'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMallocFromPoolAsync: + if (data->args.hipMallocFromPoolAsync.dev_ptr) data->args.hipMallocFromPoolAsync.dev_ptr__val = *(data->args.hipMallocFromPoolAsync.dev_ptr); + break; // hipMallocHost[('void**', 'ptr'), ('size_t', 'size')] case HIP_API_ID_hipMallocHost: if (data->args.hipMallocHost.ptr) data->args.hipMallocHost.ptr__val = *(data->args.hipMallocHost.ptr); @@ -5390,6 +5651,48 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipMemGetInfo.free) data->args.hipMemGetInfo.free__val = *(data->args.hipMemGetInfo.free); if (data->args.hipMemGetInfo.total) data->args.hipMemGetInfo.total__val = *(data->args.hipMemGetInfo.total); break; +// hipMemPoolCreate[('hipMemPool_t*', 'mem_pool'), ('const hipMemPoolProps*', 'pool_props')] + case HIP_API_ID_hipMemPoolCreate: + if (data->args.hipMemPoolCreate.mem_pool) data->args.hipMemPoolCreate.mem_pool__val = *(data->args.hipMemPoolCreate.mem_pool); + if (data->args.hipMemPoolCreate.pool_props) data->args.hipMemPoolCreate.pool_props__val = *(data->args.hipMemPoolCreate.pool_props); + break; +// hipMemPoolDestroy[('hipMemPool_t', 'mem_pool')] + case HIP_API_ID_hipMemPoolDestroy: + break; +// hipMemPoolExportPointer[('hipMemPoolPtrExportData*', 'export_data'), ('void*', 'dev_ptr')] + case HIP_API_ID_hipMemPoolExportPointer: + if (data->args.hipMemPoolExportPointer.export_data) data->args.hipMemPoolExportPointer.export_data__val = *(data->args.hipMemPoolExportPointer.export_data); + break; +// hipMemPoolExportToShareableHandle[('void*', 'shared_handle'), ('hipMemPool_t', 'mem_pool'), ('hipMemAllocationHandleType', 'handle_type'), ('unsigned int', 'flags')] + case HIP_API_ID_hipMemPoolExportToShareableHandle: + break; +// hipMemPoolGetAccess[('hipMemAccessFlags*', 'flags'), ('hipMemPool_t', 'mem_pool'), ('hipMemLocation*', 'location')] + case HIP_API_ID_hipMemPoolGetAccess: + if (data->args.hipMemPoolGetAccess.flags) data->args.hipMemPoolGetAccess.flags__val = *(data->args.hipMemPoolGetAccess.flags); + if (data->args.hipMemPoolGetAccess.location) data->args.hipMemPoolGetAccess.location__val = *(data->args.hipMemPoolGetAccess.location); + break; +// hipMemPoolGetAttribute[('hipMemPool_t', 'mem_pool'), ('hipMemPoolAttr', 'attr'), ('void*', 'value')] + case HIP_API_ID_hipMemPoolGetAttribute: + break; +// hipMemPoolImportFromShareableHandle[('hipMemPool_t*', 'mem_pool'), ('void*', 'shared_handle'), ('hipMemAllocationHandleType', 'handle_type'), ('unsigned int', 'flags')] + case HIP_API_ID_hipMemPoolImportFromShareableHandle: + if (data->args.hipMemPoolImportFromShareableHandle.mem_pool) data->args.hipMemPoolImportFromShareableHandle.mem_pool__val = *(data->args.hipMemPoolImportFromShareableHandle.mem_pool); + break; +// hipMemPoolImportPointer[('void**', 'dev_ptr'), ('hipMemPool_t', 'mem_pool'), ('hipMemPoolPtrExportData*', 'export_data')] + case HIP_API_ID_hipMemPoolImportPointer: + if (data->args.hipMemPoolImportPointer.dev_ptr) data->args.hipMemPoolImportPointer.dev_ptr__val = *(data->args.hipMemPoolImportPointer.dev_ptr); + if (data->args.hipMemPoolImportPointer.export_data) data->args.hipMemPoolImportPointer.export_data__val = *(data->args.hipMemPoolImportPointer.export_data); + break; +// hipMemPoolSetAccess[('hipMemPool_t', 'mem_pool'), ('const hipMemAccessDesc*', 'desc_list'), ('size_t', 'count')] + case HIP_API_ID_hipMemPoolSetAccess: + if (data->args.hipMemPoolSetAccess.desc_list) data->args.hipMemPoolSetAccess.desc_list__val = *(data->args.hipMemPoolSetAccess.desc_list); + break; +// hipMemPoolSetAttribute[('hipMemPool_t', 'mem_pool'), ('hipMemPoolAttr', 'attr'), ('void*', 'value')] + case HIP_API_ID_hipMemPoolSetAttribute: + break; +// hipMemPoolTrimTo[('hipMemPool_t', 'mem_pool'), ('size_t', 'min_bytes_to_hold')] + case HIP_API_ID_hipMemPoolTrimTo: + break; // hipMemPrefetchAsync[('const void*', 'dev_ptr'), ('size_t', 'count'), ('int', 'device'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemPrefetchAsync: break; @@ -6045,6 +6348,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else oss << "cacheConfig=" << data->args.hipDeviceGetCacheConfig.cacheConfig__val; oss << ")"; break; + case HIP_API_ID_hipDeviceGetDefaultMemPool: + oss << "hipDeviceGetDefaultMemPool("; + if (data->args.hipDeviceGetDefaultMemPool.mem_pool == NULL) oss << "mem_pool=NULL"; + else oss << "mem_pool=" << data->args.hipDeviceGetDefaultMemPool.mem_pool__val; + oss << ", device=" << data->args.hipDeviceGetDefaultMemPool.device; + oss << ")"; + break; case HIP_API_ID_hipDeviceGetLimit: oss << "hipDeviceGetLimit("; if (data->args.hipDeviceGetLimit.pValue == NULL) oss << "pValue=NULL"; @@ -6052,6 +6362,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", limit=" << data->args.hipDeviceGetLimit.limit; oss << ")"; break; + case HIP_API_ID_hipDeviceGetMemPool: + oss << "hipDeviceGetMemPool("; + if (data->args.hipDeviceGetMemPool.mem_pool == NULL) oss << "mem_pool=NULL"; + else oss << "mem_pool=" << data->args.hipDeviceGetMemPool.mem_pool__val; + oss << ", device=" << data->args.hipDeviceGetMemPool.device; + oss << ")"; + break; case HIP_API_ID_hipDeviceGetName: oss << "hipDeviceGetName("; if (data->args.hipDeviceGetName.name == NULL) oss << "name=NULL"; @@ -6139,6 +6456,12 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << "cacheConfig=" << data->args.hipDeviceSetCacheConfig.cacheConfig; oss << ")"; break; + case HIP_API_ID_hipDeviceSetMemPool: + oss << "hipDeviceSetMemPool("; + oss << "device=" << data->args.hipDeviceSetMemPool.device; + oss << ", mem_pool=" << data->args.hipDeviceSetMemPool.mem_pool; + oss << ")"; + break; case HIP_API_ID_hipDeviceSetSharedMemConfig: oss << "hipDeviceSetSharedMemConfig("; oss << "config=" << data->args.hipDeviceSetSharedMemConfig.config; @@ -6323,6 +6646,12 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else oss << "array=" << data->args.hipFreeArray.array__val; oss << ")"; break; + case HIP_API_ID_hipFreeAsync: + oss << "hipFreeAsync("; + oss << "dev_ptr=" << data->args.hipFreeAsync.dev_ptr; + oss << ", stream=" << data->args.hipFreeAsync.stream; + oss << ")"; + break; case HIP_API_ID_hipFreeHost: oss << "hipFreeHost("; oss << "ptr=" << data->args.hipFreeHost.ptr; @@ -7195,6 +7524,23 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", flags=" << data->args.hipMallocArray.flags; oss << ")"; break; + case HIP_API_ID_hipMallocAsync: + oss << "hipMallocAsync("; + if (data->args.hipMallocAsync.dev_ptr == NULL) oss << "dev_ptr=NULL"; + else oss << "dev_ptr=" << data->args.hipMallocAsync.dev_ptr__val; + oss << ", size=" << data->args.hipMallocAsync.size; + oss << ", stream=" << data->args.hipMallocAsync.stream; + oss << ")"; + break; + case HIP_API_ID_hipMallocFromPoolAsync: + oss << "hipMallocFromPoolAsync("; + if (data->args.hipMallocFromPoolAsync.dev_ptr == NULL) oss << "dev_ptr=NULL"; + else oss << "dev_ptr=" << data->args.hipMallocFromPoolAsync.dev_ptr__val; + oss << ", size=" << data->args.hipMallocFromPoolAsync.size; + oss << ", mem_pool=" << data->args.hipMallocFromPoolAsync.mem_pool; + oss << ", stream=" << data->args.hipMallocFromPoolAsync.stream; + oss << ")"; + break; case HIP_API_ID_hipMallocHost: oss << "hipMallocHost("; if (data->args.hipMallocHost.ptr == NULL) oss << "ptr=NULL"; @@ -7274,6 +7620,89 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else oss << ", total=" << data->args.hipMemGetInfo.total__val; oss << ")"; break; + case HIP_API_ID_hipMemPoolCreate: + oss << "hipMemPoolCreate("; + if (data->args.hipMemPoolCreate.mem_pool == NULL) oss << "mem_pool=NULL"; + else oss << "mem_pool=" << data->args.hipMemPoolCreate.mem_pool__val; + if (data->args.hipMemPoolCreate.pool_props == NULL) oss << ", pool_props=NULL"; + else oss << ", pool_props=" << data->args.hipMemPoolCreate.pool_props__val; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolDestroy: + oss << "hipMemPoolDestroy("; + oss << "mem_pool=" << data->args.hipMemPoolDestroy.mem_pool; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolExportPointer: + oss << "hipMemPoolExportPointer("; + if (data->args.hipMemPoolExportPointer.export_data == NULL) oss << "export_data=NULL"; + else oss << "export_data=" << data->args.hipMemPoolExportPointer.export_data__val; + oss << ", dev_ptr=" << data->args.hipMemPoolExportPointer.dev_ptr; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolExportToShareableHandle: + oss << "hipMemPoolExportToShareableHandle("; + oss << "shared_handle=" << data->args.hipMemPoolExportToShareableHandle.shared_handle; + oss << ", mem_pool=" << data->args.hipMemPoolExportToShareableHandle.mem_pool; + oss << ", handle_type=" << data->args.hipMemPoolExportToShareableHandle.handle_type; + oss << ", flags=" << data->args.hipMemPoolExportToShareableHandle.flags; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolGetAccess: + oss << "hipMemPoolGetAccess("; + if (data->args.hipMemPoolGetAccess.flags == NULL) oss << "flags=NULL"; + else oss << "flags=" << data->args.hipMemPoolGetAccess.flags__val; + oss << ", mem_pool=" << data->args.hipMemPoolGetAccess.mem_pool; + if (data->args.hipMemPoolGetAccess.location == NULL) oss << ", location=NULL"; + else oss << ", location=" << data->args.hipMemPoolGetAccess.location__val; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolGetAttribute: + oss << "hipMemPoolGetAttribute("; + oss << "mem_pool=" << data->args.hipMemPoolGetAttribute.mem_pool; + oss << ", attr=" << data->args.hipMemPoolGetAttribute.attr; + oss << ", value=" << data->args.hipMemPoolGetAttribute.value; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolImportFromShareableHandle: + oss << "hipMemPoolImportFromShareableHandle("; + if (data->args.hipMemPoolImportFromShareableHandle.mem_pool == NULL) oss << "mem_pool=NULL"; + else oss << "mem_pool=" << data->args.hipMemPoolImportFromShareableHandle.mem_pool__val; + oss << ", shared_handle=" << data->args.hipMemPoolImportFromShareableHandle.shared_handle; + oss << ", handle_type=" << data->args.hipMemPoolImportFromShareableHandle.handle_type; + oss << ", flags=" << data->args.hipMemPoolImportFromShareableHandle.flags; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolImportPointer: + oss << "hipMemPoolImportPointer("; + if (data->args.hipMemPoolImportPointer.dev_ptr == NULL) oss << "dev_ptr=NULL"; + else oss << "dev_ptr=" << data->args.hipMemPoolImportPointer.dev_ptr__val; + oss << ", mem_pool=" << data->args.hipMemPoolImportPointer.mem_pool; + if (data->args.hipMemPoolImportPointer.export_data == NULL) oss << ", export_data=NULL"; + else oss << ", export_data=" << data->args.hipMemPoolImportPointer.export_data__val; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolSetAccess: + oss << "hipMemPoolSetAccess("; + oss << "mem_pool=" << data->args.hipMemPoolSetAccess.mem_pool; + if (data->args.hipMemPoolSetAccess.desc_list == NULL) oss << ", desc_list=NULL"; + else oss << ", desc_list=" << data->args.hipMemPoolSetAccess.desc_list__val; + oss << ", count=" << data->args.hipMemPoolSetAccess.count; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolSetAttribute: + oss << "hipMemPoolSetAttribute("; + oss << "mem_pool=" << data->args.hipMemPoolSetAttribute.mem_pool; + oss << ", attr=" << data->args.hipMemPoolSetAttribute.attr; + oss << ", value=" << data->args.hipMemPoolSetAttribute.value; + oss << ")"; + break; + case HIP_API_ID_hipMemPoolTrimTo: + oss << "hipMemPoolTrimTo("; + oss << "mem_pool=" << data->args.hipMemPoolTrimTo.mem_pool; + oss << ", min_bytes_to_hold=" << data->args.hipMemPoolTrimTo.min_bytes_to_hold; + oss << ")"; + break; case HIP_API_ID_hipMemPrefetchAsync: oss << "hipMemPrefetchAsync("; oss << "dev_ptr=" << data->args.hipMemPrefetchAsync.dev_ptr; diff --git a/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 1babb5578e..6461e1e80a 100644 --- a/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -33,6 +33,7 @@ THE SOFTWARE. #define CUDA_10010 10010 #define CUDA_10020 10020 #define CUDA_11010 11010 +#define CUDA_11010 11020 #define CUDA_11030 11030 #define CUDA_11040 11040 @@ -408,6 +409,15 @@ typedef struct cudaPos hipPos; #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking +typedef cudaMemPool_t hipMemPool_t; +typedef enum cudaMemPoolAttr hipMemPoolAttr; +typedef struct cudaMemLocation hipMemLocation; +typedef struct cudaMemPoolProps hipMemPoolProps; +typedef struct cudaMemAccessDesc hipMemAccessDesc; +typedef enum cudaMemAccessFlags hipMemAccessFlags; +typedef enum cudaMemAllocationHandleType hipMemAllocationHandleType; +typedef struct cudaMemPoolPtrExportData hipMemPoolPtrExportData; + typedef struct cudaChannelFormatDesc hipChannelFormatDesc; typedef struct cudaResourceDesc hipResourceDesc; typedef struct cudaTextureDesc hipTextureDesc; @@ -1781,6 +1791,11 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att case hipDeviceAttributeDirectManagedMemAccessFromHost: cdattr = cudaDevAttrDirectManagedMemAccessFromHost; break; +#if CUDA_VERSION >= CUDA_11020 + case hipDeviceAttributeMemoryPoolsSupported: + cdattr = cudaDevAttrMemoryPoolsSupported; + break; +#endif // CUDA_VERSION >= CUDA_11020 default: return hipCUDAErrorTohipError(cudaErrorInvalidValue); } @@ -2284,6 +2299,100 @@ inline static hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* return hipCUDAErrorTohipError(cudaLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags)); } +#if CUDA_VERSION >= CUDA_11020 +// ========================== HIP Stream Ordered Memory Allocator ================================= +inline static hipError_t hipDeviceGetDefaultMemPool(hipMemPool_t* mem_pool, int device) { + return hipCUDAErrorTohipError(cudaDeviceGetDefaultMemPool(mem_pool, device)); +} + +inline static hipError_t hipDeviceSetMemPool(int device, hipMemPool_t mem_pool) { + return hipCUDAErrorTohipError(cudaDeviceSetMemPool(device, mem_pool)); +} + +inline static hipError_t hipDeviceGetMemPool(hipMemPool_t* mem_pool, int device) { + return hipCUDAErrorTohipError(cudaDeviceGetMemPool(mem_pool, device)); +} + +inline static hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) { + return hipCUDAErrorTohipError(cudaMallocAsync(dev_ptr, size, stream)); +} + +inline static hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) { + return hipCUDAErrorTohipError(cudaFreeAsync(dev_ptr, stream)); +} + +inline static hipError_t hipMemPoolTrimTo(hipMemPool_t mem_pool, size_t min_bytes_to_hold) { + return hipCUDAErrorTohipError(cudaMemPoolTrimTo(mem_pool, min_bytes_to_hold)); +} + +inline static hipError_t hipMemPoolSetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value) { + return hipCUDAErrorTohipError(cudaMemPoolSetAttribute(mem_pool, attr, value)); +} + +inline static hipError_t hipMemPoolGetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value) { + return hipCUDAErrorTohipError(cudaMemPoolGetAttribute(mem_pool, attr, value)); +} + +inline static hipError_t hipMemPoolSetAccess( + hipMemPool_t mem_pool, + const hipMemAccessDesc* desc_list, + size_t count) { + return hipCUDAErrorTohipError(cudaMemPoolSetAccess(mem_pool, desc_list, count)); +} + +inline static hipError_t hipMemPoolGetAccess( + hipMemAccessFlags* flags, + hipMemPool_t mem_pool, + hipMemLocation* location) { + return hipCUDAErrorTohipError(cudaMemPoolGetAccess(flags, mem_pool, location)); +} + +inline static hipError_t hipMemPoolCreate(hipMemPool_t* mem_pool, const hipMemPoolProps* pool_props) { + return hipCUDAErrorTohipError(cudaMemPoolCreate(mem_pool, pool_props)); +} + +inline static hipError_t hipMemPoolDestroy(hipMemPool_t mem_pool) { + return hipCUDAErrorTohipError(cudaMemPoolDestroy(mem_pool)); +} + +inline static hipError_t hipMallocFromPoolAsync( + void** dev_ptr, + size_t size, + hipMemPool_t mem_pool, + hipStream_t stream) { + return hipCUDAErrorTohipError(cudaMallocFromPoolAsync(dev_ptr, size, mem_pool, stream)); +} + +inline static hipError_t hipMemPoolExportToShareableHandle( + void* shared_handle, + hipMemPool_t mem_pool, + hipMemAllocationHandleType handle_type, + unsigned int flags) { + return hipCUDAErrorTohipError(cudaMemPoolExportToShareableHandle( + shared_handle, mem_pool, handle_type, flags)); +} + +inline static hipError_t hipMemPoolImportFromShareableHandle( + hipMemPool_t* mem_pool, + void* shared_handle, + hipMemAllocationHandleType handle_type, + unsigned int flags) { + return hipCUDAErrorTohipError(cudaMemPoolImportFromShareableHandle( + mem_pool, shared_handle, handle_type, flags)); +} + +inline static hipError_t hipMemPoolExportPointer(hipMemPoolPtrExportData* export_data, void* ptr) { + return hipCUDAErrorTohipError(cudaMemPoolExportPointer(export_data, ptr)); +} + +inline static hipError_t hipMemPoolImportPointer( + void** ptr, + hipMemPool_t mem_pool, + hipMemPoolPtrExportData* export_data) { + return hipCUDAErrorTohipError(cudaMemPoolImportPointer(ptr, mem_pool, export_data)); +} +#endif // CUDA_VERSION >= CUDA_11020 + #ifdef __cplusplus } #endif diff --git a/hipamd/src/CMakeLists.txt b/hipamd/src/CMakeLists.txt index a56340b520..64da097e74 100644 --- a/hipamd/src/CMakeLists.txt +++ b/hipamd/src/CMakeLists.txt @@ -112,6 +112,8 @@ target_sources(amdhip64 PRIVATE hip_hmm.cpp hip_intercept.cpp hip_memory.cpp + hip_mempool.cpp + hip_mempool_impl.cpp hip_module.cpp hip_peer.cpp hip_platform.cpp diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index b7b2bb7385..540ec0e05b 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -83,6 +83,23 @@ hipMalloc hipMalloc3D hipMalloc3DArray hipMallocManaged +hipDeviceGetDefaultMemPool +hipDeviceSetMemPool +hipDeviceGetMemPool +hipMallocAsync +hipFreeAsync +hipMemPoolTrimTo +hipMemPoolSetAttribute +hipMemPoolGetAttribute +hipMemPoolSetAccess +hipMemPoolGetAccess +hipMemPoolCreate +hipMemPoolDestroy +hipMallocFromPoolAsync +hipMemPoolExportToShareableHandle +hipMemPoolImportFromShareableHandle +hipMemPoolExportPointer +hipMemPoolImportPointer hipArrayCreate hipArray3DCreate hipArrayDestroy diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp index 509cca908d..d2f64691b7 100644 --- a/hipamd/src/hip_context.cpp +++ b/hipamd/src/hip_context.cpp @@ -63,7 +63,11 @@ bool init() { if (context && CL_SUCCESS != context->create(nullptr)) { context->release(); } else { - g_devices.push_back(new Device(context, i)); + auto device = new Device(context, i); + if ((device == nullptr) || !device->Create()) { + return false; + } + g_devices.push_back(device); } } diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index 0ed72596f7..7a2a1c0e42 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -21,6 +21,7 @@ #include #include "hip_internal.hpp" +#include "hip_mempool_impl.hpp" namespace hip { @@ -35,6 +36,58 @@ amd::HostQueue* Device::NullStream(bool skip_alloc) { return null_queue; } +// ================================================================================================ +bool Device::Create() { + // Create default memory pool + default_mem_pool_ = new MemoryPool(this); + if (default_mem_pool_ == nullptr) { + return false; + } + // Current is default pool after device creation + current_mem_pool_ = default_mem_pool_; + return true; +} + +// ================================================================================================ +void Device::AddMemoryPool(MemoryPool* pool) { + if (auto it = mem_pools_.find(pool); it == mem_pools_.end()) { + mem_pools_.insert(pool); + } +} + +// ================================================================================================ +void Device::RemoveMemoryPool(MemoryPool* pool) { + if (auto it = mem_pools_.find(pool); it != mem_pools_.end()) { + mem_pools_.erase(it); + } +} + +// ================================================================================================ +bool Device::FreeMemory(amd::Memory* memory, Stream* stream) { + // Search for memory in the entire list of pools + for (auto& it : mem_pools_) { + if (it->FreeMemory(memory, stream)) { + return true; + } + } + return false; +} + +// ================================================================================================ +void Device::ReleaseFreedMemory(Stream* stream) { + // Search for memory in the entire list of pools + for (auto& it : mem_pools_) { + it->ReleaseFreedMemory(stream); + } +} + +// ================================================================================================ +Device::~Device() { + if (default_mem_pool_ != nullptr) { + default_mem_pool_->release(); + } +} + } namespace amd { diff --git a/hipamd/src/hip_device_runtime.cpp b/hipamd/src/hip_device_runtime.cpp index 2291afa864..f094702e88 100644 --- a/hipamd/src/hip_device_runtime.cpp +++ b/hipamd/src/hip_device_runtime.cpp @@ -329,6 +329,9 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) case hipDeviceAttributeFineGrainSupport: *pi = static_cast(g_devices[device]->devices()[0]->isFineGrainSupported()); break; + case hipDeviceAttributeMemoryPoolsSupported: + *pi = HIP_MEM_POOL_SUPPORT; + break; default: HIP_RETURN(hipErrorInvalidValue); } diff --git a/hipamd/src/hip_hcc.def.in b/hipamd/src/hip_hcc.def.in index 68d6b0038c..17f0f261f4 100644 --- a/hipamd/src/hip_hcc.def.in +++ b/hipamd/src/hip_hcc.def.in @@ -83,6 +83,23 @@ hipMalloc hipMalloc3D hipMalloc3DArray hipMallocManaged +hipDeviceGetDefaultMemPool +hipDeviceSetMemPool +hipDeviceGetMemPool +hipMallocAsync +hipFreeAsync +hipMemPoolTrimTo +hipMemPoolSetAttribute +hipMemPoolGetAttribute +hipMemPoolSetAccess +hipMemPoolGetAccess +hipMemPoolCreate +hipMemPoolDestroy +hipMallocFromPoolAsync +hipMemPoolExportToShareableHandle +hipMemPoolImportFromShareableHandle +hipMemPoolExportPointer +hipMemPoolImportPointer hipArrayCreate hipArray3DCreate hipArrayDestroy diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index 982fab1445..b9a4092e9d 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -400,6 +400,23 @@ local: hip_5.1 { global: hipDeviceGetUuid; + hipDeviceGetDefaultMemPool; + hipDeviceSetMemPool; + hipDeviceGetMemPool; + hipMallocAsync; + hipFreeAsync; + hipMemPoolTrimTo; + hipMemPoolSetAttribute; + hipMemPoolGetAttribute; + hipMemPoolSetAccess; + hipMemPoolGetAccess; + hipMemPoolCreate; + hipMemPoolDestroy; + hipMallocFromPoolAsync; + hipMemPoolExportToShareableHandle; + hipMemPoolImportFromShareableHandle; + hipMemPoolExportPointer; + hipMemPoolImportPointer; local: *; } hip_5.0; diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 398e6c0651..67733e15aa 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -1,4 +1,4 @@ -/* Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. +/* Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal @@ -195,6 +195,7 @@ class accelerator_view; }; namespace hip { class Device; + class MemoryPool; class Stream { public: enum Priority : int { High = -1, Normal = 0, Low = 1 }; @@ -244,6 +245,8 @@ namespace hip { void Finish() const; /// Get device ID associated with the current stream; int DeviceId() const; + /// Get HIP device associated with the stream + Device* GetDevice() const { return device_; } /// Get device ID associated with a stream; static int DeviceId(const hipStream_t hStream); /// Returns if stream is null stream @@ -344,15 +347,23 @@ namespace hip { std::vector queues_; + MemoryPool* default_mem_pool_; + MemoryPool* current_mem_pool_; + + std::set mem_pools_; + public: Device(amd::Context* ctx, int devId): context_(ctx), - deviceId_(devId), - null_stream_(this, Stream::Priority::Normal, 0, true), + deviceId_(devId), + null_stream_(this, Stream::Priority::Normal, 0, true), flags_(hipDeviceScheduleSpin), - isActive_(false) + isActive_(false), + default_mem_pool_(nullptr), + current_mem_pool_(nullptr) { assert(ctx != nullptr); } - ~Device() {} + ~Device(); + bool Create(); amd::Context* asContext() const { return context_; } int deviceId() const { return deviceId_; } void retain() const { context_->retain(); } @@ -397,6 +408,29 @@ namespace hip { } return false; } + + /// Set the current memory pool on the device + void SetCurrentMemoryPool(MemoryPool* pool = nullptr) { + current_mem_pool_ = (pool == nullptr) ? default_mem_pool_ : pool; + } + + /// Get the current memory pool on the device + MemoryPool* GetCurrentMemoryPool() const { return current_mem_pool_; } + + /// Get the default memory pool on the device + MemoryPool* GetDefaultMemoryPool() const { return default_mem_pool_; } + + /// Add memory pool to the device + void AddMemoryPool(MemoryPool* pool); + + /// Remove memory pool from the device + void RemoveMemoryPool(MemoryPool* pool); + + /// Free memory from the device + bool FreeMemory(amd::Memory* memory, Stream* stream); + + /// Release freed memory from all pools on the current device + void ReleaseFreedMemory(Stream* stream); }; /// Current thread's device diff --git a/hipamd/src/hip_mempool.cpp b/hipamd/src/hip_mempool.cpp new file mode 100644 index 0000000000..5814c4360e --- /dev/null +++ b/hipamd/src/hip_mempool.cpp @@ -0,0 +1,229 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include "hip_mempool_impl.hpp" + +/** + * API interfaces + */ +extern hipError_t ihipFree(void* ptr); + +// ================================================================================================ +hipError_t hipDeviceGetDefaultMemPool(hipMemPool_t* mem_pool, int device) { + HIP_INIT_API(hipDeviceGetDefaultMemPool, mem_pool, device); + if ((mem_pool == nullptr) || (device >= g_devices.size())) { + HIP_RETURN(hipErrorInvalidValue); + } + *mem_pool = reinterpret_cast(g_devices[device]->GetDefaultMemoryPool()); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipDeviceSetMemPool(int device, hipMemPool_t mem_pool) { + HIP_INIT_API(hipDeviceSetMemPool, device, mem_pool); + if ((mem_pool == nullptr) || (device >= g_devices.size())) { + HIP_RETURN(hipErrorInvalidValue); + } + g_devices[device]->SetCurrentMemoryPool(reinterpret_cast(mem_pool)); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipDeviceGetMemPool(hipMemPool_t* mem_pool, int device) { + HIP_INIT_API(hipDeviceGetMemPool, mem_pool, device); + if ((mem_pool == nullptr) || (device >= g_devices.size())) { + HIP_RETURN(hipErrorInvalidValue); + } + *mem_pool = reinterpret_cast(g_devices[device]->GetCurrentMemoryPool()); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) { + HIP_INIT_API(hipMallocAsync, dev_ptr, size, stream); + if ((dev_ptr == nullptr) || (size == 0) || (stream == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + auto device = reinterpret_cast(stream)->GetDevice(); + auto mem_pool = device->GetCurrentMemoryPool(); + *dev_ptr = reinterpret_cast(mem_pool)->AllocateMemory( + size, reinterpret_cast(stream)); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) { + HIP_INIT_API(hipFreeAsync, dev_ptr, stream); + if ((dev_ptr == nullptr) || (stream == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + size_t offset = 0; + auto memory = getMemoryObject(dev_ptr, offset); + auto id = memory->getUserData().deviceId; + if (!g_devices[id]->FreeMemory(memory, reinterpret_cast(stream))) { + //! @todo It's not the most optimal logic. The current implementation has unconditional waits + HIP_RETURN(ihipFree(dev_ptr)); + } + + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolTrimTo(hipMemPool_t mem_pool, size_t min_bytes_to_hold) { + HIP_INIT_API(hipMemPoolTrimTo, mem_pool, min_bytes_to_hold); + if (mem_pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + hip::MemoryPool* hip_mem_pool = reinterpret_cast(mem_pool); + hip_mem_pool->TrimTo(min_bytes_to_hold); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolSetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value) { + HIP_INIT_API(hipMemPoolSetAttribute, mem_pool, attr, value); + if (mem_pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + auto hip_mem_pool = reinterpret_cast(mem_pool); + HIP_RETURN(hip_mem_pool->SetAttribute(attr, value)); +} + +// ================================================================================================ +hipError_t hipMemPoolGetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value) { + HIP_INIT_API(hipMemPoolGetAttribute, mem_pool, attr, value); + if (mem_pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + if (mem_pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + auto hip_mem_pool = reinterpret_cast(mem_pool); + HIP_RETURN(hip_mem_pool->GetAttribute(attr, value)); +} + +// ================================================================================================ +hipError_t hipMemPoolSetAccess( + hipMemPool_t mem_pool, + const hipMemAccessDesc* desc_list, + size_t count) { + HIP_INIT_API(hipMemPoolSetAccess, mem_pool, desc_list, count); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolGetAccess( + hipMemAccessFlags* flags, + hipMemPool_t mem_pool, + hipMemLocation* location) { + HIP_INIT_API(hipMemPoolGetAccess, flags, mem_pool, location); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolCreate(hipMemPool_t* mem_pool, const hipMemPoolProps* pool_props) { + HIP_INIT_API(hipMemPoolCreate, mem_pool, pool_props); + if (mem_pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + // Make sure the pool creation occurs on a valid device + if ((pool_props->location.type != hipMemLocationTypeDevice) || + (pool_props->location.id >= g_devices.size())) { + HIP_RETURN(hipErrorInvalidValue); + } + auto device = g_devices[pool_props->location.id]; + auto pool = new hip::MemoryPool(device); + if (pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + *mem_pool = reinterpret_cast(pool); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolDestroy(hipMemPool_t mem_pool) { + HIP_INIT_API(hipMemPoolDestroy, mem_pool); + if (mem_pool == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + hip::MemoryPool* hip_mem_pool = reinterpret_cast(mem_pool); + hip_mem_pool->ReleaseFreedMemory(); + + auto device = hip_mem_pool->Device(); + + // Force default pool if the current one is destroyed + if (hip_mem_pool == device->GetCurrentMemoryPool()) { + device->SetCurrentMemoryPool(device->GetDefaultMemoryPool()); + } + + hip_mem_pool->release(); + + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMallocFromPoolAsync( + void** dev_ptr, + size_t size, + hipMemPool_t mem_pool, + hipStream_t stream) { + HIP_INIT_API(hipMallocFromPoolAsync, dev_ptr, size, mem_pool, stream); + if ((dev_ptr == nullptr) || (size == 0) || (mem_pool == nullptr) || (stream == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + *dev_ptr = reinterpret_cast(mem_pool)->AllocateMemory( + size, reinterpret_cast(stream)); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolExportToShareableHandle( + void* shared_handle, + hipMemPool_t mem_pool, + hipMemAllocationHandleType handle_type, + unsigned int flags) { + HIP_INIT_API(hipMemPoolExportToShareableHandle, shared_handle, mem_pool, handle_type, flags); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolImportFromShareableHandle( + hipMemPool_t* mem_pool, + void* shared_handle, + hipMemAllocationHandleType handle_type, + unsigned int flags) { + HIP_INIT_API(hipMemPoolImportFromShareableHandle, mem_pool, shared_handle, handle_type, flags); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolExportPointer(hipMemPoolPtrExportData* export_data, void* ptr) { + HIP_INIT_API(hipMemPoolExportPointer, export_data, ptr); + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemPoolImportPointer( + void** ptr, + hipMemPool_t mem_pool, + hipMemPoolPtrExportData* export_data) { + HIP_INIT_API(hipMemPoolImportPointer, ptr, mem_pool, export_data); + HIP_RETURN(hipSuccess); +} diff --git a/hipamd/src/hip_mempool_impl.cpp b/hipamd/src/hip_mempool_impl.cpp new file mode 100644 index 0000000000..6d03124031 --- /dev/null +++ b/hipamd/src/hip_mempool_impl.cpp @@ -0,0 +1,299 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include "hip_mempool_impl.hpp" + +namespace hip { + +// ================================================================================================ +void Heap::AddMemory(amd::Memory* memory, hip::Stream* stream) { + allocations_.insert({memory, {stream, nullptr}}); + total_size_ += memory->getSize(); + max_total_size_ = std::max(max_total_size_, total_size_); +} + +// ================================================================================================ +void Heap::AddMemory(amd::Memory* memory, const MemoryTimestamp& ts) { + allocations_.insert({memory, ts}); + total_size_ += memory->getSize(); + max_total_size_ = std::max(max_total_size_, total_size_); +} + +// ================================================================================================ +amd::Memory* Heap::FindMemory(size_t size, hip::Stream* stream, bool opportunistic) { + amd::Memory* memory = nullptr; + for (auto it = allocations_.begin(); it != allocations_.end();) { + // Check if size can match and it's safe to use this resource + if ((it->first->getSize() >= size) && (it->second.IsSafeFind(stream, opportunistic))) { + memory = it->first; + total_size_ -= memory->getSize(); + // Remove found allocation from the map + it = allocations_.erase(it); + break; + } else { + ++it; + } + } + return memory; +} + +// ================================================================================================ +bool Heap::RemoveMemory(amd::Memory* memory, MemoryTimestamp* ts) { + if (auto it = allocations_.find(memory); it != allocations_.end()) { + if (ts != nullptr) { + // Preserve timestamp info for possible reuse later + *ts = it->second; + } else { + // Runtime will delete the timestamp object, hence make sure HIP event is released + it->second.Wait(); + it->second.SetEvent(nullptr); + } + total_size_ -= memory->getSize(); + allocations_.erase(it); + return true; + } + return false; +} + +// ================================================================================================ +std::unordered_map::iterator +Heap::EraseAllocaton(std::unordered_map::iterator& it) { + const device::Memory* dev_mem = it->first->getDeviceMemory(*device_->devices()[0]); + amd::SvmBuffer::free(it->first->getContext(), reinterpret_cast(dev_mem->virtualAddress())); + total_size_ -= it->first->getSize(); + // Clear HIP event + it->second.SetEvent(nullptr); + // Remove the allocation from the map + return allocations_.erase(it); +} + +// ================================================================================================ +bool Heap::ReleaseAllMemory(size_t min_bytes_to_hold, bool safe_release) { + for (auto it = allocations_.begin(); it != allocations_.end();) { + // Make sure the heap is smaller than the minimum value to hold + if (total_size_ <= min_bytes_to_hold) { + return true; + } + // Safe release forces unconditional wait for memory + if (safe_release) { + it->second.Wait(); + } + if (it->second.IsSafeRelease()) { + it = EraseAllocaton(it); + } else { + ++it; + } + } + return true; +} + +// ================================================================================================ +bool Heap::ReleaseAllMemory(hip::Stream* stream) { + for (auto it = allocations_.begin(); it != allocations_.end();) { + // Make sure the heap holds the minimum number of bytes + if (total_size_ <= release_threshold_) { + return true; + } + if (it->second.IsSafeRelease()) { + it = EraseAllocaton(it); + } else { + ++it; + } + } + return true; +} + +// ================================================================================================ +void* MemoryPool::AllocateMemory(size_t size, hip::Stream* stream) { + amd::ScopedLock lock(lock_pool_ops_); + + void* dev_ptr = nullptr; + amd::Memory* memory = free_heap_.FindMemory(size, stream, Opportunistic()); + if (memory == nullptr) { + amd::Context* context = device_->asContext(); + const auto& dev_info = context->devices()[0]->info(); + if (dev_info.maxMemAllocSize_ < size) { + return nullptr; + } + + dev_ptr = amd::SvmBuffer::malloc(*context, 0, size, dev_info.memBaseAddrAlign_, nullptr); + if (dev_ptr == nullptr) { + size_t free = 0, total =0; + hipError_t err = hipMemGetInfo(&free, &total); + if (err == hipSuccess) { + LogPrintfError("Allocation failed : Device memory : required :%zu | free :%zu | total :%zu \n", + size, free, total); + } + return nullptr; + } + + size_t offset = 0; + memory = getMemoryObject(dev_ptr, offset); + // Saves the current device id so that it can be accessed later + memory->getUserData().deviceId = device_->deviceId(); + } else { + free_heap_.RemoveMemory(memory); + const device::Memory* dev_mem = memory->getDeviceMemory(*device_->devices()[0]); + dev_ptr = reinterpret_cast(dev_mem->virtualAddress()); + } + // Place the allocated memory into the busy heap + busy_heap_.AddMemory(memory, stream); + + // Increment the reference counter on the pool + retain(); + + return dev_ptr; +} + +// ================================================================================================ +bool MemoryPool::FreeMemory(amd::Memory* memory, hip::Stream* stream) { + amd::ScopedLock lock(lock_pool_ops_); + + MemoryTimestamp ts; + // Remove memory object fro the busy pool + if (!busy_heap_.RemoveMemory(memory, &ts)) { + // This pool doesn't contain memory + return false; + } + // The stream of destruction is a safe stream, because the app must handle sync + ts.AddSafeStream(stream); + + // Add a marker to the stream to trace availability of this memory + Event* e = new hip::Event(0); + if (e != nullptr) { + if (hipSuccess == e->addMarker(reinterpret_cast(stream), nullptr, true)) { + ts.SetEvent(e); + } + } + free_heap_.AddMemory(memory, ts); + + // Decrement the reference counter on the pool + release(); + + return true; +} + +// ================================================================================================ +void MemoryPool::ReleaseFreedMemory(hip::Stream* stream) { + amd::ScopedLock lock(lock_pool_ops_); + + free_heap_.ReleaseAllMemory(stream); +} + +// ================================================================================================ +void MemoryPool::TrimTo(size_t min_bytes_to_hold) { + amd::ScopedLock lock(lock_pool_ops_); + + free_heap_.ReleaseAllMemory(min_bytes_to_hold); +} + +// ================================================================================================ +hipError_t MemoryPool::SetAttribute(hipMemPoolAttr attr, void* value) { + amd::ScopedLock lock(lock_pool_ops_); + uint64_t reset; + + switch (attr) { + case hipMemPoolReuseFollowEventDependencies: + // Enable/disable HIP events tracking from the app's dependencies + state_.event_dependencies_ = *reinterpret_cast(value); + break; + case hipMemPoolReuseAllowOpportunistic: + // Enable/disable HIP event check for freed memory + state_.opportunistic_ = *reinterpret_cast(value); + break; + case hipMemPoolReuseAllowInternalDependencies: + // Enable/disable internal extra dependencies introduced in runtime + state_.internal_dependencies_ = *reinterpret_cast(value); + break; + case hipMemPoolAttrReleaseThreshold: + free_heap_.SetReleaseThreshold(*reinterpret_cast(value)); + break; + case hipMemPoolAttrReservedMemCurrent: + // Should be GetAttribute only + return hipErrorInvalidValue; + break; + case hipMemPoolAttrReservedMemHigh: + reset = *reinterpret_cast(value); + // Only 0 is accepted + if (reset != 0) { + return hipErrorInvalidValue; + } + free_heap_.SetMaxTotalSize(reset); + case hipMemPoolAttrUsedMemCurrent: + // Should be GetAttribute only + return hipErrorInvalidValue; + break; + case hipMemPoolAttrUsedMemHigh: + reset = *reinterpret_cast(value); + // Only 0 is accepted + if (reset != 0) { + return hipErrorInvalidValue; + } + busy_heap_.SetMaxTotalSize(reset); + break; + default: + return hipErrorInvalidValue; + } + return hipSuccess; +} + +// ================================================================================================ +hipError_t MemoryPool::GetAttribute(hipMemPoolAttr attr, void* value) { + amd::ScopedLock lock(lock_pool_ops_); + + switch (attr) { + case hipMemPoolReuseFollowEventDependencies: + // Enable/disable HIP events tracking from the app's dependencies + *reinterpret_cast(value) = EventDependencies(); + break; + case hipMemPoolReuseAllowOpportunistic: + // Enable/disable HIP event check for freed memory + *reinterpret_cast(value) = Opportunistic(); + break; + case hipMemPoolReuseAllowInternalDependencies: + // Enable/disable internal extra dependencies introduced in runtime + *reinterpret_cast(value) = InternalDependencies(); + break; + case hipMemPoolAttrReleaseThreshold: + *reinterpret_cast(value) = free_heap_.GetReleaseThreshold(); + break; + case hipMemPoolAttrReservedMemCurrent: + // All allocate memory by the pool in OS + *reinterpret_cast(value) = busy_heap_.GetTotalSize() + free_heap_.GetTotalSize(); + break; + case hipMemPoolAttrReservedMemHigh: + // High watermark of all allocated memory in OS, since the last reset + *reinterpret_cast(value) = busy_heap_.GetTotalSize() + free_heap_.GetMaxTotalSize(); + break; + case hipMemPoolAttrUsedMemCurrent: + // Total currently used memory by the pool + *reinterpret_cast(value) = busy_heap_.GetTotalSize(); + break; + case hipMemPoolAttrUsedMemHigh: + // High watermark of all used memoryS, since the last reset + *reinterpret_cast(value) = busy_heap_.GetMaxTotalSize(); + break; + default: + return hipErrorInvalidValue; + } + return hipSuccess; +} + +} diff --git a/hipamd/src/hip_mempool_impl.hpp b/hipamd/src/hip_mempool_impl.hpp new file mode 100644 index 0000000000..6fa945a0f9 --- /dev/null +++ b/hipamd/src/hip_mempool_impl.hpp @@ -0,0 +1,212 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#pragma once + +#include +#include "hip_event.hpp" +#include "hip_internal.hpp" +#include +#include + +namespace hip { + +class Device; +class Stream; + +struct MemoryTimestamp { + MemoryTimestamp(hip::Stream* stream, hip::Event* event = nullptr): event_(event) { + safe_streams_.insert(stream); + } + MemoryTimestamp(): event_(nullptr) {} + + /// Adds a safe stream to the list of stream for possible reuse + void AddSafeStream(hip::Stream* stream) { + if (safe_streams_.find(stream) != safe_streams_.end()) { + safe_streams_.insert(stream); + } + } + /// Changes last known valid event asociated with memory + void SetEvent(hip::Event* event) { + delete event_; + event_ = event; + } + /// Wait for memory to be available + void Wait() { + if (event_ != nullptr) { + auto hip_error = event_->synchronize(); + } + } + /// Returns if memory object is safe for reuse + bool IsSafeFind(hip::Stream* stream = nullptr, bool opportunistic = true) { + bool result = false; + if (safe_streams_.find(stream) != safe_streams_.end()) { + // A safe stream doesn't require TS validation + result = true; + } else if (opportunistic && (event_ != nullptr)) { + // Check HIP event for a retired status + result = (event_->query() == hipSuccess) ? true : false; + } + return result; + } + /// Returns if memory object is safe for reuse + bool IsSafeRelease() { + bool result = true; + if (event_ != nullptr) { + // Check HIP event for a retired status + result = (event_->query() == hipSuccess) ? true : false; + } + return result; + } + + std::unordered_set safe_streams_; //!< Safe streams for memory reuse + hip::Event* event_; //!< Last known HIP event, associated with the memory object +}; + +class Heap : public amd::EmbeddedObject { +public: + Heap(hip::Device* device): + total_size_(0), max_total_size_(0), release_threshold_(0), device_(device) {} + ~Heap() {} + + /// Adds allocation into the heap on a specific stream + void AddMemory(amd::Memory* memory, hip::Stream* stream); + + /// Adds allocation into the heap with specific TS + void AddMemory(amd::Memory* memory, const MemoryTimestamp& ts); + + /// Finds memory object with the specified size + amd::Memory* FindMemory(size_t size, hip::Stream* stream, bool opportunistic); + + /// Removes allocation from the map + bool RemoveMemory(amd::Memory* memory, MemoryTimestamp* ts = nullptr); + + /// Releases all memory, until the threshold value is met + bool ReleaseAllMemory(size_t min_bytes_to_hold = std::numeric_limits::max(), bool safe_release = false); + + /// Releases all memory, safe to the provided stream, until the threshold value is met + bool ReleaseAllMemory(hip::Stream* stream); + + /// Heap doesn't have any allocations + bool IsEmpty() const { return (allocations_.size() == 0) ? true : false; } + + /// Set the memory release threshold + void SetReleaseThreshold(uint64_t value) { release_threshold_ = value; } + + /// Set the memory release threshold + uint64_t GetReleaseThreshold() const { return release_threshold_; } + + /// Get the size of all allocations in the heap + uint64_t GetTotalSize() const { return total_size_; } + + /// Get the size of all allocations in the heap + uint64_t GetMaxTotalSize() const { return max_total_size_; } + + /// Set maximum total, allocated by the heap + void SetMaxTotalSize(uint64_t value) { max_total_size_ = value; } + + std::unordered_map::iterator EraseAllocaton( + std::unordered_map::iterator& it); + +private: + Heap() = delete; + Heap(const Heap&) = delete; + Heap& operator=(const Heap&) = delete; + + std::unordered_map allocations_; //!< Map of allocations on a specific stream + uint64_t total_size_; //!< Size of all allocations in the heap + uint64_t max_total_size_; //!< Maximum heap allocation size + uint64_t release_threshold_; //!< Threshold size in bytes for memory release from heap, default 0 + + hip::Device* device_; //!< Hip device the allocations will reside +}; + +/// Allocates memory in the pool on the specified stream and places the allocation into busy_heap_ +/// @note: the logic also will look in free_heap for possible reuse. +/// hipMemPoolReuseAllowOpportunistic option will validate if HIP event, +/// associated with memory is done, then reuse can be performed. +class MemoryPool : public amd::ReferenceCountedObject { +public: + MemoryPool(hip::Device* device): + busy_heap_(device), + free_heap_(device), + lock_pool_ops_("Pool operations", true), device_(device) { + device_->AddMemoryPool(this); + state_.event_dependencies_ = 1; + state_.opportunistic_ = 1; + state_.internal_dependencies_ = 1; + } + virtual ~MemoryPool() { + assert(busy_heap_.IsEmpty() && "Can't destroy pool with busy allocations!"); + constexpr bool kSafeRelease = true; + free_heap_.ReleaseAllMemory(0, kSafeRelease); + // Remove memory pool from the list of all pool on the current device + device_->RemoveMemoryPool(this); + } + + /// The same stream can reuse memory without HIP event validation + void* AllocateMemory(size_t size, hip::Stream* stream); + + /// Frees memory by placing memory object with HIP event into free_heap_ + bool FreeMemory(amd::Memory* memory, hip::Stream* stream); + + /// Releases all allocations from free_heap_. It can be called on Stream or Device synchronization + /// @note The caller must make sure it's safe to release memory + void ReleaseFreedMemory(hip::Stream* stream = nullptr); + + /// Releases all allocations in MemoryPool + void ReleaseAllMemory(); + + /// Trims the pool until it has only min_bytes_to_hold + void TrimTo(size_t min_bytes_to_hold); + + /// Trims the pool until it has only min_bytes_to_hold + hip::Device* Device() const { return device_; } + + /// Set memory pool control attributes + hipError_t SetAttribute(hipMemPoolAttr attr, void* value); + + /// Get memory pool control attributes + hipError_t GetAttribute(hipMemPoolAttr attr, void* value); + + /// Accessors for the pool state + bool EventDependencies() const { return (state_.event_dependencies_) ? true : false; } + bool Opportunistic() const { return (state_.opportunistic_) ? true : false; } + bool InternalDependencies() const { return (state_.internal_dependencies_) ? true : false; } + +private: + MemoryPool() = delete; + MemoryPool(const MemoryPool&) = delete; + MemoryPool& operator=(const MemoryPool&) = delete; + + + Heap busy_heap_; //!< Heap of busy allocations + Heap free_heap_; //!< Heap of freed allocations + struct { + uint32_t event_dependencies_ : 1; //!< Event dependencies tracking is enabled + uint32_t opportunistic_ : 1; //!< HIP event check is enabled + uint32_t internal_dependencies_ : 1; //!< Runtime adds internal events to handle memory dependencies + } state_; + + amd::Monitor lock_pool_ops_; //!< Access to the pool must be lock protected + hip::Device* device_; //!< Hip device the heap will reside +}; + +} // Mamespace hip diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index a8b8b9138c..3daab05ea5 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -407,6 +407,12 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { // Wait for the current host queue hip::getQueue(stream)->finish(); + // Make sure runtime releases memory for all memory pools on the device, + // associated with the queue + auto hip_stream = reinterpret_cast(stream); + auto device = (hip_stream == nullptr) ? hip::getCurrentDevice() : hip_stream->GetDevice(); + device->ReleaseFreedMemory(hip_stream); + HIP_RETURN(hipSuccess); }