[SYCL][HIP] Implemented supported make_* interop functions. #10526
Conversation
Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
urDeviceCreateWithNativeHandle can never be reached in the hip backend (see comment in code) so I removed the impl I added earlier. Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
| @@ -310,7 +322,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetNativeHandle( | |||
| /// | |||
| /// \return UR_RESULT_ERROR_UNSUPPORTED_FEATURE | |||
There was a problem hiding this comment.
I missed that I should update this function descriptor. Will do this following review.
| } | ||
| // HIP uses a 32-bit int instead of an opaque pointer like other backends, | ||
| // so we need a specialization with static_cast instead of reinterpret_cast. | ||
| return static_cast<backend_return_t<backend::ext_oneapi_hip, device>>( | ||
| Obj.getNative()); | ||
| } | ||
|
|
||
| template <> | ||
| inline device make_device<backend::ext_oneapi_hip>( |
There was a problem hiding this comment.
Why is this code unique to HIP?
There was a problem hiding this comment.
There are a hierarchy of reasons. Firstly if we want to unique devices (#6055) like in the cuda backend, #7550 , then the code is unique to HIP (since the cuda implementation is in the experimental namespace). If we chose to not unique the device, then we would still need the code to be unique to hip, since, like the cuda backend, we would need to static_cast between pi_native_handle and backend_input_t<backend::ext_oneapi_cuda, device> since the default make_device uses reinterpret_cast which does not work.
| throw sycl::exception(make_error_code(errc::backend_mismatch), | ||
| "Backends mismatch"); |
There was a problem hiding this comment.
I'm not sure we can do that outside ABI breaking window.
There was a problem hiding this comment.
OK I can revert this. Thanks.
|
SYCL RT part is LGTM (I'll approve after minor revert mentioned above is pushed here). Haven't looked into anything under |
Update comments. Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
|
@sergey-semenov could you review this please? Thanks |
Co-authored-by: Sergey Semenov <[email protected]>
Signed-off-by: Jack Kirk <[email protected]>
|
@intel/llvm-gatekeepers This is now ready to be merged. Thanks |
) This PR adds missing functions in the hip backend to allow for interoperability in programs that create sycl objects from native hip objects. The new function implementations are: - `make_device` - `make_queue` - `make_event` Note that it would really make sense for intel#10491 to be merged first because this PR makes the same code change in pi2ur, for a fix that is attributed to intel#10491. --------- Signed-off-by: Jack Kirk <[email protected]>
) This PR adds missing functions in the hip backend to allow for interoperability in programs that create sycl objects from native hip objects. The new function implementations are: - `make_device` - `make_queue` - `make_event` Note that it would really make sense for intel#10491 to be merged first because this PR makes the same code change in pi2ur, for a fix that is attributed to intel#10491. --------- Signed-off-by: Jack Kirk <[email protected]>
All CI backends (cuda, hip, l0, opencl) currently support these three types of interop, that are probably the most important cases, but a complete e2e test did not previously exist. opencl is also the only backend with an existing test-e2e that calls make_* for these three sycl objects. This adds a short test-e2e corresponding to e.g. #10526. --------- Signed-off-by: JackAKirk <[email protected]>
This PR adds missing functions in the hip backend to allow for interoperability in programs that create sycl objects from native hip objects. The new function implementations are:
make_devicemake_queuemake_eventNote that it would really make sense for #10491 to be merged first because this PR makes the same code change in pi2ur, for a fix that is attributed to #10491.