diff --git a/include/ur_api.h b/include/ur_api.h index 12eb2dd643..fff58cc7cc 100644 --- a/include/ur_api.h +++ b/include/ur_api.h @@ -5278,6 +5278,8 @@ typedef struct ur_kernel_native_properties_t { /// - The application may call this function from simultaneous threads for /// the same context. /// - The implementation of this function should be thread-safe. +/// - The implementation may require a valid program handle to return the +/// native kernel handle /// /// @returns /// - ::UR_RESULT_SUCCESS @@ -5286,7 +5288,7 @@ typedef struct ur_kernel_native_properties_t { /// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == hContext` -/// + `NULL == hProgram` +/// + If `hProgram == NULL` and the implementation requires a valid program. /// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER /// + `NULL == phKernel` /// - ::UR_RESULT_ERROR_UNSUPPORTED_FEATURE @@ -5295,7 +5297,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelCreateWithNativeHandle( ur_native_handle_t hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object - ur_program_handle_t hProgram, ///< [in] handle of the program associated with the kernel + ur_program_handle_t hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t *pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t *phKernel ///< [out] pointer to the handle of the kernel object created. ); diff --git a/scripts/core/kernel.yml b/scripts/core/kernel.yml index e657d806be..4815c8413a 100644 --- a/scripts/core/kernel.yml +++ b/scripts/core/kernel.yml @@ -513,6 +513,7 @@ details: - "Creates runtime kernel handle from native driver kernel handle." - "The application may call this function from simultaneous threads for the same context." - "The implementation of this function should be thread-safe." + - "The implementation may require a valid program handle to return the native kernel handle" params: - type: $x_native_handle_t name: hNativeKernel @@ -523,7 +524,7 @@ params: desc: "[in] handle of the context object" - type: $x_program_handle_t name: hProgram - desc: "[in] handle of the program associated with the kernel" + desc: "[in][optional] handle of the program associated with the kernel" - type: "const $x_kernel_native_properties_t*" name: pProperties desc: "[in][optional] pointer to native kernel properties struct" @@ -534,6 +535,8 @@ params: returns: - $X_RESULT_ERROR_UNSUPPORTED_FEATURE: - "If the adapter has no underlying equivalent handle." + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "If `hProgram == NULL` and the implementation requires a valid program." --- #-------------------------------------------------------------------------- type: function desc: "Get the suggested local work size for a kernel." diff --git a/source/adapters/level_zero/kernel.cpp b/source/adapters/level_zero/kernel.cpp index af0dfdd0d5..cb020395ed 100644 --- a/source/adapters/level_zero/kernel.cpp +++ b/source/adapters/level_zero/kernel.cpp @@ -1142,6 +1142,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urKernelCreateWithNativeHandle( ur_kernel_handle_t * RetKernel ///< [out] pointer to the handle of the kernel object created. ) { + if (!Program) { + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } ze_kernel_handle_t ZeKernel = ur_cast(NativeKernel); ur_kernel_handle_t_ *Kernel = nullptr; try { diff --git a/source/adapters/mock/ur_mockddi.cpp b/source/adapters/mock/ur_mockddi.cpp index c204cc248e..bd024f90f7 100644 --- a/source/adapters/mock/ur_mockddi.cpp +++ b/source/adapters/mock/ur_mockddi.cpp @@ -4489,7 +4489,7 @@ __urdlllocal ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object ur_program_handle_t - hProgram, ///< [in] handle of the program associated with the kernel + hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t * pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t diff --git a/source/loader/layers/tracing/ur_trcddi.cpp b/source/loader/layers/tracing/ur_trcddi.cpp index 4da6da9081..d9d35f7a1b 100644 --- a/source/loader/layers/tracing/ur_trcddi.cpp +++ b/source/loader/layers/tracing/ur_trcddi.cpp @@ -3428,7 +3428,7 @@ __urdlllocal ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object ur_program_handle_t - hProgram, ///< [in] handle of the program associated with the kernel + hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t * pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t diff --git a/source/loader/layers/validation/ur_valddi.cpp b/source/loader/layers/validation/ur_valddi.cpp index 38e0776650..9ba1839c9f 100644 --- a/source/loader/layers/validation/ur_valddi.cpp +++ b/source/loader/layers/validation/ur_valddi.cpp @@ -3855,7 +3855,7 @@ __urdlllocal ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object ur_program_handle_t - hProgram, ///< [in] handle of the program associated with the kernel + hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t * pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t @@ -3873,10 +3873,6 @@ __urdlllocal ur_result_t UR_APICALL urKernelCreateWithNativeHandle( return UR_RESULT_ERROR_INVALID_NULL_HANDLE; } - if (NULL == hProgram) { - return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - } - if (NULL == phKernel) { return UR_RESULT_ERROR_INVALID_NULL_POINTER; } diff --git a/source/loader/ur_ldrddi.cpp b/source/loader/ur_ldrddi.cpp index f1d65dd67c..1eca3ecc2b 100644 --- a/source/loader/ur_ldrddi.cpp +++ b/source/loader/ur_ldrddi.cpp @@ -3592,7 +3592,7 @@ __urdlllocal ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object ur_program_handle_t - hProgram, ///< [in] handle of the program associated with the kernel + hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t * pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t @@ -3614,7 +3614,9 @@ __urdlllocal ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hContext = reinterpret_cast(hContext)->handle; // convert loader handle to platform handle - hProgram = reinterpret_cast(hProgram)->handle; + hProgram = (hProgram) + ? reinterpret_cast(hProgram)->handle + : nullptr; // forward to device-platform result = pfnCreateWithNativeHandle(hNativeKernel, hContext, hProgram, diff --git a/source/loader/ur_libapi.cpp b/source/loader/ur_libapi.cpp index f594aaf481..a52d4678a6 100644 --- a/source/loader/ur_libapi.cpp +++ b/source/loader/ur_libapi.cpp @@ -4139,6 +4139,8 @@ ur_result_t UR_APICALL urKernelGetNativeHandle( /// - The application may call this function from simultaneous threads for /// the same context. /// - The implementation of this function should be thread-safe. +/// - The implementation may require a valid program handle to return the +/// native kernel handle /// /// @returns /// - ::UR_RESULT_SUCCESS @@ -4147,7 +4149,7 @@ ur_result_t UR_APICALL urKernelGetNativeHandle( /// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == hContext` -/// + `NULL == hProgram` +/// + If `hProgram == NULL` and the implementation requires a valid program. /// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER /// + `NULL == phKernel` /// - ::UR_RESULT_ERROR_UNSUPPORTED_FEATURE @@ -4157,7 +4159,7 @@ ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object ur_program_handle_t - hProgram, ///< [in] handle of the program associated with the kernel + hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t * pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t diff --git a/source/ur_api.cpp b/source/ur_api.cpp index 366232e6ee..2b4f98c76e 100644 --- a/source/ur_api.cpp +++ b/source/ur_api.cpp @@ -3516,6 +3516,8 @@ ur_result_t UR_APICALL urKernelGetNativeHandle( /// - The application may call this function from simultaneous threads for /// the same context. /// - The implementation of this function should be thread-safe. +/// - The implementation may require a valid program handle to return the +/// native kernel handle /// /// @returns /// - ::UR_RESULT_SUCCESS @@ -3524,7 +3526,7 @@ ur_result_t UR_APICALL urKernelGetNativeHandle( /// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC /// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE /// + `NULL == hContext` -/// + `NULL == hProgram` +/// + If `hProgram == NULL` and the implementation requires a valid program. /// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER /// + `NULL == phKernel` /// - ::UR_RESULT_ERROR_UNSUPPORTED_FEATURE @@ -3534,7 +3536,7 @@ ur_result_t UR_APICALL urKernelCreateWithNativeHandle( hNativeKernel, ///< [in][nocheck] the native handle of the kernel. ur_context_handle_t hContext, ///< [in] handle of the context object ur_program_handle_t - hProgram, ///< [in] handle of the program associated with the kernel + hProgram, ///< [in][optional] handle of the program associated with the kernel const ur_kernel_native_properties_t * pProperties, ///< [in][optional] pointer to native kernel properties struct ur_kernel_handle_t diff --git a/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp b/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp index 62c667b242..6ee49dbbfb 100644 --- a/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp +++ b/test/adapters/level_zero/urKernelCreateWithNativeHandle.cpp @@ -61,3 +61,46 @@ TEST_P(urLevelZeroKernelNativeHandleTest, OwnedHandleRelease) { ASSERT_SUCCESS(urKernelRelease(kernel)); ASSERT_SUCCESS(urProgramRelease(program)); } + +TEST_P(urLevelZeroKernelNativeHandleTest, NullProgram) { + ze_context_handle_t native_context; + urContextGetNativeHandle(context, (ur_native_handle_t *)&native_context); + + ze_device_handle_t native_device; + urDeviceGetNativeHandle(device, (ur_native_handle_t *)&native_device); + + std::shared_ptr> il_binary; + uur::KernelsEnvironment::instance->LoadSource("foo", il_binary); + + auto kernel_name = + uur::KernelsEnvironment::instance->GetEntryPointNames("foo")[0]; + + ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC}; + moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + moduleDesc.inputSize = il_binary->size(); + moduleDesc.pInputModule = + reinterpret_cast(il_binary->data()); + moduleDesc.pBuildFlags = ""; + ze_module_handle_t module; + + ASSERT_EQ(zeModuleCreate(native_context, native_device, &moduleDesc, + &module, NULL), + ZE_RESULT_SUCCESS); + + ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC}; + kernelDesc.pKernelName = kernel_name.c_str(); + + ze_kernel_handle_t native_kernel; + + ASSERT_EQ(zeKernelCreate(module, &kernelDesc, &native_kernel), + ZE_RESULT_SUCCESS); + + ur_kernel_native_properties_t kprops = { + UR_STRUCTURE_TYPE_KERNEL_NATIVE_PROPERTIES, nullptr, 1}; + + ur_kernel_handle_t kernel; + EXPECT_EQ(urKernelCreateWithNativeHandle((ur_native_handle_t)native_kernel, + context, nullptr, &kprops, + &kernel), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} diff --git a/test/conformance/kernel/kernel_adapter_native_cpu.match b/test/conformance/kernel/kernel_adapter_native_cpu.match index 0c10ec3ee4..4d3b506fcf 100644 --- a/test/conformance/kernel/kernel_adapter_native_cpu.match +++ b/test/conformance/kernel/kernel_adapter_native_cpu.match @@ -5,7 +5,6 @@ urKernelCreateTest.InvalidNullPointerKernel/SYCL_NATIVE_CPU___SYCL_Native_CPU__{ urKernelCreateTest.InvalidKernelName/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} urKernelCreateWithNativeHandleTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} urKernelCreateWithNativeHandleTest.InvalidNullHandleContext/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} -urKernelCreateWithNativeHandleTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} urKernelCreateWithNativeHandleTest.InvalidNullPointerNativeKernel/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}} urKernelGetGroupInfoTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}__UR_KERNEL_GROUP_INFO_GLOBAL_WORK_SIZE urKernelGetGroupInfoTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}__UR_KERNEL_GROUP_INFO_WORK_GROUP_SIZE diff --git a/test/conformance/kernel/urKernelCreateWithNativeHandle.cpp b/test/conformance/kernel/urKernelCreateWithNativeHandle.cpp index 4f7dae4598..8640463334 100644 --- a/test/conformance/kernel/urKernelCreateWithNativeHandle.cpp +++ b/test/conformance/kernel/urKernelCreateWithNativeHandle.cpp @@ -50,13 +50,6 @@ TEST_P(urKernelCreateWithNativeHandleTest, InvalidNullHandleContext) { &properties, &native_kernel)); } -TEST_P(urKernelCreateWithNativeHandleTest, InvalidNullHandleProgram) { - ASSERT_EQ_RESULT( - UR_RESULT_ERROR_INVALID_NULL_HANDLE, - urKernelCreateWithNativeHandle(native_kernel_handle, context, nullptr, - &properties, &native_kernel)); -} - TEST_P(urKernelCreateWithNativeHandleTest, InvalidNullPointerNativeKernel) { ASSERT_EQ_RESULT(UR_RESULT_ERROR_INVALID_NULL_POINTER, urKernelCreateWithNativeHandle(native_kernel_handle,