PyTorch xpu support (machine learning ops in sycl) - #7529
Open
ssheorey wants to merge 98 commits into
Open
Conversation
… Registration, RGBDOdometry, TransformationConverter Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
…L declaration Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
… to SYCL kernel files Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
…l-kernels-for-cuda
…e coverage Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
…ead of plain global atomics Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
…/Open3D into copilot/add-sycl-kernels-for-cuda
…als, RotatePoints, RotateNormals Co-authored-by: ssheorey <41028320+ssheorey@users.noreply.github.com>
…l-kernels-for-cuda
- Implemented SYCL device support in ImageImpl.h and PointCloud.h for various functions including ToSYCL, ClipTransformSYCL, and PyrDownDepthSYCL. - Added corresponding SYCL implementations in ImageSYCL.cpp and PointCloudSYCL.cpp. - Updated PointCloud.cpp and PointCloudImpl.h to handle SYCL devices in functions like Unproject, Project, and GetPointMaskWithinAABB. - Modified tests to accommodate SYCL devices, ensuring compatibility and skipping unsupported tests where necessary. - Introduced new test cases for SYCL in Python tests for nearest neighbor search and registration.
…rnels Implement SYCL custom kernels for TriangleMesh normals/areas and VoxelBlockGrid touch operations. Stub out remaining VoxelBlockGrid kernels due to core::HashMap limitations on SYCL, and add corresponding C++ and Python tests.
FP64 check missing hash function improved object copy fix early return bug in IndexAddContiguousSYCL use single workgroup launch to avoid atomics in ArgReduce (Check and revert) custom kernel for merging top-k results in KnnIndex (k<256) fix nanoflann call missing sort option Add RGB2Gray with tensor ops (for SYCL)
…L kernel invocations to use direct queue parallel_for calls. Add tests for nearest nbr and hashmap
SYCL nearest-neighbor search - Refactor tiled KNN/radius/hybrid into KnnSearchOpsSYCL.cpp + KnnSearchSYCLImpl.h so AddMM stays in the driver while top-K, count, and gather kernels are reusable and documented (small-k fused heap path vs legacy select/merge for large k). - Add configurable tile_bytes on KnnIndex/FixedRadiusIndex (defaults in NeighborSearchCommon.h) because integrated vs discrete Intel GPUs need different distance-tile sizes to stay cache-friendly without blowing memory. - Implement fused UpdateTopKFromTile for k ≤ 512, per-query threshold handling for radius/hybrid (radius² − |q|²), and finalize/clamp rules (C1/C4) so distances are non-negative and ties break by index like CPU/CUDA. - Extend C++/Python SYCL NNS tests (parity, coincident query, tie-break, radius, hybrid) to lock in correctness after the algorithm rewrite. SYCL hash map backend - Pack slot state, buf_index, and fingerprint into one uint64 per bucket to cut probe traffic and skip key-buffer loads on fingerprint mismatch. - Use power-of-two buckets with HashMix (fmix64) so probing uses masks instead of 64-bit modulo on GPU, and reserve/rehash when tombstones fill the table (GetNonEmptyCount + HashMap::Insert/Activate checks), not only live size. - Harden Insert for Intel Xe L1 coherence (seq_cst fences, LOCKED slots, restart-on-LOCKED instead of subgroup spin) to avoid stale keys and hangs. - Vectorize value copies via SYCLBlockCopyDispatch; improve GetActiveIndices with work-group scan + one atomic per group instead of per-slot atomics. - SYCLHashDeviceLookup uses plain loads when the table is read-only during raycast-style kernels. Core SYCL utilities - Add SYCLBlockCopyDispatch.h and use it in CopySYCL for object dtypes so copies use wide vector loads/stores instead of per-element queue.memcpy. Build and tooling (SYCL-without-CUDA / local dev) - Gate OPEN3D_CUDA_COMPILER_* defines and CompilerInfo CUDA strings on BUILD_CUDA_MODULE so SYCL-only builds do not reference undefined CUDA macros. - Add ENABLE_SANITIZER CMake option and wire -fsanitize into Open3D when set. - Comment out optional EGL/X11 linking block in cpp/open3d/CMakeLists.txt (local build adjustment—confirm this is intended before upstreaming).
Some optimizations. (e.g. restrict) lock free hashmap insert: write buffer then CAS design is correct and fast, but leaves holes in the data buffer. Center data before Knn, if using expanded L2 distance formula (p^2+q^2-2pq) to prevent cancellation bug in indexer: TensorIterator::GetPtr() incorrect for non-contiguous. Added Knn search benchmark
- SYCLContext: process-wide static singleton (was thread_local) to avoid per-thread SYCL contexts/USM mismatches; cache all device properties in one place (SYCLContext::Impl) and expose via GetDeviceProperties(). - Add SYCL launch helpers (SYCLPreferredWorkGroupSize, SYCLNdRange1D, group-reduction helpers) and use nd_range<1> + sycl::reduce_over_group across ParallelFor, elementwise, and reduction kernels instead of flat parallel_for(n) / per-output kernel launches. - ReductionSYCL: on-device GetInputPtrDevice() enables one kernel with one work-group per output for multi-output reductions (incl. arg-reductions). - Registration/RGBD odometry/SLAC kernels: accumulate AtA/Atb/residual in SLM per work-group (restrict-qualified pointers) instead of global atomics. - Rename SYCLBlockCopyDispatch.h -> BlockCopyDispatch.h, generalize the vectorized object-copy dispatch (up to 64-byte blocks) shared by the hash map and tensor copy paths; align CUDA hashmap Dispatch.h divisors to match. - Build: set -fsycl-max-parallel-link-jobs, prefer lld linker when available. - Update ParallelFor/Reduction benchmarks and Linalg/Tensor tests for the refactored SYCLContext API.
Add core::sy::IsCPUDevice() to detect the SYCL CPU fallback device and use it to throw clear errors from LeastSquaresSYCL (gels_batch), the SYCL hash map, and FixedRadiusSearch/HybridSearch, which are broken on SYCL CPU. Skip the corresponding C++ tests (HashMap, NNS, VoxelBlockGrid, Registration, Feature, PointCloud) and opt affected Python tests out of the SYCL CPU fallback via list_devices(also_sycl_cpu=False).
…ync everywhere possible.
… regression - cutlass::arch::global_load/global_store in the vendored sycl-tla only implemented the IEEE-fp32 GEMM epilogue's device read/write via CUDA PTX, guarded by a macro that's never defined on genuine SPIR-V/Intel targets. The epilogue silently no-op'd on real XPU hardware, producing all-zero GEMM output. Add a SYCL-native fallback and fold it into the existing sycl_tla ieee-gemm patch. This was the root cause of the test_sparseconv/test_knn_search/test_nms ml_ops failures on XPU. - Fix test_sparseconv_allow_tf32 comparing two differently-initialized SparseConv layers instead of toggling allow_tf32 on a single layer. - Remove leftover debug instrumentation from ContinuousConvSYCL.h. - Restore the global index offset in CUDA KnnSearchOps.cu's multi-batch combine step, lost during the ready_event/user_stream refactor; without it, batches after the first returned point indices local to their own slice instead of into the full points tensor. - Document verification results and known pre-existing/out-of-scope test failures in docs/dev/pr_7529_flagged_issues.md. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
# Conflicts: # .github/workflows/ubuntu-sycl.yml # cpp/open3d/core/hashmap/SYCL/SYCLHashBackend.h
Add TBB-parallelized CPU implementations of three_nn, three_interpolate,
three_interpolate_grad, ball_query, and furthest_point_sampling so these
ops work on CPU-only (e.g. macOS) builds instead of throwing
TORCH_CHECK(false, ...).
- New BallQueryKernelCPU.cpp, InterpolateKernelCPU.cpp,
SamplingKernelCPU.cpp with CPU launchers, parallelized with
tbb::parallel_for (flattened iteration spaces to avoid nested
parallelism; scatter-add grads partitioned by (batch, channel) slice
to avoid data races without atomics).
- Launcher functions use a _cpu suffix to avoid multiple-definition
link errors now that CPU kernel files are always compiled alongside
the CUDA .cu variants (matches the existing convention used by
misc/*.cpp ops of always compiling the CPU kernel and gating only
the CUDA/SYCL variants).
- CMakeLists.txt: removed the conditional guard around the CPU kernel
sources; they now build unconditionally like other ops.
- Fixed a stray unbalanced #endif in BallQueryOps.cpp.
- Enabled the corresponding CPU test paths in
python/test/ml_ops/{test_query_pts,test_sampling,test_three_interp,
test_three_nn}.py, adjusting a couple of assertions to tolerate
legitimate floating-point tie/reduction-order differences between
the CPU (serial) and CUDA (parallel-tree) implementations.
- python/open3d/ml/torch/__init__.py: fix Windows DLL search paths for
the SYCL/XPU runtime, verified against actual PyPI wheel layouts
(intel-sycl-rt and friends install to <sys.prefix>/Library/bin; the
torch+xpu wheel bundles its own DLLs under torch/lib).
Replace hardcoded torch/torchvision cu126 index-url installs with the pinned requirements-torch-cuda.txt from Open3D-ML, matching the pattern already used for the CPU and XPU cases in both windows.yml and util/ci_utils.sh's test_wheel().
Isolate sharded gtest temp files via per-shard TMPDIR, restrict PointNet ml_ops tests to GPU frameworks plus torch CPU, always register trilinear devoxelize torch ops for CPU wheels, and widen Windows DLL paths when generating torch op wrappers. Co-authored-by: Cursor <cursoragent@cursor.com>
Escape TMPDIR paths so the host shell does not expand shard variables under set -u, and broaden Windows DLL discovery when loading open3d_torch_ops for torch op wrapper generation. Co-authored-by: Cursor <cursoragent@cursor.com>
…oad. Set per-shard OPEN3D_DATA_ROOT alongside TMPDIR in SYCL docker tests so parallel shards do not race on shared download files, and add Open3D.dll to the DLL search path when generating torch op wrappers on Windows. Co-authored-by: Cursor <cursoragent@cursor.com>
CMake treats semicolons in custom COMMAND as argument separators; pass Open3D.dll location via a separate --open3d_lib_dir flag instead. Co-authored-by: Cursor <cursoragent@cursor.com>
Dataset gtests use ~/open3d_data; parallel shards still raced when only OPEN3D_DATA_ROOT was set. Use per-shard HOME and TMPDIR subdirs. Co-authored-by: Cursor <cursoragent@cursor.com>
Remove SYCL gtest sharding and the duplicate bare docker gtest step. Add docker_test.sh phases (cpp, lib, python, all) and split Ubuntu, OpenBLAS, and SYCL workflows so C++ runs once per config and pytest runs once per Python inside the CI image. Run build-wheel jobs even when build-lib fails; document the layout in workflows README. Co-authored-by: Cursor <cursoragent@cursor.com>
…sts. Restore and extend the CPU/CUDA/SYCL routing guide with Sphinx and API links; clean up maintainer comments in SYCL contrib and conv paths; expand PyTorch allow_tf32 coverage in existing ml_ops tests. Co-authored-by: Cursor <cursoragent@cursor.com>
|
Check out this pull request on See visual diffs & provide feedback on Jupyter Notebooks. Powered by ReviewNB |
Keep allow_tf32 on torch-only paths so TensorFlow gradient tests are unchanged, and align continuous conv TF32 smoke with matching in/out channel counts for transpose. Co-authored-by: Cursor <cursoragent@cursor.com>
Use forward output features for transpose allow_tf32 checks; narrow infeat scope and drop unused RNG in NMS test. Co-authored-by: Cursor <cursoragent@cursor.com>
Register Intel pip runtime paths (*.data/Library/bin), torch/lib for XPU builds, and PATH for transitive deps so pybind loads in CI wheel tests. Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Type
Motivation and Context
Checklist:
python util/check_style.py --applyto apply Open3D code styleto my code.
updated accordingly.
results (e.g. screenshots or numbers) here.
Description