Merge pull request elalish#126 from pca006132/dynamic-backend

dynamic backend implementation

.github/workflows/manifold.yml

@@ -14,7 +14,8 @@ jobs:
  build:
  strategy:
  matrix:
- backend: [CPP, OMP, TBB, CUDA]
+ cuda_support: [ON, OFF]
+ parallel_backend: [NONE, OMP, TBB]
  runs-on: ubuntu-20.04
  container:
  image: docker://nvidia/cuda:11.6.0-devel-ubuntu20.04
@@ -33,9 +34,10 @@ jobs:
  patch third_party/assimp/code/CMakeLists.txt < assimp.diff
  mkdir build
  cd build
- cmake -DCMAKE_BUILD_TYPE=Release -DTHRUST_BACKEND=${{matrix.backend}} .. && make
- - name: Test ${{matrix.backend}}
- if: matrix.backend != 'CUDA'
+ cmake -DCMAKE_BUILD_TYPE=Release -DMANIFOLD_PAR=${{matrix.parallel_backend}} -DMANIFOLD_USE_CUDA=${{matrix.cuda_support}} .. && make
+ - name: Test ${{matrix.parallel_backend}} with CUDA ${{matrix.cuda_support}}
+ # note that the test for CUDA backend does not really test CUDA, as we
+ # don't have CUDA GPU on GitHub Action
  run: |
  cd build/test
  ./manifold_test
@@ -74,8 +76,8 @@ jobs:
  build_windows:
  strategy:
  matrix:
- # backend: [CUDA, CPP]
- backend: [CPP]
+ parallel_backend: [NONE]
+ cuda_support: [OFF]
  max-parallel: 1
  runs-on: windows-2019
  steps:
@@ -107,11 +109,10 @@ jobs:
  - name: Build ${{matrix.backend}}
  shell: powershell
  run: |
- cmake . -DCMAKE_BUILD_TYPE=Release -B build -DTHRUST_BACKEND=${{matrix.backend}} -A x64
+ cmake . -DCMAKE_BUILD_TYPE=Release -B build -DMANIFOLD_PAR=${{matrix.parallel_backend}} -DMANIFOLD_USE_CUDA=${{matrix.cuda_support}} -A x64
  cd build
  cmake --build . --target ALL_BUILD --config Release
- - name: Test ${{matrix.backend}}
- if: matrix.backend != 'CUDA'
+ - name: Test ${{matrix.parallel_backend}} with CUDA ${{matrix.cuda_support}}
  shell: bash
  run: |
  cd build/test
@@ -120,7 +121,7 @@ jobs:
  build_nix:
  strategy:
  matrix:
- variant: [cpp, omp, cuda, js, tbb]
+ variant: [none, omp, tbb, none-cuda, omp-cuda, tbb-cuda, js]
  runs-on: ubuntu-latest
  steps:
  - uses: actions/checkout@v2.4.0

CMakeLists.txt

@@ -9,13 +9,14 @@ set(CMAKE_VERBOSE_MAKEFILE ON)
 # )
 # endif()
 
-set(THRUST_BACKEND "CUDA" CACHE STRING
- "Thrust backend, either \"CUDA\", \"OMP\" or \"CPP\"(single-threaded)"
-)
+option(MANIFOLD_USE_CUDA on)
+set(MANIFOLD_PAR "TBB" CACHE STRING "Parallel backend, either \"TBB\" or \"OpenMP\" or \"NONE\"")
+
 if(EMSCRIPTEN)
  message("Building for Emscripten")
  set(CMAKE_EXE_LINKER_FLAGS ${CMAKE_EXE_LINKER_FLAGS} -sALLOW_MEMORY_GROWTH=1)
- set(THRUST_BACKEND "CPP")
+ set(MANIFOLD_USE_CUDA OFF)
+ set(MANIFOLD_PAR "NONE")
 endif()
 
 option(BUILD_TEST_CGAL off)
@@ -26,8 +27,6 @@ set(ASSIMP_INC_DIR
  ${CMAKE_BINARY_DIR}/third_party/assimp/include)
 set(GLM_INC_DIR ${PROJECT_SOURCE_DIR}/third_party/glm)
 
-set_property(CACHE THRUST_BACKEND PROPERTY STRINGS CUDA OMP CPP)
-
 option(ASSIMP_FAST_BUILD "build ASSIMP just for tests" ON)
 if(ASSIMP_FAST_BUILD)
  option(ASSIMP_INSTALL FALSE)
