diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index b89099ea0d..07d58599cf 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1866,8 +1866,8 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, } hc::accelerator acc; - hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); - hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo dstPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo srcPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); bool dstTracked = getTailoredPtrInfo(&dstPtrInfo, dst, sizeBytes); bool srcTracked = getTailoredPtrInfo(&srcPtrInfo, src, sizeBytes); @@ -1902,7 +1902,7 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, } void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes) { - hc::AmPointerInfo ptrInfo(NULL, ptr, sizeBytes, acc, true, false); + hc::AmPointerInfo ptrInfo(NULL, ptr, ptr, sizeBytes, acc, true, false); hc::am_memtracker_add(ptr, ptrInfo); } @@ -1926,7 +1926,7 @@ void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* { // TODO - review - this looks broken , should not be adding pointers to tracker dynamically: if(kind == hipMemcpyHostToDevice) { - hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo srcPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); if(srcTracked) { addSymbolPtrToTracker(acc, dst, sizeBytes); @@ -1938,7 +1938,7 @@ void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* } } if(kind == hipMemcpyDeviceToHost) { - hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo dstPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); if(dstTracked) { addSymbolPtrToTracker(acc, src, sizeBytes); @@ -1977,8 +1977,8 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes } else { hc::accelerator acc; - hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); - hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo dstPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo srcPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); bool dstTracked = getTailoredPtrInfo(&dstPtrInfo, dst, sizeBytes); bool srcTracked = getTailoredPtrInfo(&srcPtrInfo, src, sizeBytes); diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 9bb31d2d78..bf51748070 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -101,14 +101,18 @@ int sharePtr(void *ptr, ihipCtx_t *ctx, bool shareWithAll, unsigned hipFlags) // Allocate a new pointer with am_alloc and share with all valid peers. // Returns null-ptr if a memory error occurs (either allocation or sharing) -void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool shareWithAll, unsigned amFlags, unsigned hipFlags) +void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool shareWithAll, unsigned amFlags, unsigned hipFlags, size_t alignment) { void *ptr = nullptr; auto device = ctx->getWriteableDevice(); - ptr = hc::am_alloc(sizeBytes, device->_acc, amFlags); + if (alignment != 0) { + ptr = hc::am_aligned_alloc(sizeBytes, device->_acc, amFlags, alignment); + } else { + ptr = hc::am_alloc(sizeBytes, device->_acc, amFlags); + } tprintf(DB_MEM, " alloc %s ptr:%p-%p size:%zu on dev:%d\n", msg, ptr, static_cast(ptr)+sizeBytes, sizeBytes, device->_deviceId); @@ -145,7 +149,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void hipError_t e = hipSuccess; hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if (status == AM_SUCCESS) { @@ -195,7 +199,7 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi e = hipErrorInvalidValue; } else { hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPointer); if (status == AM_SUCCESS) { *devicePointer = static_cast(amPointerInfo._devicePointer) + (static_cast(hostPointer) - static_cast(amPointerInfo._hostPointer)) ; @@ -226,7 +230,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) } else { auto device = ctx->getWriteableDevice(); - *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false/*shareWithAll*/, 0/*amFlags*/, 0/*hipFlags*/); + *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false/*shareWithAll*/, 0/*amFlags*/, 0/*hipFlags*/, 0); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; @@ -294,7 +298,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) *ptr = hip_internal::allocAndSharePtr((amFlags & amHostCoherent) ? "finegrained_host":"pinned_host", - sizeBytes, ctx, (trueFlags & hipHostMallocPortable) /*shareWithAll*/, amFlags, flags); + sizeBytes, ctx, (trueFlags & hipHostMallocPortable) /*shareWithAll*/, amFlags, flags, 0); if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; @@ -340,10 +344,29 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height //err = hipMalloc(ptr, (*pitch)*height); if (ctx) { - auto device = ctx->getWriteableDevice(); + hc::accelerator acc = ctx->getDevice()->_acc; + hsa_agent_t* agent =static_cast(acc.get_hsa_agent()); + + size_t allocGranularity = 0; + hsa_amd_memory_pool_t *allocRegion = static_cast(acc.get_hsa_am_region()); + hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); + + hsa_ext_image_descriptor_t imageDescriptor; + imageDescriptor.width = *pitch; + imageDescriptor.height = height; + imageDescriptor.depth = 0; + imageDescriptor.array_size = 0; + imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; + imageDescriptor.format.channel_order = HSA_EXT_IMAGE_CHANNEL_ORDER_R; + imageDescriptor.format.channel_type = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; + + hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + hsa_ext_image_data_info_t imageInfo; + hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); + size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; const unsigned am_flags = 0; - *ptr = hip_internal::allocAndSharePtr("device_pitch", sizeBytes, ctx, false/*shareWithAll*/, am_flags, 0); + *ptr = hip_internal::allocAndSharePtr("device_pitch", sizeBytes, ctx, false/*shareWithAll*/, am_flags, 0, alignment); if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; @@ -363,6 +386,11 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel return cd; } +extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, + enum hipTextureReadMode readMode, + hsa_ext_image_channel_order_t& channelOrder, + hsa_ext_image_channel_type_t& channelType); + hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { @@ -403,7 +431,43 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, hip_status = hipErrorUnknown; break; } - *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0); + hc::accelerator acc = ctx->getDevice()->_acc; + hsa_agent_t* agent =static_cast(acc.get_hsa_agent()); + + size_t allocGranularity = 0; + hsa_amd_memory_pool_t *allocRegion = static_cast(acc.get_hsa_am_region()); + hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); + + hsa_ext_image_descriptor_t imageDescriptor; + + imageDescriptor.width = width; + imageDescriptor.height = height; + imageDescriptor.depth = 0; + imageDescriptor.array_size = 0; + switch (flags) { + case hipArrayLayered: + case hipArrayCubemap: + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + default: + imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; + break; + } + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType); + imageDescriptor.format.channel_order = channelOrder; + imageDescriptor.format.channel_type = channelType; + + hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + hsa_ext_image_data_info_t imageInfo; + hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); + size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; + + *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment); if (size && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; } @@ -457,7 +521,47 @@ hipError_t hipMalloc3DArray(hipArray_t *array, hip_status = hipErrorUnknown; break; } - *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0); + + hc::accelerator acc = ctx->getDevice()->_acc; + hsa_agent_t* agent =static_cast(acc.get_hsa_agent()); + + size_t allocGranularity = 0; + hsa_amd_memory_pool_t *allocRegion = static_cast(acc.get_hsa_am_region()); + hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity); + + hsa_ext_image_descriptor_t imageDescriptor; + imageDescriptor.width = extent.width; + imageDescriptor.height = extent.height; + imageDescriptor.depth = 0; + imageDescriptor.array_size = 0; + switch (flags) { + case hipArrayLayered: + imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA; + imageDescriptor.array_size = extent.depth; + break; + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + case hipArrayDefault: + assert(0); + break; + case hipArrayCubemap: + default: + imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; + imageDescriptor.depth = extent.depth; + break; + } + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType); + imageDescriptor.format.channel_order = channelOrder; + imageDescriptor.format.channel_type = channelType; + + hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; + hsa_ext_image_data_info_t imageInfo; + hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo); + size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; + + *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0, alignment); if (size && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; @@ -467,7 +571,6 @@ hipError_t hipMalloc3DArray(hipArray_t *array, hip_status = hipErrorMemoryAllocation; } - //return ihipLogStatus(hip_status); return hip_status; } @@ -478,7 +581,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) hipError_t hip_status = hipSuccess; hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); if(status == AM_SUCCESS){ *flagsPtr = amPointerInfo._appAllocationFlags; @@ -509,7 +612,7 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) } hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t am_status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); if(am_status == AM_SUCCESS){ @@ -1309,7 +1412,7 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) if(ptr != nullptr && size != nullptr){ hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if(status == AM_SUCCESS){ *size = amPointerInfo._sizeBytes; @@ -1334,7 +1437,7 @@ hipError_t hipFree(void* ptr) if (ptr) { hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if(status == AM_SUCCESS){ if(amPointerInfo._hostPointer == NULL){ @@ -1362,7 +1465,7 @@ hipError_t hipHostFree(void* ptr) hipError_t hipStatus = hipErrorInvalidValue; if (ptr) { hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if(status == AM_SUCCESS){ if(amPointerInfo._hostPointer == ptr){ @@ -1396,7 +1499,7 @@ hipError_t hipFreeArray(hipArray* array) if(array->data) { hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data); if(status == AM_SUCCESS){ if(amPointerInfo._hostPointer == NULL){ @@ -1414,7 +1517,7 @@ hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDevi HIP_INIT_API ( pbase , psize , dptr ); hipError_t hipStatus = hipSuccess; hc::accelerator acc; - hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 ); + hc::AmPointerInfo amPointerInfo( NULL , NULL , NULL, 0 , acc , 0 , 0 ); am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , dptr ); if (status == AM_SUCCESS) { *pbase = amPointerInfo._devicePointer; @@ -1437,7 +1540,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){ if((handle == NULL) || (devPtr == NULL)) { hipStatus = hipErrorInvalidResourceHandle; } else { - hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 ); + hc::AmPointerInfo amPointerInfo( NULL , NULL , NULL, 0 , acc , 0 , 0 ); am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , devPtr ); if (status == AM_SUCCESS) { psize = (size_t)amPointerInfo._sizeBytes;