diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index d2f81c8b7b..9cb59f14ea 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -66,24 +66,24 @@ typedef enum hipMemcpyKind { #define HIP_LIBRARY_PATCH_LEVEL PATCH_LEVEL // hipTextureAddressMode -#define hipTextureAddressMode cudaTextureAddressMode +typedef enum cudaTextureAddressMode hipTextureAddressMode; #define hipAddressModeWrap cudaAddressModeWrap #define hipAddressModeClamp cudaAddressModeClamp #define hipAddressModeMirror cudaAddressModeMirror #define hipAddressModeBorder cudaAddressModeBorder // hipTextureFilterMode -#define hipTextureFilterMode cudaTextureFilterMode +typedef enum cudaTextureFilterMode hipTextureFilterMode; #define hipFilterModePoint cudaFilterModePoint #define hipFilterModeLinear cudaFilterModeLinear // hipTextureReadMode -#define hipTextureReadMode cudaTextureReadMode +typedef enum cudaTextureReadMode hipTextureReadMode; #define hipReadModeElementType cudaReadModeElementType #define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat // hipChannelFormatKind -#define hipChannelFormatKind cudaChannelFormatKind +typedef enum cudaChannelFormatKind hipChannelFormatKind; #define hipChannelFormatKindSigned cudaChannelFormatKindSigned #define hipChannelFormatKindUnsigned cudaChannelFormatKindUnsigned #define hipChannelFormatKindFloat cudaChannelFormatKindFloat @@ -167,7 +167,7 @@ typedef cudaIpcMemHandle_t hipIpcMemHandle_t; typedef enum cudaLimit hipLimit_t; typedef enum cudaFuncCache hipFuncCache_t; typedef CUcontext hipCtx_t; -typedef cudaSharedMemConfig hipSharedMemConfig; +typedef enum cudaSharedMemConfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; typedef CUjit_option hipJitOption; typedef CUdevice hipDevice_t; @@ -177,8 +177,8 @@ typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_t; typedef struct cudaArray* hipArray_const_t; -typedef cudaFuncAttributes hipFuncAttributes; -typedef CUfunction_attribute hipFunction_attribute; +typedef struct cudaFuncAttributes hipFuncAttributes; +#define hipFunction_attribute CUfunction_attribute #define hip_Memcpy2D CUDA_MEMCPY2D #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault @@ -196,8 +196,8 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipTextureType3D cudaTextureType3D #define hipDeviceMapHost cudaDeviceMapHost -#define hipExtent cudaExtent -#define hipPitchedPtr cudaPitchedPtr +typedef struct cudaExtent hipExtent; +typedef struct cudaPitchedPtr hipPitchedPtr; #define make_hipExtent make_cudaExtent #define make_hipPos make_cudaPos #define make_hipPitchedPtr make_cudaPitchedPtr @@ -205,10 +205,10 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking -#define hipChannelFormatDesc cudaChannelFormatDesc -#define hipResourceDesc cudaResourceDesc -#define hipTextureDesc cudaTextureDesc -#define hipResourceViewDesc cudaResourceViewDesc +typedef struct cudaChannelFormatDesc hipChannelFormatDesc; +typedef struct cudaResourceDesc hipResourceDesc; +typedef struct cudaTextureDesc hipTextureDesc; +typedef struct cudaResourceViewDesc hipResourceViewDesc; // adding code for hipmemSharedConfig #define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte @@ -394,7 +394,7 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki } } -inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode( +inline static enum cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode( hipTextureAddressMode kind) { switch (kind) { case hipAddressModeWrap: @@ -410,7 +410,7 @@ inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMo } } -inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode( +inline static enum cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode( hipTextureFilterMode kind) { switch (kind) { case hipFilterModePoint: @@ -422,7 +422,7 @@ inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode( } } -inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) { +inline static enum cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) { switch (kind) { case hipReadModeElementType: return cudaReadModeElementType; @@ -433,7 +433,7 @@ inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTex } } -inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind( +inline static enum cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind( hipChannelFormatKind kind) { switch (kind) { case hipChannelFormatKindSigned: @@ -510,14 +510,14 @@ inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags)); } -inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, +inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } -inline static hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, - struct hipExtent extent, unsigned int flags) { +inline static hipError_t hipMalloc3DArray(hipArray** array, const hipChannelFormatDesc* desc, + hipExtent extent, unsigned int flags) { return hipCUDAErrorTohipError(cudaMalloc3DArray(array, desc, extent, flags)); } @@ -1048,7 +1048,6 @@ inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attribut return err; } - inline static hipError_t hipMemGetInfo(size_t* free, size_t* total) { return hipCUDAErrorTohipError(cudaMemGetInfo(free, total)); } @@ -1073,7 +1072,6 @@ inline static hipError_t hipEventDestroy(hipEvent_t event) { return hipCUDAErrorTohipError(cudaEventDestroy(event)); } - inline static hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) { return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags)); } @@ -1360,79 +1358,15 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gri kernelParams, extra)); } - inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) { return hipCUDAErrorTohipError(cudaFuncSetCacheConfig(func, cacheConfig)); } -#ifdef __cplusplus -} -#endif - -#ifdef __CUDACC__ - -template -inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func, - size_t dynamicSMemSize = 0, - int blockSizeLimit = 0) { - cudaError_t cerror; - cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit); - return hipCUDAErrorTohipError(cerror); -} - -template -inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, - const void* devPtr, size_t size = UINT_MAX) { - return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); -} - -template -inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, - const void* devPtr, const struct hipChannelFormatDesc& desc, - size_t size = UINT_MAX) { +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)); } -template -inline static hipError_t hipUnbindTexture(struct texture* tex) { - return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); -} - -template -inline static hipError_t hipUnbindTexture(struct texture &tex) { - return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); -} - -inline static hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, - const hipChannelFormatDesc* desc, size_t size = UINT_MAX){ - return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); -} - -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, - 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(); -} - inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { return cudaCreateChannelDesc(x, y, z, w, hipChannelFormatKindToCudaChannelFormatKind(f)); @@ -1464,7 +1398,7 @@ inline static hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDe return hipCUDAErrorTohipError(cudaGetTextureObjectResourceDesc( pResDesc, textureObject)); } -inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref) +inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const struct textureReference* texref) { return hipCUDAErrorTohipError(cudaGetTextureAlignmentOffset(offset,texref)); } @@ -1473,6 +1407,70 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_ { return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array)); } + + +#ifdef __cplusplus +} +#endif + +#ifdef __CUDACC__ + +template +inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func, + size_t dynamicSMemSize = 0, + int blockSizeLimit = 0) { + cudaError_t cerror; + cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit); + return hipCUDAErrorTohipError(cerror); +} + +template +inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, + const void* devPtr, size_t size = UINT_MAX) { + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); +} + +template +inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, + const void* devPtr, const hipChannelFormatDesc& desc, + size_t size = UINT_MAX) { + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); +} + +template +inline static hipError_t hipUnbindTexture(struct texture* tex) { + return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); +} + +template +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) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); +} + +template +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) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); +} + +template +inline static hipChannelFormatDesc hipCreateChannelDesc() { + return cudaCreateChannelDesc(); +} #endif //__CUDACC__ #endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H diff --git a/projects/hip/tests/src/gcc/hipMalloc.c b/projects/hip/tests/src/gcc/hipMalloc.c index e3c7a66919..19e90d5222 100644 --- a/projects/hip/tests/src/gcc/hipMalloc.c +++ b/projects/hip/tests/src/gcc/hipMalloc.c @@ -18,8 +18,10 @@ * */ /* HIT_START - * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t EXCLUDE_HIP_PLATFORM nvcc - * TEST: %t EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVCC__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM hcc + * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc + * TEST: hipMalloc_nv EXCLUDE_HIP_PLATFORM hcc + * TEST: hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc * HIT_END */ diff --git a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp index b6c016ecc0..4580220d1d 100644 --- a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp +++ b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp @@ -53,7 +53,7 @@ int runTest(int argc, char** argv) { hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - struct hipResourceDesc resDesc; + hipResourceDesc resDesc; memset(&resDesc, 0, sizeof(resDesc)); resDesc.resType = hipResourceTypeArray; resDesc.res.array.array = hipArray; @@ -62,7 +62,7 @@ int runTest(int argc, char** argv) { hipCreateSurfaceObject(&surfaceObject, &resDesc); hipMallocArray(&hipOutArray, &channelDesc, width, height); - struct hipResourceDesc resOutDesc; + hipResourceDesc resOutDesc; memset(&resOutDesc, 0, sizeof(resOutDesc)); resOutDesc.resType = hipResourceTypeArray; resOutDesc.res.array.array = hipOutArray; diff --git a/projects/hip/tests/src/texture/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp index 148c966e80..649c748fd2 100644 --- a/projects/hip/tests/src/texture/hipTextureObj2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureObj2D.cpp @@ -53,13 +53,13 @@ int runTest(int argc, char** argv) { hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - struct hipResourceDesc resDesc; + hipResourceDesc resDesc; memset(&resDesc, 0, sizeof(resDesc)); resDesc.resType = hipResourceTypeArray; resDesc.res.array.array = hipArray; // Specify texture object parameters - struct hipTextureDesc texDesc; + hipTextureDesc texDesc; memset(&texDesc, 0, sizeof(texDesc)); texDesc.addressMode[0] = hipAddressModeWrap; texDesc.addressMode[1] = hipAddressModeWrap;