Add SYCL kernels for t/pipelines/kernel (Feature, FillInLinearSystem, Registration, RGBDOdometry, TransformationConverter)

[Copilot]

Ports all CUDA kernels in cpp/open3d/t/pipelines/kernel/ to SYCL 2020, enabling ICP registration, RGBD odometry, SLAC optimization, and FPFH feature computation on Intel GPUs via oneAPI.

Type

• New feature (non-breaking change which adds functionality). Resolves #

Motivation and Context

Open3D had no SYCL equivalents for the CUDA kernels in t/pipelines/kernel/, leaving Intel GPU users unable to run these pipelines on SYCL devices.

Checklist:

• I have run python util/check_style.py --apply to apply Open3D code style
  to my code.
• This PR changes Open3D behavior or adds new functionality.
  • Both C++ (Doxygen) and Python (Sphinx / Google style) documentation is
    updated accordingly.
  • I have added or updated C++ and / or Python unit tests OR included test
    results (e.g. screenshots or numbers) here.
• I will follow up and update the code if CI fails.
• For fork PRs, I have selected Allow edits from maintainers.

Description

New SYCL implementation files

New file	CUDA equivalent	Key primitive
FeatureSYCL.cpp	FeatureCUDA.cu	queue.parallel_for
FillInLinearSystemSYCL.cpp	FillInLinearSystemCUDA.cu	sycl::atomic_ref
RegistrationSYCL.cpp	RegistrationCUDA.cu	sycl::atomic_ref (replaces CUB BlockReduce)
RGBDOdometrySYCL.cpp	RGBDOdometryCUDA.cu	sycl::atomic_ref (replaces CUB BlockReduce)
TransformationConverterSYCL.cpp	TransformationConverter.cu	queue.single_task

All reductions use sycl::atomic_ref<T, memory_order::acq_rel, memory_scope::device>.

Dispatch pattern

Each *CPU/CUDA dispatch site gains an else if (device.IsSYCL()) branch under #ifdef BUILD_SYCL_MODULE, matching the existing CUDA pattern exactly:

if (source_points.IsCPU()) {
    ComputePosePointToPlaneCPU(...);
} else if (source_points.IsCUDA()) {
    CUDA_CALL(ComputePosePointToPlaneCUDA, ...);
} else if (source_points.IsSYCL()) {
#ifdef BUILD_SYCL_MODULE
    ComputePosePointToPlaneSYCL(...);
#else
    utility::LogError("Not compiled with SYCL, but SYCL device is used.");
#endif
}

Build system

• All five SYCL files are listed under if(BUILD_SYCL_MODULE) in CMakeLists.txt using open3d_sycl_target_sources (not bare target_sources) so they receive -fsycl -fsycl-targets=... per-file compile options — the CI failure from missing this was caught and fixed in the same PR.

Code reuse

• Feature / FillInLinearSystem: Added OPEN3D_SKIP_FPFH_MAIN / OPEN3D_SKIP_FILL_IN_LS_MAIN guards in the shared *Impl.h headers. SYCL files define these macros before inclusion to reuse only the helper templates, then define their own SYCL main functions.
• ComputeRtPointToPoint: No custom SYCL kernel — implemented via existing tensor ops (IndexGet, SVD, Matmul) which already dispatch to SYCL.

---

💡 You can make Copilot smarter by setting up custom instructions, customizing its development environment and configuring Model Context Protocol (MCP) servers. Learn more Copilot coding agent tips in the docs.

[update-docs]

Thanks for submitting this pull request! The maintainers of this repository would appreciate if you could update the CHANGELOG.md based on your changes.