From 94bfd485ef92305a8569ab421bfbfc2abdfa4d70 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Tue, 28 Jul 2020 08:16:58 -0400 Subject: [PATCH] Cuda 11 changes: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaPointerAttributes.html\#structcudaPointerAttributes Change-Id: I8a5389b88df286043c365a734983a4c5de352102 --- .../include/hip/hcc_detail/hip_runtime_api.h | 51 +++++++++----- .../include/hip/nvcc_detail/hip_runtime_api.h | 66 ++++++++++++------- 2 files changed, 80 insertions(+), 37 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 3c8a775289..853c3ce6a8 100755 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -2173,6 +2173,7 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, * hipMemcpyAsync */ +DEPRECATED(DEPRECATED_MSG) hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind); @@ -2191,6 +2192,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, * hipMemcpyAsync */ +DEPRECATED(DEPRECATED_MSG) hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind); @@ -3531,6 +3533,7 @@ hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 void** args, size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent, hipEvent_t stopEvent, int flags); +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture( size_t* offset, const textureReference* tex, @@ -3538,6 +3541,7 @@ hipError_t hipBindTexture( const hipChannelFormatDesc* desc, size_t size __dparm(UINT_MAX)); +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture2D( size_t* offset, const textureReference* tex, @@ -3547,6 +3551,7 @@ hipError_t hipBindTexture2D( size_t height, size_t pitch); +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTextureToArray( const textureReference* tex, hipArray_const_t array, @@ -3557,6 +3562,7 @@ hipError_t hipBindTextureToMipmappedArray( hipMipmappedArray_const_t mipmappedArray, const hipChannelFormatDesc* desc); +DEPRECATED(DEPRECATED_MSG) hipError_t hipGetTextureAlignmentOffset( size_t* offset, const textureReference* texref); @@ -3565,6 +3571,7 @@ hipError_t hipGetTextureReference( const textureReference** texref, const void* symbol); +DEPRECATED(DEPRECATED_MSG) hipError_t hipUnbindTexture(const textureReference* tex); hipError_t hipCreateTextureObject( @@ -3834,6 +3841,7 @@ inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( class TlsData; #if !__HIP_ROCclr__ +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, const hipChannelFormatDesc* desc, size_t size = UINT_MAX); #endif @@ -3861,6 +3869,7 @@ hipError_t ihipBindTextureImpl(TlsData *tls, int dim, enum hipTextureReadMode re **/ #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture(size_t* offset, struct texture& tex, const void* devPtr, const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) { return ihipBindTextureImpl(nullptr, dim, readMode, offset, devPtr, &desc, size, &tex); @@ -3883,6 +3892,7 @@ hipError_t hipBindTexture(size_t* offset, struct texture& tex, **/ #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture(size_t* offset, struct texture& tex, const void* devPtr, size_t size = UINT_MAX) { return ihipBindTextureImpl(nullptr, dim, readMode, offset, devPtr, &(tex.channelDesc), size, &tex); @@ -3891,6 +3901,7 @@ hipError_t hipBindTexture(size_t* offset, struct texture& tex, // C API #if !__HIP_ROCclr__ +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* devPtr, const hipChannelFormatDesc* desc, size_t width, size_t height, size_t pitch); @@ -3904,6 +3915,7 @@ hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture2D(size_t* offset, struct texture& tex, const void* devPtr, size_t width, size_t height, size_t pitch) { return ihipBindTexture2DImpl(dim, readMode, offset, devPtr, &(tex.channelDesc), width, height, @@ -3913,6 +3925,7 @@ hipError_t hipBindTexture2D(size_t* offset, struct texture& te #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTexture2D(size_t* offset, struct texture& tex, const void* devPtr, const struct hipChannelFormatDesc& desc, size_t width, size_t height, size_t pitch) { @@ -3922,6 +3935,7 @@ hipError_t hipBindTexture2D(size_t* offset, struct texture& te // C API #if !__HIP_ROCclr__ +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, const hipChannelFormatDesc* desc); #endif @@ -3935,6 +3949,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls, int dim, enum hipTextureRead #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTextureToArray(struct texture& tex, hipArray_const_t array) { return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, tex.channelDesc, &tex); } @@ -3942,6 +3957,7 @@ hipError_t hipBindTextureToArray(struct texture& tex, hipArray #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipBindTextureToArray(struct texture& tex, hipArray_const_t array, const struct hipChannelFormatDesc& desc) { return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, desc, &tex); @@ -3950,6 +3966,7 @@ hipError_t hipBindTextureToArray(struct texture& tex, hipArray #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) inline static hipError_t hipBindTextureToArray(struct texture *tex, hipArray_const_t array, const struct hipChannelFormatDesc* desc) { @@ -4019,6 +4036,7 @@ inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchPara * @return #hipSuccess **/ #if !__HIP_ROCclr__ +DEPRECATED(DEPRECATED_MSG) hipError_t hipUnbindTexture(const textureReference* tex); #endif @@ -4028,6 +4046,7 @@ extern hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) #if !__HIP_ROCclr__ template +DEPRECATED(DEPRECATED_MSG) hipError_t hipUnbindTexture(struct texture& tex) { return ihipUnbindTextureImpl(tex.textureObject); } @@ -4035,7 +4054,10 @@ hipError_t hipUnbindTexture(struct texture& tex) { #if !__HIP_ROCclr__ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array); + +DEPRECATED(DEPRECATED_MSG) hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref); + hipError_t hipGetTextureReference(const textureReference** texref, const void* symbol); hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, @@ -4078,28 +4100,23 @@ hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipReso hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject); #if __HIP_ROCclr__ -template -static inline hipError_t hipBindTexture( - size_t *offset, - const struct texture &tex, - const void *devPtr, - size_t size = UINT_MAX) -{ +template +DEPRECATED(DEPRECATED_MSG) +static inline hipError_t hipBindTexture(size_t* offset, const struct texture& tex, + const void* devPtr, size_t size = UINT_MAX) { return hipBindTexture(offset, &tex, devPtr, &tex.channelDesc, size); } -template -static inline hipError_t hipBindTexture( - size_t *offset, - const struct texture &tex, - const void *devPtr, - const struct hipChannelFormatDesc &desc, - size_t size = UINT_MAX) -{ +template +DEPRECATED(DEPRECATED_MSG) +static inline hipError_t + hipBindTexture(size_t* offset, const struct texture& tex, const void* devPtr, + const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) { return hipBindTexture(offset, &tex, devPtr, &desc, size); } template +DEPRECATED(DEPRECATED_MSG) static inline hipError_t hipBindTexture2D( size_t *offset, const struct texture &tex, @@ -4112,6 +4129,7 @@ static inline hipError_t hipBindTexture2D( } template +DEPRECATED(DEPRECATED_MSG) static inline hipError_t hipBindTexture2D( size_t *offset, const struct texture &tex, @@ -4125,6 +4143,7 @@ static inline hipError_t hipBindTexture2D( } template +DEPRECATED(DEPRECATED_MSG) static inline hipError_t hipBindTextureToArray( const struct texture &tex, hipArray_const_t array) @@ -4135,6 +4154,7 @@ static inline hipError_t hipBindTextureToArray( } template +DEPRECATED(DEPRECATED_MSG) static inline hipError_t hipBindTextureToArray( const struct texture &tex, hipArray_const_t array, @@ -4168,6 +4188,7 @@ static inline hipError_t hipBindTextureToMipmappedArray( } template +DEPRECATED(DEPRECATED_MSG) static inline hipError_t hipUnbindTexture( const struct texture &tex) { diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 3744c6740c..f9a2992cd1 100755 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -37,6 +37,18 @@ extern "C" { #define __dparm(x) #endif +// Add Deprecated Support for CUDA Mapped HIP APIs +#if defined(__DOXYGEN_ONLY__) || defined(HIP_ENABLE_DEPRECATED) +#define __HIP_DEPRECATED +#elif defined(_MSC_VER) +#define __HIP_DEPRECATED __declspec(deprecated) +#elif defined(__GNUC__) +#define __HIP_DEPRECATED __attribute__((deprecated)) +#else +#define __HIP_DEPRECATED +#endif + + // TODO -move to include/hip_runtime_api.h as a common implementation. /** * Memory copy types @@ -963,14 +975,16 @@ inline static hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_ height, hipMemcpyKindToCudaMemcpyKind(kind))); } -inline static hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, - const void* src, size_t count, hipMemcpyKind kind) { +__HIP_DEPRECATED inline static hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, + size_t hOffset, const void* src, + size_t count, hipMemcpyKind kind) { return hipCUDAErrorTohipError( cudaMemcpyToArray(dst, wOffset, hOffset, src, count, hipMemcpyKindToCudaMemcpyKind(kind))); } -inline static hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, - size_t hOffset, size_t count, hipMemcpyKind kind) { +__HIP_DEPRECATED inline static hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, + size_t wOffset, size_t hOffset, + size_t count, hipMemcpyKind kind) { return hipCUDAErrorTohipError(cudaMemcpyFromArray(dst, srcArray, wOffset, hOffset, count, hipMemcpyKindToCudaMemcpyKind(kind))); } @@ -1353,7 +1367,12 @@ inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attribut struct cudaPointerAttributes cPA; hipError_t err = hipCUDAErrorTohipError(cudaPointerGetAttributes(&cPA, ptr)); if (err == hipSuccess) { - switch (cPA.memoryType) { +#if (CUDART_VERSION >= 11000) + auto memType = cPA.type; +#else + unsigned memType = cPA.memoryType; // No auto because cuda 10.2 doesnt force c++11 +#endif + switch (memType) { case cudaMemoryTypeDevice: attributes->memoryType = hipMemoryTypeDevice; break; @@ -1695,14 +1714,17 @@ inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t return hipCUDAErrorTohipError(cudaFuncSetCacheConfig(func, cacheConfig)); } -inline static hipError_t hipBindTexture(size_t* offset, struct textureReference* tex, const void* devPtr, - const hipChannelFormatDesc* desc, size_t size __dparm(UINT_MAX)){ +__HIP_DEPRECATED inline static hipError_t hipBindTexture(size_t* offset, + struct textureReference* tex, + const void* devPtr, + const hipChannelFormatDesc* desc, + size_t size __dparm(UINT_MAX)) { return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); } -inline static hipError_t hipBindTexture2D(size_t* offset, struct textureReference* tex, const void* devPtr, - const hipChannelFormatDesc* desc, size_t width, size_t height, - size_t pitch) { +__HIP_DEPRECATED inline static hipError_t hipBindTexture2D( + size_t* offset, struct textureReference* tex, const void* devPtr, + const hipChannelFormatDesc* desc, size_t width, size_t height, size_t pitch) { return hipCUDAErrorTohipError(cudaBindTexture2D(offset, tex, devPtr, desc, width, height, pitch)); } @@ -1737,8 +1759,8 @@ inline static hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDe return hipCUDAErrorTohipError(cudaGetTextureObjectResourceDesc( pResDesc, textureObject)); } -inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const struct textureReference* texref) -{ +__HIP_DEPRECATED inline static hipError_t hipGetTextureAlignmentOffset( + size_t* offset, const struct textureReference* texref) { return hipCUDAErrorTohipError(cudaGetTextureAlignmentOffset(offset,texref)); } @@ -1811,32 +1833,32 @@ inline static hipError_t hipBindTexture(size_t* offset, struct texture -inline static hipError_t hipUnbindTexture(struct texture* tex) { +__HIP_DEPRECATED inline static hipError_t hipUnbindTexture(struct texture* tex) { return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); } template -inline static hipError_t hipUnbindTexture(struct texture &tex) { +__HIP_DEPRECATED inline static hipError_t hipUnbindTexture(struct texture& tex) { return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); } template -inline static hipError_t hipBindTextureToArray(struct texture& tex, - hipArray_const_t array, - const hipChannelFormatDesc& desc) { +__HIP_DEPRECATED inline static hipError_t hipBindTextureToArray( + struct texture& tex, hipArray_const_t array, + const hipChannelFormatDesc& desc) { return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); } template -inline static hipError_t hipBindTextureToArray(struct texture *tex, - hipArray_const_t array, - const hipChannelFormatDesc* desc) { +__HIP_DEPRECATED inline static hipError_t hipBindTextureToArray( + struct texture* tex, hipArray_const_t array, + const hipChannelFormatDesc* desc) { return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); } template -inline static hipError_t hipBindTextureToArray(struct texture& tex, - hipArray_const_t array) { +__HIP_DEPRECATED inline static hipError_t hipBindTextureToArray( + struct texture& tex, hipArray_const_t array) { return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); }