From dd373854f38f8b59d25947f7dd0b3f514d699203 Mon Sep 17 00:00:00 2001 From: sdashmiz Date: Tue, 25 Jul 2023 14:29:18 -0400 Subject: [PATCH] SWDEV-332969 - [ABI Break]Substitute hipArray* with hipArray_t - hipArray will be an internal struct from rocm6.0 Signed-off-by: sdashmiz Change-Id: Icf97fe96b87be8532098cd7f9ceaad099f99c9b9 [ROCm/hipother commit: 452c1060f3efffa77b902cca4710e7418fc003c6] --- .../nvidia_detail/nvidia_hip_runtime_api.h | 292 +++++++++++++++--- 1 file changed, 257 insertions(+), 35 deletions(-) diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 89ae35428c..ac387be5d5 100644 --- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -254,7 +254,6 @@ inline static CUresourcetype hipResourcetype_enumToCUresourcetype( #define hipStreamPerThread ((cudaStream_t)2) #define hipTexRef CUtexref -#define hiparray CUarray typedef CUmipmappedArray hipmipmappedArray; typedef cudaMipmappedArray_t hipMipmappedArray_t; @@ -438,7 +437,6 @@ typedef enum cudaDeviceP2PAttr hipDeviceP2PAttr; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_t; typedef struct cudaArray* hipArray_const_t; typedef struct cudaFuncAttributes hipFuncAttributes; @@ -446,8 +444,209 @@ typedef struct cudaLaunchParams hipLaunchParams; typedef CUDA_LAUNCH_PARAMS hipFunctionLaunchParams; #define hipFunction_attribute CUfunction_attribute #define hipPointer_attribute CUpointer_attribute -#define hip_Memcpy2D CUDA_MEMCPY2D -#define HIP_MEMCPY3D CUDA_MEMCPY3D + +typedef struct HIP_RESOURCE_DESC_st +{ + hipResourcetype resType; /**< Resource type */ + union { + struct { + hipArray_t hArray; /**< HIP array */ + } array; + struct { + hipMipmappedArray_t hMipmappedArray; /**< HIP mipmapped array */ + } mipmap; + struct { + hipDeviceptr_t devPtr; /**< Device pointer */ + hipArray_Format format; /**< Array format */ + unsigned int numChannels; /**< Channels per array element */ + size_t sizeInBytes; /**< Size in bytes */ + } linear; + struct { + hipDeviceptr_t devPtr; /**< Device pointer */ + hipArray_Format format; /**< Array format */ + unsigned int numChannels; /**< Channels per array element */ + size_t width; /**< Width of the array in elements */ + size_t height; /**< Height of the array in elements */ + size_t pitchInBytes; /**< Pitch between two rows in bytes */ + } pitch2D; + struct { + int reserved[32]; + } reserved; + } res; + unsigned int flags; /**< Flags (must be zero) */ +} HIP_RESOURCE_DESC; + +static inline CUDA_RESOURCE_DESC* hipResourceDesTocudaResourceDes(const HIP_RESOURCE_DESC* p){ + CUDA_RESOURCE_DESC a; + switch (p->resType) { + case HIP_RESOURCE_TYPE_ARRAY: + a.resType = CU_RESOURCE_TYPE_ARRAY; + case HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY: + a.resType = CU_RESOURCE_TYPE_MIPMAPPED_ARRAY; + case HIP_RESOURCE_TYPE_LINEAR: + a.resType = CU_RESOURCE_TYPE_LINEAR; + case HIP_RESOURCE_TYPE_PITCH2D: + a.resType = CU_RESOURCE_TYPE_PITCH2D; + default: + a.resType = CU_RESOURCE_TYPE_ARRAY; + } + a.res.array.hArray = (CUarray)p->res.array.hArray; + a.res.mipmap.hMipmappedArray = (CUmipmappedArray)p->res.mipmap.hMipmappedArray; + a.res.linear.devPtr = p->res.linear.devPtr; + a.res.linear.format = p->res.linear.format; + a.res.linear.numChannels = p->res.linear.numChannels; + a.res.linear.sizeInBytes = p->res.linear.sizeInBytes; + a.res.pitch2D.devPtr = p->res.pitch2D.devPtr; + a.res.pitch2D.numChannels = p->res.pitch2D.numChannels; + a.res.pitch2D.format = p->res.pitch2D.format; + a.res.pitch2D.width = p->res.pitch2D.width; + a.res.pitch2D.height = p->res.pitch2D.height; + a.res.pitch2D.pitchInBytes = p->res.pitch2D.pitchInBytes; + a.flags = p->flags; + return &a; +} + +typedef struct hip_Memcpy2D { + size_t srcXInBytes; + size_t srcY; + hipMemoryType srcMemoryType; + const void* srcHost; + hipDeviceptr_t srcDevice; + hipArray_t srcArray; + size_t srcPitch; + size_t dstXInBytes; + size_t dstY; + hipMemoryType dstMemoryType; + void* dstHost; + hipDeviceptr_t dstDevice; + hipArray_t dstArray; + size_t dstPitch; + size_t WidthInBytes; + size_t Height; +} hip_Memcpy2D; + +typedef struct HIP_MEMCPY3D { + unsigned int srcXInBytes; + unsigned int srcY; + unsigned int srcZ; + unsigned int srcLOD; + hipMemoryType srcMemoryType; + const void* srcHost; + hipDeviceptr_t srcDevice; + hipArray_t srcArray; + unsigned int srcPitch; + unsigned int srcHeight; + unsigned int dstXInBytes; + unsigned int dstY; + unsigned int dstZ; + unsigned int dstLOD; + hipMemoryType dstMemoryType; + void* dstHost; + hipDeviceptr_t dstDevice; + hipArray_t dstArray; + unsigned int dstPitch; + unsigned int dstHeight; + unsigned int WidthInBytes; + unsigned int Height; + unsigned int Depth; +} HIP_MEMCPY3D; + +static inline void hipMemcpy3DTocudaMemcpy3D(CUDA_MEMCPY3D &a, const HIP_MEMCPY3D* p){ + a.srcXInBytes = (size_t)p->srcXInBytes; + a.srcY = (size_t)p->srcY; + a.srcZ = (size_t)p->srcZ; + a.srcLOD = (size_t)p->srcLOD; + switch (p->srcMemoryType) { + case hipMemoryTypeHost: + a.srcMemoryType = CU_MEMORYTYPE_HOST; + break; + case hipMemoryTypeDevice: + a.srcMemoryType = CU_MEMORYTYPE_DEVICE; + break; + case hipMemoryTypeArray: + a.srcMemoryType = CU_MEMORYTYPE_ARRAY; + break; + default: + a.srcMemoryType = CU_MEMORYTYPE_UNIFIED; + } + a.srcHost = p->srcHost; + a.srcDevice =(CUdeviceptr)p->srcDevice; + a.srcArray = (CUarray)p->srcArray; + a.reserved0 = nullptr; + a.srcPitch = (size_t)p->srcPitch; + a.srcHeight = (size_t)p->srcHeight; + a.dstXInBytes = (size_t)p->dstXInBytes; + a.dstY = (size_t)p->dstY; + a.dstZ = (size_t)p->dstZ; + a.dstLOD = (size_t)p->dstLOD; + switch (p->dstMemoryType) { + case hipMemoryTypeHost: + a.dstMemoryType = CU_MEMORYTYPE_HOST; + break; + case hipMemoryTypeDevice: + a.dstMemoryType = CU_MEMORYTYPE_DEVICE; + break; + case hipMemoryTypeArray: + a.dstMemoryType = CU_MEMORYTYPE_ARRAY; + break; + default: + a.dstMemoryType = CU_MEMORYTYPE_UNIFIED; + } + a.dstHost = p->dstHost; + a.dstDevice = (CUdeviceptr)p->dstDevice; + a.dstArray = (CUarray)p->dstArray; + a.reserved1 = nullptr; + a.dstPitch = (size_t)p->dstPitch; + a.dstHeight = (size_t)p->dstHeight; + a.WidthInBytes = (size_t)p->WidthInBytes; + a.Height = (size_t)p->Height; + a.Depth = (size_t)p->Depth; +} + +static inline void hipMemcpy2DTocudaMemcpy2D(CUDA_MEMCPY2D &a, const hip_Memcpy2D* p){ + a.srcXInBytes = (size_t)p->srcXInBytes; + a.srcY = (size_t)p->srcY; + switch (p->srcMemoryType) { + case hipMemoryTypeHost: + a.srcMemoryType = CU_MEMORYTYPE_HOST; + break; + case hipMemoryTypeDevice: + a.srcMemoryType = CU_MEMORYTYPE_DEVICE; + break; + case hipMemoryTypeArray: + a.srcMemoryType = CU_MEMORYTYPE_ARRAY; + break; + default: + a.srcMemoryType = CU_MEMORYTYPE_UNIFIED; + } + a.srcHost = p->srcHost; + a.srcDevice = (CUdeviceptr)p->srcDevice; + a.srcArray = (CUarray)p->srcArray; + a.srcPitch = (size_t)p->srcPitch; + a.dstXInBytes = (size_t)p->dstXInBytes; + a.dstY = (size_t)p->dstY; + switch (p->dstMemoryType) { + case hipMemoryTypeHost: + a.dstMemoryType = CU_MEMORYTYPE_HOST; + break; + case hipMemoryTypeDevice: + a.dstMemoryType = CU_MEMORYTYPE_DEVICE; + break; + case hipMemoryTypeArray: + a.dstMemoryType = CU_MEMORYTYPE_ARRAY; + break; + default: + a.dstMemoryType = CU_MEMORYTYPE_UNIFIED; + } + a.dstHost = p->dstHost; + a.dstDevice = (CUdeviceptr)p->dstDevice; + a.dstArray = (CUarray)p->dstArray; + a.dstPitch = (size_t)p->dstPitch; + a.WidthInBytes = (size_t)p->WidthInBytes; + a.Height = (size_t)p->Height; +} + + #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault #define hipArrayLayered cudaArrayLayered @@ -507,7 +706,6 @@ typedef struct cudaChannelFormatDesc hipChannelFormatDesc; typedef struct cudaResourceDesc hipResourceDesc; typedef struct cudaTextureDesc hipTextureDesc; typedef struct cudaResourceViewDesc hipResourceViewDesc; -typedef CUDA_RESOURCE_DESC HIP_RESOURCE_DESC; typedef CUDA_TEXTURE_DESC HIP_TEXTURE_DESC; typedef CUDA_RESOURCE_VIEW_DESC HIP_RESOURCE_VIEW_DESC; // adding code for hipmemSharedConfig @@ -1568,18 +1766,18 @@ 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 hipChannelFormatDesc* desc, +inline static hipError_t hipMallocArray(hipArray_t* array, const hipChannelFormatDesc* desc, size_t width, size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } -inline static hipError_t hipMalloc3DArray(hipArray** array, const hipChannelFormatDesc* desc, +inline static hipError_t hipMalloc3DArray(hipArray_t* array, const hipChannelFormatDesc* desc, hipExtent extent, unsigned int flags) { return hipCUDAErrorTohipError(cudaMalloc3DArray(array, desc, extent, flags)); } -inline static hipError_t hipFreeArray(hipArray* array) { +inline static hipError_t hipFreeArray(hipArray_t array) { return hipCUDAErrorTohipError(cudaFreeArray(array)); } @@ -1593,7 +1791,7 @@ inline static hipError_t hipMipmappedArrayDestroy(hipmipmappedArray hMipmappedAr return hipCUResultTohipError(cuMipmappedArrayDestroy(hMipmappedArray)); } -inline static hipError_t hipMipmappedArrayGetLevel(hiparray* pLevelArray, +inline static hipError_t hipMipmappedArrayGetLevel(hipArray_t* pLevelArray, hipmipmappedArray hMipMappedArray, unsigned int level) { return hipCUResultTohipError(cuMipmappedArrayGetLevel((CUarray*)pLevelArray, hMipMappedArray, level)); @@ -1763,11 +1961,23 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, } inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - return hipCUResultTohipError(cuMemcpy2D(pCopy)); + if(pCopy == nullptr) { + return hipCUResultTohipError(cuMemcpy2D(nullptr)); + } else { + CUDA_MEMCPY2D cudaCopy = {0}; + hipMemcpy2DTocudaMemcpy2D(cudaCopy, pCopy); + return hipCUResultTohipError(cuMemcpy2D((const CUDA_MEMCPY2D*)&cudaCopy)); + } } inline static hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream __dparm(0)) { - return hipCUResultTohipError(cuMemcpy2DAsync(pCopy, stream)); + if(pCopy == nullptr) { + return hipCUResultTohipError(cuMemcpy2DAsync(nullptr, stream)); + } else { + CUDA_MEMCPY2D cudaCopy = {0}; + hipMemcpy2DTocudaMemcpy2D(cudaCopy, pCopy); + return hipCUResultTohipError(cuMemcpy2DAsync((const CUDA_MEMCPY2D*)&cudaCopy, stream)); + } } inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) { @@ -1778,12 +1988,24 @@ inline static hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms *p, hipS return hipCUDAErrorTohipError(cudaMemcpy3DAsync(p, stream)); } -inline static hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy) { - return hipCUResultTohipError(cuMemcpy3D(pCopy)); +inline static hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pcopy) { + if(pcopy == nullptr) { + return hipCUResultTohipError(cuMemcpy3D(nullptr)); + } else { + CUDA_MEMCPY3D cudaCopy = {0}; + hipMemcpy3DTocudaMemcpy3D(cudaCopy, pcopy); + return hipCUResultTohipError(cuMemcpy3D((const CUDA_MEMCPY3D*)&cudaCopy)); + } } -inline static hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream) { - return hipCUResultTohipError(cuMemcpy3DAsync(pCopy, stream)); +inline static hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D *pcopy, hipStream_t stream) { + if(pcopy == nullptr) { + return hipCUResultTohipError(cuMemcpy3DAsync(nullptr, stream)); + } else { + CUDA_MEMCPY3D cudaCopy = {0}; + hipMemcpy3DTocudaMemcpy3D(cudaCopy, pcopy); + return hipCUResultTohipError(cuMemcpy3DAsync((const CUDA_MEMCPY3D*)&cudaCopy, stream)); + } } inline static hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, @@ -1793,7 +2015,7 @@ inline static hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* kind, stream)); } -inline static hipError_t hipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray* src, +inline static hipError_t hipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind) { return hipCUDAErrorTohipError(cudaMemcpy2DFromArray(dst, dpitch, src, wOffset, hOffset, width, @@ -1801,7 +2023,7 @@ inline static hipError_t hipMemcpy2DFromArray(void* dst, size_t dpitch, hipArray kind)); } -inline static hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray* src, +inline static hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hipArray_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { @@ -1811,14 +2033,14 @@ inline static hipError_t hipMemcpy2DFromArrayAsync(void* dst, size_t dpitch, hip stream)); } -inline static hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, +inline static hipError_t hipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { return hipCUDAErrorTohipError(cudaMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind)); } -inline static hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, +inline static hipError_t hipMemcpy2DToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { @@ -1828,7 +2050,7 @@ inline static hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, stream)); } -__HIP_DEPRECATED inline static hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, +__HIP_DEPRECATED inline static hipError_t hipMemcpyToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { return hipCUDAErrorTohipError( @@ -1842,12 +2064,12 @@ __HIP_DEPRECATED inline static hipError_t hipMemcpyFromArray(void* dst, hipArray kind)); } -inline static hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, +inline static hipError_t hipMemcpyAtoH(void* dst, hipArray_t srcArray, size_t srcOffset, size_t count) { return hipCUResultTohipError(cuMemcpyAtoH(dst, (CUarray)srcArray, srcOffset, count)); } -inline static hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, +inline static hipError_t hipMemcpyHtoA(hipArray_t dstArray, size_t dstOffset, const void* srcHost, size_t count) { return hipCUResultTohipError(cuMemcpyHtoA((CUarray)dstArray, dstOffset, srcHost, count)); } @@ -3241,7 +3463,7 @@ inline static hipError_t hipTexObjectCreate(hipTextureObject_t* pTexObject, const HIP_RESOURCE_DESC* pResDesc, const HIP_TEXTURE_DESC* pTexDesc, const HIP_RESOURCE_VIEW_DESC* pResViewDesc) { - return hipCUResultTohipError(cuTexObjectCreate((CUtexObject*)pTexObject, pResDesc, pTexDesc, pResViewDesc)); + return hipCUResultTohipError(cuTexObjectCreate((CUtexObject*)pTexObject,(CUDA_RESOURCE_DESC*)pResDesc, pTexDesc, pResViewDesc)); } inline static hipError_t hipTexObjectDestroy(hipTextureObject_t texObject) { @@ -3249,7 +3471,7 @@ inline static hipError_t hipTexObjectDestroy(hipTextureObject_t texObject) { } inline static hipError_t hipTexObjectGetResourceDesc(HIP_RESOURCE_DESC* pResDesc, hipTextureObject_t texObject) { - return hipCUResultTohipError(cuTexObjectGetResourceDesc(pResDesc, (CUtexObject)texObject)); + return hipCUResultTohipError(cuTexObjectGetResourceDesc((CUDA_RESOURCE_DESC*)pResDesc, (CUtexObject)texObject)); } inline static hipError_t hipTexObjectGetResourceViewDesc(HIP_RESOURCE_VIEW_DESC* pResViewDesc, hipTextureObject_t texObject) { @@ -3284,35 +3506,35 @@ __HIP_DEPRECATED inline static hipError_t hipTexRefSetFlags(hipTexRef hTexRef, u return hipCUResultTohipError(cuTexRefSetFlags(hTexRef,Flags)); } -__HIP_DEPRECATED inline static hipError_t hipTexRefSetArray(hipTexRef hTexRef, hiparray hArray, unsigned int Flags){ - return hipCUResultTohipError(cuTexRefSetArray(hTexRef,hArray,Flags)); +__HIP_DEPRECATED inline static hipError_t hipTexRefSetArray(hipTexRef hTexRef, hipArray_t hArray, unsigned int Flags){ + return hipCUResultTohipError(cuTexRefSetArray(hTexRef,(CUarray)hArray,Flags)); } -inline static hipError_t hipArrayCreate(hiparray* pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray){ - return hipCUResultTohipError(cuArrayCreate(pHandle, pAllocateArray)); +inline static hipError_t hipArrayCreate(hipArray_t* pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray){ + return hipCUResultTohipError(cuArrayCreate((CUarray*)pHandle, pAllocateArray)); } -inline static hipError_t hipArrayDestroy(hiparray hArray){ - return hipCUResultTohipError(cuArrayDestroy(hArray)); +inline static hipError_t hipArrayDestroy(hipArray_t hArray){ + return hipCUResultTohipError(cuArrayDestroy((CUarray)hArray)); } -inline static hipError_t hipArray3DCreate(hiparray* pHandle, +inline static hipError_t hipArray3DCreate(hipArray_t* pHandle, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray){ - return hipCUResultTohipError(cuArray3DCreate(pHandle, pAllocateArray)); + return hipCUResultTohipError(cuArray3DCreate((CUarray*)pHandle, pAllocateArray)); } inline static hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc, hipExtent* extent, - unsigned int* flags, hipArray* array) { + unsigned int* flags, hipArray_t array) { return hipCUDAErrorTohipError(cudaArrayGetInfo(desc, extent, flags, array)); } inline static hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor, - hipArray* array) { + hipArray_t array) { return hipCUResultTohipError(cuArrayGetDescriptor(pArrayDescriptor, (CUarray)array)); } inline static hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor, - hipArray* array) { + hipArray_t array) { return hipCUResultTohipError(cuArray3DGetDescriptor(pArrayDescriptor, (CUarray)array)); }