From fcfa0d073352d3fad813214c8ae2b0df8beb8dd3 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/clr commit: c089ef4a2b7d2c6530595c8f8c722521d514fc83] --- .../hip/amd_detail/amd_hip_runtime_pt_api.h | 4 +- .../include/hip/amd_detail/hip_prof_str.h | 159 ++++------ .../nvidia_detail/nvidia_hip_runtime_api.h | 292 +++++++++++++++--- projects/clr/hipamd/src/hip_formatting.hpp | 35 --- projects/clr/hipamd/src/hip_graph.cpp | 6 +- projects/clr/hipamd/src/hip_graph_capture.hpp | 6 +- projects/clr/hipamd/src/hip_graph_helper.hpp | 10 +- projects/clr/hipamd/src/hip_internal.hpp | 14 + projects/clr/hipamd/src/hip_memory.cpp | 80 ++--- 9 files changed, 389 insertions(+), 217 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h index a597dbd03b..070c58bd31 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h @@ -136,14 +136,14 @@ hipError_t hipMemcpyToSymbolAsync_spt(const void* symbol, const void* src, size_ hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t src, size_t wOffsetSrc, size_t hOffset, size_t count, hipMemcpyKind kind); -hipError_t hipMemcpy2DToArray_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, +hipError_t hipMemcpy2DToArray_spt(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind); hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream); -hipError_t hipMemcpy2DToArrayAsync_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, +hipError_t hipMemcpy2DToArrayAsync_spt(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); diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 1bdd5ff70c..4753028e00 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -1173,32 +1173,29 @@ typedef struct hip_api_data_s { hipStream_t stream; } __hipPushCallConfiguration; struct { - hipArray** array; - hipArray* array__val; + hipArray_t* array; + hipArray_t array__val; const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray; HIP_ARRAY3D_DESCRIPTOR pAllocateArray__val; } hipArray3DCreate; struct { HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor; HIP_ARRAY3D_DESCRIPTOR pArrayDescriptor__val; - hipArray* array; - hipArray array__val; + hipArray_t array; } hipArray3DGetDescriptor; struct { - hipArray** pHandle; - hipArray* pHandle__val; + hipArray_t* pHandle; + hipArray_t pHandle__val; const HIP_ARRAY_DESCRIPTOR* pAllocateArray; HIP_ARRAY_DESCRIPTOR pAllocateArray__val; } hipArrayCreate; struct { - hipArray* array; - hipArray array__val; + hipArray_t array; } hipArrayDestroy; struct { HIP_ARRAY_DESCRIPTOR* pArrayDescriptor; HIP_ARRAY_DESCRIPTOR pArrayDescriptor__val; - hipArray* array; - hipArray array__val; + hipArray_t array; } hipArrayGetDescriptor; struct { hipChannelFormatDesc* desc; @@ -1207,8 +1204,7 @@ typedef struct hip_api_data_s { hipExtent extent__val; unsigned int* flags; unsigned int flags__val; - hipArray* array; - hipArray array__val; + hipArray_t array; } hipArrayGetInfo; struct { int* device; @@ -1565,8 +1561,7 @@ typedef struct hip_api_data_s { void* ptr; } hipFree; struct { - hipArray* array; - hipArray array__val; + hipArray_t array; } hipFreeArray; struct { void* dev_ptr; @@ -2282,8 +2277,8 @@ typedef struct hip_api_data_s { unsigned int flags; } hipMalloc3DArray; struct { - hipArray** array; - hipArray* array__val; + hipArray_t* array; + hipArray_t array__val; const hipChannelFormatDesc* desc; hipChannelFormatDesc desc__val; size_t width; @@ -2584,8 +2579,7 @@ typedef struct hip_api_data_s { hipStream_t stream; } hipMemcpy2DFromArrayAsync; struct { - hipArray* dst; - hipArray dst__val; + hipArray_t dst; size_t wOffset; size_t hOffset; const void* src; @@ -2595,8 +2589,7 @@ typedef struct hip_api_data_s { hipMemcpyKind kind; } hipMemcpy2DToArray; struct { - hipArray* dst; - hipArray dst__val; + hipArray_t dst; size_t wOffset; size_t hOffset; const void* src; @@ -2624,8 +2617,7 @@ typedef struct hip_api_data_s { } hipMemcpyAsync; struct { void* dst; - hipArray* srcArray; - hipArray srcArray__val; + hipArray_t srcArray; size_t srcOffset; size_t count; } hipMemcpyAtoH; @@ -2675,8 +2667,7 @@ typedef struct hip_api_data_s { hipStream_t stream; } hipMemcpyFromSymbolAsync; struct { - hipArray* dstArray; - hipArray dstArray__val; + hipArray_t dstArray; size_t dstOffset; const void* srcHost; size_t count; @@ -2717,8 +2708,7 @@ typedef struct hip_api_data_s { hipStream_t stream; } hipMemcpyPeerAsync; struct { - hipArray* dst; - hipArray dst__val; + hipArray_t dst; size_t wOffset; size_t hOffset; const void* src; @@ -3294,36 +3284,36 @@ typedef struct hip_api_data_s { cb_data.args.__hipPushCallConfiguration.sharedMem = (size_t)sharedMem; \ cb_data.args.__hipPushCallConfiguration.stream = (hipStream_t)stream; \ }; -// hipArray3DCreate[('hipArray**', 'array'), ('const HIP_ARRAY3D_DESCRIPTOR*', 'pAllocateArray')] +// hipArray3DCreate[('hipArray_t*', 'array'), ('const HIP_ARRAY3D_DESCRIPTOR*', 'pAllocateArray')] #define INIT_hipArray3DCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipArray3DCreate.array = (hipArray**)array; \ + cb_data.args.hipArray3DCreate.array = (hipArray_t*)array; \ cb_data.args.hipArray3DCreate.pAllocateArray = (const HIP_ARRAY3D_DESCRIPTOR*)pAllocateArray; \ }; -// hipArray3DGetDescriptor[('HIP_ARRAY3D_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')] +// hipArray3DGetDescriptor[('HIP_ARRAY3D_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray_t', 'array')] #define INIT_hipArray3DGetDescriptor_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipArray3DGetDescriptor.pArrayDescriptor = (HIP_ARRAY3D_DESCRIPTOR*)pArrayDescriptor; \ - cb_data.args.hipArray3DGetDescriptor.array = (hipArray*)array; \ + cb_data.args.hipArray3DGetDescriptor.array = (hipArray_t)array; \ }; -// hipArrayCreate[('hipArray**', 'pHandle'), ('const HIP_ARRAY_DESCRIPTOR*', 'pAllocateArray')] +// hipArrayCreate[('hipArray_t*', 'pHandle'), ('const HIP_ARRAY_DESCRIPTOR*', 'pAllocateArray')] #define INIT_hipArrayCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipArrayCreate.pHandle = (hipArray**)array; \ + cb_data.args.hipArrayCreate.pHandle = (hipArray_t*)array; \ cb_data.args.hipArrayCreate.pAllocateArray = (const HIP_ARRAY_DESCRIPTOR*)pAllocateArray; \ }; -// hipArrayDestroy[('hipArray*', 'array')] +// hipArrayDestroy[('hipArray_t', 'array')] #define INIT_hipArrayDestroy_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipArrayDestroy.array = (hipArray*)array; \ + cb_data.args.hipArrayDestroy.array = (hipArray_t)array; \ }; -// hipArrayGetDescriptor[('HIP_ARRAY_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')] +// hipArrayGetDescriptor[('HIP_ARRAY_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray_t', 'array')] #define INIT_hipArrayGetDescriptor_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipArrayGetDescriptor.pArrayDescriptor = (HIP_ARRAY_DESCRIPTOR*)pArrayDescriptor; \ - cb_data.args.hipArrayGetDescriptor.array = (hipArray*)array; \ + cb_data.args.hipArrayGetDescriptor.array = (hipArray_t)array; \ }; -// hipArrayGetInfo[('hipChannelFormatDesc*', 'desc'), ('hipExtent*', 'extent'), ('unsigned int*', 'flags'), ('hipArray*', 'array')] +// hipArrayGetInfo[('hipChannelFormatDesc*', 'desc'), ('hipExtent*', 'extent'), ('unsigned int*', 'flags'), ('hipArray_t', 'array')] #define INIT_hipArrayGetInfo_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipArrayGetInfo.desc = (hipChannelFormatDesc*)desc; \ cb_data.args.hipArrayGetInfo.extent = (hipExtent*)extent; \ cb_data.args.hipArrayGetInfo.flags = (unsigned int*)flags; \ - cb_data.args.hipArrayGetInfo.array = (hipArray*)array; \ + cb_data.args.hipArrayGetInfo.array = (hipArray_t)array; \ }; // hipChooseDevice[('int*', 'device'), ('const hipDeviceProp_tR0600*', 'prop')] #define INIT_hipChooseDevice_CB_ARGS_DATA(cb_data) { \ @@ -3711,9 +3701,9 @@ typedef struct hip_api_data_s { #define INIT_hipFree_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipFree.ptr = (void*)ptr; \ }; -// hipFreeArray[('hipArray*', 'array')] +// hipFreeArray[('hipArray_t', 'array')] #define INIT_hipFreeArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipFreeArray.array = (hipArray*)array; \ + cb_data.args.hipFreeArray.array = (hipArray_t)array; \ }; // hipFreeAsync[('void*', 'dev_ptr'), ('hipStream_t', 'stream')] #define INIT_hipFreeAsync_CB_ARGS_DATA(cb_data) { \ @@ -4434,9 +4424,9 @@ typedef struct hip_api_data_s { cb_data.args.hipMalloc3DArray.extent = (hipExtent)extent; \ cb_data.args.hipMalloc3DArray.flags = (unsigned int)flags; \ }; -// hipMallocArray[('hipArray**', 'array'), ('const hipChannelFormatDesc*', 'desc'), ('size_t', 'width'), ('size_t', 'height'), ('unsigned int', 'flags')] +// hipMallocArray[('hipArray_t*', 'array'), ('const hipChannelFormatDesc*', 'desc'), ('size_t', 'width'), ('size_t', 'height'), ('unsigned int', 'flags')] #define INIT_hipMallocArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMallocArray.array = (hipArray**)array; \ + cb_data.args.hipMallocArray.array = (hipArray_t*)array; \ cb_data.args.hipMallocArray.desc = (const hipChannelFormatDesc*)desc; \ cb_data.args.hipMallocArray.width = (size_t)width; \ cb_data.args.hipMallocArray.height = (size_t)height; \ @@ -4740,9 +4730,9 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpy2DFromArrayAsync.kind = (hipMemcpyKind)kind; \ cb_data.args.hipMemcpy2DFromArrayAsync.stream = (hipStream_t)stream; \ }; -// hipMemcpy2DToArray[('hipArray*', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] +// hipMemcpy2DToArray[('hipArray_t', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] #define INIT_hipMemcpy2DToArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy2DToArray.dst = (hipArray*)dst; \ + cb_data.args.hipMemcpy2DToArray.dst = (hipArray_t)dst; \ cb_data.args.hipMemcpy2DToArray.wOffset = (size_t)wOffset; \ cb_data.args.hipMemcpy2DToArray.hOffset = (size_t)hOffset; \ cb_data.args.hipMemcpy2DToArray.src = (const void*)src; \ @@ -4751,9 +4741,9 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpy2DToArray.height = (size_t)height; \ cb_data.args.hipMemcpy2DToArray.kind = (hipMemcpyKind)kind; \ }; -// hipMemcpy2DToArrayAsync[('hipArray*', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] +// 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')] #define INIT_hipMemcpy2DToArrayAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy2DToArrayAsync.dst = (hipArray*)dst; \ + cb_data.args.hipMemcpy2DToArrayAsync.dst = (hipArray_t)dst; \ cb_data.args.hipMemcpy2DToArrayAsync.wOffset = (size_t)wOffset; \ cb_data.args.hipMemcpy2DToArrayAsync.hOffset = (size_t)hOffset; \ cb_data.args.hipMemcpy2DToArrayAsync.src = (const void*)src; \ @@ -4780,10 +4770,10 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyAsync.kind = (hipMemcpyKind)kind; \ cb_data.args.hipMemcpyAsync.stream = (hipStream_t)stream; \ }; -// hipMemcpyAtoH[('void*', 'dst'), ('hipArray*', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'count')] +// hipMemcpyAtoH[('void*', 'dst'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'count')] #define INIT_hipMemcpyAtoH_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpyAtoH.dst = (void*)dstHost; \ - cb_data.args.hipMemcpyAtoH.srcArray = (hipArray*)srcArray; \ + cb_data.args.hipMemcpyAtoH.srcArray = (hipArray_t)srcArray; \ cb_data.args.hipMemcpyAtoH.srcOffset = (size_t)srcOffset; \ cb_data.args.hipMemcpyAtoH.count = (size_t)ByteCount; \ }; @@ -4839,9 +4829,9 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyFromSymbolAsync.kind = (hipMemcpyKind)kind; \ cb_data.args.hipMemcpyFromSymbolAsync.stream = (hipStream_t)stream; \ }; -// hipMemcpyHtoA[('hipArray*', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'count')] +// hipMemcpyHtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'count')] #define INIT_hipMemcpyHtoA_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyHtoA.dstArray = (hipArray*)dstArray; \ + cb_data.args.hipMemcpyHtoA.dstArray = (hipArray_t)dstArray; \ cb_data.args.hipMemcpyHtoA.dstOffset = (size_t)dstOffset; \ cb_data.args.hipMemcpyHtoA.srcHost = (const void*)srcHost; \ cb_data.args.hipMemcpyHtoA.count = (size_t)ByteCount; \ @@ -4885,9 +4875,9 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyPeerAsync.sizeBytes = (size_t)sizeBytes; \ cb_data.args.hipMemcpyPeerAsync.stream = (hipStream_t)stream; \ }; -// hipMemcpyToArray[('hipArray*', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'count'), ('hipMemcpyKind', 'kind')] +// hipMemcpyToArray[('hipArray_t', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'count'), ('hipMemcpyKind', 'kind')] #define INIT_hipMemcpyToArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyToArray.dst = (hipArray*)dst; \ + cb_data.args.hipMemcpyToArray.dst = (hipArray_t)dst; \ cb_data.args.hipMemcpyToArray.wOffset = (size_t)wOffset; \ cb_data.args.hipMemcpyToArray.hOffset = (size_t)hOffset; \ cb_data.args.hipMemcpyToArray.src = (const void*)src; \ @@ -5542,36 +5532,32 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // __hipPushCallConfiguration[('dim3', 'gridDim'), ('dim3', 'blockDim'), ('size_t', 'sharedMem'), ('hipStream_t', 'stream')] case HIP_API_ID___hipPushCallConfiguration: break; -// hipArray3DCreate[('hipArray**', 'array'), ('const HIP_ARRAY3D_DESCRIPTOR*', 'pAllocateArray')] +// hipArray3DCreate[('hipArray_t*', 'array'), ('const HIP_ARRAY3D_DESCRIPTOR*', 'pAllocateArray')] case HIP_API_ID_hipArray3DCreate: if (data->args.hipArray3DCreate.array) data->args.hipArray3DCreate.array__val = *(data->args.hipArray3DCreate.array); if (data->args.hipArray3DCreate.pAllocateArray) data->args.hipArray3DCreate.pAllocateArray__val = *(data->args.hipArray3DCreate.pAllocateArray); break; -// hipArray3DGetDescriptor[('HIP_ARRAY3D_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')] +// hipArray3DGetDescriptor[('HIP_ARRAY3D_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray_t', 'array')] case HIP_API_ID_hipArray3DGetDescriptor: if (data->args.hipArray3DGetDescriptor.pArrayDescriptor) data->args.hipArray3DGetDescriptor.pArrayDescriptor__val = *(data->args.hipArray3DGetDescriptor.pArrayDescriptor); - if (data->args.hipArray3DGetDescriptor.array) data->args.hipArray3DGetDescriptor.array__val = *(data->args.hipArray3DGetDescriptor.array); break; -// hipArrayCreate[('hipArray**', 'pHandle'), ('const HIP_ARRAY_DESCRIPTOR*', 'pAllocateArray')] +// hipArrayCreate[('hipArray_t*', 'pHandle'), ('const HIP_ARRAY_DESCRIPTOR*', 'pAllocateArray')] case HIP_API_ID_hipArrayCreate: if (data->args.hipArrayCreate.pHandle) data->args.hipArrayCreate.pHandle__val = *(data->args.hipArrayCreate.pHandle); if (data->args.hipArrayCreate.pAllocateArray) data->args.hipArrayCreate.pAllocateArray__val = *(data->args.hipArrayCreate.pAllocateArray); break; -// hipArrayDestroy[('hipArray*', 'array')] +// hipArrayDestroy[('hipArray_t', 'array')] case HIP_API_ID_hipArrayDestroy: - if (data->args.hipArrayDestroy.array) data->args.hipArrayDestroy.array__val = *(data->args.hipArrayDestroy.array); break; -// hipArrayGetDescriptor[('HIP_ARRAY_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')] +// hipArrayGetDescriptor[('HIP_ARRAY_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray_t', 'array')] case HIP_API_ID_hipArrayGetDescriptor: if (data->args.hipArrayGetDescriptor.pArrayDescriptor) data->args.hipArrayGetDescriptor.pArrayDescriptor__val = *(data->args.hipArrayGetDescriptor.pArrayDescriptor); - if (data->args.hipArrayGetDescriptor.array) data->args.hipArrayGetDescriptor.array__val = *(data->args.hipArrayGetDescriptor.array); break; -// hipArrayGetInfo[('hipChannelFormatDesc*', 'desc'), ('hipExtent*', 'extent'), ('unsigned int*', 'flags'), ('hipArray*', 'array')] +// hipArrayGetInfo[('hipChannelFormatDesc*', 'desc'), ('hipExtent*', 'extent'), ('unsigned int*', 'flags'), ('hipArray_t', 'array')] case HIP_API_ID_hipArrayGetInfo: if (data->args.hipArrayGetInfo.desc) data->args.hipArrayGetInfo.desc__val = *(data->args.hipArrayGetInfo.desc); if (data->args.hipArrayGetInfo.extent) data->args.hipArrayGetInfo.extent__val = *(data->args.hipArrayGetInfo.extent); if (data->args.hipArrayGetInfo.flags) data->args.hipArrayGetInfo.flags__val = *(data->args.hipArrayGetInfo.flags); - if (data->args.hipArrayGetInfo.array) data->args.hipArrayGetInfo.array__val = *(data->args.hipArrayGetInfo.array); break; // hipChooseDevice[('int*', 'device'), ('const hipDeviceProp_tR0600*', 'prop')] case HIP_API_ID_hipChooseDevice: @@ -5857,9 +5843,8 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipFree[('void*', 'ptr')] case HIP_API_ID_hipFree: break; -// hipFreeArray[('hipArray*', 'array')] +// hipFreeArray[('hipArray_t', 'array')] case HIP_API_ID_hipFreeArray: - if (data->args.hipFreeArray.array) data->args.hipFreeArray.array__val = *(data->args.hipFreeArray.array); break; // hipFreeAsync[('void*', 'dev_ptr'), ('hipStream_t', 'stream')] case HIP_API_ID_hipFreeAsync: @@ -6331,7 +6316,7 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipMalloc3DArray.array) data->args.hipMalloc3DArray.array__val = *(data->args.hipMalloc3DArray.array); if (data->args.hipMalloc3DArray.desc) data->args.hipMalloc3DArray.desc__val = *(data->args.hipMalloc3DArray.desc); break; -// hipMallocArray[('hipArray**', 'array'), ('const hipChannelFormatDesc*', 'desc'), ('size_t', 'width'), ('size_t', 'height'), ('unsigned int', 'flags')] +// hipMallocArray[('hipArray_t*', 'array'), ('const hipChannelFormatDesc*', 'desc'), ('size_t', 'width'), ('size_t', 'height'), ('unsigned int', 'flags')] case HIP_API_ID_hipMallocArray: if (data->args.hipMallocArray.array) data->args.hipMallocArray.array__val = *(data->args.hipMallocArray.array); if (data->args.hipMallocArray.desc) data->args.hipMallocArray.desc__val = *(data->args.hipMallocArray.desc); @@ -6511,13 +6496,11 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // 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')] case HIP_API_ID_hipMemcpy2DFromArrayAsync: break; -// hipMemcpy2DToArray[('hipArray*', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] +// hipMemcpy2DToArray[('hipArray_t', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind')] case HIP_API_ID_hipMemcpy2DToArray: - if (data->args.hipMemcpy2DToArray.dst) data->args.hipMemcpy2DToArray.dst__val = *(data->args.hipMemcpy2DToArray.dst); break; -// hipMemcpy2DToArrayAsync[('hipArray*', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'spitch'), ('size_t', 'width'), ('size_t', 'height'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] +// 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')] case HIP_API_ID_hipMemcpy2DToArrayAsync: - if (data->args.hipMemcpy2DToArrayAsync.dst) data->args.hipMemcpy2DToArrayAsync.dst__val = *(data->args.hipMemcpy2DToArrayAsync.dst); break; // hipMemcpy3D[('const hipMemcpy3DParms*', 'p')] case HIP_API_ID_hipMemcpy3D: @@ -6530,9 +6513,8 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpyAsync[('void*', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpyAsync: break; -// hipMemcpyAtoH[('void*', 'dst'), ('hipArray*', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'count')] +// hipMemcpyAtoH[('void*', 'dst'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'count')] case HIP_API_ID_hipMemcpyAtoH: - if (data->args.hipMemcpyAtoH.srcArray) data->args.hipMemcpyAtoH.srcArray__val = *(data->args.hipMemcpyAtoH.srcArray); break; // hipMemcpyDtoD[('hipDeviceptr_t', 'dst'), ('hipDeviceptr_t', 'src'), ('size_t', 'sizeBytes')] case HIP_API_ID_hipMemcpyDtoD: @@ -6555,9 +6537,8 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpyFromSymbolAsync[('void*', 'dst'), ('const void*', 'symbol'), ('size_t', 'sizeBytes'), ('size_t', 'offset'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpyFromSymbolAsync: break; -// hipMemcpyHtoA[('hipArray*', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'count')] +// hipMemcpyHtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'count')] case HIP_API_ID_hipMemcpyHtoA: - if (data->args.hipMemcpyHtoA.dstArray) data->args.hipMemcpyHtoA.dstArray__val = *(data->args.hipMemcpyHtoA.dstArray); break; // hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')] case HIP_API_ID_hipMemcpyHtoD: @@ -6579,9 +6560,8 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpyPeerAsync[('void*', 'dst'), ('int', 'dstDeviceId'), ('const void*', 'src'), ('int', 'srcDevice'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpyPeerAsync: break; -// hipMemcpyToArray[('hipArray*', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'count'), ('hipMemcpyKind', 'kind')] +// hipMemcpyToArray[('hipArray_t', 'dst'), ('size_t', 'wOffset'), ('size_t', 'hOffset'), ('const void*', 'src'), ('size_t', 'count'), ('hipMemcpyKind', 'kind')] case HIP_API_ID_hipMemcpyToArray: - if (data->args.hipMemcpyToArray.dst) data->args.hipMemcpyToArray.dst__val = *(data->args.hipMemcpyToArray.dst); break; // hipMemcpyToSymbol[('const void*', 'symbol'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('size_t', 'offset'), ('hipMemcpyKind', 'kind')] case HIP_API_ID_hipMemcpyToSymbol: @@ -6979,8 +6959,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << "hipArray3DGetDescriptor("; if (data->args.hipArray3DGetDescriptor.pArrayDescriptor == NULL) oss << "pArrayDescriptor=NULL"; else { oss << "pArrayDescriptor="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.pArrayDescriptor__val); } - if (data->args.hipArray3DGetDescriptor.array == NULL) oss << ", array=NULL"; - else { oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.array__val); } + oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.array); oss << ")"; break; case HIP_API_ID_hipArrayCreate: @@ -6993,16 +6972,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipArrayDestroy: oss << "hipArrayDestroy("; - if (data->args.hipArrayDestroy.array == NULL) oss << "array=NULL"; - else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayDestroy.array__val); } + oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayDestroy.array); oss << ")"; break; case HIP_API_ID_hipArrayGetDescriptor: oss << "hipArrayGetDescriptor("; if (data->args.hipArrayGetDescriptor.pArrayDescriptor == NULL) oss << "pArrayDescriptor=NULL"; else { oss << "pArrayDescriptor="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetDescriptor.pArrayDescriptor__val); } - if (data->args.hipArrayGetDescriptor.array == NULL) oss << ", array=NULL"; - else { oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetDescriptor.array__val); } + oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetDescriptor.array); oss << ")"; break; case HIP_API_ID_hipArrayGetInfo: @@ -7014,7 +6991,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da if (data->args.hipArrayGetInfo.flags == NULL) oss << ", flags=NULL"; else { oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.flags__val); } if (data->args.hipArrayGetInfo.array == NULL) oss << ", array=NULL"; - else { oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.array__val); } + oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.array); oss << ")"; break; case HIP_API_ID_hipChooseDevice: @@ -7537,8 +7514,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipFreeArray: oss << "hipFreeArray("; - if (data->args.hipFreeArray.array == NULL) oss << "array=NULL"; - else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipFreeArray.array__val); } + oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipFreeArray.array); oss << ")"; break; case HIP_API_ID_hipFreeAsync: @@ -8890,8 +8866,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipMemcpy2DToArray: oss << "hipMemcpy2DToArray("; - if (data->args.hipMemcpy2DToArray.dst == NULL) oss << "dst=NULL"; - else { oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArray.dst__val); } + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArray.dst); oss << ", wOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArray.wOffset); oss << ", hOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArray.hOffset); oss << ", src="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArray.src); @@ -8903,8 +8878,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipMemcpy2DToArrayAsync: oss << "hipMemcpy2DToArrayAsync("; - if (data->args.hipMemcpy2DToArrayAsync.dst == NULL) oss << "dst=NULL"; - else { oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArrayAsync.dst__val); } + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArrayAsync.dst); oss << ", wOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArrayAsync.wOffset); oss << ", hOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArrayAsync.hOffset); oss << ", src="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy2DToArrayAsync.src); @@ -8940,8 +8914,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da case HIP_API_ID_hipMemcpyAtoH: oss << "hipMemcpyAtoH("; oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.dst); - if (data->args.hipMemcpyAtoH.srcArray == NULL) oss << ", srcArray=NULL"; - else { oss << ", srcArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.srcArray__val); } + oss << ", srcArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.srcArray); oss << ", srcOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.srcOffset); oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoH.count); oss << ")"; @@ -9007,8 +8980,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipMemcpyHtoA: oss << "hipMemcpyHtoA("; - if (data->args.hipMemcpyHtoA.dstArray == NULL) oss << "dstArray=NULL"; - else { oss << "dstArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoA.dstArray__val); } + oss << "dstArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoA.dstArray); oss << ", dstOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoA.dstOffset); oss << ", srcHost="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoA.srcHost); oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyHtoA.count); @@ -9063,8 +9035,7 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da break; case HIP_API_ID_hipMemcpyToArray: oss << "hipMemcpyToArray("; - if (data->args.hipMemcpyToArray.dst == NULL) oss << "dst=NULL"; - else { oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyToArray.dst__val); } + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyToArray.dst); oss << ", wOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyToArray.wOffset); oss << ", hOffset="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyToArray.hOffset); oss << ", src="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyToArray.src); diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 89ae35428c..ac387be5d5 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/clr/hipamd/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)); } diff --git a/projects/clr/hipamd/src/hip_formatting.hpp b/projects/clr/hipamd/src/hip_formatting.hpp index 00a27476be..c4e874bd02 100644 --- a/projects/clr/hipamd/src/hip_formatting.hpp +++ b/projects/clr/hipamd/src/hip_formatting.hpp @@ -683,41 +683,6 @@ inline std::ostream& operator<<(std::ostream& os, const hipResourceDesc* s) { return os; } - -inline std::ostream& operator<<(std::ostream& os, const hipArray& s) { - os << '{' - << s.data - << ',' - << s.desc - << ',' - << s.type - << ',' - << s.width - << ',' - << s.height - << ',' - << s.depth - << ',' - << s.Format - << ',' - << s.NumChannels - << ',' - << s.isDrv - << ',' - << s.textureType - << '}'; - return os; -} - -inline std::ostream& operator<<(std::ostream& os, const hipArray* s) { - if (s) { - os << *s; - } else { - os << "nullptr"; - } - return os; -} - inline std::ostream& operator<<(std::ostream& os, const textureReference& s) { os << '{' << s.normalized diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 95aa925f61..d87b19e40b 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -394,7 +394,7 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz return hipSuccess; } -hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, size_t& wOffset, +hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray_t& dst, size_t& wOffset, size_t& hOffset, const void*& src, size_t& spitch, size_t& width, size_t& height, hipMemcpyKind& kind) { ClPrint(amd::LOG_INFO, amd::LOG_API, @@ -477,7 +477,7 @@ hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D* return hipSuccess; } -hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArray*& srcArray, +hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArray_t& srcArray, size_t& srcOffset, size_t& ByteCount) { ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node MemcpyParam2D on stream : %p", stream); @@ -505,7 +505,7 @@ hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArr return hipSuccess; } -hipError_t capturehipMemcpyHtoAAsync(hipStream_t& stream, hipArray*& dstArray, size_t& dstOffset, +hipError_t capturehipMemcpyHtoAAsync(hipStream_t& stream, hipArray_t& dstArray, size_t& dstOffset, const void*& srcHost, size_t& ByteCount) { ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] Current capture node MemcpyParam2D on stream : %p", stream); diff --git a/projects/clr/hipamd/src/hip_graph_capture.hpp b/projects/clr/hipamd/src/hip_graph_capture.hpp index 9506ae99cd..77a27a5dcd 100644 --- a/projects/clr/hipamd/src/hip_graph_capture.hpp +++ b/projects/clr/hipamd/src/hip_graph_capture.hpp @@ -52,14 +52,14 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz size_t& hOffsetSrc, size_t& width, size_t& height, hipMemcpyKind& kind); -hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray*& dst, size_t& wOffset, +hipError_t capturehipMemcpy2DToArrayAsync(hipStream_t& stream, hipArray_t& dst, size_t& wOffset, size_t& hOffset, const void*& src, size_t& spitch, size_t& width, size_t& height, hipMemcpyKind& kind); -hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArray*& srcArray, +hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArray_t& srcArray, size_t& srcOffset, size_t& ByteCount); -hipError_t capturehipMemcpyHtoAAsync(hipStream_t& stream, hipArray*& dstArray, size_t& dstOffset, +hipError_t capturehipMemcpyHtoAAsync(hipStream_t& stream, hipArray_t& dstArray, size_t& dstOffset, const void*& srcHost, size_t& ByteCount); hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*& p); diff --git a/projects/clr/hipamd/src/hip_graph_helper.hpp b/projects/clr/hipamd/src/hip_graph_helper.hpp index a1d82edb1b..b195267395 100644 --- a/projects/clr/hipamd/src/hip_graph_helper.hpp +++ b/projects/clr/hipamd/src/hip_graph_helper.hpp @@ -65,13 +65,13 @@ hipError_t ihipMemset3DCommand(std::vector& commands, hipPitchedP hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes, size_t offset, size_t& sym_size, hipDeviceptr_t& device_ptr); -hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyAtoDValidate(hipArray_t srcArray, void* dstDevice, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t dstRowPitch, size_t dstSlicePitch, amd::Memory*& dstMemory, amd::Image*& srcImage, amd::BufferRect& srcRect, amd::BufferRect& dstRect); -hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray* dstArray, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage, amd::Memory*& srcMemory, amd::BufferRect& dstRect, @@ -98,17 +98,17 @@ hipError_t ihipMemcpyHtoDValidate(const void* srcHost, void* dstDevice, amd::Coo amd::BufferRect& srcRect, amd::BufferRect& dstRect); -hipError_t ihipMemcpyAtoAValidate(hipArray* srcArray, hipArray* dstArray, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyAtoAValidate(hipArray_t srcArray, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, amd::Image*& srcImage, amd::Image*& dstImage); -hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray* dstArray, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage, amd::BufferRect& srcRect); -hipError_t ihipMemcpyAtoHValidate(hipArray* srcArray, void* dstHost, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t dstRowPitch, size_t dstSlicePitch, amd::Image*& srcImage, amd::BufferRect& dstRect); diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 4c34d5baa7..81b91a7aa1 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/hipamd/src/hip_internal.hpp @@ -67,6 +67,20 @@ typedef struct ihipIpcMemHandle_st { char reserved[IHIP_IPC_MEM_RESERVED_SIZE]; } ihipIpcMemHandle_t; +typedef struct hipArray { + void* data; // FIXME: generalize this + struct hipChannelFormatDesc desc; + unsigned int type; + unsigned int width; + unsigned int height; + unsigned int depth; + enum hipArray_Format Format; + unsigned int NumChannels; + bool isDrv; + unsigned int textureType; + unsigned int flags; +}hipArray; + #define IHIP_IPC_EVENT_HANDLE_SIZE 32 #define IHIP_IPC_EVENT_RESERVED_SIZE LP64_SWITCH(28,24) typedef struct ihipIpcEventHandle_st { diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 233fc1e0ae..71445219c4 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -711,7 +711,7 @@ hipError_t hipHostFree(void* ptr) { HIP_RETURN(ihipFree(ptr)); } -hipError_t ihipArrayDestroy(hipArray* array) { +hipError_t ihipArrayDestroy(hipArray_t array) { if (array == nullptr) { return hipErrorInvalidValue; } @@ -737,7 +737,7 @@ hipError_t ihipArrayDestroy(hipArray* array) { return hipSuccess; } -hipError_t hipFreeArray(hipArray* array) { +hipError_t hipFreeArray(hipArray_t array) { HIP_INIT_API(hipFreeArray, array); CHECK_STREAM_CAPTURE_SUPPORTED(); HIP_RETURN(ihipArrayDestroy(array)); @@ -1033,7 +1033,7 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder, return image; } -hipError_t ihipArrayCreate(hipArray** array, +hipError_t ihipArrayCreate(hipArray_t* array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray, unsigned int numMipmapLevels) { if (array == nullptr) { @@ -1096,7 +1096,7 @@ hipError_t ihipArrayCreate(hipArray** array, return hipSuccess; } -hipError_t hipArrayCreate(hipArray** array, +hipError_t hipArrayCreate(hipArray_t* array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_API(hipArrayCreate, array, pAllocateArray); if (pAllocateArray == nullptr) { @@ -1114,7 +1114,7 @@ hipError_t hipArrayCreate(hipArray** array, } -hipError_t hipMallocArray(hipArray** array, +hipError_t hipMallocArray(hipArray_t* array, const hipChannelFormatDesc* desc, size_t width, size_t height, @@ -1136,7 +1136,7 @@ hipError_t hipMallocArray(hipArray** array, HIP_RETURN(ihipArrayCreate(array, &allocateArray, 0 /* numMipLevels */)); } -hipError_t hipArray3DCreate(hipArray** array, +hipError_t hipArray3DCreate(hipArray_t* array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray) { HIP_INIT_API(hipArray3DCreate, array, pAllocateArray); CHECK_STREAM_CAPTURE_SUPPORTED(); @@ -1522,7 +1522,7 @@ hipError_t hipMemcpyDtoHAsync(void* dstHost, hipDeviceptr_t srcDevice, size_t By ihipMemcpy(dstHost, srcDevice, ByteCount, kind, *hip_stream, true)); } -hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyAtoDValidate(hipArray_t srcArray, void* dstDevice, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t dstRowPitch, size_t dstSlicePitch, amd::Memory*& dstMemory, amd::Image*& srcImage, @@ -1564,7 +1564,7 @@ hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coor return hipSuccess; } -hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, void* dstDevice, +hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray_t srcArray, void* dstDevice, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch, hip::Stream* stream) { @@ -1595,7 +1595,7 @@ hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, voi return hipSuccess; } -hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray* dstArray, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage, amd::Memory*& srcMemory, amd::BufferRect& dstRect, @@ -1637,7 +1637,7 @@ hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray* dstArray, amd::Coor return hipSuccess; } -hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArray* dstArray, +hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArray_t dstArray, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch, hip::Stream* stream) { @@ -1906,7 +1906,7 @@ hipError_t ihipMemcpyHtoH(const void* srcHost, void* dstHost, amd::Coord3D srcOr return hipSuccess; } -hipError_t ihipMemcpyAtoAValidate(hipArray* srcArray, hipArray* dstArray, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyAtoAValidate(hipArray_t srcArray, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, amd::Image*& srcImage, amd::Image*& dstImage) { if (dstArray == nullptr || srcArray == nullptr) { @@ -1937,7 +1937,7 @@ hipError_t ihipMemcpyAtoAValidate(hipArray* srcArray, hipArray* dstArray, amd::C return hipSuccess; } -hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hipArray* dstArray, +hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray_t srcArray, hipArray_t dstArray, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, hip::Stream* stream) { amd::Image* srcImage; @@ -1965,7 +1965,7 @@ hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hip return hipSuccess; } -hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray* dstArray, +hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage, @@ -1995,7 +1995,7 @@ hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray* dstArray, return hipSuccess; } -hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hipArray* dstArray, +hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hipArray_t dstArray, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch, hip::Stream* stream, bool isAsync = false) { @@ -2026,7 +2026,7 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi return hipSuccess; } -hipError_t ihipMemcpyAtoHValidate(hipArray* srcArray, void* dstHost, amd::Coord3D& srcOrigin, +hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t dstRowPitch, size_t dstSlicePitch, amd::Image*& srcImage, amd::BufferRect& dstRect) { @@ -2056,7 +2056,7 @@ hipError_t ihipMemcpyAtoHValidate(hipArray* srcArray, void* dstHost, amd::Coord3 return hipSuccess; } -hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, void* dstHost, +hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray_t srcArray, void* dstHost, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch, hip::Stream* stream, bool isAsync = false) { @@ -2481,7 +2481,7 @@ hipError_t ihipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, c return ihipMemcpyParam2D(&desc, stream, isAsync); } -hipError_t hipMemcpy2DToArray_common(hipArray* dst, size_t wOffset, size_t hOffset, +hipError_t hipMemcpy2DToArray_common(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=nullptr, bool isAsync = false) { @@ -2499,20 +2499,20 @@ hipError_t hipMemcpy2DToArray_common(hipArray* dst, size_t wOffset, size_t hOffs return ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, isAsync); } -hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { +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) { HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind); CHECK_STREAM_CAPTURING(); HIP_RETURN_DURATION(hipMemcpy2DToArray_common(dst, wOffset, hOffset, src, spitch, width, height, kind)); } -hipError_t hipMemcpy2DToArray_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { +hipError_t hipMemcpy2DToArray_spt(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind); CHECK_STREAM_CAPTURING(); HIP_RETURN_DURATION(hipMemcpy2DToArray_common(dst, wOffset, hOffset, src, spitch, width, height, kind, getPerThreadDefaultStream())); } -hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { +hipError_t hipMemcpyToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t count, hipMemcpyKind kind) { HIP_INIT_API(hipMemcpyToArray, dst, wOffset, hOffset, src, count, kind); CHECK_STREAM_CAPTURING(); if (dst == nullptr) { @@ -2582,7 +2582,7 @@ hipError_t hipMemcpyFromArray_spt(void* dst, hipArray_const_t src, size_t wOffse getPerThreadDefaultStream())); } -hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin, +hipError_t ihipMemcpyAtoD(hipArray_t srcArray, void* dstDevice, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch, hipStream_t stream, bool isAsync = false) { amd::Command* command; @@ -2596,7 +2596,7 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcO if (status != hipSuccess) return status; return ihipMemcpyCmdEnqueue(command, isAsync); } -hipError_t ihipMemcpyDtoA(void* srcDevice, hipArray* dstArray, amd::Coord3D srcOrigin, +hipError_t ihipMemcpyDtoA(void* srcDevice, hipArray_t dstArray, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch, hipStream_t stream, bool isAsync = false) { amd::Command* command; @@ -2655,7 +2655,7 @@ hipError_t ihipMemcpyHtoD(const void* srcHost, void* dstDevice, amd::Coord3D src if (status != hipSuccess) return status; return ihipMemcpyCmdEnqueue(command, isAsync); } -hipError_t ihipMemcpyAtoA(hipArray* srcArray, hipArray* dstArray, amd::Coord3D srcOrigin, +hipError_t ihipMemcpyAtoA(hipArray_t srcArray, hipArray_t dstArray, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, hipStream_t stream, bool isAsync = false) { amd::Command* command; @@ -2668,7 +2668,7 @@ hipError_t ihipMemcpyAtoA(hipArray* srcArray, hipArray* dstArray, amd::Coord3D s if (status != hipSuccess) return status; return ihipMemcpyCmdEnqueue(command, isAsync); } -hipError_t ihipMemcpyHtoA(const void* srcHost, hipArray* dstArray, amd::Coord3D srcOrigin, +hipError_t ihipMemcpyHtoA(const void* srcHost, hipArray_t dstArray, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch, hipStream_t stream, bool isAsync = false) { amd::Command* command; @@ -2682,7 +2682,7 @@ hipError_t ihipMemcpyHtoA(const void* srcHost, hipArray* dstArray, amd::Coord3D if (status != hipSuccess) return status; return ihipMemcpyCmdEnqueue(command, isAsync); } -hipError_t ihipMemcpyAtoH(hipArray* srcArray, void* dstHost, amd::Coord3D srcOrigin, +hipError_t ihipMemcpyAtoH(hipArray_t srcArray, void* dstHost, amd::Coord3D srcOrigin, amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch, hipStream_t stream, bool isAsync = false) { amd::Command* command; @@ -2697,7 +2697,7 @@ hipError_t ihipMemcpyAtoH(hipArray* srcArray, void* dstHost, amd::Coord3D srcOri return ihipMemcpyCmdEnqueue(command, isAsync); } -hipError_t hipMemcpyHtoA(hipArray* dstArray, +hipError_t hipMemcpyHtoA(hipArray_t dstArray, size_t dstOffset, const void* srcHost, size_t ByteCount) { @@ -2707,7 +2707,7 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, } hipError_t hipMemcpyAtoH(void* dstHost, - hipArray* srcArray, + hipArray_t srcArray, size_t srcOffset, size_t ByteCount) { HIP_INIT_API(hipMemcpyAtoH, dstHost, srcArray, srcOffset, ByteCount); @@ -3674,14 +3674,14 @@ hipError_t hipDrvPointerGetAttributes(unsigned int numAttributes, hipPointer_att } // ================================================================================================ -hipError_t hipArrayDestroy(hipArray* array) { +hipError_t hipArrayDestroy(hipArray_t array) { HIP_INIT_API(hipArrayDestroy, array); CHECK_STREAM_CAPTURE_SUPPORTED(); HIP_RETURN(ihipArrayDestroy(array)); } hipError_t ihipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* desc, - hipArray* array) { + hipArray_t array) { { amd::ScopedLock lock(hip::hipArraySetLock); if (hip::hipArraySet.find(array) == hip::hipArraySet.end()) { @@ -3702,7 +3702,7 @@ hipError_t ihipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* desc, hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc, hipExtent* extent, unsigned int* flags, - hipArray* array) { + hipArray_t array) { HIP_INIT_API(hipArrayGetInfo, desc, extent, flags, array); CHECK_STREAM_CAPTURE_SUPPORTED(); @@ -3739,7 +3739,7 @@ hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc, } hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor, - hipArray* array) { + hipArray_t array) { HIP_INIT_API(hipArrayGetDescriptor, pArrayDescriptor, array); CHECK_STREAM_CAPTURE_SUPPORTED(); @@ -3766,7 +3766,7 @@ hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor, } hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor, - hipArray* array) { + hipArray_t array) { HIP_INIT_API(hipArray3DGetDescriptor, pArrayDescriptor, array); CHECK_STREAM_CAPTURE_SUPPORTED(); @@ -3875,14 +3875,14 @@ hipError_t hipMemcpy2DFromArrayAsync_spt(void* dst, size_t dpitch, hipArray_cons HIP_RETURN_DURATION(hipMemcpy2DFromArray_common(dst, dpitch, src, wOffsetSrc, hOffsetSrc, width, height, kind, stream, true)); } -hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { +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) { HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind, stream); STREAM_CAPTURE(hipMemcpy2DToArrayAsync, stream, dst, wOffset, hOffset, src, spitch, width, height, kind); HIP_RETURN_DURATION(hipMemcpy2DToArray_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true)); } -hipError_t hipMemcpy2DToArrayAsync_spt(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { +hipError_t hipMemcpy2DToArrayAsync_spt(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) { HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind, stream); PER_THREAD_DEFAULT_STREAM(stream); STREAM_CAPTURE(hipMemcpy2DToArrayAsync, stream, dst, wOffset, hOffset, src, spitch, width, height, @@ -3890,9 +3890,9 @@ hipError_t hipMemcpy2DToArrayAsync_spt(hipArray* dst, size_t wOffset, size_t hOf HIP_RETURN_DURATION(hipMemcpy2DToArray_common(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true)); } -hipError_t hipMemcpyAtoA(hipArray* dstArray, +hipError_t hipMemcpyAtoA(hipArray_t dstArray, size_t dstOffset, - hipArray* srcArray, + hipArray_t srcArray, size_t srcOffset, size_t ByteCount) { HIP_INIT_API(hipMemcpyAtoA, dstArray, dstOffset, srcArray, srcOffset, ByteCount); @@ -3901,7 +3901,7 @@ hipError_t hipMemcpyAtoA(hipArray* dstArray, } hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, - hipArray* srcArray, + hipArray_t srcArray, size_t srcOffset, size_t ByteCount) { HIP_INIT_API(hipMemcpyAtoD, dstDevice, srcArray, srcOffset, ByteCount); @@ -3910,7 +3910,7 @@ hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, } hipError_t hipMemcpyAtoHAsync(void* dstHost, - hipArray* srcArray, + hipArray_t srcArray, size_t srcOffset, size_t ByteCount, hipStream_t stream) { @@ -3919,7 +3919,7 @@ hipError_t hipMemcpyAtoHAsync(void* dstHost, HIP_RETURN_DURATION(ihipMemcpyAtoH(srcArray, dstHost, {srcOffset, 0, 0}, {0, 0, 0}, {ByteCount, 1, 1}, 0, 0, stream, true)); } -hipError_t hipMemcpyDtoA(hipArray* dstArray, +hipError_t hipMemcpyDtoA(hipArray_t dstArray, size_t dstOffset, hipDeviceptr_t srcDevice, size_t ByteCount) { @@ -3928,7 +3928,7 @@ hipError_t hipMemcpyDtoA(hipArray* dstArray, HIP_RETURN_DURATION(ihipMemcpyDtoA(srcDevice, dstArray, {0, 0, 0}, {dstOffset, 0, 0}, {ByteCount, 1, 1}, 0, 0, nullptr)); } -hipError_t hipMemcpyHtoAAsync(hipArray* dstArray, +hipError_t hipMemcpyHtoAAsync(hipArray_t dstArray, size_t dstOffset, const void* srcHost, size_t ByteCount,