AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività

Fixed hipMallocArray for 1D cases

Sfoglia il codice sorgente 
[ROCm/hip commit: 2a915b2790]
This commit is contained in:
Rahul Garg
2017-09-12 21:52:11 +05:30
parent f66a03fa9c
commit 3178dc04cd
1 ha cambiato i file con 69 aggiunte e 63 eliminazioni
Scarica il file Patch Scarica il file Diff Expand all files Collapse all files
projects/hip/src/hip_memory.cpp
+69 -63
Vedi File
@@ -407,83 +407,89 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
HIP_INIT_SPECIAL_API((TRACE_MEM), array, desc, width, height, flags);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if(width > 0) {
auto ctx = ihipGetTlsDefaultCtx();
auto ctx = ihipGetTlsDefaultCtx();
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = flags;
array[0]->width = width;
array[0]->height = height;
array[0]->depth = 1;
array[0]->desc = *desc;
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->type = flags;
array[0]->width = width;
array[0]->height = height;
array[0]->depth = 1;
array[0]->desc = *desc;
void ** ptr = &array[0]->data;
void ** ptr = &array[0]->data;
if (ctx) {
const unsigned am_flags = 0;
size_t size = width;
if(height > 0) {
size = size * height;
}
if (ctx) {
const unsigned am_flags = 0;
const size_t size = width*height;
size_t allocSize = 0;
switch(desc->f) {
case hipChannelFormatKindSigned:
allocSize = size * sizeof(int);
break;
case hipChannelFormatKindUnsigned:
allocSize = size * sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
allocSize = size * sizeof(float);
break;
case hipChannelFormatKindNone:
allocSize = size * sizeof(size_t);
break;
default:
hip_status = hipErrorUnknown;
break;
}
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
size_t allocSize = 0;
switch(desc->f) {
case hipChannelFormatKindSigned:
allocSize = size * sizeof(int);
break;
case hipChannelFormatKindUnsigned:
allocSize = size * sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
allocSize = size * sizeof(float);
break;
case hipChannelFormatKindNone:
allocSize = size * sizeof(size_t);
size_t allocGranularity = 0;
hsa_amd_memory_pool_t *allocRegion = static_cast<hsa_amd_memory_pool_t*>(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:
hip_status = hipErrorUnknown;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
break;
}
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
}
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;
size_t allocGranularity = 0;
hsa_amd_memory_pool_t *allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
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;
hsa_ext_image_descriptor_t imageDescriptor;
*ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment);
if (size && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
}
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)) {
} else {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorMemoryAllocation;
hip_status = hipErrorInvalidValue;
}
return ihipLogStatus(hip_status);