diff --git a/hipamd/include/hip/amd_detail/hip_runtime_api.h b/hipamd/include/hip/amd_detail/hip_runtime_api.h index eb5890759c..8ddd7fd1e3 100644 --- a/hipamd/include/hip/amd_detail/hip_runtime_api.h +++ b/hipamd/include/hip/amd_detail/hip_runtime_api.h @@ -3508,6 +3508,14 @@ hipError_t hipLaunchKernel(const void* function_address, void** args, size_t sharedMemBytes __dparm(0), hipStream_t stream __dparm(0)); +/** + * Copies memory for 2D arrays. + * + * @param pCopy - Parameters for the memory copy + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy); //TODO: Move this to hip_ext.h hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks, diff --git a/hipamd/rocclr/hip_hcc.def.in b/hipamd/rocclr/hip_hcc.def.in index 65d15b3c6c..595b0ae226 100755 --- a/hipamd/rocclr/hip_hcc.def.in +++ b/hipamd/rocclr/hip_hcc.def.in @@ -265,6 +265,7 @@ hipExtStreamCreateWithCUMask hipStreamGetPriority hipMemcpy2DFromArray hipMemcpy2DFromArrayAsync +hipDrvMemcpy2DUnaligned hipMemcpyAtoH hipMemcpyHtoA hipMemcpyParam2DAsync diff --git a/hipamd/rocclr/hip_hcc.map.in b/hipamd/rocclr/hip_hcc.map.in index e0b4893e67..2a513d5ebe 100755 --- a/hipamd/rocclr/hip_hcc.map.in +++ b/hipamd/rocclr/hip_hcc.map.in @@ -96,6 +96,7 @@ global: hipMemcpy2D; hipMemcpy2DAsync; hipMemcpy2DToArray; + hipDrvMemcpy2DUnaligned; hipMemcpy3D; hipMemcpy3DAsync; hipDrvMemcpy3D; diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp index 190d7c6d99..3404c55622 100755 --- a/hipamd/rocclr/hip_memory.cpp +++ b/hipamd/rocclr/hip_memory.cpp @@ -2392,3 +2392,12 @@ hipError_t hipFreeHost(void *ptr) { HIP_RETURN(ihipFree(ptr)); } + +hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy) { + HIP_INIT_API(hipDrvMemcpy2DUnaligned, pCopy); + + HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*pCopy); + + HIP_RETURN(ihipMemcpyParam3D(&desc, nullptr)); +} +