From 90190185ee30c45f972c601149ddeafe983f16f2 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Tue, 20 Jun 2023 14:41:25 +0100 Subject: [PATCH 1/5] [EXP][CMD-BUFFER] Command-buffer additional function definitions for buffer read and write --- include/ur.py | 40 +++ include/ur_api.h | 224 +++++++++++++++ include/ur_ddi.h | 62 ++++ scripts/core/EXP-COMMAND-BUFFER.rst | 27 +- scripts/core/exp-command-buffer.yml | 186 ++++++++++++ scripts/core/registry.yml | 12 + source/adapters/null/ur_nullddi.cpp | 177 ++++++++++++ source/common/ur_params.hpp | 271 ++++++++++++++++++ source/loader/layers/tracing/ur_trcddi.cpp | 247 ++++++++++++++++ source/loader/layers/validation/ur_valddi.cpp | 255 ++++++++++++++++ source/loader/ur_ldrddi.cpp | 213 ++++++++++++++ source/loader/ur_libapi.cpp | 226 +++++++++++++++ source/ur_api.cpp | 186 ++++++++++++ 13 files changed, 2120 insertions(+), 6 deletions(-) diff --git a/include/ur.py b/include/ur.py index fcee6474e2..229cd870af 100644 --- a/include/ur.py +++ b/include/ur.py @@ -1981,6 +1981,10 @@ class ur_function_v(IntEnum): USM_P2P_ENABLE_PEER_ACCESS_EXP = 165 ## Enumerator for ::urUsmP2PEnablePeerAccessExp USM_P2P_DISABLE_PEER_ACCESS_EXP = 166 ## Enumerator for ::urUsmP2PDisablePeerAccessExp USM_P2P_PEER_ACCESS_GET_INFO_EXP = 167 ## Enumerator for ::urUsmP2PPeerAccessGetInfoExp + COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP = 168 ## Enumerator for ::urCommandBufferAppendMembufferWriteExp + COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP = 169 ## Enumerator for ::urCommandBufferAppendMembufferReadExp + COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP = 170## Enumerator for ::urCommandBufferAppendMembufferWriteRectExp + COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP = 171 ## Enumerator for ::urCommandBufferAppendMembufferReadRectExp class ur_function_t(c_int): def __str__(self): @@ -3289,6 +3293,20 @@ class ur_usm_exp_dditable_t(Structure): else: _urCommandBufferAppendMembufferCopyExp_t = CFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_mem_handle_t, c_size_t, c_size_t, c_size_t, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) +############################################################################### +## @brief Function-pointer for urCommandBufferAppendMembufferWriteExp +if __use_win_types: + _urCommandBufferAppendMembufferWriteExp_t = WINFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) +else: + _urCommandBufferAppendMembufferWriteExp_t = CFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) + +############################################################################### +## @brief Function-pointer for urCommandBufferAppendMembufferReadExp +if __use_win_types: + _urCommandBufferAppendMembufferReadExp_t = WINFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) +else: + _urCommandBufferAppendMembufferReadExp_t = CFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) + ############################################################################### ## @brief Function-pointer for urCommandBufferAppendMembufferCopyRectExp if __use_win_types: @@ -3296,6 +3314,20 @@ class ur_usm_exp_dditable_t(Structure): else: _urCommandBufferAppendMembufferCopyRectExp_t = CFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_mem_handle_t, ur_rect_offset_t, ur_rect_offset_t, ur_rect_region_t, c_size_t, c_size_t, c_size_t, c_size_t, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) +############################################################################### +## @brief Function-pointer for urCommandBufferAppendMembufferWriteRectExp +if __use_win_types: + _urCommandBufferAppendMembufferWriteRectExp_t = WINFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_rect_offset_t, ur_rect_offset_t, ur_rect_region_t, c_size_t, c_size_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) +else: + _urCommandBufferAppendMembufferWriteRectExp_t = CFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_rect_offset_t, ur_rect_offset_t, ur_rect_region_t, c_size_t, c_size_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) + +############################################################################### +## @brief Function-pointer for urCommandBufferAppendMembufferReadRectExp +if __use_win_types: + _urCommandBufferAppendMembufferReadRectExp_t = WINFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_rect_offset_t, ur_rect_offset_t, ur_rect_region_t, c_size_t, c_size_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) +else: + _urCommandBufferAppendMembufferReadRectExp_t = CFUNCTYPE( ur_result_t, ur_exp_command_buffer_handle_t, ur_mem_handle_t, ur_rect_offset_t, ur_rect_offset_t, ur_rect_region_t, c_size_t, c_size_t, c_size_t, c_size_t, c_void_p, c_ulong, POINTER(ur_exp_command_buffer_sync_point_t), POINTER(ur_exp_command_buffer_sync_point_t) ) + ############################################################################### ## @brief Function-pointer for urCommandBufferEnqueueExp if __use_win_types: @@ -3315,7 +3347,11 @@ class ur_command_buffer_exp_dditable_t(Structure): ("pfnAppendKernelLaunchExp", c_void_p), ## _urCommandBufferAppendKernelLaunchExp_t ("pfnAppendMemcpyUSMExp", c_void_p), ## _urCommandBufferAppendMemcpyUSMExp_t ("pfnAppendMembufferCopyExp", c_void_p), ## _urCommandBufferAppendMembufferCopyExp_t + ("pfnAppendMembufferWriteExp", c_void_p), ## _urCommandBufferAppendMembufferWriteExp_t + ("pfnAppendMembufferReadExp", c_void_p), ## _urCommandBufferAppendMembufferReadExp_t ("pfnAppendMembufferCopyRectExp", c_void_p), ## _urCommandBufferAppendMembufferCopyRectExp_t + ("pfnAppendMembufferWriteRectExp", c_void_p), ## _urCommandBufferAppendMembufferWriteRectExp_t + ("pfnAppendMembufferReadRectExp", c_void_p), ## _urCommandBufferAppendMembufferReadRectExp_t ("pfnEnqueueExp", c_void_p) ## _urCommandBufferEnqueueExp_t ] @@ -3817,7 +3853,11 @@ def __init__(self, version : ur_api_version_t): self.urCommandBufferAppendKernelLaunchExp = _urCommandBufferAppendKernelLaunchExp_t(self.__dditable.CommandBufferExp.pfnAppendKernelLaunchExp) self.urCommandBufferAppendMemcpyUSMExp = _urCommandBufferAppendMemcpyUSMExp_t(self.__dditable.CommandBufferExp.pfnAppendMemcpyUSMExp) self.urCommandBufferAppendMembufferCopyExp = _urCommandBufferAppendMembufferCopyExp_t(self.__dditable.CommandBufferExp.pfnAppendMembufferCopyExp) + self.urCommandBufferAppendMembufferWriteExp = _urCommandBufferAppendMembufferWriteExp_t(self.__dditable.CommandBufferExp.pfnAppendMembufferWriteExp) + self.urCommandBufferAppendMembufferReadExp = _urCommandBufferAppendMembufferReadExp_t(self.__dditable.CommandBufferExp.pfnAppendMembufferReadExp) self.urCommandBufferAppendMembufferCopyRectExp = _urCommandBufferAppendMembufferCopyRectExp_t(self.__dditable.CommandBufferExp.pfnAppendMembufferCopyRectExp) + self.urCommandBufferAppendMembufferWriteRectExp = _urCommandBufferAppendMembufferWriteRectExp_t(self.__dditable.CommandBufferExp.pfnAppendMembufferWriteRectExp) + self.urCommandBufferAppendMembufferReadRectExp = _urCommandBufferAppendMembufferReadRectExp_t(self.__dditable.CommandBufferExp.pfnAppendMembufferReadRectExp) self.urCommandBufferEnqueueExp = _urCommandBufferEnqueueExp_t(self.__dditable.CommandBufferExp.pfnEnqueueExp) # call driver to get function pointers diff --git a/include/ur_api.h b/include/ur_api.h index 309a309627..cb47830b47 100644 --- a/include/ur_api.h +++ b/include/ur_api.h @@ -5277,6 +5277,10 @@ typedef enum ur_function_t { UR_FUNCTION_USM_P2P_ENABLE_PEER_ACCESS_EXP = 165, ///< Enumerator for ::urUsmP2PEnablePeerAccessExp UR_FUNCTION_USM_P2P_DISABLE_PEER_ACCESS_EXP = 166, ///< Enumerator for ::urUsmP2PDisablePeerAccessExp UR_FUNCTION_USM_P2P_PEER_ACCESS_GET_INFO_EXP = 167, ///< Enumerator for ::urUsmP2PPeerAccessGetInfoExp + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP = 168, ///< Enumerator for ::urCommandBufferAppendMembufferWriteExp + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP = 169, ///< Enumerator for ::urCommandBufferAppendMembufferReadExp + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP = 170, ///< Enumerator for ::urCommandBufferAppendMembufferWriteRectExp + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP = 171, ///< Enumerator for ::urCommandBufferAppendMembufferReadRectExp /// @cond UR_FUNCTION_FORCE_UINT32 = 0x7fffffff /// @endcond @@ -7322,6 +7326,72 @@ urCommandBufferAppendMembufferCopyExp( ur_exp_command_buffer_sync_point_t *pSyncPoint ///< [out][optional] sync point associated with this command ); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a memory write command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pSrc` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL +urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void *pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t *pSyncPoint ///< [out][optional] sync point associated with this command +); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a memory read command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pDst` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL +urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t *pSyncPoint ///< [out][optional] sync point associated with this command +); + /////////////////////////////////////////////////////////////////////////////// /// @brief Append a rectangular memory copy command to a command-buffer object /// @@ -7359,6 +7429,87 @@ urCommandBufferAppendMembufferCopyRectExp( ur_exp_command_buffer_sync_point_t *pSyncPoint ///< [out][optional] sync point associated with this command ); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a rectangular memory write command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pSrc` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL +urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void *pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t *pSyncPoint ///< [out][optional] sync point associated with this command +); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a rectangular memory read command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pDst` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL +urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t *pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t *pSyncPoint ///< [out][optional] sync point associated with this command +); + /////////////////////////////////////////////////////////////////////////////// /// @brief Submit a command-buffer for execution on a queue. /// @@ -7494,6 +7645,7 @@ typedef enum ur_exp_peer_info_t { /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -7535,6 +7687,7 @@ urUsmP2PEnablePeerAccessExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -7561,6 +7714,7 @@ urUsmP2PDisablePeerAccessExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -9192,6 +9346,36 @@ typedef struct ur_command_buffer_append_membuffer_copy_exp_params_t { ur_exp_command_buffer_sync_point_t **ppSyncPoint; } ur_command_buffer_append_membuffer_copy_exp_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urCommandBufferAppendMembufferWriteExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_command_buffer_append_membuffer_write_exp_params_t { + ur_exp_command_buffer_handle_t *phCommandBuffer; + ur_mem_handle_t *phBuffer; + size_t *poffset; + size_t *psize; + const void **ppSrc; + uint32_t *pnumSyncPointsInWaitList; + const ur_exp_command_buffer_sync_point_t **ppSyncPointWaitList; + ur_exp_command_buffer_sync_point_t **ppSyncPoint; +} ur_command_buffer_append_membuffer_write_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urCommandBufferAppendMembufferReadExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_command_buffer_append_membuffer_read_exp_params_t { + ur_exp_command_buffer_handle_t *phCommandBuffer; + ur_mem_handle_t *phBuffer; + size_t *poffset; + size_t *psize; + void **ppDst; + uint32_t *pnumSyncPointsInWaitList; + const ur_exp_command_buffer_sync_point_t **ppSyncPointWaitList; + ur_exp_command_buffer_sync_point_t **ppSyncPoint; +} ur_command_buffer_append_membuffer_read_exp_params_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief Function parameters for urCommandBufferAppendMembufferCopyRectExp /// @details Each entry is a pointer to the parameter passed to the function; @@ -9212,6 +9396,46 @@ typedef struct ur_command_buffer_append_membuffer_copy_rect_exp_params_t { ur_exp_command_buffer_sync_point_t **ppSyncPoint; } ur_command_buffer_append_membuffer_copy_rect_exp_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urCommandBufferAppendMembufferWriteRectExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_command_buffer_append_membuffer_write_rect_exp_params_t { + ur_exp_command_buffer_handle_t *phCommandBuffer; + ur_mem_handle_t *phBuffer; + ur_rect_offset_t *pbufferOffset; + ur_rect_offset_t *phostOffset; + ur_rect_region_t *pregion; + size_t *pbufferRowPitch; + size_t *pbufferSlicePitch; + size_t *phostRowPitch; + size_t *phostSlicePitch; + void **ppSrc; + uint32_t *pnumSyncPointsInWaitList; + const ur_exp_command_buffer_sync_point_t **ppSyncPointWaitList; + ur_exp_command_buffer_sync_point_t **ppSyncPoint; +} ur_command_buffer_append_membuffer_write_rect_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urCommandBufferAppendMembufferReadRectExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_command_buffer_append_membuffer_read_rect_exp_params_t { + ur_exp_command_buffer_handle_t *phCommandBuffer; + ur_mem_handle_t *phBuffer; + ur_rect_offset_t *pbufferOffset; + ur_rect_offset_t *phostOffset; + ur_rect_region_t *pregion; + size_t *pbufferRowPitch; + size_t *pbufferSlicePitch; + size_t *phostRowPitch; + size_t *phostSlicePitch; + void **ppDst; + uint32_t *pnumSyncPointsInWaitList; + const ur_exp_command_buffer_sync_point_t **ppSyncPointWaitList; + ur_exp_command_buffer_sync_point_t **ppSyncPoint; +} ur_command_buffer_append_membuffer_read_rect_exp_params_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief Function parameters for urCommandBufferEnqueueExp /// @details Each entry is a pointer to the parameter passed to the function; diff --git a/include/ur_ddi.h b/include/ur_ddi.h index 795edc605b..6ba809041e 100644 --- a/include/ur_ddi.h +++ b/include/ur_ddi.h @@ -1670,6 +1670,30 @@ typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferCopyExp_t)( const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urCommandBufferAppendMembufferWriteExp +typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferWriteExp_t)( + ur_exp_command_buffer_handle_t, + ur_mem_handle_t, + size_t, + size_t, + const void *, + uint32_t, + const ur_exp_command_buffer_sync_point_t *, + ur_exp_command_buffer_sync_point_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urCommandBufferAppendMembufferReadExp +typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferReadExp_t)( + ur_exp_command_buffer_handle_t, + ur_mem_handle_t, + size_t, + size_t, + void *, + uint32_t, + const ur_exp_command_buffer_sync_point_t *, + ur_exp_command_buffer_sync_point_t *); + /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urCommandBufferAppendMembufferCopyRectExp typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferCopyRectExp_t)( @@ -1687,6 +1711,40 @@ typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferCopyRectExp_t) const ur_exp_command_buffer_sync_point_t *, ur_exp_command_buffer_sync_point_t *); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urCommandBufferAppendMembufferWriteRectExp +typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferWriteRectExp_t)( + ur_exp_command_buffer_handle_t, + ur_mem_handle_t, + ur_rect_offset_t, + ur_rect_offset_t, + ur_rect_region_t, + size_t, + size_t, + size_t, + size_t, + void *, + uint32_t, + const ur_exp_command_buffer_sync_point_t *, + ur_exp_command_buffer_sync_point_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urCommandBufferAppendMembufferReadRectExp +typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferAppendMembufferReadRectExp_t)( + ur_exp_command_buffer_handle_t, + ur_mem_handle_t, + ur_rect_offset_t, + ur_rect_offset_t, + ur_rect_region_t, + size_t, + size_t, + size_t, + size_t, + void *, + uint32_t, + const ur_exp_command_buffer_sync_point_t *, + ur_exp_command_buffer_sync_point_t *); + /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urCommandBufferEnqueueExp typedef ur_result_t(UR_APICALL *ur_pfnCommandBufferEnqueueExp_t)( @@ -1706,7 +1764,11 @@ typedef struct ur_command_buffer_exp_dditable_t { ur_pfnCommandBufferAppendKernelLaunchExp_t pfnAppendKernelLaunchExp; ur_pfnCommandBufferAppendMemcpyUSMExp_t pfnAppendMemcpyUSMExp; ur_pfnCommandBufferAppendMembufferCopyExp_t pfnAppendMembufferCopyExp; + ur_pfnCommandBufferAppendMembufferWriteExp_t pfnAppendMembufferWriteExp; + ur_pfnCommandBufferAppendMembufferReadExp_t pfnAppendMembufferReadExp; ur_pfnCommandBufferAppendMembufferCopyRectExp_t pfnAppendMembufferCopyRectExp; + ur_pfnCommandBufferAppendMembufferWriteRectExp_t pfnAppendMembufferWriteRectExp; + ur_pfnCommandBufferAppendMembufferReadRectExp_t pfnAppendMembufferReadRectExp; ur_pfnCommandBufferEnqueueExp_t pfnEnqueueExp; } ur_command_buffer_exp_dditable_t; diff --git a/scripts/core/EXP-COMMAND-BUFFER.rst b/scripts/core/EXP-COMMAND-BUFFER.rst index 2f61388339..46aa160630 100644 --- a/scripts/core/EXP-COMMAND-BUFFER.rst +++ b/scripts/core/EXP-COMMAND-BUFFER.rst @@ -94,7 +94,11 @@ Currently only the following commands are supported: * ${x}CommandBufferAppendMemcpyUSMExp * ${x}CommandBufferAppendMembufferCopyExp * ${x}CommandBufferAppendMembufferCopyRectExp - +* ${x}CommandBufferAppendMembufferReadExp +* ${x}CommandBufferAppendMembufferReadRectExp +* ${x}CommandBufferAppendMembufferWriteExp +* ${x}CommandBufferAppendMembufferWriteRectExp + It is planned to eventually support any command type from the Core API which can actually be appended to the equiavalent adapter native constructs. @@ -161,6 +165,11 @@ Enums * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMCPY_USM_EXP * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_COPY_EXP * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_COPY_RECT_EXP + * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP + * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP + * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP + * ${X}_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP + Types @@ -180,19 +189,25 @@ Functions * ${x}CommandBufferAppendMemcpyUSMExp * ${x}CommandBufferAppendMembufferCopyExp * ${x}CommandBufferAppendMembufferCopyRectExp +* ${x}CommandBufferAppendMembufferReadExp +* ${x}CommandBufferAppendMembufferReadRectExp +* ${x}CommandBufferAppendMembufferWriteExp +* ${x}CommandBufferAppendMembufferWriteRectExp * ${x}CommandBufferEnqueueExp Changelog -------------------------------------------------------------------------------- -+-----------+------------------------+ -| Revision | Changes | -+===========+========================+ -| 1.0 | Initial Draft | -+-----------+------------------------+ ++-----------+-------------------------------------------------------+ +| Revision | Changes | ++===========+=======================================================+ +| 1.0 | Initial Draft | +| 1.1 | add function definitions for buffer read and write | ++-----------+-------------------------------------------------------+ Contributors -------------------------------------------------------------------------------- * Ben Tracy `ben.tracy@codeplay.com `_ * Ewan Crawford `ewan@codeplay.com `_ +* Maxime France-Pillois `maxime.francepillois@codeplay.com `_ diff --git a/scripts/core/exp-command-buffer.yml b/scripts/core/exp-command-buffer.yml index 0ad073bfcc..e8c5417831 100644 --- a/scripts/core/exp-command-buffer.yml +++ b/scripts/core/exp-command-buffer.yml @@ -259,6 +259,84 @@ returns: - $X_RESULT_ERROR_OUT_OF_RESOURCES --- #-------------------------------------------------------------------------- type: function +desc: "Append a memory write command to a command-buffer object" +class: $xCommandBuffer +name: AppendMembufferWriteExp +params: + - type: $x_exp_command_buffer_handle_t + name: hCommandBuffer + desc: "[in] handle of the command-buffer object." + - type: $x_mem_handle_t + name: hBuffer + desc: "[in] handle of the buffer object." + - type: "size_t" + name: offset + desc: "[in] offset in bytes in the buffer object." + - type: "size_t" + name: size + desc: "[in] size in bytes of data being written." + - type: "const void*" + name: pSrc + desc: "[in] pointer to host memory where data is to be written from." + - type: uint32_t + name: numSyncPointsInWaitList + desc: "[in] The number of sync points in the provided dependency list." + - type: "const $x_exp_command_buffer_sync_point_t*" + name: pSyncPointWaitList + desc: "[in][optional] A list of sync points that this command depends on." + - type: "$x_exp_command_buffer_sync_point_t*" + name: pSyncPoint + desc: "[out][optional] sync point associated with this command" +returns: + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP: + - "`pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0`" + - "`pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0`" + - $X_RESULT_ERROR_INVALID_MEM_OBJECT + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Append a memory read command to a command-buffer object" +class: $xCommandBuffer +name: AppendMembufferReadExp +params: + - type: $x_exp_command_buffer_handle_t + name: hCommandBuffer + desc: "[in] handle of the command-buffer object." + - type: $x_mem_handle_t + name: hBuffer + desc: "[in] handle of the buffer object." + - type: "size_t" + name: offset + desc: "[in] offset in bytes in the buffer object." + - type: "size_t" + name: size + desc: "[in] size in bytes of data being written." + - type: "void*" + name: pDst + desc: "[in] pointer to host memory where data is to be written to." + - type: uint32_t + name: numSyncPointsInWaitList + desc: "[in] The number of sync points in the provided dependency list." + - type: "const $x_exp_command_buffer_sync_point_t*" + name: pSyncPointWaitList + desc: "[in][optional] A list of sync points that this command depends on." + - type: "$x_exp_command_buffer_sync_point_t*" + name: pSyncPoint + desc: "[out][optional] sync point associated with this command" +returns: + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP: + - "`pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0`" + - "`pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0`" + - $X_RESULT_ERROR_INVALID_MEM_OBJECT + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function desc: "Append a rectangular memory copy command to a command-buffer object" class: $xCommandBuffer name: AppendMembufferCopyRectExp @@ -313,6 +391,114 @@ returns: - $X_RESULT_ERROR_OUT_OF_RESOURCES --- #-------------------------------------------------------------------------- type: function +desc: "Append a rectangular memory write command to a command-buffer object" +class: $xCommandBuffer +name: AppendMembufferWriteRectExp +params: + - type: $x_exp_command_buffer_handle_t + name: hCommandBuffer + desc: "[in] handle of the command-buffer object." + - type: $x_mem_handle_t + name: hBuffer + desc: "[in] handle of the buffer object." + - type: $x_rect_offset_t + name: bufferOffset + desc: "[in] 3D offset in the buffer." + - type: $x_rect_offset_t + name: hostOffset + desc: "[in] 3D offset in the host region." + - type: $x_rect_region_t + name: region + desc: "[in] 3D rectangular region descriptor: width, height, depth." + - type: "size_t" + name: bufferRowPitch + desc: "[in] length of each row in bytes in the buffer object." + - type: "size_t" + name: bufferSlicePitch + desc: "[in] length of each 2D slice in bytes in the buffer object being written." + - type: "size_t" + name: hostRowPitch + desc: "[in] length of each row in bytes in the host memory region pointed to by pSrc." + - type: "size_t" + name: hostSlicePitch + desc: "[in] length of each 2D slice in bytes in the host memory region pointed to by pSrc." + - type: "void*" + name: pSrc + desc: "[in] pointer to host memory where data is to be written from." + - type: uint32_t + name: numSyncPointsInWaitList + desc: "[in] The number of sync points in the provided dependency list." + - type: "const $x_exp_command_buffer_sync_point_t*" + name: pSyncPointWaitList + desc: "[in][optional] A list of sync points that this command depends on." + - type: $x_exp_command_buffer_sync_point_t* + name: pSyncPoint + desc: "[out][optional] sync point associated with this command" +returns: + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP: + - "`pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0`" + - "`pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0`" + - $X_RESULT_ERROR_INVALID_MEM_OBJECT + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Append a rectangular memory read command to a command-buffer object" +class: $xCommandBuffer +name: AppendMembufferReadRectExp +params: + - type: $x_exp_command_buffer_handle_t + name: hCommandBuffer + desc: "[in] handle of the command-buffer object." + - type: $x_mem_handle_t + name: hBuffer + desc: "[in] handle of the buffer object." + - type: $x_rect_offset_t + name: bufferOffset + desc: "[in] 3D offset in the buffer." + - type: $x_rect_offset_t + name: hostOffset + desc: "[in] 3D offset in the host region." + - type: $x_rect_region_t + name: region + desc: "[in] 3D rectangular region descriptor: width, height, depth." + - type: "size_t" + name: bufferRowPitch + desc: "[in] length of each row in bytes in the buffer object." + - type: "size_t" + name: bufferSlicePitch + desc: "[in] length of each 2D slice in bytes in the buffer object being read." + - type: "size_t" + name: hostRowPitch + desc: "[in] length of each row in bytes in the host memory region pointed to by pDst." + - type: "size_t" + name: hostSlicePitch + desc: "[in] length of each 2D slice in bytes in the host memory region pointed to by pDst." + - type: "void*" + name: pDst + desc: "[in] pointer to host memory where data is to be read into." + - type: uint32_t + name: numSyncPointsInWaitList + desc: "[in] The number of sync points in the provided dependency list." + - type: "const $x_exp_command_buffer_sync_point_t*" + name: pSyncPointWaitList + desc: "[in][optional] A list of sync points that this command depends on." + - type: $x_exp_command_buffer_sync_point_t* + name: pSyncPoint + desc: "[out][optional] sync point associated with this command" +returns: + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP + - $X_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP: + - "`pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0`" + - "`pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0`" + - $X_RESULT_ERROR_INVALID_MEM_OBJECT + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function desc: "Submit a command-buffer for execution on a queue." class: $xCommandBuffer name: EnqueueExp diff --git a/scripts/core/registry.yml b/scripts/core/registry.yml index 0aa2703187..e29203afdc 100644 --- a/scripts/core/registry.yml +++ b/scripts/core/registry.yml @@ -490,3 +490,15 @@ etors: - name: USM_P2P_PEER_ACCESS_GET_INFO_EXP desc: Enumerator for $xUsmP2PPeerAccessGetInfoExp value: '167' +- name: COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP + desc: Enumerator for $xCommandBufferAppendMembufferWriteExp + value: '168' +- name: COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP + desc: Enumerator for $xCommandBufferAppendMembufferReadExp + value: '169' +- name: COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP + desc: Enumerator for $xCommandBufferAppendMembufferWriteRectExp + value: '170' +- name: COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP + desc: Enumerator for $xCommandBufferAppendMembufferReadRectExp + value: '171' diff --git a/source/adapters/null/ur_nullddi.cpp b/source/adapters/null/ur_nullddi.cpp index a029e4ff1e..9b3dddd191 100644 --- a/source/adapters/null/ur_nullddi.cpp +++ b/source/adapters/null/ur_nullddi.cpp @@ -4493,6 +4493,75 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + ur_result_t result = UR_RESULT_SUCCESS; + + // if the driver has created a custom function, then call it instead of using the generic path + auto pfnAppendMembufferWriteExp = + d_context.urDdiTable.CommandBufferExp.pfnAppendMembufferWriteExp; + if (nullptr != pfnAppendMembufferWriteExp) { + result = pfnAppendMembufferWriteExp(hCommandBuffer, hBuffer, offset, + size, pSrc, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + } else { + // generic implementation + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + ur_result_t result = UR_RESULT_SUCCESS; + + // if the driver has created a custom function, then call it instead of using the generic path + auto pfnAppendMembufferReadExp = + d_context.urDdiTable.CommandBufferExp.pfnAppendMembufferReadExp; + if (nullptr != pfnAppendMembufferReadExp) { + result = pfnAppendMembufferReadExp(hCommandBuffer, hBuffer, offset, + size, pDst, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + } else { + // generic implementation + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferAppendMembufferCopyRectExp __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( @@ -4536,6 +4605,102 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + ur_result_t result = UR_RESULT_SUCCESS; + + // if the driver has created a custom function, then call it instead of using the generic path + auto pfnAppendMembufferWriteRectExp = + d_context.urDdiTable.CommandBufferExp.pfnAppendMembufferWriteRectExp; + if (nullptr != pfnAppendMembufferWriteRectExp) { + result = pfnAppendMembufferWriteRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, + pSrc, numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + } else { + // generic implementation + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + ur_result_t result = UR_RESULT_SUCCESS; + + // if the driver has created a custom function, then call it instead of using the generic path + auto pfnAppendMembufferReadRectExp = + d_context.urDdiTable.CommandBufferExp.pfnAppendMembufferReadRectExp; + if (nullptr != pfnAppendMembufferReadRectExp) { + result = pfnAppendMembufferReadRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, + pDst, numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + } else { + // generic implementation + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferEnqueueExp __urdlllocal ur_result_t UR_APICALL urCommandBufferEnqueueExp( @@ -4845,9 +5010,21 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetCommandBufferExpProcAddrTable( pDdiTable->pfnAppendMembufferCopyExp = driver::urCommandBufferAppendMembufferCopyExp; + pDdiTable->pfnAppendMembufferWriteExp = + driver::urCommandBufferAppendMembufferWriteExp; + + pDdiTable->pfnAppendMembufferReadExp = + driver::urCommandBufferAppendMembufferReadExp; + pDdiTable->pfnAppendMembufferCopyRectExp = driver::urCommandBufferAppendMembufferCopyRectExp; + pDdiTable->pfnAppendMembufferWriteRectExp = + driver::urCommandBufferAppendMembufferWriteRectExp; + + pDdiTable->pfnAppendMembufferReadRectExp = + driver::urCommandBufferAppendMembufferReadRectExp; + pDdiTable->pfnEnqueueExp = driver::urCommandBufferEnqueueExp; return result; diff --git a/source/common/ur_params.hpp b/source/common/ur_params.hpp index 27ecf34a62..ffb78e36e3 100644 --- a/source/common/ur_params.hpp +++ b/source/common/ur_params.hpp @@ -9072,6 +9072,22 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_function_t value) { case UR_FUNCTION_USM_P2P_PEER_ACCESS_GET_INFO_EXP: os << "UR_FUNCTION_USM_P2P_PEER_ACCESS_GET_INFO_EXP"; break; + + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP: + os << "UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP"; + break; + + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP: + os << "UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP"; + break; + + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP: + os << "UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP"; + break; + + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP: + os << "UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP"; + break; default: os << "unknown enumerator"; break; @@ -10067,6 +10083,99 @@ inline std::ostream &operator<<( return os; } +inline std::ostream & +operator<<(std::ostream &os, + const struct ur_command_buffer_append_membuffer_write_exp_params_t + *params) { + + os << ".hCommandBuffer = "; + + ur_params::serializePtr(os, *(params->phCommandBuffer)); + + os << ", "; + os << ".hBuffer = "; + + ur_params::serializePtr(os, *(params->phBuffer)); + + os << ", "; + os << ".offset = "; + + os << *(params->poffset); + + os << ", "; + os << ".size = "; + + os << *(params->psize); + + os << ", "; + os << ".pSrc = "; + + ur_params::serializePtr(os, *(params->ppSrc)); + + os << ", "; + os << ".numSyncPointsInWaitList = "; + + os << *(params->pnumSyncPointsInWaitList); + + os << ", "; + os << ".pSyncPointWaitList = "; + + ur_params::serializePtr(os, *(params->ppSyncPointWaitList)); + + os << ", "; + os << ".pSyncPoint = "; + + ur_params::serializePtr(os, *(params->ppSyncPoint)); + + return os; +} + +inline std::ostream &operator<<( + std::ostream &os, + const struct ur_command_buffer_append_membuffer_read_exp_params_t *params) { + + os << ".hCommandBuffer = "; + + ur_params::serializePtr(os, *(params->phCommandBuffer)); + + os << ", "; + os << ".hBuffer = "; + + ur_params::serializePtr(os, *(params->phBuffer)); + + os << ", "; + os << ".offset = "; + + os << *(params->poffset); + + os << ", "; + os << ".size = "; + + os << *(params->psize); + + os << ", "; + os << ".pDst = "; + + ur_params::serializePtr(os, *(params->ppDst)); + + os << ", "; + os << ".numSyncPointsInWaitList = "; + + os << *(params->pnumSyncPointsInWaitList); + + os << ", "; + os << ".pSyncPointWaitList = "; + + ur_params::serializePtr(os, *(params->ppSyncPointWaitList)); + + os << ", "; + os << ".pSyncPoint = "; + + ur_params::serializePtr(os, *(params->ppSyncPoint)); + + return os; +} + inline std::ostream &operator<<( std::ostream &os, const struct ur_command_buffer_append_membuffer_copy_rect_exp_params_t @@ -10139,6 +10248,150 @@ inline std::ostream &operator<<( return os; } +inline std::ostream &operator<<( + std::ostream &os, + const struct ur_command_buffer_append_membuffer_write_rect_exp_params_t + *params) { + + os << ".hCommandBuffer = "; + + ur_params::serializePtr(os, *(params->phCommandBuffer)); + + os << ", "; + os << ".hBuffer = "; + + ur_params::serializePtr(os, *(params->phBuffer)); + + os << ", "; + os << ".bufferOffset = "; + + os << *(params->pbufferOffset); + + os << ", "; + os << ".hostOffset = "; + + os << *(params->phostOffset); + + os << ", "; + os << ".region = "; + + os << *(params->pregion); + + os << ", "; + os << ".bufferRowPitch = "; + + os << *(params->pbufferRowPitch); + + os << ", "; + os << ".bufferSlicePitch = "; + + os << *(params->pbufferSlicePitch); + + os << ", "; + os << ".hostRowPitch = "; + + os << *(params->phostRowPitch); + + os << ", "; + os << ".hostSlicePitch = "; + + os << *(params->phostSlicePitch); + + os << ", "; + os << ".pSrc = "; + + ur_params::serializePtr(os, *(params->ppSrc)); + + os << ", "; + os << ".numSyncPointsInWaitList = "; + + os << *(params->pnumSyncPointsInWaitList); + + os << ", "; + os << ".pSyncPointWaitList = "; + + ur_params::serializePtr(os, *(params->ppSyncPointWaitList)); + + os << ", "; + os << ".pSyncPoint = "; + + ur_params::serializePtr(os, *(params->ppSyncPoint)); + + return os; +} + +inline std::ostream &operator<<( + std::ostream &os, + const struct ur_command_buffer_append_membuffer_read_rect_exp_params_t + *params) { + + os << ".hCommandBuffer = "; + + ur_params::serializePtr(os, *(params->phCommandBuffer)); + + os << ", "; + os << ".hBuffer = "; + + ur_params::serializePtr(os, *(params->phBuffer)); + + os << ", "; + os << ".bufferOffset = "; + + os << *(params->pbufferOffset); + + os << ", "; + os << ".hostOffset = "; + + os << *(params->phostOffset); + + os << ", "; + os << ".region = "; + + os << *(params->pregion); + + os << ", "; + os << ".bufferRowPitch = "; + + os << *(params->pbufferRowPitch); + + os << ", "; + os << ".bufferSlicePitch = "; + + os << *(params->pbufferSlicePitch); + + os << ", "; + os << ".hostRowPitch = "; + + os << *(params->phostRowPitch); + + os << ", "; + os << ".hostSlicePitch = "; + + os << *(params->phostSlicePitch); + + os << ", "; + os << ".pDst = "; + + ur_params::serializePtr(os, *(params->ppDst)); + + os << ", "; + os << ".numSyncPointsInWaitList = "; + + os << *(params->pnumSyncPointsInWaitList); + + os << ", "; + os << ".pSyncPointWaitList = "; + + ur_params::serializePtr(os, *(params->ppSyncPointWaitList)); + + os << ", "; + os << ".pSyncPoint = "; + + ur_params::serializePtr(os, *(params->ppSyncPoint)); + + return os; +} + inline std::ostream & operator<<(std::ostream &os, const struct ur_command_buffer_enqueue_exp_params_t *params) { @@ -14321,11 +14574,29 @@ inline int serializeFunctionParams(std::ostream &os, uint32_t function, os << (const struct ur_command_buffer_append_membuffer_copy_exp_params_t *)params; } break; + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP: { + os << (const struct + ur_command_buffer_append_membuffer_write_exp_params_t *)params; + } break; + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP: { + os << (const struct ur_command_buffer_append_membuffer_read_exp_params_t + *)params; + } break; case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_COPY_RECT_EXP: { os << (const struct ur_command_buffer_append_membuffer_copy_rect_exp_params_t *) params; } break; + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP: { + os << (const struct + ur_command_buffer_append_membuffer_write_rect_exp_params_t *) + params; + } break; + case UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP: { + os << (const struct + ur_command_buffer_append_membuffer_read_rect_exp_params_t *) + params; + } break; case UR_FUNCTION_COMMAND_BUFFER_ENQUEUE_EXP: { os << (const struct ur_command_buffer_enqueue_exp_params_t *)params; } break; diff --git a/source/loader/layers/tracing/ur_trcddi.cpp b/source/loader/layers/tracing/ur_trcddi.cpp index a7d43afe5e..4f380b4f50 100644 --- a/source/loader/layers/tracing/ur_trcddi.cpp +++ b/source/loader/layers/tracing/ur_trcddi.cpp @@ -5140,6 +5140,101 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferWriteExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferWriteExp; + + if (nullptr == pfnAppendMembufferWriteExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + ur_command_buffer_append_membuffer_write_exp_params_t params = { + &hCommandBuffer, + &hBuffer, + &offset, + &size, + &pSrc, + &numSyncPointsInWaitList, + &pSyncPointWaitList, + &pSyncPoint}; + uint64_t instance = context.notify_begin( + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP, + "urCommandBufferAppendMembufferWriteExp", ¶ms); + + ur_result_t result = pfnAppendMembufferWriteExp( + hCommandBuffer, hBuffer, offset, size, pSrc, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + + context.notify_end(UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_EXP, + "urCommandBufferAppendMembufferWriteExp", ¶ms, + &result, instance); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferReadExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferReadExp; + + if (nullptr == pfnAppendMembufferReadExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + ur_command_buffer_append_membuffer_read_exp_params_t params = { + &hCommandBuffer, + &hBuffer, + &offset, + &size, + &pDst, + &numSyncPointsInWaitList, + &pSyncPointWaitList, + &pSyncPoint}; + uint64_t instance = context.notify_begin( + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP, + "urCommandBufferAppendMembufferReadExp", ¶ms); + + ur_result_t result = pfnAppendMembufferReadExp( + hCommandBuffer, hBuffer, offset, size, pDst, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + + context.notify_end(UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_EXP, + "urCommandBufferAppendMembufferReadExp", ¶ms, + &result, instance); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferAppendMembufferCopyRectExp __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( @@ -5202,6 +5297,140 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferWriteRectExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferWriteRectExp; + + if (nullptr == pfnAppendMembufferWriteRectExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + ur_command_buffer_append_membuffer_write_rect_exp_params_t params = { + &hCommandBuffer, + &hBuffer, + &bufferOffset, + &hostOffset, + ®ion, + &bufferRowPitch, + &bufferSlicePitch, + &hostRowPitch, + &hostSlicePitch, + &pSrc, + &numSyncPointsInWaitList, + &pSyncPointWaitList, + &pSyncPoint}; + uint64_t instance = context.notify_begin( + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP, + "urCommandBufferAppendMembufferWriteRectExp", ¶ms); + + ur_result_t result = pfnAppendMembufferWriteRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pSrc, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + + context.notify_end( + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_WRITE_RECT_EXP, + "urCommandBufferAppendMembufferWriteRectExp", ¶ms, &result, + instance); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferReadRectExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferReadRectExp; + + if (nullptr == pfnAppendMembufferReadRectExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + ur_command_buffer_append_membuffer_read_rect_exp_params_t params = { + &hCommandBuffer, + &hBuffer, + &bufferOffset, + &hostOffset, + ®ion, + &bufferRowPitch, + &bufferSlicePitch, + &hostRowPitch, + &hostSlicePitch, + &pDst, + &numSyncPointsInWaitList, + &pSyncPointWaitList, + &pSyncPoint}; + uint64_t instance = context.notify_begin( + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP, + "urCommandBufferAppendMembufferReadRectExp", ¶ms); + + ur_result_t result = pfnAppendMembufferReadRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pDst, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + + context.notify_end( + UR_FUNCTION_COMMAND_BUFFER_APPEND_MEMBUFFER_READ_RECT_EXP, + "urCommandBufferAppendMembufferReadRectExp", ¶ms, &result, + instance); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferEnqueueExp __urdlllocal ur_result_t UR_APICALL urCommandBufferEnqueueExp( @@ -5581,11 +5810,29 @@ __urdlllocal ur_result_t UR_APICALL urGetCommandBufferExpProcAddrTable( pDdiTable->pfnAppendMembufferCopyExp = ur_tracing_layer::urCommandBufferAppendMembufferCopyExp; + dditable.pfnAppendMembufferWriteExp = pDdiTable->pfnAppendMembufferWriteExp; + pDdiTable->pfnAppendMembufferWriteExp = + ur_tracing_layer::urCommandBufferAppendMembufferWriteExp; + + dditable.pfnAppendMembufferReadExp = pDdiTable->pfnAppendMembufferReadExp; + pDdiTable->pfnAppendMembufferReadExp = + ur_tracing_layer::urCommandBufferAppendMembufferReadExp; + dditable.pfnAppendMembufferCopyRectExp = pDdiTable->pfnAppendMembufferCopyRectExp; pDdiTable->pfnAppendMembufferCopyRectExp = ur_tracing_layer::urCommandBufferAppendMembufferCopyRectExp; + dditable.pfnAppendMembufferWriteRectExp = + pDdiTable->pfnAppendMembufferWriteRectExp; + pDdiTable->pfnAppendMembufferWriteRectExp = + ur_tracing_layer::urCommandBufferAppendMembufferWriteRectExp; + + dditable.pfnAppendMembufferReadRectExp = + pDdiTable->pfnAppendMembufferReadRectExp; + pDdiTable->pfnAppendMembufferReadRectExp = + ur_tracing_layer::urCommandBufferAppendMembufferReadRectExp; + dditable.pfnEnqueueExp = pDdiTable->pfnEnqueueExp; pDdiTable->pfnEnqueueExp = ur_tracing_layer::urCommandBufferEnqueueExp; diff --git a/source/loader/layers/validation/ur_valddi.cpp b/source/loader/layers/validation/ur_valddi.cpp index 6084954b58..1f9ecdf443 100644 --- a/source/loader/layers/validation/ur_valddi.cpp +++ b/source/loader/layers/validation/ur_valddi.cpp @@ -6401,6 +6401,111 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferWriteExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferWriteExp; + + if (nullptr == pfnAppendMembufferWriteExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + if (context.enableParameterValidation) { + if (NULL == hCommandBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == hBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == pSrc) { + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + } + + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + + if (pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + } + + ur_result_t result = pfnAppendMembufferWriteExp( + hCommandBuffer, hBuffer, offset, size, pSrc, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferReadExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferReadExp; + + if (nullptr == pfnAppendMembufferReadExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + if (context.enableParameterValidation) { + if (NULL == hCommandBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == hBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == pDst) { + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + } + + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + + if (pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + } + + ur_result_t result = pfnAppendMembufferReadExp( + hCommandBuffer, hBuffer, offset, size, pDst, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferAppendMembufferCopyRectExp __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( @@ -6462,6 +6567,138 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferWriteRectExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferWriteRectExp; + + if (nullptr == pfnAppendMembufferWriteRectExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + if (context.enableParameterValidation) { + if (NULL == hCommandBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == hBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == pSrc) { + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + } + + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + + if (pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + } + + ur_result_t result = pfnAppendMembufferWriteRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pSrc, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + auto pfnAppendMembufferReadRectExp = + context.urDdiTable.CommandBufferExp.pfnAppendMembufferReadRectExp; + + if (nullptr == pfnAppendMembufferReadRectExp) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + + if (context.enableParameterValidation) { + if (NULL == hCommandBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == hBuffer) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (NULL == pDst) { + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + } + + if (pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + + if (pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0) { + return UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP; + } + } + + ur_result_t result = pfnAppendMembufferReadRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pDst, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferEnqueueExp __urdlllocal ur_result_t UR_APICALL urCommandBufferEnqueueExp( @@ -6878,11 +7115,29 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetCommandBufferExpProcAddrTable( pDdiTable->pfnAppendMembufferCopyExp = ur_validation_layer::urCommandBufferAppendMembufferCopyExp; + dditable.pfnAppendMembufferWriteExp = pDdiTable->pfnAppendMembufferWriteExp; + pDdiTable->pfnAppendMembufferWriteExp = + ur_validation_layer::urCommandBufferAppendMembufferWriteExp; + + dditable.pfnAppendMembufferReadExp = pDdiTable->pfnAppendMembufferReadExp; + pDdiTable->pfnAppendMembufferReadExp = + ur_validation_layer::urCommandBufferAppendMembufferReadExp; + dditable.pfnAppendMembufferCopyRectExp = pDdiTable->pfnAppendMembufferCopyRectExp; pDdiTable->pfnAppendMembufferCopyRectExp = ur_validation_layer::urCommandBufferAppendMembufferCopyRectExp; + dditable.pfnAppendMembufferWriteRectExp = + pDdiTable->pfnAppendMembufferWriteRectExp; + pDdiTable->pfnAppendMembufferWriteRectExp = + ur_validation_layer::urCommandBufferAppendMembufferWriteRectExp; + + dditable.pfnAppendMembufferReadRectExp = + pDdiTable->pfnAppendMembufferReadRectExp; + pDdiTable->pfnAppendMembufferReadRectExp = + ur_validation_layer::urCommandBufferAppendMembufferReadRectExp; + dditable.pfnEnqueueExp = pDdiTable->pfnEnqueueExp; pDdiTable->pfnEnqueueExp = ur_validation_layer::urCommandBufferEnqueueExp; diff --git a/source/loader/ur_ldrddi.cpp b/source/loader/ur_ldrddi.cpp index 144d812382..3cb1885039 100644 --- a/source/loader/ur_ldrddi.cpp +++ b/source/loader/ur_ldrddi.cpp @@ -6249,6 +6249,95 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + + // extract platform's function pointer table + auto dditable = + reinterpret_cast(hCommandBuffer) + ->dditable; + auto pfnAppendMembufferWriteExp = + dditable->ur.CommandBufferExp.pfnAppendMembufferWriteExp; + if (nullptr == pfnAppendMembufferWriteExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + // convert loader handle to platform handle + hCommandBuffer = + reinterpret_cast(hCommandBuffer) + ->handle; + + // convert loader handle to platform handle + hBuffer = reinterpret_cast(hBuffer)->handle; + + // forward to device-platform + result = pfnAppendMembufferWriteExp(hCommandBuffer, hBuffer, offset, size, + pSrc, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + + // extract platform's function pointer table + auto dditable = + reinterpret_cast(hCommandBuffer) + ->dditable; + auto pfnAppendMembufferReadExp = + dditable->ur.CommandBufferExp.pfnAppendMembufferReadExp; + if (nullptr == pfnAppendMembufferReadExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + // convert loader handle to platform handle + hCommandBuffer = + reinterpret_cast(hCommandBuffer) + ->handle; + + // convert loader handle to platform handle + hBuffer = reinterpret_cast(hBuffer)->handle; + + // forward to device-platform + result = pfnAppendMembufferReadExp(hCommandBuffer, hBuffer, offset, size, + pDst, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferAppendMembufferCopyRectExp __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( @@ -6305,6 +6394,122 @@ __urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferWriteRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + + // extract platform's function pointer table + auto dditable = + reinterpret_cast(hCommandBuffer) + ->dditable; + auto pfnAppendMembufferWriteRectExp = + dditable->ur.CommandBufferExp.pfnAppendMembufferWriteRectExp; + if (nullptr == pfnAppendMembufferWriteRectExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + // convert loader handle to platform handle + hCommandBuffer = + reinterpret_cast(hCommandBuffer) + ->handle; + + // convert loader handle to platform handle + hBuffer = reinterpret_cast(hBuffer)->handle; + + // forward to device-platform + result = pfnAppendMembufferWriteRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pSrc, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urCommandBufferAppendMembufferReadRectExp +__urdlllocal ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + + // extract platform's function pointer table + auto dditable = + reinterpret_cast(hCommandBuffer) + ->dditable; + auto pfnAppendMembufferReadRectExp = + dditable->ur.CommandBufferExp.pfnAppendMembufferReadRectExp; + if (nullptr == pfnAppendMembufferReadRectExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + // convert loader handle to platform handle + hCommandBuffer = + reinterpret_cast(hCommandBuffer) + ->handle; + + // convert loader handle to platform handle + hBuffer = reinterpret_cast(hBuffer)->handle; + + // forward to device-platform + result = pfnAppendMembufferReadRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pDst, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urCommandBufferEnqueueExp __urdlllocal ur_result_t UR_APICALL urCommandBufferEnqueueExp( @@ -6727,8 +6932,16 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetCommandBufferExpProcAddrTable( ur_loader::urCommandBufferAppendMemcpyUSMExp; pDdiTable->pfnAppendMembufferCopyExp = ur_loader::urCommandBufferAppendMembufferCopyExp; + pDdiTable->pfnAppendMembufferWriteExp = + ur_loader::urCommandBufferAppendMembufferWriteExp; + pDdiTable->pfnAppendMembufferReadExp = + ur_loader::urCommandBufferAppendMembufferReadExp; pDdiTable->pfnAppendMembufferCopyRectExp = ur_loader::urCommandBufferAppendMembufferCopyRectExp; + pDdiTable->pfnAppendMembufferWriteRectExp = + ur_loader::urCommandBufferAppendMembufferWriteRectExp; + pDdiTable->pfnAppendMembufferReadRectExp = + ur_loader::urCommandBufferAppendMembufferReadRectExp; pDdiTable->pfnEnqueueExp = ur_loader::urCommandBufferEnqueueExp; } else { // return pointers directly to platform's DDIs diff --git a/source/loader/ur_libapi.cpp b/source/loader/ur_libapi.cpp index 664e7e3d1a..e5beb86b87 100644 --- a/source/loader/ur_libapi.cpp +++ b/source/loader/ur_libapi.cpp @@ -6779,6 +6779,103 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a memory write command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pSrc` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + auto pfnAppendMembufferWriteExp = + ur_lib::context->urDdiTable.CommandBufferExp.pfnAppendMembufferWriteExp; + if (nullptr == pfnAppendMembufferWriteExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + return pfnAppendMembufferWriteExp(hCommandBuffer, hBuffer, offset, size, + pSrc, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a memory read command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pDst` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + auto pfnAppendMembufferReadExp = + ur_lib::context->urDdiTable.CommandBufferExp.pfnAppendMembufferReadExp; + if (nullptr == pfnAppendMembufferReadExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + return pfnAppendMembufferReadExp(hCommandBuffer, hBuffer, offset, size, + pDst, numSyncPointsInWaitList, + pSyncPointWaitList, pSyncPoint); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Append a rectangular memory copy command to a command-buffer object /// @@ -6836,6 +6933,132 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a rectangular memory write command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pSrc` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + auto pfnAppendMembufferWriteRectExp = + ur_lib::context->urDdiTable.CommandBufferExp + .pfnAppendMembufferWriteRectExp; + if (nullptr == pfnAppendMembufferWriteRectExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + return pfnAppendMembufferWriteRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pSrc, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a rectangular memory read command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pDst` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command + ) try { + auto pfnAppendMembufferReadRectExp = + ur_lib::context->urDdiTable.CommandBufferExp + .pfnAppendMembufferReadRectExp; + if (nullptr == pfnAppendMembufferReadRectExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + return pfnAppendMembufferReadRectExp( + hCommandBuffer, hBuffer, bufferOffset, hostOffset, region, + bufferRowPitch, bufferSlicePitch, hostRowPitch, hostSlicePitch, pDst, + numSyncPointsInWaitList, pSyncPointWaitList, pSyncPoint); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Submit a command-buffer for execution on a queue. /// @@ -6974,6 +7197,7 @@ ur_result_t UR_APICALL urUSMReleaseExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -7025,6 +7249,7 @@ ur_result_t UR_APICALL urUsmP2PEnablePeerAccessExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -7061,6 +7286,7 @@ ur_result_t UR_APICALL urUsmP2PDisablePeerAccessExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` diff --git a/source/ur_api.cpp b/source/ur_api.cpp index b44a830b3c..ae93fd7f94 100644 --- a/source/ur_api.cpp +++ b/source/ur_api.cpp @@ -5691,6 +5691,85 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a memory write command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pSrc` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + const void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a memory read command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pDst` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + size_t offset, ///< [in] offset in bytes in the buffer object. + size_t size, ///< [in] size in bytes of data being written. + void *pDst, ///< [in] pointer to host memory where data is to be written to. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Append a rectangular memory copy command to a command-buffer object /// @@ -5737,6 +5816,110 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a rectangular memory write command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pSrc` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being + ///< written. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pSrc. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pSrc. + void * + pSrc, ///< [in] pointer to host memory where data is to be written from. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Append a rectangular memory read command to a command-buffer object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hCommandBuffer` +/// + `NULL == hBuffer` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pDst` +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_EXP +/// - ::UR_RESULT_ERROR_INVALID_COMMAND_BUFFER_SYNC_POINT_WAIT_LIST_EXP +/// + `pSyncPointWaitList == NULL && numSyncPointsInWaitList > 0` +/// + `pSyncPointWaitList != NULL && numSyncPointsInWaitList == 0` +/// - ::UR_RESULT_ERROR_INVALID_MEM_OBJECT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + ur_exp_command_buffer_handle_t + hCommandBuffer, ///< [in] handle of the command-buffer object. + ur_mem_handle_t hBuffer, ///< [in] handle of the buffer object. + ur_rect_offset_t bufferOffset, ///< [in] 3D offset in the buffer. + ur_rect_offset_t hostOffset, ///< [in] 3D offset in the host region. + ur_rect_region_t + region, ///< [in] 3D rectangular region descriptor: width, height, depth. + size_t + bufferRowPitch, ///< [in] length of each row in bytes in the buffer object. + size_t + bufferSlicePitch, ///< [in] length of each 2D slice in bytes in the buffer object being read. + size_t + hostRowPitch, ///< [in] length of each row in bytes in the host memory region pointed to + ///< by pDst. + size_t + hostSlicePitch, ///< [in] length of each 2D slice in bytes in the host memory region + ///< pointed to by pDst. + void *pDst, ///< [in] pointer to host memory where data is to be read into. + uint32_t + numSyncPointsInWaitList, ///< [in] The number of sync points in the provided dependency list. + const ur_exp_command_buffer_sync_point_t * + pSyncPointWaitList, ///< [in][optional] A list of sync points that this command depends on. + ur_exp_command_buffer_sync_point_t + *pSyncPoint ///< [out][optional] sync point associated with this command +) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Submit a command-buffer for execution on a queue. /// @@ -5855,6 +6038,7 @@ ur_result_t UR_APICALL urUSMReleaseExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -5899,6 +6083,7 @@ ur_result_t UR_APICALL urUsmP2PEnablePeerAccessExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` @@ -5928,6 +6113,7 @@ ur_result_t UR_APICALL urUsmP2PDisablePeerAccessExp( /// - ::UR_RESULT_SUCCESS /// - ::UR_RESULT_ERROR_UNINITIALIZED /// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == commandDevice` /// + `NULL == peerDevice` From fb460fcb68dba32b93ad5eea593a62b474bf9eb8 Mon Sep 17 00:00:00 2001 From: Krzysztof Swiecicki Date: Fri, 23 Jun 2023 13:03:02 +0200 Subject: [PATCH 2/5] Add coverity scan fix to search order test --- test/loader/adapter_registry/search_order.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/loader/adapter_registry/search_order.cpp b/test/loader/adapter_registry/search_order.cpp index b946b2164f..9264bafe44 100644 --- a/test/loader/adapter_registry/search_order.cpp +++ b/test/loader/adapter_registry/search_order.cpp @@ -11,7 +11,7 @@ void assertRegistryPathSequence(std::vector testAdapterPaths, static size_t assertIndex = 0; auto pathIt = std::find_if(testAdapterPaths.cbegin(), - testAdapterPaths.cend(), predicate); + testAdapterPaths.cend(), std::move(predicate)); size_t index = std::distance(testAdapterPaths.cbegin(), pathIt); ASSERT_EQ(index, assertIndex++); } From 2b50550aee0e2a8a5319a6159bbc2705328e372b Mon Sep 17 00:00:00 2001 From: Petr Vesely Date: Wed, 12 Jul 2023 11:57:37 +0100 Subject: [PATCH 3/5] [UR] update adapter to use newer sycl plugin version --- cmake/helpers.cmake | 4 +- ...move-sycl-namespaces-from-ur-adapter.patch | 390 +++++++++++++----- source/adapters/CMakeLists.txt | 9 +- 3 files changed, 301 insertions(+), 102 deletions(-) diff --git a/cmake/helpers.cmake b/cmake/helpers.cmake index a4a3a70e78..0892fecdb4 100644 --- a/cmake/helpers.cmake +++ b/cmake/helpers.cmake @@ -63,7 +63,9 @@ function(FetchSource GIT_REPOSITORY GIT_TAG GIT_DIR DEST) message(STATUS "Fetching sparse source ${GIT_DIR} from ${GIT_REPOSITORY} ${GIT_TAG}") IF(NOT EXISTS ${DEST}) file(MAKE_DIRECTORY ${DEST}) - execute_process(COMMAND git init -b main + execute_process(COMMAND git init + WORKING_DIRECTORY ${DEST}) + execute_process(COMMAND git checkout -b main WORKING_DIRECTORY ${DEST}) execute_process(COMMAND git remote add origin ${GIT_REPOSITORY} WORKING_DIRECTORY ${DEST}) diff --git a/source/adapters/0001-SYCL-CUDA-remove-sycl-namespaces-from-ur-adapter.patch b/source/adapters/0001-SYCL-CUDA-remove-sycl-namespaces-from-ur-adapter.patch index e7214e94b2..8153e1cb85 100644 --- a/source/adapters/0001-SYCL-CUDA-remove-sycl-namespaces-from-ur-adapter.patch +++ b/source/adapters/0001-SYCL-CUDA-remove-sycl-namespaces-from-ur-adapter.patch @@ -1,23 +1,174 @@ -From e30434a4a9e893f24e0bb18051576f297d1f4f08 Mon Sep 17 00:00:00 2001 +From fd78871a6bd2ff41ff37b8bd786c17f59911c677 Mon Sep 17 00:00:00 2001 From: pbalcer -Date: Thu, 29 Jun 2023 14:26:26 +0200 -Subject: [PATCH] [SYCL][CUDA] remove sycl namespaces from ur adapter +Date: Thu, 29 Jun 2023 15:11:43 +0200 +Subject: [PATCH] [SYCL][CUDA] remove sycl dependencies from cuda ur adapter +This was preventing out-of-tree build of the adapter for standalone +use with unified runtime. + +Signed-off-by: Piotr Balcer --- + .../ur/adapters/cuda/command_buffer.cpp | 52 ++--- .../ur/adapters/cuda/common.cpp | 6 +- .../ur/adapters/cuda/common.hpp | 5 - .../ur/adapters/cuda/context.cpp | 2 +- - .../ur/adapters/cuda/device.cpp | 170 +++++++++--------- + .../ur/adapters/cuda/device.cpp | 209 +++++++++--------- .../ur/adapters/cuda/enqueue.cpp | 2 +- - .../ur/adapters/cuda/event.cpp | 12 +- - .../ur/adapters/cuda/kernel.cpp | 26 +-- - .../ur/adapters/cuda/memory.cpp | 4 +- + .../ur/adapters/cuda/event.cpp | 17 +- + .../ur/adapters/cuda/kernel.cpp | 42 ++-- + .../ur/adapters/cuda/memory.cpp | 5 +- .../ur/adapters/cuda/queue.cpp | 2 +- .../ur/adapters/cuda/sampler.cpp | 2 +- - 10 files changed, 113 insertions(+), 118 deletions(-) + 11 files changed, 167 insertions(+), 177 deletions(-) +diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/command_buffer.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/command_buffer.cpp +index c83e9e732303..57956cb64a67 100644 +--- a/sycl/plugins/unified_runtime/ur/adapters/cuda/command_buffer.cpp ++++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/command_buffer.cpp +@@ -19,8 +19,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferCreateExp( + (void)hDevice; + (void)pCommandBufferDesc; + (void)phCommandBuffer; +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -28,8 +28,8 @@ UR_APIEXPORT ur_result_t UR_APICALL + urCommandBufferRetainExp(ur_exp_command_buffer_handle_t hCommandBuffer) { + (void)hCommandBuffer; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -37,8 +37,8 @@ UR_APIEXPORT ur_result_t UR_APICALL + urCommandBufferReleaseExp(ur_exp_command_buffer_handle_t hCommandBuffer) { + (void)hCommandBuffer; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -46,8 +46,8 @@ UR_APIEXPORT ur_result_t UR_APICALL + urCommandBufferFinalizeExp(ur_exp_command_buffer_handle_t hCommandBuffer) { + (void)hCommandBuffer; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -68,8 +68,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendKernelLaunchExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -86,8 +86,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendMemcpyUSMExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -107,8 +107,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -134,8 +134,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferAppendMembufferCopyRectExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -155,8 +155,8 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -175,8 +175,8 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferReadExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -203,8 +203,8 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferWriteRectExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -232,8 +232,8 @@ ur_result_t UR_APICALL urCommandBufferAppendMembufferReadRectExp( + (void)pSyncPointWaitList; + (void)pSyncPoint; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } + +@@ -247,7 +247,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferEnqueueExp( + (void)phEventWaitList; + (void)phEvent; + +- sycl::detail::ur::die("Experimental Command-buffer feature is not " +- "implemented for CUDA adapter."); ++ detail::ur::die("Experimental Command-buffer feature is not " ++ "implemented for CUDA adapter."); + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + } diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/common.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/common.cpp -index 86975e509..83264160e 100644 +index 86975e509725..83264160e700 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/common.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/common.cpp @@ -72,17 +72,17 @@ std::string getCudaVersionString() { @@ -42,7 +193,7 @@ index 86975e509..83264160e 100644 } diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/common.hpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/common.hpp -index 5cfa60901..82b38c10d 100644 +index 5cfa609018b2..82b38c10d449 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/common.hpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/common.hpp @@ -8,7 +8,6 @@ @@ -69,7 +220,7 @@ index 5cfa60901..82b38c10d 100644 -} // __SYCL_INLINE_VER_NAMESPACE(_V1) -} // namespace sycl diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/context.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/context.cpp -index 74a32bdac..2b621383d 100644 +index 74a32bdac274..2b621383da09 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/context.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/context.cpp @@ -66,7 +66,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urContextGetInfo( @@ -82,7 +233,7 @@ index 74a32bdac..2b621383d 100644 CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hContext->getDevice()->get()) == CUDA_SUCCESS); diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/device.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/device.cpp -index 24f9d52a0..c6b6bc07e 100644 +index 52d4e3badc8f..a81599d629a7 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/device.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/device.cpp @@ -15,7 +15,7 @@ @@ -177,7 +328,7 @@ index 24f9d52a0..c6b6bc07e 100644 return ReturnValue(size_t(MaxWorkGroupSize)); } -@@ -172,12 +172,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -172,14 +172,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_MAX_NUM_SUB_GROUPS: { // Number of sub-groups = max block size / warp size + possible remainder int MaxThreads = 0; @@ -188,10 +339,14 @@ index 24f9d52a0..c6b6bc07e 100644 hDevice->get()) == CUDA_SUCCESS); int WarpSize = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, +- hDevice->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetAttribute(&WarpSize, ++ CU_DEVICE_ATTRIBUTE_WARP_SIZE, ++ hDevice->get()) == CUDA_SUCCESS); int MaxWarps = (MaxThreads + WarpSize - 1) / WarpSize; + return ReturnValue(MaxWarps); + } @@ -187,7 +187,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // Volta provides independent thread scheduling // TODO: Revisit for previous generation GPUs @@ -228,37 +383,43 @@ index 24f9d52a0..c6b6bc07e 100644 cuDeviceGetAttribute(&Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hDevice->get()) == CUDA_SUCCESS); -@@ -266,7 +266,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -266,18 +266,18 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_SUB_GROUP_SIZES_INTEL: { // NVIDIA devices only support one sub-group size (the warp size) int WarpSize = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, +- hDevice->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetAttribute(&WarpSize, ++ CU_DEVICE_ATTRIBUTE_WARP_SIZE, ++ hDevice->get()) == CUDA_SUCCESS); size_t Sizes[1] = {static_cast(WarpSize)}; -@@ -274,10 +274,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, + return ReturnValue(Sizes, 1); } case UR_DEVICE_INFO_MAX_CLOCK_FREQUENCY: { int ClockFreq = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&ClockFreq, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&ClockFreq, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, +- hDevice->get()) == CUDA_SUCCESS); - sycl::detail::ur::assertion(ClockFreq >= 0); ++ detail::ur::assertion(cuDeviceGetAttribute(&ClockFreq, ++ CU_DEVICE_ATTRIBUTE_CLOCK_RATE, ++ hDevice->get()) == CUDA_SUCCESS); + detail::ur::assertion(ClockFreq >= 0); return ReturnValue(static_cast(ClockFreq) / 1000u); } case UR_DEVICE_INFO_ADDRESS_BITS: { -@@ -292,7 +292,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -292,8 +292,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, // CL_DEVICE_TYPE_CUSTOM. size_t Global = 0; - sycl::detail::ur::assertion(cuDeviceTotalMem(&Global, hDevice->get()) == +- CUDA_SUCCESS); + detail::ur::assertion(cuDeviceTotalMem(&Global, hDevice->get()) == - CUDA_SUCCESS); ++ CUDA_SUCCESS); auto QuarterGlobal = static_cast(Global / 4u); + @@ -308,7 +308,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, if (std::getenv("SYCL_PI_CUDA_ENABLE_IMAGE_SUPPORT") != nullptr) { Enabled = true; @@ -426,8 +587,9 @@ index 24f9d52a0..c6b6bc07e 100644 size_t Bytes = 0; // Runtime API has easy access to this value, driver API info is scarse. - sycl::detail::ur::assertion(cuDeviceTotalMem(&Bytes, hDevice->get()) == +- CUDA_SUCCESS); + detail::ur::assertion(cuDeviceTotalMem(&Bytes, hDevice->get()) == - CUDA_SUCCESS); ++ CUDA_SUCCESS); return ReturnValue(uint64_t{Bytes}); } case UR_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE: { @@ -458,9 +620,11 @@ index 24f9d52a0..c6b6bc07e 100644 case UR_DEVICE_INFO_ERROR_CORRECTION_SUPPORT: { int ECCEnabled = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&ECCEnabled, CU_DEVICE_ATTRIBUTE_ECC_ENABLED, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&ECCEnabled, CU_DEVICE_ATTRIBUTE_ECC_ENABLED, +- hDevice->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetAttribute(&ECCEnabled, ++ CU_DEVICE_ATTRIBUTE_ECC_ENABLED, ++ hDevice->get()) == CUDA_SUCCESS); - sycl::detail::ur::assertion((ECCEnabled == 0) | (ECCEnabled == 1)); + detail::ur::assertion((ECCEnabled == 0) | (ECCEnabled == 1)); @@ -470,25 +634,30 @@ index 24f9d52a0..c6b6bc07e 100644 case UR_DEVICE_INFO_HOST_UNIFIED_MEMORY: { int IsIntegrated = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&IsIntegrated, CU_DEVICE_ATTRIBUTE_INTEGRATED, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&IsIntegrated, CU_DEVICE_ATTRIBUTE_INTEGRATED, +- hDevice->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetAttribute(&IsIntegrated, ++ CU_DEVICE_ATTRIBUTE_INTEGRATED, ++ hDevice->get()) == CUDA_SUCCESS); - sycl::detail::ur::assertion((IsIntegrated == 0) | (IsIntegrated == 1)); + detail::ur::assertion((IsIntegrated == 0) | (IsIntegrated == 1)); auto result = static_cast(IsIntegrated); return ReturnValue(result); } -@@ -620,7 +620,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -620,9 +620,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_NAME: { static constexpr size_t MaxDeviceNameLength = 256u; char Name[MaxDeviceNameLength]; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetName(Name, MaxDeviceNameLength, hDevice->get()) == - CUDA_SUCCESS); +- cuDeviceGetName(Name, MaxDeviceNameLength, hDevice->get()) == +- CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetName(Name, MaxDeviceNameLength, ++ hDevice->get()) == CUDA_SUCCESS); return ReturnValue(Name, strlen(Name) + 1); -@@ -641,13 +641,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, + } + case UR_DEVICE_INFO_VENDOR: { +@@ -641,13 +640,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_VERSION: { std::stringstream SS; int Major; @@ -504,7 +673,7 @@ index 24f9d52a0..c6b6bc07e 100644 cuDeviceGetAttribute(&Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hDevice->get()) == CUDA_SUCCESS); -@@ -666,11 +666,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -666,11 +665,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, int Major = 0; int Minor = 0; @@ -518,14 +687,16 @@ index 24f9d52a0..c6b6bc07e 100644 cuDeviceGetAttribute(&Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hDevice->get()) == CUDA_SUCCESS); -@@ -847,27 +847,27 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -847,27 +846,27 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_GLOBAL_MEM_FREE: { size_t FreeMemory = 0; size_t TotalMemory = 0; - sycl::detail::ur::assertion(cuMemGetInfo(&FreeMemory, &TotalMemory) == +- CUDA_SUCCESS, +- "failed cuMemGetInfo() API."); + detail::ur::assertion(cuMemGetInfo(&FreeMemory, &TotalMemory) == - CUDA_SUCCESS, - "failed cuMemGetInfo() API."); ++ CUDA_SUCCESS, ++ "failed cuMemGetInfo() API."); return ReturnValue(FreeMemory); } case UR_DEVICE_INFO_MEMORY_CLOCK_RATE: { @@ -551,7 +722,7 @@ index 24f9d52a0..c6b6bc07e 100644 return ReturnValue(Value); } case UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES: { -@@ -875,10 +875,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -875,20 +874,20 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_DEVICE_ID: { int Value = 0; @@ -564,20 +735,21 @@ index 24f9d52a0..c6b6bc07e 100644 return ReturnValue(Value); } case UR_DEVICE_INFO_UUID: { -@@ -888,10 +888,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, - int Minor = DriverVersion % 1000 / 10; CUuuid UUID; - if ((Major > 11) || (Major == 11 && Minor >= 4)) { -- sycl::detail::ur::assertion(cuDeviceGetUuid_v2(&UUID, hDevice->get()) == -+ detail::ur::assertion(cuDeviceGetUuid_v2(&UUID, hDevice->get()) == - CUDA_SUCCESS); - } else { -- sycl::detail::ur::assertion(cuDeviceGetUuid(&UUID, hDevice->get()) == -+ detail::ur::assertion(cuDeviceGetUuid(&UUID, hDevice->get()) == - CUDA_SUCCESS); - } + #if (CUDA_VERSION >= 11040) +- sycl::detail::ur::assertion(cuDeviceGetUuid_v2(&UUID, hDevice->get()) == +- CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetUuid_v2(&UUID, hDevice->get()) == ++ CUDA_SUCCESS); + #else +- sycl::detail::ur::assertion(cuDeviceGetUuid(&UUID, hDevice->get()) == +- CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetUuid(&UUID, hDevice->get()) == ++ CUDA_SUCCESS); + #endif std::array Name; -@@ -900,13 +900,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, + std::copy(UUID.bytes, UUID.bytes + 16, Name.begin()); +@@ -896,13 +895,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_MAX_MEMORY_BANDWIDTH: { int Major = 0; @@ -593,7 +765,7 @@ index 24f9d52a0..c6b6bc07e 100644 cuDeviceGetAttribute(&Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hDevice->get()) == CUDA_SUCCESS); -@@ -922,7 +922,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -918,7 +917,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } else if (IsOrinAGX) { MemoryClockKHz = 3200000; } else { @@ -602,7 +774,7 @@ index 24f9d52a0..c6b6bc07e 100644 cuDeviceGetAttribute(&MemoryClockKHz, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, hDevice->get()) == CUDA_SUCCESS); -@@ -932,7 +932,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -928,7 +927,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, if (IsOrinAGX) { MemoryBusWidth = 256; } else { @@ -611,7 +783,7 @@ index 24f9d52a0..c6b6bc07e 100644 cuDeviceGetAttribute(&MemoryBusWidth, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, hDevice->get()) == CUDA_SUCCESS); -@@ -977,7 +977,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -973,7 +972,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, &MaxRegisters, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK, hDevice->get())); @@ -620,25 +792,27 @@ index 24f9d52a0..c6b6bc07e 100644 return ReturnValue(static_cast(MaxRegisters)); } -@@ -988,11 +988,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, +@@ -984,12 +983,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_PCI_ADDRESS: { constexpr size_t AddressBufferSize = 13; char AddressBuffer[AddressBufferSize]; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, hDevice->get()) == - CUDA_SUCCESS); +- cuDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, hDevice->get()) == +- CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetPCIBusId(AddressBuffer, AddressBufferSize, ++ hDevice->get()) == CUDA_SUCCESS); // CUDA API (8.x - 12.1) guarantees 12 bytes + \0 are written - sycl::detail::ur::assertion(strnlen(AddressBuffer, AddressBufferSize) == -+ detail::ur::assertion(strnlen(AddressBuffer, AddressBufferSize) == - 12); +- 12); ++ detail::ur::assertion(strnlen(AddressBuffer, AddressBufferSize) == 12); return ReturnValue(AddressBuffer, strnlen(AddressBuffer, AddressBufferSize - 1) + 1); + } diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/enqueue.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/enqueue.cpp -index 52c4c3895..55c56aee2 100644 +index 1cfc5cc40a4a..792f69092682 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/enqueue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/enqueue.cpp -@@ -806,7 +806,7 @@ static size_t imageElementByteSize(CUDA_ARRAY_DESCRIPTOR ArrayDesc) { +@@ -794,7 +794,7 @@ static size_t imageElementByteSize(CUDA_ARRAY_DESCRIPTOR ArrayDesc) { case CU_AD_FORMAT_FLOAT: return 4; default: @@ -648,7 +822,7 @@ index 52c4c3895..55c56aee2 100644 } } diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/event.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/event.cpp -index 8916197b7..9d86189b9 100644 +index 8916197b73f1..066c0498f1d0 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/event.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/event.cpp @@ -119,7 +119,7 @@ ur_result_t ur_event_handle_t_::record() { @@ -687,26 +861,31 @@ index 8916197b7..9d86189b9 100644 return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -@@ -254,7 +254,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { +@@ -254,8 +254,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) { const auto RefCount = hEvent->incrementReferenceCount(); - sycl::detail::ur::assertion( -+ detail::ur::assertion( - RefCount != 0, "Reference count overflow detected in urEventRetain."); +- RefCount != 0, "Reference count overflow detected in urEventRetain."); ++ detail::ur::assertion(RefCount != 0, ++ "Reference count overflow detected in urEventRetain."); return UR_RESULT_SUCCESS; -@@ -265,7 +265,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRelease(ur_event_handle_t hEvent) { + } +@@ -265,9 +265,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRelease(ur_event_handle_t hEvent) { // double delete or someone is messing with the ref count. // either way, cannot safely proceed. - sycl::detail::ur::assertion( -+ detail::ur::assertion( - hEvent->getReferenceCount() != 0, - "Reference count overflow detected in urEventRelease."); +- hEvent->getReferenceCount() != 0, +- "Reference count overflow detected in urEventRelease."); ++ detail::ur::assertion(hEvent->getReferenceCount() != 0, ++ "Reference count overflow detected in urEventRelease."); + // decrement ref count. If it is 0, delete the event. + if (hEvent->decrementReferenceCount() == 0) { diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/kernel.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/kernel.cpp -index 358f59c49..cae080401 100644 +index 358f59c499e1..7d46ce039bab 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/kernel.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/kernel.cpp @@ -73,24 +73,24 @@ urKernelGetGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, @@ -758,43 +937,55 @@ index 358f59c49..cae080401 100644 cuFuncGetAttribute(&Bytes, CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, hKernel->get()) == CUDA_SUCCESS); return ReturnValue(uint64_t(Bytes)); -@@ -130,7 +130,7 @@ urKernelGetGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, +@@ -130,17 +130,17 @@ urKernelGetGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, case UR_KERNEL_GROUP_INFO_PREFERRED_WORK_GROUP_SIZE_MULTIPLE: { // Work groups should be multiples of the warp size int WarpSize = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, +- hDevice->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetAttribute(&WarpSize, ++ CU_DEVICE_ATTRIBUTE_WARP_SIZE, ++ hDevice->get()) == CUDA_SUCCESS); return ReturnValue(static_cast(WarpSize)); -@@ -138,7 +138,7 @@ urKernelGetGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, + } case UR_KERNEL_GROUP_INFO_PRIVATE_MEM_SIZE: { // OpenCL PRIVATE == CUDA LOCAL int Bytes = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuFuncGetAttribute(&Bytes, CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, - hKernel->get()) == CUDA_SUCCESS); +- cuFuncGetAttribute(&Bytes, CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, +- hKernel->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuFuncGetAttribute(&Bytes, ++ CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, ++ hKernel->get()) == CUDA_SUCCESS); return ReturnValue(uint64_t(Bytes)); -@@ -231,7 +231,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetInfo(ur_kernel_handle_t hKernel, + } + default: +@@ -231,9 +231,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelGetInfo(ur_kernel_handle_t hKernel, return ReturnValue(""); case UR_KERNEL_INFO_NUM_REGS: { int NumRegs = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuFuncGetAttribute(&NumRegs, CU_FUNC_ATTRIBUTE_NUM_REGS, - hKernel->get()) == CUDA_SUCCESS); +- cuFuncGetAttribute(&NumRegs, CU_FUNC_ATTRIBUTE_NUM_REGS, +- hKernel->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuFuncGetAttribute(&NumRegs, ++ CU_FUNC_ATTRIBUTE_NUM_REGS, ++ hKernel->get()) == CUDA_SUCCESS); return ReturnValue(static_cast(NumRegs)); -@@ -254,7 +254,7 @@ urKernelGetSubGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, + } + default: +@@ -254,15 +254,15 @@ urKernelGetSubGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, case UR_KERNEL_SUB_GROUP_INFO_MAX_SUB_GROUP_SIZE: { // Sub-group size is equivalent to warp size int WarpSize = 0; - sycl::detail::ur::assertion( -+ detail::ur::assertion( - cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, - hDevice->get()) == CUDA_SUCCESS); +- cuDeviceGetAttribute(&WarpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, +- hDevice->get()) == CUDA_SUCCESS); ++ detail::ur::assertion(cuDeviceGetAttribute(&WarpSize, ++ CU_DEVICE_ATTRIBUTE_WARP_SIZE, ++ hDevice->get()) == CUDA_SUCCESS); return ReturnValue(static_cast(WarpSize)); -@@ -262,7 +262,7 @@ urKernelGetSubGroupInfo(ur_kernel_handle_t hKernel, ur_device_handle_t hDevice, + } case UR_KERNEL_SUB_GROUP_INFO_MAX_NUM_SUB_GROUPS: { // Number of sub-groups = max block size / warp size + possible remainder int MaxThreads = 0; @@ -804,19 +995,20 @@ index 358f59c49..cae080401 100644 hKernel->get()) == CUDA_SUCCESS); int WarpSize = 0; diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/memory.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/memory.cpp -index b19acea31..ecf840330 100644 +index b19acea3159f..f0c276579476 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/memory.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/memory.cpp -@@ -162,7 +162,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemRelease(ur_mem_handle_t hMem) { +@@ -162,8 +162,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemRelease(ur_mem_handle_t hMem) { // error for which it is unclear if the function that reported it succeeded // or not. Either way, the state of the program is compromised and likely // unrecoverable. - sycl::detail::ur::die( -+ detail::ur::die( - "Unrecoverable program state reached in urMemRelease"); +- "Unrecoverable program state reached in urMemRelease"); ++ detail::ur::die("Unrecoverable program state reached in urMemRelease"); } -@@ -331,7 +331,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreate( + return UR_RESULT_SUCCESS; +@@ -331,7 +330,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageCreate( PixelTypeSizeBytes = 4; break; default: @@ -826,7 +1018,7 @@ index b19acea31..ecf840330 100644 } diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp -index 05443eeed..32391fec5 100644 +index 05443eeed89d..32391fec5c13 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp @@ -265,7 +265,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle( @@ -839,7 +1031,7 @@ index 05443eeed..32391fec5 100644 std::vector ComputeCuStreams(1, CuStream); std::vector TransferCuStreams(0); diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/sampler.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/sampler.cpp -index 36ec89fb9..836e47f98 100644 +index 36ec89fb9da3..836e47f988e5 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/sampler.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/cuda/sampler.cpp @@ -73,7 +73,7 @@ urSamplerRelease(ur_sampler_handle_t hSampler) { diff --git a/source/adapters/CMakeLists.txt b/source/adapters/CMakeLists.txt index d9d4d29156..9a817f4abe 100644 --- a/source/adapters/CMakeLists.txt +++ b/source/adapters/CMakeLists.txt @@ -9,9 +9,14 @@ add_subdirectory(null) if(UR_BUILD_ADAPTER_L0 OR UR_BUILD_ADAPTER_CUDA) # fetch adapter sources from SYCL set(SYCL_ADAPTER_DIR "${CMAKE_CURRENT_SOURCE_DIR}/external") - FetchSource(https://github.com/intel/llvm.git sycl-nightly/20230628 "sycl/plugins/unified_runtime/ur" ${SYCL_ADAPTER_DIR}) + FetchSource(https://github.com/intel/llvm.git sycl-nightly/20230706 "sycl/plugins/unified_runtime/ur" ${SYCL_ADAPTER_DIR}) - execute_process(COMMAND git apply --quiet ../0001-SYCL-CUDA-remove-sycl-namespaces-from-ur-adapter.patch + get_program_version_major_minor(git GIT_VERSION) + set(GIT_QUIET_OPTION "") + if(GIT_VERSION VERSION_GREATER_EQUAL "3.35.0") + set(GIT_QUIET_OPTION "--quiet") + endif() + execute_process(COMMAND git apply ${GIT_QUIET_OPTION} ../0001-SYCL-CUDA-remove-sycl-namespaces-from-ur-adapter.patch WORKING_DIRECTORY ${SYCL_ADAPTER_DIR}) endif() From 7c2db62b18ea170fcf13fbca1cdc1ded92a81c11 Mon Sep 17 00:00:00 2001 From: Petr Vesely Date: Wed, 12 Jul 2023 12:09:19 +0100 Subject: [PATCH 4/5] [UR] Fix L0 plugin file names --- source/adapters/level_zero/CMakeLists.txt | 46 +++++++++++------------ 1 file changed, 23 insertions(+), 23 deletions(-) diff --git a/source/adapters/level_zero/CMakeLists.txt b/source/adapters/level_zero/CMakeLists.txt index e1d483cdd4..e9b4aaa547 100644 --- a/source/adapters/level_zero/CMakeLists.txt +++ b/source/adapters/level_zero/CMakeLists.txt @@ -73,30 +73,30 @@ target_include_directories(LevelZeroLoader-Headers add_library(${TARGET_NAME} SHARED - ${L0_DIR}/ur_loader_interface.cpp - ${L0_DIR}/ur_level_zero_common.hpp - ${L0_DIR}/ur_level_zero_context.hpp - ${L0_DIR}/ur_level_zero_device.hpp - ${L0_DIR}/ur_level_zero_event.hpp - ${L0_DIR}/ur_level_zero_usm.hpp - ${L0_DIR}/ur_level_zero_mem.hpp - ${L0_DIR}/ur_level_zero_kernel.hpp - ${L0_DIR}/ur_level_zero_platform.hpp - ${L0_DIR}/ur_level_zero_program.hpp - ${L0_DIR}/ur_level_zero_queue.hpp - ${L0_DIR}/ur_level_zero_sampler.hpp + ${L0_DIR}/ur_interface_loader.cpp + ${L0_DIR}/common.hpp + ${L0_DIR}/context.hpp + ${L0_DIR}/device.hpp + ${L0_DIR}/event.hpp + ${L0_DIR}/usm.hpp + ${L0_DIR}/memory.hpp + ${L0_DIR}/kernel.hpp + ${L0_DIR}/platform.hpp + ${L0_DIR}/program.hpp + ${L0_DIR}/queue.hpp + ${L0_DIR}/sampler.hpp ${L0_DIR}/ur_level_zero.cpp - ${L0_DIR}/ur_level_zero_common.cpp - ${L0_DIR}/ur_level_zero_context.cpp - ${L0_DIR}/ur_level_zero_device.cpp - ${L0_DIR}/ur_level_zero_event.cpp - ${L0_DIR}/ur_level_zero_usm.cpp - ${L0_DIR}/ur_level_zero_mem.cpp - ${L0_DIR}/ur_level_zero_kernel.cpp - ${L0_DIR}/ur_level_zero_platform.cpp - ${L0_DIR}/ur_level_zero_program.cpp - ${L0_DIR}/ur_level_zero_queue.cpp - ${L0_DIR}/ur_level_zero_sampler.cpp + ${L0_DIR}/common.cpp + ${L0_DIR}/context.cpp + ${L0_DIR}/device.cpp + ${L0_DIR}/event.cpp + ${L0_DIR}/usm.cpp + ${L0_DIR}/memory.cpp + ${L0_DIR}/kernel.cpp + ${L0_DIR}/platform.cpp + ${L0_DIR}/program.cpp + ${L0_DIR}/queue.cpp + ${L0_DIR}/sampler.cpp ${L0_DIR}/../../ur.cpp ${L0_DIR}/../../usm_allocator.cpp ${L0_DIR}/../../usm_allocator.hpp From 43f421691c8d7b2e5cb3466498e09428122ab885 Mon Sep 17 00:00:00 2001 From: Petr Vesely Date: Wed, 12 Jul 2023 13:25:16 +0100 Subject: [PATCH 5/5] [UR] Bump L0 Loader version --- source/adapters/level_zero/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/source/adapters/level_zero/CMakeLists.txt b/source/adapters/level_zero/CMakeLists.txt index e9b4aaa547..52d22c3ea6 100644 --- a/source/adapters/level_zero/CMakeLists.txt +++ b/source/adapters/level_zero/CMakeLists.txt @@ -11,7 +11,7 @@ if (NOT DEFINED LEVEL_ZERO_LIBRARY OR NOT DEFINED LEVEL_ZERO_INCLUDE_DIR) message(STATUS "Download Level Zero loader and headers from github.com") set(LEVEL_ZERO_LOADER_REPO "https://github.com/oneapi-src/level-zero.git") - set(LEVEL_ZERO_LOADER_TAG v1.8.8) + set(LEVEL_ZERO_LOADER_TAG v1.11.0) # Disable due to a bug https://github.com/oneapi-src/level-zero/issues/104 set(CMAKE_INCLUDE_CURRENT_DIR OFF)