Merge branch 'ss/sycl-tensor' of github.com:isl-org/Open3D into ss/sy…

…cl-tensor

.github/workflows/ubuntu-sycl.yml

@@ -49,6 +49,23 @@ jobs:
             docker/docker_test.sh sycl-static
           fi
 
+      - name: Upload wheel to GitHub artifacts
+        if: ${{ matrix.BUILD_SHARED_LIBS == 'ON' }}
+        uses: actions/upload-artifact@v4
+        with:
+          name: open3d-sycl-linux-python-wheel
+          path: open3d-*.whl
+          if-no-files-found: error
+      - name: Update devel release
+        if: ${{ github.ref == 'refs/heads/main' }}
+        env:
+          GH_TOKEN: ${{ github.token }}
+        run: |
+          if [ ${{ matrix.BUILD_SHARED_LIBS }} == 'ON' ] ; then
+            gh release upload main-devel open3d-*.whl --clobber
+          fi
+          gh release view main-devel
+
       - name: GCloud CLI auth
         if: ${{ github.ref == 'refs/heads/main' }}
         uses: 'google-github-actions/auth@v2'

3rdparty/find_dependencies.cmake

@@ -1538,8 +1538,11 @@ if(BUILD_SYCL_MODULE)
     target_link_libraries(3rdparty_sycl INTERFACE
         $<$<AND:$<CXX_COMPILER_ID:IntelLLVM>,$<NOT:$<LINK_LANGUAGE:ISPC>>>:sycl>)
     target_link_options(3rdparty_sycl INTERFACE
-        $<$<AND:$<CXX_COMPILER_ID:IntelLLVM>,$<NOT:$<LINK_LANGUAGE:ISPC>>>:-fsycl -fsycl-targets=intel_gpu_acm_g10>)
-        # $<$<AND:$<CXX_COMPILER_ID:IntelLLVM>,$<NOT:$<LINK_LANGUAGE:ISPC>>>:-fsycl -fsycl-targets=spir64,spir64_gen>)
+        $<$<AND:$<CXX_COMPILER_ID:IntelLLVM>,$<NOT:$<LINK_LANGUAGE:ISPC>>>:-fsycl -fsycl-targets=${OPEN3D_SYCL_TARGETS}>)
+    if (OPEN3D_SYCL_TARGET_BACKEND_OPTIONS)
+        target_link_options(3rdparty_sycl INTERFACE
+            $<$<AND:$<CXX_COMPILER_ID:IntelLLVM>,$<NOT:$<LINK_LANGUAGE:ISPC>>>:-Xs ${OPEN3D_SYCL_TARGET_BACKEND_OPTIONS}>)
+    endif()
     if(NOT BUILD_SHARED_LIBS OR arg_PUBLIC)
         install(TARGETS 3rdparty_sycl EXPORT Open3DTargets)
     endif()

CMakeLists.txt

@@ -63,6 +63,16 @@ else()
     option(STATIC_WINDOWS_RUNTIME "Use static (MT/MTd) Windows runtime"      ON )
 endif()
 option(BUILD_SYCL_MODULE          "Build SYCL module with Intel oneAPI"      OFF)
+if(BUILD_SYCL_MODULE)
+    set(OPEN3D_SYCL_TARGETS "spir64" CACHE STRING 
+    "SYCL targets: spir64 for JIT, or another for AOT compilation. See https://github.com/intel/llvm/blob/sycl/sycl/doc/UsersManual.md."
+)
+    set(OPEN3D_SYCL_TARGET_BACKEND_OPTIONS "" CACHE STRING 
+    "SYCL target backend options, e.g. to compile for a specific device. See https://github.com/intel/llvm/blob/sycl/sycl/doc/UsersManual.md."
+)
+    set(BUILD_ISPC_MODULE OFF CACHE BOOL "Build the ISPC module" FORCE)
+    set(BUILD_CUDA_MODULE OFF CACHE BOOL "Build the CUDA module" FORCE)
+endif()
 option(GLIBCXX_USE_CXX11_ABI      "Set -D_GLIBCXX_USE_CXX11_ABI=1"           ON )
 option(ENABLE_SYCL_UNIFIED_SHARED_MEMORY "Enable SYCL unified shared memory" OFF)
 if(BUILD_GUI AND (WIN32 OR UNIX AND NOT LINUX_AARCH64 AND NOT APPLE_AARCH64))

cmake/Open3DSYCLTargetSources.cmake

