From 4264a4ce563f9e6b6f27b44583ce1dc31284cb40 Mon Sep 17 00:00:00 2001 From: Weixing Zhang Date: Thu, 3 Aug 2017 00:29:10 -0400 Subject: [PATCH 1/6] [HIP Texture] The GPU virtual address for texture memory needs to be aligned. In hcc_am, a bigger buffer will be allocated for alignment purpose and _unalignedDevicePointer is added in struct AmPointerInfo for original allocated address. --- src/hip_hcc.cpp | 14 ++--- src/hip_memory.cpp | 141 +++++++++++++++++++++++++++++++++++++++------ 2 files changed, 129 insertions(+), 26 deletions(-) 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; From 9a04c5cdd419b30b68970f0164d579eed627bd3b Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 10 Aug 2017 17:58:36 +0000 Subject: [PATCH 2/6] add C++11 compilation flags and minor bug fixes --- tests/src/deviceLib/hip_threadfence_system.cpp | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/tests/src/deviceLib/hip_threadfence_system.cpp b/tests/src/deviceLib/hip_threadfence_system.cpp index 673cb4507a..be59d56ab1 100644 --- a/tests/src/deviceLib/hip_threadfence_system.cpp +++ b/tests/src/deviceLib/hip_threadfence_system.cpp @@ -21,11 +21,13 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../test_common.cpp -std=c++11 * RUN: %t * HIT_END */ +#include +#include #include #include #include @@ -33,15 +35,14 @@ THE SOFTWARE. #include "hip/device_functions.h" #include "test_common.h" - #define HIP_ASSERT(x) (assert((x)==hipSuccess)) -__host__ void fence_system() { - std::atomic_thread_fence(std::memory_order_seq_cst); -} - -__device__ void fence_system() { +__host__ __device__ void fence_system() { +#ifdef __HIP_DEVICE_COMPILE__ __threadfence_system(); +#else + std::atomic_thread_fence(std::memory_order_seq_cst); +#endif } __host__ __device__ void round_robin(const int id, const int num_dev, const int num_iter, volatile int* data, volatile int* flag) { From 4bcb3372b52f150cfead9ca9c82be54a8b8bcbc8 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 11 Aug 2017 14:49:17 +0530 Subject: [PATCH 3/6] [dtests] Fix hip_threadfence_system test's BUILD command Change-Id: I037c523369e4bddd7a386ffe117fe97a178900fd --- tests/src/deviceLib/hip_threadfence_system.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/src/deviceLib/hip_threadfence_system.cpp b/tests/src/deviceLib/hip_threadfence_system.cpp index be59d56ab1..31f6db1e17 100644 --- a/tests/src/deviceLib/hip_threadfence_system.cpp +++ b/tests/src/deviceLib/hip_threadfence_system.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp -std=c++11 + * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11 * RUN: %t * HIT_END */ From ec3db6791e01b76cae504cc2d99f7fb6c05b860e Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 11 Aug 2017 15:57:20 +0530 Subject: [PATCH 4/6] [ci] Continue building remaining dtests on failure Change-Id: Iab7fd53517e0b66dd5f71efb9e65bd2e9c91355b --- Jenkinsfile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Jenkinsfile b/Jenkinsfile index 526a825c65..8aba810dc0 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -118,7 +118,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf set -x cd ${build_dir_rel} make install -j\$(nproc) - make build_tests -j\$(nproc) + make build_tests -i -j\$(nproc) make test """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file From e40047f2bf698515e4b75e21d81cd0d57c462f81 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 15 Aug 2017 15:51:38 +0530 Subject: [PATCH 5/6] [texture] guard new HCC APIs under workweek Change-Id: I4f60a64fb0b0496ca1eb01ffe6ddda121c25d976 --- src/hip_hcc.cpp | 22 ++++++++++++++++++++++ src/hip_memory.cpp | 45 ++++++++++++++++++++++++++++++++++++++++++++- 2 files changed, 66 insertions(+), 1 deletion(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 07d58599cf..f7082e7e9e 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -1866,8 +1866,13 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, } hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo dstPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); hc::AmPointerInfo srcPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); +#endif bool dstTracked = getTailoredPtrInfo(&dstPtrInfo, dst, sizeBytes); bool srcTracked = getTailoredPtrInfo(&srcPtrInfo, src, sizeBytes); @@ -1902,7 +1907,11 @@ 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) { +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo ptrInfo(NULL, ptr, ptr, sizeBytes, acc, true, false); +#else + hc::AmPointerInfo ptrInfo(NULL, ptr, sizeBytes, acc, true, false); +#endif hc::am_memtracker_add(ptr, ptrInfo); } @@ -1926,7 +1935,11 @@ 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) { +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo srcPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); +#endif bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS); if(srcTracked) { addSymbolPtrToTracker(acc, dst, sizeBytes); @@ -1938,7 +1951,11 @@ void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* } } if(kind == hipMemcpyDeviceToHost) { +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo dstPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); +#endif bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS); if(dstTracked) { addSymbolPtrToTracker(acc, src, sizeBytes); @@ -1977,8 +1994,13 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes } else { hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo dstPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); hc::AmPointerInfo srcPtrInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0); + hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0); +#endif 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 bf51748070..6db3d82376 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -108,9 +108,12 @@ void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool auto device = ctx->getWriteableDevice(); +#if (__hcc_workweek__ >= 17332) if (alignment != 0) { ptr = hc::am_aligned_alloc(sizeBytes, device->_acc, amFlags, alignment); - } else { + } else +#endif + { ptr = hc::am_alloc(sizeBytes, device->_acc, amFlags); } tprintf(DB_MEM, " alloc %s ptr:%p-%p size:%zu on dev:%d\n", @@ -149,7 +152,11 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void hipError_t e = hipSuccess; hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if (status == AM_SUCCESS) { @@ -199,7 +206,11 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi e = hipErrorInvalidValue; } else { hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif 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)) ; @@ -581,7 +592,11 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) hipError_t hip_status = hipSuccess; hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); if(status == AM_SUCCESS){ *flagsPtr = amPointerInfo._appAllocationFlags; @@ -612,7 +627,11 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) } hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t am_status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); if(am_status == AM_SUCCESS){ @@ -1412,7 +1431,11 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) if(ptr != nullptr && size != nullptr){ hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if(status == AM_SUCCESS){ *size = amPointerInfo._sizeBytes; @@ -1437,7 +1460,11 @@ hipError_t hipFree(void* ptr) if (ptr) { hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if(status == AM_SUCCESS){ if(amPointerInfo._hostPointer == NULL){ @@ -1465,7 +1492,11 @@ hipError_t hipHostFree(void* ptr) hipError_t hipStatus = hipErrorInvalidValue; if (ptr) { hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); if(status == AM_SUCCESS){ if(amPointerInfo._hostPointer == ptr){ @@ -1499,7 +1530,11 @@ hipError_t hipFreeArray(hipArray* array) if(array->data) { hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data); if(status == AM_SUCCESS){ if(amPointerInfo._hostPointer == NULL){ @@ -1517,7 +1552,11 @@ hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDevi HIP_INIT_API ( pbase , psize , dptr ); hipError_t hipStatus = hipSuccess; hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo( NULL , NULL , NULL, 0 , acc , 0 , 0 ); +#else + hc::AmPointerInfo amPointerInfo( NULL , NULL, 0 , acc , 0 , 0 ); +#endif am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , dptr ); if (status == AM_SUCCESS) { *pbase = amPointerInfo._devicePointer; @@ -1540,7 +1579,11 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){ if((handle == NULL) || (devPtr == NULL)) { hipStatus = hipErrorInvalidResourceHandle; } else { +#if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo( NULL , NULL , NULL, 0 , acc , 0 , 0 ); +#else + hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 ); +#endif am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , devPtr ); if (status == AM_SUCCESS) { psize = (size_t)amPointerInfo._sizeBytes; From c11b8737b92ab9cb11b1ecce09e4574fd5ffa97b Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 16 Aug 2017 10:57:07 +0530 Subject: [PATCH 6/6] [ci] Archive packages based on hcc-ctu instead of hcc-1.6 Change-Id: Ie1ea6411ed19aae87b979fa001db3e8927fc11f6 --- Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 8aba810dc0..af80a2c18b 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -140,8 +140,8 @@ def docker_build_inside_image( def build_image, String inside_args, String platf // No matter the base platform, all packages have the same name // Only upload 1 set of packages, so we don't have a race condition uploading packages - // I arbitrarily pick hcc-1.6 as the most stable - if( platform.toLowerCase( ).startsWith( 'hcc-1.6' ) ) + // Using hcc-ctu since that is what most people want + if( platform.toLowerCase( ).startsWith( 'hcc-ctu' ) ) { archiveArtifacts artifacts: "${build_dir_rel}/*.deb", fingerprint: true archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true