@@ -39,39 +38,49 @@ if(ASSIMP_FAST_BUILD)
  endforeach()
 endif()
 
-if(MSVC)
- set(MANIFOLD_FLAGS)
-else()
- set(MANIFOLD_FLAGS -Werror -Wall -Wno-sign-compare -Wno-unused)
-endif()
 set(MANIFOLD_NVCC_FLAGS -Xcudafe --diag_suppress=esa_on_defaulted_function_ignored --extended-lambda)
 set(MANIFOLD_NVCC_RELEASE_FLAGS -O3)
 set(MANIFOLD_NVCC_DEBUG_FLAGS -G)
 
 set(THRUST_INC_DIR ${PROJECT_SOURCE_DIR}/third_party/thrust)
 
-message("Using ${THRUST_BACKEND} backend")
-if (THRUST_BACKEND STREQUAL "CUDA")
- # we cannot set THRUST_INC_DIR when building with CUDA, otherwise the
- # compiler will not use the system CUDA headers which causes incompatibility
- enable_language(CUDA)
- # clear THRUST_INC_DIR, we use the one from nvcc
- set(THRUST_INC_DIR "")
-elseif (THRUST_BACKEND STREQUAL "OMP")
+message("CUDA Support: ${MANIFOLD_USE_CUDA}")
+message("Parallel Backend: ${MANIFOLD_PAR}")
+
+if (MANIFOLD_PAR STREQUAL "OMP")
  find_package(OpenMP REQUIRED)
  set(MANIFOLD_INCLUDE OpenMP::OpenMP_CXX)
- set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_OMP -fopenmp)
-elseif (THRUST_BACKEND STREQUAL "CPP")
- set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CPP)
-elseif (THRUST_BACKEND STREQUAL "TBB")
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -fopenmp -DMANIFOLD_PAR='O')
+elseif (MANIFOLD_PAR STREQUAL "TBB")
  find_package(PkgConfig REQUIRED)
  pkg_check_modules(TBB REQUIRED tbb)
-
  set(MANIFOLD_INCLUDE ${TBB_LINK_LIBRARIES})
- set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_TBB)
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -DMANIFOLD_PAR='T')
+elseif (MANIFOLD_PAR STREQUAL "NONE")
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS})
+ set(MANIFOLD_PAR "CPP")
 else ()
- message(FATAL_ERROR "Invalid value for THRUST_BACKEND: ${THRUST_BACKEND}. "
- "Should be one of \"CUDA\", \"OMP\", \"CPP\" or \"TBB\"")
+ message(FATAL_ERROR "Invalid value for MANIFOLD_PAR: ${MANIFOLD_PAR}. "
+ "Should be one of \"TBB\", \"OMP\", \"CPP\"")
+endif()
+
+if (MANIFOLD_USE_CUDA)
+ enable_language(CUDA)
+ find_package(CUDA REQUIRED)
+ # we cannot set THRUST_INC_DIR when building with CUDA, otherwise the
+ # compiler will not use the system CUDA headers which causes incompatibility
+ # clear THRUST_INC_DIR, we use the one from nvcc
+ set(THRUST_INC_DIR "")
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -DMANIFOLD_USE_CUDA)
+ set(MANIFOLD_NVCC_FLAGS ${MANIFOLD_NVCC_FLAGS} ${MANIFOLD_FLAGS})
+else ()
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_${MANIFOLD_PAR})
+endif()
+
+if(MSVC)
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS})
+else()
+ set(MANIFOLD_FLAGS ${MANIFOLD_FLAGS} -Werror -Wall -Wno-sign-compare -Wno-unused)
 endif()
 
 add_subdirectory(utilities)

collider/CMakeLists.txt

@@ -2,7 +2,7 @@ project (collider)
 
 add_library(${PROJECT_NAME} src/collider.cpp)
 