@@ -32,8 +32,7 @@ function(open3d_sycl_target_sources target)
         if(BUILD_SYCL_MODULE)
             foreach(sycl_file IN LISTS arg_UNPARSED_ARGUMENTS)
                 set_source_files_properties(${sycl_file} PROPERTIES
-                    COMPILE_OPTIONS "-fsycl;-fsycl-targets=intel_gpu_acm_g10")
-                    #COMPILE_OPTIONS "-fsycl;-fsycl-targets=spir64,spir64_gen")
+                    COMPILE_OPTIONS "-fsycl;-fsycl-targets=${OPEN3D_SYCL_TARGETS}")
                 if(arg_VERBOSE)
                     message(STATUS "open3d_sycl_target_sources(${target}): marked ${sycl_file} as SYCL code")
                 endif()

cpp/open3d/core/Indexer.h

@@ -34,11 +34,11 @@ class Indexer;
 class IndexerIterator;
 
 // Maximum number of dimensions of TensorRef.
-static constexpr int64_t MAX_DIMS = 4;
+static constexpr int64_t MAX_DIMS = 5;
 
 // Maximum number of inputs of an op.
 // MAX_INPUTS shall be >= MAX_DIMS to support advanced indexing.
-static constexpr int64_t MAX_INPUTS = 4;
+static constexpr int64_t MAX_INPUTS = 5;
 
 // Maximum number of outputs of an op. This number can be increased when
 // necessary.
