Merge pull request #1791 from aarongreig/aaron/makeKernelNativeProgra…

…mOpt

Make urKernelCreateWithNativeHandle's program param optional.

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.
 );

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."

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<ze_kernel_handle_t>(NativeKernel);
  ur_kernel_handle_t_ *Kernel = nullptr;
  try {

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

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

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;
  }

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<ur_context_object_t *>(hContext)->handle;
 
  // convert loader handle to platform handle
- hProgram = reinterpret_cast<ur_program_object_t *>(hProgram)->handle;
+ hProgram = (hProgram)
+ ? reinterpret_cast<ur_program_object_t *>(hProgram)->handle
+ : nullptr;
 
  // forward to device-platform
  result = pfnCreateWithNativeHandle(hNativeKernel, hContext, hProgram,

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

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

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<std::vector<char>> 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<const uint8_t *>(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);
+}

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

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,