-if(THRUST_BACKEND STREQUAL "CUDA")
+if(MANIFOLD_USE_CUDA)
  set_source_files_properties(src/collider.cpp PROPERTIES LANGUAGE CUDA)
  set_property(TARGET ${PROJECT_NAME} PROPERTY CUDA_ARCHITECTURES 61)
  target_compile_options(${PROJECT_NAME} 

collider/src/collider.cpp

@@ -14,6 +14,7 @@
 
 #include "collider.h"
 #include "utils.h"
+#include "par.h"
 
 #ifdef _MSC_VER
 #include <intrin.h>
@@ -252,7 +253,7 @@ Collider::Collider(const VecDH<Box>& leafBB,
  nodeParent_.resize(num_nodes, -1);
  internalChildren_.resize(leafBB.size() - 1, thrust::make_pair(-1, -1));
  // organize tree
- thrust::for_each_n(thrust::device, countAt(0), NumInternal(),
+ for_each_n(autoPolicy(NumInternal()), countAt(0), NumInternal(),
  CreateRadixTree({nodeParent_.ptrD(),
  internalChildren_.ptrD(), leafMorton}));
  UpdateBoxes(leafBB);
@@ -273,8 +274,8 @@ SparseIndices Collider::Collisions(const VecDH<T>& querriesIn) const {
  // scalar number of overlaps found
  VecDH<int> nOverlapsD(1, 0);
  // calculate Bounding Box overlaps
- thrust::for_each_n(
- thrust::device, zip(querriesIn.cbegin(), countAt(0)), querriesIn.size(),
+ for_each_n(
+ autoPolicy(querriesIn.size()), zip(querriesIn.cbegin(), countAt(0)), querriesIn.size(),
  FindCollisions<T>({querryTri.ptrDpq(), nOverlapsD.ptrD(), maxOverlaps,
  nodeBBox_.ptrD(), internalChildren_.ptrD()}));
  nOverlaps = nOverlapsD[0];
@@ -305,13 +306,13 @@ void Collider::UpdateBoxes(const VecDH<Box>& leafBB) {
  // copy in leaf node Boxs
  strided_range<VecDH<Box>::Iter> leaves(nodeBBox_.begin(), nodeBBox_.end(),
  2);
- thrust::copy(thrust::device, leafBB.cbegin(), leafBB.cend(), leaves.begin());
+ auto policy = autoPolicy(NumInternal());
+ copy(policy, leafBB.cbegin(), leafBB.cend(), leaves.begin());
  // create global counters
- VecDH<int> counter(NumInternal());
- thrust::fill(thrust::device, counter.begin(), counter.end(), 0);
+ VecDH<int> counter(NumInternal(), 0);
  // kernel over leaves to save internal Boxs
- thrust::for_each_n(
- thrust::device, countAt(0), NumLeaves(),
+ for_each_n(
+ policy, countAt(0), NumLeaves(),
  BuildInternalBoxes({nodeBBox_.ptrD(), counter.ptrD(), nodeParent_.ptrD(),
  internalChildren_.ptrD()}));
 }
@@ -330,7 +331,7 @@ bool Collider::Transform(glm::mat4x3 transform) {
  if (count != 2) axisAligned = false;
  }
  if (axisAligned) {
- thrust::for_each(thrust::device, nodeBBox_.begin(), nodeBBox_.end(),
+ for_each(autoPolicy(nodeBBox_.size()), nodeBBox_.begin(), nodeBBox_.end(),
  TransformBox({transform}));
  }
  return axisAligned;

flake.nix

@@ -10,33 +10,65 @@
  inherit system;
  config.allowUnfree = true;
  };
- manifold = { backend ? "CPP", doCheck ? true, build-tools ? [ ], runtime ? [ ] }: pkgs.stdenv.mkDerivation {
- inherit doCheck;
- pname = "manifold-${backend}";
- version = "beta";
- src = self;
- patches = [ ./assimp.diff ];
- nativeBuildInputs = (with pkgs; [ cmake python38 ]) ++ build-tools;
- buildInputs = runtime;
- cmakeFlags = [ "-DTHRUST_BACKEND=${backend}" ];
- checkPhase = ''
- cd test
- ./manifold_test
- cd ../
- '';
- installPhase = ''
- mkdir -p $out
- cp manifold/libmanifold.a $out/
- cp meshIO/libmeshIO.a $out/
- cp tools/loadMesh $out
- cp tools/perfTest $out
- '';
- };
+ manifold =
+ { parallel-backend ? "none"
+ , cuda-support ? false
+ , doCheck ? true
+ , build-tools ? [ ]
+ , ...
+ }: pkgs.stdenv.mkDerivation {
+ inherit doCheck;
+ pname =
+ if cuda-support then
+ "manifold-${parallel-backend}-cuda"
+ else
+ "manifold-${parallel-backend}";
+ version = "beta";
+ src = self;
+ patches = [ ./assimp.diff ];
+ nativeBuildInputs = (with pkgs; [ cmake python38 ]) ++ build-tools ++
+ (if cuda-support then [ pkgs.cudatoolkit_11_5 pkgs.addOpenGLRunpath ] else [ ]);
+ cmakeFlags = [
+ "-DMANIFOLD_PAR=${pkgs.lib.strings.toUpper parallel-backend}"
+ "-DMANIFOLD_USE_CUDA=${if cuda-support then "ON" else "OFF"}"
+ ];
+ checkPhase = ''
+ cd test
+ ./manifold_test
+ cd ../
+ '';
+ installPhase = ''
+ mkdir -p $out
+ cp manifold/libmanifold.a $out/
+ cp meshIO/libmeshIO.a $out/
+ cp tools/loadMesh $out
+ cp tools/perfTest $out
+ '';
+ };
+ parallelBackends = [
+ { parallel-backend = "none"; }
+ {
+ parallel-backend = "omp";
+ build-tools = [ pkgs.llvmPackages_13.openmp ];
+ }
+ {
+ parallel-backend = "tbb";
+ build-tools = with pkgs; [ tbb pkg-config ];
+ }
+ ];
+ buildMatrix = with pkgs; with lib; lists.flatten (map
+ (env: map
+ (x: x // env)
+ parallelBackends) [
+ { cuda-support = false; }
+ {
+ cuda-support = true;
+ }
+ ]);
  devShell = { additional ? [ ] }: pkgs.mkShell {
  buildInputs = with pkgs; [
  cmake
- ccls
- llvmPackages.openmp
+ llvmPackages_13.openmp
  clang-tools
  clang_13
  emscripten
@@ -45,48 +77,45 @@
  };
  in
  {
- packages.manifold-cpp = manifold { };
- packages.manifold-omp = manifold { backend = "OMP"; runtime = [ pkgs.llvmPackages.openmp ]; };
- packages.manifold-tbb = manifold { backend = "TBB"; runtime = [ pkgs.tbb pkgs.pkg-config ]; };
- packages.manifold-cuda = manifold {
- backend = "CUDA";
- runtime = [
- pkgs.cudaPackages.cudatoolkit_11
- ];
- doCheck = false;
- };
- packages.manifold-js = pkgs.buildEmscriptenPackage {
- name = "manifold-js";
- version = "beta";
- src = self;
- patches = [ ./assimp.diff ];
- nativeBuildInputs = (with pkgs; [ cmake python38 ]);
- buildInputs = [ pkgs.nodejs ];
- configurePhase = ''
- mkdir build
- cd build
- emcmake cmake -DCMAKE_BUILD_TYPE=Release ..
- '';
- buildPhase = ''
- emmake make
- '';
- checkPhase = ''
- cd test
- node manifold_test.js
- cd ../
- '';
- installPhase = ''
- mkdir -p $out
- cd tools
- cp *.js $out/
- cp *.wasm $out/
- '';
+ packages = (builtins.listToAttrs
+ (map
+ (x: {
+ name = "manifold-" + x.parallel-backend + (if
+ x.cuda-support then "-cuda" else "");
+ value = manifold x;
+ })
+ buildMatrix)) // {
+ manifold-js = pkgs.buildEmscriptenPackage {
+ name = "manifold-js";
+ version = "beta";
+ src = self;
+ patches = [ ./assimp.diff ];
+ nativeBuildInputs = (with pkgs; [ cmake python38 ]);
+ buildInputs = [ pkgs.nodejs ];
+ configurePhase = ''
+ mkdir build
+ cd build
+ emcmake cmake -DCMAKE_BUILD_TYPE=Release ..
+ '';
+ buildPhase = ''
+ emmake make
+ '';
+ checkPhase = ''
+ cd test
+ node manifold_test.js
+ cd ../
+ '';
+ installPhase = ''
+ mkdir -p $out
+ cd tools
+ cp *.js $out/
+ cp *.wasm $out/
+ '';
+ };
  };
  devShell = devShell { };
  devShells.cuda = devShell {
- additional = [
- pkgs.cudaPackages.cudatoolkit_11
- ];
+ additional = [ pkgs.cudatoolkit_11_5 ];
  };
  }
  );