@@ -638,7 +638,7 @@ class Indexer {
 class IndexerIterator {
 public:
     struct Iterator {
-        Iterator(){};
+        Iterator() {};
         Iterator(const Indexer& indexer);
         Iterator(Iterator&& other) = default;
 

cpp/open3d/core/kernel/NonZeroSYCL.cpp

@@ -16,10 +16,6 @@
 
 namespace open3d {
 namespace core {
-/// Maximum number of dimensions of TensorRef, rounded up to the power of 2 for
-/// sycl::vec support.
-static constexpr int64_t MAX_DIMS_POW2 = 4;
-static_assert(MAX_DIMS_POW2 >= MAX_DIMS, "MAX_DIMS_POW2 too small.");
 namespace kernel {
 
 Tensor NonZeroSYCL(const Tensor& src) {
@@ -48,8 +44,13 @@ Tensor NonZeroSYCL(const Tensor& src) {
     // Transform flattened indices to indices in each dimension.
     const auto num_dims = src.NumDims();
     SizeVector shape = src.GetShape();
-    sycl::vec<int64_t, MAX_DIMS_POW2> shape_vec;  // device copyable
-    OPEN3D_ASSERT(shape.size() <= MAX_DIMS_POW2 && "Too many dimensions.");
+    // MAX_DIMS: Maximum number of dimensions of TensorRef, defined in
+    // Indexer.h.
+    sycl::marray<int64_t, MAX_DIMS> shape_vec;  // device copyable
+    if (shape.size() <= MAX_DIMS) {
+        utility::LogError("Too many dimensions: {} > MAX_DIMS={}.",
+                          shape.size(), MAX_DIMS);
+    }
     for (auto k = 0; k < num_dims; ++k) shape_vec[k] = shape[k];
     Tensor result({num_dims, static_cast<int64_t>(num_non_zeros)}, Int64,
                   device);

docker/Dockerfile.ci

@@ -87,7 +87,11 @@ RUN apt-get update && apt-get install -y \
     libxmlsec1-dev \
     libffi-dev \
     liblzma-dev \
+ && if [ "${BUILD_SYCL_MODULE}" = "ON" ]; then \
+        apt-get install g++-11; \
+    fi \
  && rm -rf /var/lib/apt/lists/*
+# OneDPL TBB backend requires libstdc++ >= v11
 
 # pyenv
 # The pyenv python paths are used during docker run, in this way docker run
@@ -238,6 +242,7 @@ RUN ccache -s \
  && tar -caf /${CCACHE_TAR_NAME}.tar.xz ${CCACHE_DIR_NAME} \
  && if [ "${PACKAGE}" = "ON" ]; then mv /root/Open3D/build/package/open3d-devel*.tar.xz /; fi \
  && if [ "${PACKAGE}" = "VIEWER" ]; then mv /root/Open3D/build/package-Open3DViewer-deb/open3d-viewer-*-Linux.deb /; fi \
+ && if [ "${BUILD_SYCL_MODULE}" = "ON"  && "${BUILD_SHARED_LIBS}" = "ON"]; then mv /root/Open3D/build/lib/python_package/pip_package/open3d-*.whl /; fi \
  && ls -alh /
 
 RUN echo "Docker build done."

docker/docker_build.sh

@@ -451,13 +451,17 @@ sycl-shared_export_env() {
     export BASE_IMAGE=intel/oneapi-basekit:2024.1.0-devel-ubuntu20.04
     export DEVELOPER_BUILD=ON
     export CCACHE_TAR_NAME=open3d-ci-sycl
-    export PYTHON_VERSION=3.8
+    export PYTHON_VERSION=3.10
     export BUILD_SHARED_LIBS=ON
     export BUILD_CUDA_MODULE=OFF
-    export BUILD_TENSORFLOW_OPS=OFF
-    export BUILD_PYTORCH_OPS=OFF
+    export BUILD_TENSORFLOW_OPS=ON
+    export BUILD_PYTORCH_OPS=ON
     export PACKAGE=OFF
     export BUILD_SYCL_MODULE=ON
+
+    export IGC_EnableDPEmulation=1          # Enable float64 emulation during compilation
+    export SYCL_CACHE_PERSISTENT=1          # Cache SYCL kernel binaries.
+    export OverrideDefaultFP64Settings=1    # Enable double precision emulation at runtime.
 }
 
 sycl-static_export_env() {
@@ -468,13 +472,17 @@ sycl-static_export_env() {
     export BASE_IMAGE=intel/oneapi-basekit:2024.1.0-devel-ubuntu20.04
     export DEVELOPER_BUILD=ON
     export CCACHE_TAR_NAME=open3d-ci-sycl
-    export PYTHON_VERSION=3.8
+    export PYTHON_VERSION=3.10
     export BUILD_SHARED_LIBS=OFF
     export BUILD_CUDA_MODULE=OFF
     export BUILD_TENSORFLOW_OPS=OFF
     export BUILD_PYTORCH_OPS=OFF
     export PACKAGE=OFF
     export BUILD_SYCL_MODULE=ON
+
+    export IGC_EnableDPEmulation=1          # Enable float64 emulation during compilation
+    export SYCL_CACHE_PERSISTENT=1          # Cache SYCL kernel binaries.
+    export OverrideDefaultFP64Settings=1    # Enable double precision emulation at runtime.
 }
 
 function main() {

util/ci_utils.sh

@@ -21,13 +21,15 @@ fi
 BUILD_TENSORFLOW_OPS=${BUILD_TENSORFLOW_OPS:-ON}
 BUILD_PYTORCH_OPS=${BUILD_PYTORCH_OPS:-ON}
 LOW_MEM_USAGE=${LOW_MEM_USAGE:-OFF}
+BUILD_SYCL_MODULE=${BUILD_SYCL_MODULE:-OFF}
 
 # Dependency versions:
 # CUDA: see docker/docker_build.sh
 # ML
 TENSORFLOW_VER="2.16.2"
 TORCH_VER="2.2.2"
 TORCH_REPO_URL="https://download.pytorch.org/whl/torch/"
+TORCH_CXX11_URL="https://download.pytorch.org/whl/"
 # Python
 PIP_VER="23.2.1"
 WHEEL_VER="0.38.4"
@@ -77,9 +79,12 @@ install_python_dependencies() {
         python -m pip install -U "$TF_ARCH_NAME"=="$TENSORFLOW_VER" # ML/requirements-tensorflow.txt
     fi
     if [ "$BUILD_PYTORCH_OPS" == "ON" ]; then # ML/requirements-torch.txt
-        if [[ "$OSTYPE" == "linux-gnu"* ]]; then
-            python -m pip install -U "${TORCH_GLNX}" -f "$TORCH_REPO_URL" tensorboard
-
+        if [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "OFF" ]]; then
+            python -m pip install -U "${TORCH_GLNX}" -f "$TORCH_REPO_URL" 
+            python -m pip install tensorboard
+        elif [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "ON" ]]; then
+            python -m pip install -U "${TORCH_GLNX}.cxx11.abi" -i "$TORCH_CXX11_URL"
+            python -m pip install tensorboard
         elif [[ "$OSTYPE" == "darwin"* ]]; then
             python -m pip install -U torch=="$TORCH_VER" -f "$TORCH_REPO_URL" tensorboard
         else