From da95b80a6cfd92f68b750bb6559fa44fbec08f80 Mon Sep 17 00:00:00 2001 From: Sameer Sheorey Date: Mon, 30 Dec 2024 10:25:57 +0100 Subject: [PATCH 1/3] Build and upload SYCL wheel --- .github/workflows/ubuntu-sycl.yml | 17 +++++++++++++++++ 3rdparty/find_dependencies.cmake | 7 +++++-- CMakeLists.txt | 10 ++++++++++ cmake/Open3DSYCLTargetSources.cmake | 3 +-- docker/Dockerfile.ci | 1 + docker/docker_build.sh | 16 ++++++++++++---- util/ci_utils.sh | 8 ++++++-- 7 files changed, 52 insertions(+), 10 deletions(-) diff --git a/.github/workflows/ubuntu-sycl.yml b/.github/workflows/ubuntu-sycl.yml index daad144ce08..46c274925b0 100644 --- a/.github/workflows/ubuntu-sycl.yml +++ b/.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' diff --git a/3rdparty/find_dependencies.cmake b/3rdparty/find_dependencies.cmake index 62ea33ee1fe..f4362c08c77 100644 --- a/3rdparty/find_dependencies.cmake +++ b/3rdparty/find_dependencies.cmake @@ -1538,8 +1538,11 @@ if(BUILD_SYCL_MODULE) target_link_libraries(3rdparty_sycl INTERFACE $<$,$>>:sycl>) target_link_options(3rdparty_sycl INTERFACE - $<$,$>>:-fsycl -fsycl-targets=intel_gpu_acm_g10>) - # $<$,$>>:-fsycl -fsycl-targets=spir64,spir64_gen>) + $<$,$>>:-fsycl -fsycl-targets=${OPEN3D_SYCL_TARGETS}>) + if (OPEN3D_SYCL_TARGET_BACKEND_OPTIONS) + target_link_options(3rdparty_sycl INTERFACE + $<$,$>>:-Xs ${OPEN3D_SYCL_TARGET_BACKEND_OPTIONS}>) + endif() if(NOT BUILD_SHARED_LIBS OR arg_PUBLIC) install(TARGETS 3rdparty_sycl EXPORT Open3DTargets) endif() diff --git a/CMakeLists.txt b/CMakeLists.txt index 635e43e6abf..cc2cace76a3 100644 --- a/CMakeLists.txt +++ b/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)) diff --git a/cmake/Open3DSYCLTargetSources.cmake b/cmake/Open3DSYCLTargetSources.cmake index baa075f38e7..3314cb24303 100644 --- a/cmake/Open3DSYCLTargetSources.cmake +++ b/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() diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 189fa6c50b0..1dc37163c84 100755 --- a/docker/Dockerfile.ci +++ b/docker/Dockerfile.ci @@ -238,6 +238,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." diff --git a/docker/docker_build.sh b/docker/docker_build.sh index e6c22bb10df..bbbeebfc40d 100755 --- a/docker/docker_build.sh +++ b/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() { diff --git a/util/ci_utils.sh b/util/ci_utils.sh index 1adb4b7d8c5..68687127fc3 100644 --- a/util/ci_utils.sh +++ b/util/ci_utils.sh @@ -21,6 +21,7 @@ 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 @@ -28,6 +29,7 @@ LOW_MEM_USAGE=${LOW_MEM_USAGE:-OFF} 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,11 @@ 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 + if [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "OFF" ]]; then python -m pip install -U "${TORCH_GLNX}" -f "$TORCH_REPO_URL" tensorboard - + elif [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "ON" ]]; then + python -m pip install -U "${TORCH_GLNX}+cpu.cxx11.abi" -i "$TORCH_CXX11_URL" + python -m pip install -U tensorboard elif [[ "$OSTYPE" == "darwin"* ]]; then python -m pip install -U torch=="$TORCH_VER" -f "$TORCH_REPO_URL" tensorboard else From 973670ee92268ac55cabde5a8eb7db6e0f3d578d Mon Sep 17 00:00:00 2001 From: Sameer Sheorey Date: Mon, 30 Dec 2024 22:47:08 +0100 Subject: [PATCH 2/3] MAX_DIMS=5 for VoxelBlockGrid, CI error fixes. --- cpp/open3d/core/Indexer.h | 6 +++--- cpp/open3d/core/kernel/NonZeroSYCL.cpp | 13 +++++++------ util/ci_utils.sh | 2 +- 3 files changed, 11 insertions(+), 10 deletions(-) diff --git a/cpp/open3d/core/Indexer.h b/cpp/open3d/core/Indexer.h index 1f6627e9d1a..76f99acba5d 100644 --- a/cpp/open3d/core/Indexer.h +++ b/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; diff --git a/cpp/open3d/core/kernel/NonZeroSYCL.cpp b/cpp/open3d/core/kernel/NonZeroSYCL.cpp index 9a9795e9f6a..6a1a4817b7f 100644 --- a/cpp/open3d/core/kernel/NonZeroSYCL.cpp +++ b/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 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 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(num_non_zeros)}, Int64, device); diff --git a/util/ci_utils.sh b/util/ci_utils.sh index 68687127fc3..4387e4e10a4 100644 --- a/util/ci_utils.sh +++ b/util/ci_utils.sh @@ -82,7 +82,7 @@ install_python_dependencies() { if [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "OFF" ]]; then python -m pip install -U "${TORCH_GLNX}" -f "$TORCH_REPO_URL" tensorboard elif [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "ON" ]]; then - python -m pip install -U "${TORCH_GLNX}+cpu.cxx11.abi" -i "$TORCH_CXX11_URL" + python -m pip install -U "${TORCH_GLNX}.cxx11.abi" -i "$TORCH_CXX11_URL" python -m pip install -U tensorboard elif [[ "$OSTYPE" == "darwin"* ]]; then python -m pip install -U torch=="$TORCH_VER" -f "$TORCH_REPO_URL" tensorboard From 061059ddba79a7bbdbe368caabb376ef6984fc08 Mon Sep 17 00:00:00 2001 From: Sameer Sheorey Date: Mon, 30 Dec 2024 23:53:14 +0100 Subject: [PATCH 3/3] Install g++-11 for OneDPL TBB backend --- docker/Dockerfile.ci | 4 ++++ util/ci_utils.sh | 5 +++-- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 1dc37163c84..2f4eff65273 100755 --- a/docker/Dockerfile.ci +++ b/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 diff --git a/util/ci_utils.sh b/util/ci_utils.sh index 4387e4e10a4..f7c5c73521b 100644 --- a/util/ci_utils.sh +++ b/util/ci_utils.sh @@ -80,10 +80,11 @@ install_python_dependencies() { fi if [ "$BUILD_PYTORCH_OPS" == "ON" ]; then # ML/requirements-torch.txt if [[ "$OSTYPE" == "linux-gnu"* && "$BUILD_SYCL_MODULE" == "OFF" ]]; then - python -m pip install -U "${TORCH_GLNX}" -f "$TORCH_REPO_URL" tensorboard + 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 -U tensorboard + python -m pip install tensorboard elif [[ "$OSTYPE" == "darwin"* ]]; then python -m pip install -U torch=="$TORCH_VER" -f "$TORCH_REPO_URL" tensorboard else