From c1cb518025bbc30f9d97df87e543823c02aa2c43 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sat, 23 Sep 2017 05:23:24 +0530 Subject: [PATCH 1/4] Fix texture support on HIP/NVCC path [ROCm/clr commit: e6725e2b7a65a16ffb97bf0cc44c9a3acecd778c] --- .../include/hip/nvcc_detail/hip_runtime_api.h | 174 ++++++++++++++++-- 1 file changed, 162 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 1103bb2274..6459839180 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -50,15 +50,64 @@ hipMemcpyHostToHost ,hipMemcpyDeviceToDevice ,hipMemcpyDefault } hipMemcpyKind ; +#if 0 +typedef enum hipTextureAddressMode { + hipAddressModeWrap = 0, + hipAddressModeClamp = 1, + hipAddressModeMirror = 2, + hipAddressModeBorder = 3 +}hipTextureAddressMode; +#endif +#define hipTextureAddressMode cudaTextureAddressMode +#define hipAddressModeWrap cudaAddressModeWrap +#define hipAddressModeClamp cudaAddressModeClamp +#define hipAddressModeMirror cudaAddressModeMirror +#define hipAddressModeBorder cudaAddressModeBorder +#if 0 +typedef enum hipTextureFilterMode { + hipFilterModePoint = 0, + hipFilterModeLinear = 1 +}hipTextureFilterMode; + +#endif +#if 1 +#define hipTextureFilterMode cudaTextureFilterMode +#define hipFilterModePoint cudaFilterModePoint +#define hipFilterModeLinear cudaFilterModeLinear +#endif +#if 0 +typedef enum hipTextureReadMode { + hipReadModeElementType = 0, + hipReadModeNormalizedFloat = 1 +}hipTextureReadMode; +#endif +#define hipTextureReadMode cudaTextureReadMode +#define hipReadModeElementType cudaReadModeElementType +#define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat + +typedef enum hipChannelFormatKind { + hipChannelFormatKindSigned = 0, + hipChannelFormatKindUnsigned = 1, + hipChannelFormatKindFloat = 2, + hipChannelFormatKindNone = 3 +}hipChannelFormatKind; +#if 0 +typedef enum hipResourceType { + hipResourceTypeArray = 0, + hipResourceTypeMipmappedArray = 1, + hipResourceTypeLinear = 2, + hipResourceTypePitch2D = 3 +}hipResourceType; +#endif +#define hipResourceType cudaResourceType +#define hipResourceTypeArray cudaResourceTypeArray +#define hipResourceTypeMipmappedArray cudaResourceTypeMipmappedArray +#define hipResourceTypeLinear cudaResourceTypeLinear +#define hipResourceTypePitch2D cudaResourceTypePitch2D +// // hipErrorNoDevice. -/*typedef enum hipTextureFilterMode -{ - hipFilterModePoint = cudaFilterModePoint, ///< Point filter mode. -//! @warning cudaFilterModeLinear is not supported. -} hipTextureFilterMode;*/ -#define hipFilterModePoint cudaFilterModePoint //! Flags that can be used with hipEventCreateWithFlags: #define hipEventDefault cudaEventDefault @@ -119,10 +168,16 @@ typedef CUdevice hipDevice_t; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -typedef enum cudaChannelFormatKind hipChannelFormatKind; -typedef struct cudaChannelFormatDesc hipChannelFormatDesc; -typedef enum cudaTextureReadMode hipTextureReadMode; +//typedef enum cudaChannelFormatKind hipChannelFormatKind; +//typedef struct cudaChannelFormatDesc hipChannelFormatDesc; typedef struct cudaArray hipArray; +typedef struct cudaArray* hipArray_const_t; +//typedef cudaArray_const_t hipArray; +#define hipArrayDefault cudaArrayDefault + +typedef cudaTextureObject_t hipTextureObject_t; +#define hipTextureType2D cudaTextureType2D; +#define hipDeviceMapHost cudaDeviceMapHost // Flags that can be used with hipStreamCreateWithFlags #define hipStreamDefault cudaStreamDefault @@ -130,7 +185,9 @@ typedef struct cudaArray hipArray; //typedef cudaChannelFormatDesc hipChannelFormatDesc; #define hipChannelFormatDesc cudaChannelFormatDesc - +#define hipResourceDesc cudaResourceDesc +#define hipTextureDesc cudaTextureDesc +#define hipResourceViewDesc cudaResourceViewDesc //adding code for hipmemSharedConfig #define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte @@ -214,6 +271,58 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki } } +inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(hipTextureAddressMode kind) { + switch(kind) { + case hipAddressModeWrap: + return cudaAddressModeWrap; + case hipAddressModeClamp: + return cudaAddressModeClamp; + case hipAddressModeMirror: + return cudaAddressModeMirror; + case hipAddressModeBorder: + return cudaAddressModeBorder; + default: + return cudaAddressModeWrap; +} +} + +inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(hipTextureFilterMode kind) { + switch(kind) { + case hipFilterModePoint: + return cudaFilterModePoint; + case hipFilterModeLinear: + return cudaFilterModeLinear; + default: + return cudaFilterModePoint; +} +} + +inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) { + switch(kind) { + case hipReadModeElementType: + return cudaReadModeElementType; + case hipReadModeNormalizedFloat: + return cudaReadModeNormalizedFloat; + default: + return cudaReadModeElementType; +} +} + +inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(hipChannelFormatKind kind) { + switch(kind) { + case hipChannelFormatKindSigned : + return cudaChannelFormatKindSigned ; + case hipChannelFormatKindUnsigned : + return cudaChannelFormatKindUnsigned ; + case hipChannelFormatKindFloat : + return cudaChannelFormatKindFloat ; + case hipChannelFormatKindNone : + return cudaChannelFormatKindNone ; + default: + return cudaChannelFormatKindNone ; +} +} + /** * Stream CallBack struct */ @@ -262,9 +371,17 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { - return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); +#if __cplusplus +inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, + size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) { + return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } +#else +inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, + size_t width, size_t height, unsigned int flags) { + return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); +} +#endif inline static hipError_t hipFreeArray(hipArray* array) { return hipCUDAErrorTohipError(cudaFreeArray(array)); @@ -1007,11 +1124,44 @@ 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 struct hipChannelFormatDesc& desc) +{ + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array,desc)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture &tex, + hipArray_const_t array) +{ + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); +} + template inline static hipChannelFormatDesc hipCreateChannelDesc() { return cudaCreateChannelDesc(); } + +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) +{ + return cudaCreateChannelDesc(x , y , z , w, hipChannelFormatKindToCudaChannelFormatKind(f)); +} + +inline static hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, + const hipResourceDesc* pResDesc, + const hipTextureDesc* pTexDesc, + const hipResourceViewDesc* pResViewDesc) +{ + return hipCUDAErrorTohipError(cudaCreateTextureObject(pTexObject, pResDesc, pTexDesc, pResViewDesc)); +} + +inline static hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject) +{ + return hipCUDAErrorTohipError(cudaDestroyTextureObject(textureObject)); +} #endif //__CUDACC__ #endif //HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H From 59fce18f03cb5ada2a90b39620a2c8540b129c6f Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 25 Sep 2017 17:32:56 +0530 Subject: [PATCH 2/4] Fixed build issue [ROCm/clr commit: d92cfa44d24396da4b7c5243a10a0118c172913e] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 6459839180..d8d0bc2c62 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -1145,7 +1145,7 @@ inline static hipChannelFormatDesc hipCreateChannelDesc() return cudaCreateChannelDesc(); } -hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) +inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { return cudaCreateChannelDesc(x , y , z , w, hipChannelFormatKindToCudaChannelFormatKind(f)); } From f32bb0e089a15ba79301491a6f26b8626021247d Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 2 Oct 2017 10:57:25 +0530 Subject: [PATCH 3/4] Fixed review comments [ROCm/clr commit: 79a943e6544ca5c4415aa61d6bca307385a3a100] --- .../include/hip/nvcc_detail/hip_runtime_api.h | 46 ++++--------------- 1 file changed, 9 insertions(+), 37 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index d8d0bc2c62..245327a041 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -50,38 +50,20 @@ hipMemcpyHostToHost ,hipMemcpyDeviceToDevice ,hipMemcpyDefault } hipMemcpyKind ; -#if 0 -typedef enum hipTextureAddressMode { - hipAddressModeWrap = 0, - hipAddressModeClamp = 1, - hipAddressModeMirror = 2, - hipAddressModeBorder = 3 -}hipTextureAddressMode; -#endif + +//hipTextureAddressMode #define hipTextureAddressMode cudaTextureAddressMode #define hipAddressModeWrap cudaAddressModeWrap #define hipAddressModeClamp cudaAddressModeClamp #define hipAddressModeMirror cudaAddressModeMirror #define hipAddressModeBorder cudaAddressModeBorder -#if 0 -typedef enum hipTextureFilterMode { - hipFilterModePoint = 0, - hipFilterModeLinear = 1 -}hipTextureFilterMode; - -#endif -#if 1 +//hipTextureFilterMode #define hipTextureFilterMode cudaTextureFilterMode #define hipFilterModePoint cudaFilterModePoint #define hipFilterModeLinear cudaFilterModeLinear -#endif -#if 0 -typedef enum hipTextureReadMode { - hipReadModeElementType = 0, - hipReadModeNormalizedFloat = 1 -}hipTextureReadMode; -#endif + +//hipTextureReadMode #define hipTextureReadMode cudaTextureReadMode #define hipReadModeElementType cudaReadModeElementType #define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat @@ -92,14 +74,8 @@ typedef enum hipChannelFormatKind { hipChannelFormatKindFloat = 2, hipChannelFormatKindNone = 3 }hipChannelFormatKind; -#if 0 -typedef enum hipResourceType { - hipResourceTypeArray = 0, - hipResourceTypeMipmappedArray = 1, - hipResourceTypeLinear = 2, - hipResourceTypePitch2D = 3 -}hipResourceType; -#endif + +//hipResourceType #define hipResourceType cudaResourceType #define hipResourceTypeArray cudaResourceTypeArray #define hipResourceTypeMipmappedArray cudaResourceTypeMipmappedArray @@ -168,11 +144,8 @@ typedef CUdevice hipDevice_t; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -//typedef enum cudaChannelFormatKind hipChannelFormatKind; -//typedef struct cudaChannelFormatDesc hipChannelFormatDesc; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; -//typedef cudaArray_const_t hipArray; #define hipArrayDefault cudaArrayDefault typedef cudaTextureObject_t hipTextureObject_t; @@ -183,7 +156,6 @@ typedef cudaTextureObject_t hipTextureObject_t; #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking -//typedef cudaChannelFormatDesc hipChannelFormatDesc; #define hipChannelFormatDesc cudaChannelFormatDesc #define hipResourceDesc cudaResourceDesc #define hipTextureDesc cudaTextureDesc @@ -371,14 +343,14 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -#if __cplusplus +#if 0 //__cplusplus inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } #else inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, - size_t width, size_t height, unsigned int flags) { + size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } #endif From bc675b15b8256a34c5ea1ac787ab71fdf256177c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 2 Oct 2017 17:29:08 +0530 Subject: [PATCH 4/4] Fixed review comment [ROCm/clr commit: 71a5a0cb4286a3b8b17ebb0fa3206d6dda4591e1] --- .../clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 7 ------- 1 file changed, 7 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 245327a041..870dcc5dae 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -343,17 +343,10 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -#if 0 //__cplusplus -inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, - size_t width, size_t height = 0, unsigned int flags = hipArrayDefault) { - return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); -} -#else inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } -#endif inline static hipError_t hipFreeArray(hipArray* array) { return hipCUDAErrorTohipError(cudaFreeArray(array));