From 8def98b74e4e4e3e1bcb418c9d2d40863fdf3d6b Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 4 Mar 2020 14:17:07 +0530 Subject: [PATCH 1/5] Expose support for non-returning atomic FADD Change-Id: If5359488324477315a9bd4f308a75f606c065b39 [ROCm/clr commit: 48f5524fdcca2e15f0475f5f93b5ee987f2c1371] --- .../hipamd/include/hip/hcc_detail/device_library_decls.h | 1 + projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h index ac35823cd2..182565ad61 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h @@ -72,6 +72,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank( extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void); extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void); +extern "C" __device__ void __ockl_global_atomic_add_f32(__attribute__((address_space(1))) float*, float); // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h index 263f639e96..7ccfa6b43e 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h @@ -73,6 +73,14 @@ float atomicAdd(float* address, float val) return __uint_as_float(r); } + +__device__ +inline +void atomicAddNoRet(float* address, float val) +{ + __ockl_global_atomic_add_f32((__attribute__((address_space(1))) float*)address, val); +} + __device__ inline double atomicAdd(double* address, double val) From 988b489efd9b2b836fc80e2c0f89eb3d22dc2664 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe Date: Fri, 6 Mar 2020 11:47:21 +0530 Subject: [PATCH 2/5] separate printf declaration for vdi/clang There are now two implementations of printf in HIP: 1. The implemenation for HCC is controlled by the HC_FEATURE_PRINTF macro, and it works only with the HCC compiler used in combination with the HCC runtime. 2. The implementation for hip-clang requires the VDI runtime, and is always enabled with that combination. [ROCm/clr commit: 2c3dfdda4176f3c04e9d70d61dca252ca130ffef] --- projects/clr/hipamd/bin/hipcc | 1 + .../hipamd/include/hip/hcc_detail/device_functions.h | 9 +++++++++ .../clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 10 +++++++--- 3 files changed, 17 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 06a3f9e385..debc1d46c9 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -133,6 +133,7 @@ if (defined $HIP_RUNTIME and $HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) { } else { $HIP_VDI_HOME = $HIP_PATH; # use HIP_PATH } + $HIPCXXFLAGS .= "-D__HIP_VDI__"; } if (defined $HIP_VDI_HOME) { diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index 7096841da8..e3544a491e 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -33,6 +33,15 @@ THE SOFTWARE. #include #include #include + +#if __HIP_CLANG_ONLY__ +#if __HIP_VDI__ +extern "C" __device__ int printf(const char *fmt, ...); +#else +static inline __device__ void printf(const char* format, All... all) {} +#endif +#endif + /* Integer Intrinsics */ diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index e86611b213..c1ad5b2fe5 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -311,15 +311,19 @@ extern "C" __device__ void* __hip_free(void* ptr); static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); } static inline __device__ void* free(void* ptr) { return __hip_free(ptr); } -#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF) +// Declare printf only for the HCC compiler. hip-clang is handled in +// device_functions.h +#if __HCC_ACCELERATOR__ +#if HC_FEATURE_PRINTF template static inline __device__ void printf(const char* format, All... all) { hc::printf(format, all...); } -#elif defined(__HCC_ACCELERATOR__) || __HIP__ +#else template static inline __device__ void printf(const char* format, All... all) {} -#endif +#endif // HC_FEATURE_PRINTF +#endif // __HCC_ACCELERATOR__ #endif //__HCC_OR_HIP_CLANG__ From 2b17cd51892d46ab4e62b0d62d8d406601662c24 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 10 Mar 2020 18:04:01 +0300 Subject: [PATCH 3/5] [HIP][doc] Update README.md [ROCm/clr commit: 2bb166503ffe18430d5bdd14d1462929bdcd3f7a] --- projects/clr/hipamd/README.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md index 2bffd12162..c2e2a7a456 100644 --- a/projects/clr/hipamd/README.md +++ b/projects/clr/hipamd/README.md @@ -1,6 +1,7 @@ ## What is this repository for? ### -HIP allows developers to convert CUDA code to portable C++. The same source code can be compiled to run on NVIDIA or AMD GPUs. +**HIP is a C++ Runtime API and Kernel Language that allows developers to create portable applications for AMD and NVIDIA GPUs from single source code.** + Key features include: * HIP is very thin and has little or no performance impact over coding directly in CUDA or hcc "HC" mode. From cf2629a95cd00152ec65dc3a42126ce15ab3a12f Mon Sep 17 00:00:00 2001 From: Nick Curtis Date: Tue, 10 Mar 2020 12:04:05 -0500 Subject: [PATCH 4/5] Fix incorrect shfl_xor for Windows copy/paste error, need __shfl_xor w/ lane_mask [ROCm/clr commit: a9af4755cb46aa554690a0a6ef5fd59b53cb968f] --- projects/clr/hipamd/include/hip/hcc_detail/device_functions.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index e3544a491e..6e6756fd9c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -557,7 +557,7 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize) return tmp1; #else static_assert(sizeof(long) == sizeof(int), ""); - return static_cast(__shfl_down(static_cast(var), lane_delta, width)); + return static_cast(__shfl_xor(static_cast(var), lane_mask, width)); #endif } __device__ From 195abb8a1897ba3a1180a73ce7e9ab93cd1c788f Mon Sep 17 00:00:00 2001 From: srinivamd <52507740+srinivamd@users.noreply.github.com> Date: Wed, 11 Mar 2020 02:02:54 -0700 Subject: [PATCH 5/5] return hipSuccess when count is zero (#1900) [ROCm/clr commit: 32960ab8f95ca250f268b138a444d2e0f9bed3fc] --- projects/clr/hipamd/src/hip_memory.cpp | 44 ++++++++++++++++++-------- 1 file changed, 31 insertions(+), 13 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 8159f22a97..4a126532f4 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -478,6 +478,10 @@ void* allocAndSharePtr(const char* msg, size_t sizeBytes, ihipCtx_t* ctx, bool s hipError_t ihipHostMalloc(TlsData *tls, void** ptr, size_t sizeBytes, unsigned int flags) { hipError_t hip_status = hipSuccess; + if (sizeBytes == 0) { + return hipSuccess; + } + if (HIP_SYNC_HOST_ALLOC) { hipDeviceSynchronize(); } @@ -485,10 +489,6 @@ hipError_t ihipHostMalloc(TlsData *tls, void** ptr, size_t sizeBytes, unsigned i auto ctx = ihipGetTlsDefaultCtx(); if ((ctx == nullptr) || (ptr == nullptr)) { hip_status = hipErrorInvalidValue; - } - else if (sizeBytes == 0) { - hip_status = hipSuccess; - // TODO - should size of 0 return err or be siliently ignored? } else { unsigned trueFlags = flags; if (flags == hipHostMallocDefault) { @@ -673,14 +673,15 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) { HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; + if (sizeBytes == 0) { + if (ptr) *ptr = NULL; + return ihipLogStatus(hipSuccess); + } + auto ctx = ihipGetTlsDefaultCtx(); // return NULL pointer when malloc size is 0 if ( nullptr == ctx || nullptr == ptr) { hip_status = hipErrorInvalidValue; - } - else if (sizeBytes == 0) { - *ptr = NULL; - hip_status = hipSuccess; } else { auto device = ctx->getWriteableDevice(); *ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/, @@ -700,14 +701,15 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag HIP_SET_DEVICE(); #if (__hcc_workweek__ >= 19115) + if (sizeBytes == 0) { + if (ptr) *ptr = NULL; + return ihipLogStatus(hipSuccess); + } + hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); - // return NULL pointer when malloc size is 0 - if (sizeBytes == 0) { - *ptr = NULL; - hip_status = hipSuccess; - } else if ((ctx == nullptr) || (ptr == nullptr)) { + if ((ctx == nullptr) || (ptr == nullptr)) { hip_status = hipErrorInvalidValue; } else { unsigned amFlags = 0; @@ -736,6 +738,9 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); HIP_SET_DEVICE(); + if (sizeBytes == 0) { + return ihipLogStatus(hipSuccess); + } hipError_t hip_status = hipSuccess; hip_status = hip_internal::ihipHostMalloc(tls, ptr, sizeBytes, flags); return ihipLogStatus(hip_status); @@ -744,6 +749,9 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) { HIP_INIT_SPECIAL_API(hipMallocManaged, (TRACE_MEM), devPtr, size, flags); HIP_SET_DEVICE(); + if (size == 0) { + return ihipLogStatus(hipSuccess); + } hipError_t hip_status = hipSuccess; if(flags != hipMemAttachGlobal) hip_status = hipErrorInvalidValue; @@ -1224,6 +1232,7 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); + if (count == 0) return ihipLogStatus(hipSuccess); if (dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } @@ -1246,6 +1255,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); + if (count == 0) return ihipLogStatus(hipSuccess); if (src == nullptr || dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } @@ -1269,6 +1279,7 @@ hipError_t hipMemcpyToSymbolAsync(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); + if (count == 0) return ihipLogStatus(hipSuccess); if (dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } @@ -1301,6 +1312,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, src); + if (count == 0) return ihipLogStatus(hipSuccess); if (src == nullptr || dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } @@ -1592,6 +1604,7 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo srcXoffset = p->srcPos.x; srcYoffset = p->srcPos.y; srcZoffset = p->srcPos.z; + if (copyWidth == 0) return hipSuccess; if (p->dstArray != nullptr) { if ((p->dstArray->isDrv == true) ||( p->dstPtr.ptr!= nullptr)){ return hipErrorInvalidValue; @@ -1933,6 +1946,7 @@ hipError_t getLockedPointer(void *hostPtr, size_t dataLen, void **devicePtrPtr) // TODO - review and optimize hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { + if (height == 0 || width == 0) return hipSuccess; if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue; hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); @@ -1989,6 +2003,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, hipError_t ihipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { + if (height == 0 || width == 0) return hipSuccess; if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue; hipError_t e = hipSuccess; int isLockedOrD2D = 0; @@ -2043,6 +2058,7 @@ hipError_t ihip2dOffsetMemcpy(void* dst, size_t dpitch, const void* src, size_t size_t height, size_t srcXOffsetInBytes, size_t srcYOffset, size_t dstXOffsetInBytes, size_t dstYOffset,hipMemcpyKind kind, hipStream_t stream, bool isAsync) { + if (height == 0 || width == 0) return hipSuccess; if((spitch < width + srcXOffsetInBytes) || (srcYOffset >= height)){ return hipErrorInvalidValue; } else if((dpitch < width + dstXOffsetInBytes) || (dstYOffset >= height)){ @@ -2061,6 +2077,7 @@ hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool if (pCopy == nullptr) { return hipErrorInvalidValue; } + if (pCopy->Height == 0 || pCopy->WidthInBytes == 0) return hipSuccess; void* dst; const void* src; size_t spitch = pCopy->srcPitch; size_t dpitch = pCopy->dstPitch; @@ -2140,6 +2157,7 @@ hipError_t hipMemcpy2DFromArray( void* dst, size_t dpitch, hipArray_const_t src, hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream ){ HIP_INIT_SPECIAL_API(hipMemcpy2DFromArrayAsync, (TRACE_MCMD), dst, dpitch, src, wOffset, hOffset, width, height, kind, stream); size_t byteSize; + if (height == 0 || width == 0) return ihipLogStatus(hipSuccess); if(src) { switch (src->desc.f) { case hipChannelFormatKindSigned: