diff --git a/CMakePresets.json b/CMakePresets.json
index 29dd152ed01..270433b7103 100644
--- a/CMakePresets.json
+++ b/CMakePresets.json
@@ -167,6 +167,14 @@
"type": "BOOL",
"value": "ON"
},
+ "SOFA_FETCH_SOFADISTANCEGRID": {
+ "type": "BOOL",
+ "value": "ON"
+ },
+ "PLUGIN_SOFADISTANCEGRID": {
+ "type": "BOOL",
+ "value": "ON"
+ },
"PLUGIN_SOFADISTANCEGRID_CUDA": {
"type": "BOOL",
"value": "ON"
@@ -248,10 +256,6 @@
"type": "BOOL",
"value": "ON"
},
- "PLUGIN_SOFADISTANCEGRID": {
- "type": "BOOL",
- "value": "ON"
- },
"PLUGIN_SOFAEULERIANFLUID": {
"type": "BOOL",
"value": "ON"
diff --git a/applications/plugins/CMakeLists.txt b/applications/plugins/CMakeLists.txt
index 736fbf72423..17f28750931 100644
--- a/applications/plugins/CMakeLists.txt
+++ b/applications/plugins/CMakeLists.txt
@@ -72,5 +72,5 @@ else()
message("Sofa.GL not found; disabling VolumetricRendering plugin")
endif()
-sofa_add_subdirectory(plugin SofaDistanceGrid SofaDistanceGrid) # Also defines SofaDistanceGrid.CUDA
+sofa_add_external(plugin SofaDistanceGrid GIT_REF master GIT_REPOSITORY https://www.github.com/sofa-framework/SofaDistanceGrid.git) # Also defines SofaDistanceGrid.CUDA
sofa_add_subdirectory(plugin SofaImplicitField SofaImplicitField)
diff --git a/applications/plugins/SofaDistanceGrid/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/CMakeLists.txt
deleted file mode 100644
index 643ddab08b1..00000000000
--- a/applications/plugins/SofaDistanceGrid/CMakeLists.txt
+++ /dev/null
@@ -1,94 +0,0 @@
-cmake_minimum_required(VERSION 3.22)
-project(SofaDistanceGrid LANGUAGES CXX)
-
-option(SOFADISTANCEGRID_USE_MINIFLOWVR "Build and use the miniFlowVR library with DistanceGrid." OFF)
-if(SOFADISTANCEGRID_USE_MINIFLOWVR)
- add_subdirectory(extlibs/miniFlowVR)
- set_target_properties(miniFlowVR PROPERTIES DEBUG_POSTFIX "_d")
- #set_target_properties(miniFlowVR PROPERTIES FOLDER "SofaExtlibs")
-endif()
-
-option(SOFADISTANCEGRID_USE_MESHCONV "Build meshconv, mainly used to generate DistanceGrid file. Needs miniFlowVR." OFF)
-if(SOFADISTANCEGRID_USE_MESHCONV)
- find_package(MiniFlowVR QUIET)
- if(MiniFlowVR_FOUND)
- sofa_add_subdirectory(application applications/meshconv meshconv OFF)
- else()
- message(WARNING "meshconv needs miniFlowVR; enable SOFADISTANCEGRID_USE_MINIFLOWVR to build it.")
- endif()
-endif()
-
-set(SOFADISTANCEGRID_SRC "src/${PROJECT_NAME}")
-
-set(HEADER_FILES
- ${SOFADISTANCEGRID_SRC}/config.h.in
- ${SOFADISTANCEGRID_SRC}/initSofaDistanceGrid.h
- ${SOFADISTANCEGRID_SRC}/DistanceGrid.h
- ${SOFADISTANCEGRID_SRC}/components/collision/FFDDistanceGridDiscreteIntersection.h
- ${SOFADISTANCEGRID_SRC}/components/collision/FFDDistanceGridDiscreteIntersection.inl
- ${SOFADISTANCEGRID_SRC}/components/collision/RigidDistanceGridDiscreteIntersection.h
- ${SOFADISTANCEGRID_SRC}/components/collision/RigidDistanceGridDiscreteIntersection.inl
- ${SOFADISTANCEGRID_SRC}/components/collision/DistanceGridCollisionModel.h
- ${SOFADISTANCEGRID_SRC}/components/forcefield/DistanceGridForceField.h
- ${SOFADISTANCEGRID_SRC}/components/forcefield/DistanceGridForceField.inl
-)
-
-set(SOURCE_FILES
- ${SOFADISTANCEGRID_SRC}/initSofaDistanceGrid.cpp
- ${SOFADISTANCEGRID_SRC}/DistanceGrid.cpp
- ${SOFADISTANCEGRID_SRC}/RegisterModelToCollisionFactory.cpp
- ${SOFADISTANCEGRID_SRC}/components/forcefield/DistanceGridForceField.cpp
- ${SOFADISTANCEGRID_SRC}/components/collision/FFDDistanceGridDiscreteIntersection.cpp
- ${SOFADISTANCEGRID_SRC}/components/collision/RayDistanceGridContact.cpp
- ${SOFADISTANCEGRID_SRC}/components/collision/RigidDistanceGridDiscreteIntersection.cpp
- ${SOFADISTANCEGRID_SRC}/components/collision/DistanceGridCollisionModel.cpp
- ${SOFADISTANCEGRID_SRC}/components/collision/BarycentricPenalityContact_DistanceGrid.cpp
-)
-
-set(EXTRA_FILES
- README.md
- examples/README.md
- examples/DistanceGridForceField_liver.scn
- examples/FFDDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn
- examples/FFDDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn
- examples/RigidDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn
- examples/RigidDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn
- #examples/RigidDistanceGridCollisionModel_skull_FreeMotionAnimationLoop.scn
- )
-
-
-find_package(MiniFlowVR QUIET)
-sofa_find_package(Sofa.Core REQUIRED)
-sofa_find_package(Sofa.Component.Collision REQUIRED)
-sofa_find_package(Sofa.GL QUIET)
-
-add_library(${PROJECT_NAME} SHARED ${HEADER_FILES} ${SOURCE_FILES} ${EXTRA_FILES})
-target_link_libraries(${PROJECT_NAME} PUBLIC Sofa.Core)
-target_link_libraries(${PROJECT_NAME} PUBLIC Sofa.Component.Collision)
-
-if(Sofa.GL_FOUND)
- target_link_libraries(${PROJECT_NAME} PUBLIC Sofa.GL)
-endif()
-
-if(MiniFlowVR_FOUND)
- target_link_libraries(${PROJECT_NAME} PRIVATE miniFlowVR) # Private because not exported in API
- message(STATUS "SofaDistanceGrid: MiniFlowVR found.")
-else()
- message(STATUS "SofaDistanceGrid: MiniFlowVR not found, some features will not be compiled.")
-endif()
-
-if(SOFA_BUILD_TESTS)
- add_subdirectory(SofaDistanceGrid_test)
-endif()
-
-## Install rules for the library and headers; CMake package configurations files
-sofa_create_package_with_targets(
- PACKAGE_NAME ${PROJECT_NAME}
- PACKAGE_VERSION ${Sofa_VERSION}
- TARGETS ${PROJECT_NAME} AUTO_SET_TARGET_PROPERTIES
- INCLUDE_SOURCE_DIR "src"
- INCLUDE_INSTALL_DIR "${PROJECT_NAME}"
- RELOCATABLE "plugins"
- )
-
-sofa_add_subdirectory(plugin extensions/CUDA SofaDistanceGrid.CUDA)
diff --git a/applications/plugins/SofaDistanceGrid/README.md b/applications/plugins/SofaDistanceGrid/README.md
deleted file mode 100644
index e289f6ec213..00000000000
--- a/applications/plugins/SofaDistanceGrid/README.md
+++ /dev/null
@@ -1,3 +0,0 @@
-= The SofaDistanceGrid plugin.
-
-
diff --git a/applications/plugins/SofaDistanceGrid/SofaDistanceGridConfig.cmake.in b/applications/plugins/SofaDistanceGrid/SofaDistanceGridConfig.cmake.in
deleted file mode 100644
index d2e411e2f59..00000000000
--- a/applications/plugins/SofaDistanceGrid/SofaDistanceGridConfig.cmake.in
+++ /dev/null
@@ -1,27 +0,0 @@
-# CMake package configuration file for the SofaDistanceGrid plugin
-
-@PACKAGE_GUARD@
-@PACKAGE_INIT@
-
-find_package(Sofa.Core QUIET REQUIRED)
-find_package(Sofa.Component.Collision QUIET REQUIRED)
-
-set(SOFADISTANCEGRID_HAVE_MINIFLOWVR @SOFADISTANCEGRID_HAVE_MINIFLOWVR@)
-set(SOFADISTANCEGRID_HAVE_SOFA_GL @SOFADISTANCEGRID_HAVE_SOFA_GL@)
-
-if(SOFADISTANCEGRID_HAVE_MINIFLOWVR)
- find_package(MiniFlowVR QUIET REQUIRED)
-endif()
-
-if(SOFADISTANCEGRID_HAVE_SOFA_GL)
- find_package(Sofa.GL QUIET REQUIRED)
-endif()
-
-if(NOT TARGET SofaDistanceGrid)
- include("${CMAKE_CURRENT_LIST_DIR}/SofaDistanceGridTargets.cmake")
-endif()
-
-check_required_components(SofaDistanceGrid)
-set(SofaDistanceGrid_LIBRARIES SOFADISTANCEGRID)
-set(SofaDistanceGrid_INCLUDE_DIRS @PACKAGE_SOFADISTANCEGRID_INCLUDE_DIR@ ${SOFADISTANCEGRID_INCLUDE_DIR})
-
diff --git a/applications/plugins/SofaDistanceGrid/SofaDistanceGrid_test/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/SofaDistanceGrid_test/CMakeLists.txt
deleted file mode 100644
index 5fc39e20534..00000000000
--- a/applications/plugins/SofaDistanceGrid/SofaDistanceGrid_test/CMakeLists.txt
+++ /dev/null
@@ -1,12 +0,0 @@
-cmake_minimum_required(VERSION 3.22)
-
-project(SofaDistanceGrid_test)
-
-set(SOURCE_FILES
- DistanceGrid_test.cpp
-)
-
-add_executable(${PROJECT_NAME} ${SOURCE_FILES})
-target_link_libraries(${PROJECT_NAME} Sofa.Testing SofaDistanceGrid)
-
-add_test(NAME ${PROJECT_NAME} COMMAND ${PROJECT_NAME})
diff --git a/applications/plugins/SofaDistanceGrid/SofaDistanceGrid_test/DistanceGrid_test.cpp b/applications/plugins/SofaDistanceGrid/SofaDistanceGrid_test/DistanceGrid_test.cpp
deleted file mode 100644
index de3d9ebe118..00000000000
--- a/applications/plugins/SofaDistanceGrid/SofaDistanceGrid_test/DistanceGrid_test.cpp
+++ /dev/null
@@ -1,102 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include
-#include
-
-#include
-using sofa::component::container::DistanceGrid ;
-
-namespace sofa
-{
-namespace component
-{
-namespace container
-{
-namespace _distancegrid_
-{
-using sofa::type::Vec3 ;
-
-struct DistanceGrid_test : public sofa::testing::NumericTest
-{
- void chekcValidConstructorsCube(){
- EXPECT_MSG_NOEMIT(Warning, Error) ;
-
- DistanceGrid grid(10, 10, 10,
- DistanceGrid::Coord(-1,-1,-1),
- DistanceGrid::Coord(1.0,1.0,1.0)) ;
-
- EXPECT_EQ(grid.getNx(), 10) ;
- EXPECT_EQ(grid.getNy(), 10) ;
- EXPECT_EQ(grid.getNz(), 10) ;
-
- EXPECT_FALSE(grid.inBBox(type::Vec3(-2, 0, 0), 0.0f)) ;
- EXPECT_FALSE(grid.inBBox(type::Vec3( 0,-2, 0), 0.0f)) ;
- EXPECT_FALSE(grid.inBBox(type::Vec3( 0, 0,-2), 0.0f)) ;
- EXPECT_FALSE(grid.inBBox(type::Vec3( 2, 0, 0), 0.0f)) ;
- EXPECT_FALSE(grid.inBBox(type::Vec3( 0, 2, 0), 0.0f)) ;
- EXPECT_FALSE(grid.inBBox(type::Vec3( 0, 0, 2), 0.0f)) ;
-
- EXPECT_EQ(grid.size(), 10*10*10);
-
- //todo(dmarchal:2017-05-02) This "isCube" & "getCubeDim" stuff is ugly as hell !
- EXPECT_FALSE(grid.isCube());
- }
-
- void checInvalidConstructorsCube(int x, int y, int z,
- float mx, float my, float mz,
- float ex, float ey, float ez){
- std::cout << "x-y-z:" << x << ", " << y << ", " << z << std::endl ;
- std::cout << "mx-my-mz:" << mx << ", " << my << ", " << mz << std::endl ;
- std::cout << "ex-ey-ez:" << ex << ", " << ey << ", " << ez << std::endl ;
-
- DistanceGrid grid(x, y, z,
- DistanceGrid::Coord(mx,my,mz),
- DistanceGrid::Coord(ex,ey,ez)) ;
- }
-};
-
-TEST_F(DistanceGrid_test, chekcValidConstructorsCube) {
- ASSERT_NO_THROW(this->chekcValidConstructorsCube()) ;
-}
-
-TEST_F(DistanceGrid_test, chekcInvalidConstructorsCube) {
- std::vector< std::vector< float >> values = {
- {-10, 10, 10, -1,-1,-1, 1, 1,1},
- { 10,-10, 10, -1,-1,-1, 1, 1,1},
- { 10, 10,-10, -1,-1,-1, 1, 1,1},
- { 10, 10, 0, -1,-1,-1, 1, 1,1},
- { 10, 0, 10, -1,-1,-1, 1, 1,1},
- { 0, 10, 0, -1,-1,-1, 1, 1,1},
- { 0, 10, 0, -1, 1,-1, 1,-1,1} };
- for(auto& v : values ){
- ASSERT_NO_THROW( this->checInvalidConstructorsCube(
- (int)v[0], (int)v[1], (int)v[2],
- v[3], v[4], v[5],
- v[6], v[7], v[8])) ;
- }
-}
-
-
-} // __distance_grid__
-} // container
-} // component
-} // sofa
diff --git a/applications/plugins/SofaDistanceGrid/applications/meshconv/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/applications/meshconv/CMakeLists.txt
deleted file mode 100644
index 7b73ed3a729..00000000000
--- a/applications/plugins/SofaDistanceGrid/applications/meshconv/CMakeLists.txt
+++ /dev/null
@@ -1,9 +0,0 @@
-cmake_minimum_required(VERSION 3.22)
-project(meshconv)
-
-find_package(MiniFlowVR REQUIRED)
-
-add_executable(${PROJECT_NAME} meshconv.cpp tesselate.cpp)
-
-target_link_libraries(${PROJECT_NAME} miniFlowVR)
-
diff --git a/applications/plugins/SofaDistanceGrid/applications/meshconv/meshconv.cpp b/applications/plugins/SofaDistanceGrid/applications/meshconv/meshconv.cpp
deleted file mode 100644
index 3d5c6b76a8e..00000000000
--- a/applications/plugins/SofaDistanceGrid/applications/meshconv/meshconv.cpp
+++ /dev/null
@@ -1,293 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU General Public License as published by the Free *
-* Software Foundation; either version 2 of the License, or (at your option) *
-* any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License for *
-* more details. *
-* *
-* You should have received a copy of the GNU General Public License along *
-* with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-/******* COPYRIGHT ************************************************
-* *
-* FlowVR Render *
-* Parallel Rendering Modules *
-* *
-*-----------------------------------------------------------------*
-* COPYRIGHT (C) 2005 by *
-* Laboratoire Informatique et Distribution (UMR5132) and *
-* INRIA Project MOVI. ALL RIGHTS RESERVED. *
-* *
-* This source is covered by the GNU GPL, please refer to the *
-* COPYING file for further information. *
-* *
-*-----------------------------------------------------------------*
-* *
-* Original Contributors: *
-* Jeremie Allard, *
-* Clement Menier. *
-* *
-*******************************************************************
-* *
-* File: ./src/utils/meshconv.cpp *
-* *
-* Contacts: *
-* *
-******************************************************************/
-#include
-#include
-#include
-
-#include
-#include
-
-using namespace flowvr::render;
-
-extern void tesselateMesh(Mesh& obj, int rec=1, bool onSphere=false);
-
-int main(int argc, char** argv)
-{
- bool normalize = false;
- bool flip = false;
- bool wnormals = false;
- bool rnormals = false;
- bool rtexcoords = false;
- bool closemesh = false;
- float closedist = 0.0f;
- float mergedist = -1.0f;
- Vec3f translation;
- Vec3f rotation;
- Vec3f scale(1,1,1);
- float dilate = 0.0f;
- int tesselate = 0;
- bool sphere = false;
- bool dist = false;
- int res = 16;
- int rx=0,ry=0,rz=0;
- float border = 0.25f;
- float vsize = 0.0f;
- ftl::CmdLine cmd("Usage: meshconv [options] mesh.input [mesh.output]");
- cmd.opt("normalize",'n',"transform points so that the center is at <0,0,0> and the max coodinate is 1",&normalize);
- cmd.opt("flip",'f',"flip normals",&flip);
- cmd.opt("wnormals",'N',"force writing of normals",&wnormals);
- cmd.opt("rnormals",'O',"remove normals",&rnormals);
- cmd.opt("rtexcoords",'T',"remove texcoords",&rtexcoords);
- cmd.opt("close",'c',"close mesh",&closemesh);
- cmd.opt("close2",'C',"close mesh creating intermediate vertices no further appart than given dist",&closedist);
- cmd.opt("merge",'m',"merge vertices closer than the given distance",&mergedist);
- cmd.opt("translate",'t',"translate the mesh",&translation);
- cmd.opt("rotate",'r',"rotate the mesh using euler angles in degree",&rotation);
- cmd.opt("scale",'s',"scale the mesh using 3 coefficients",&scale);
- cmd.opt("dilate",'d',"dilate (i.e. displace vertices of the given distance along their normals)",&dilate);
- cmd.opt("tesselate",'a',"tesselate (split each edge in 2 resursivly n times)",&tesselate);
- cmd.opt("sphere",'S',"consider the mesh as a sphere for tesselation",&sphere);
- cmd.opt("dist",'D',"compute distance field",&dist);
- cmd.opt("res",'R',"resolution of distance field",&res);
- cmd.opt("rx",'X',"X resolution of distance field",&rx);
- cmd.opt("ry",'Y',"Y resolution of distance field",&ry);
- cmd.opt("rz",'Z',"Z resolution of distance field",&rz);
- cmd.opt("vsize",'V',"size of each voxel in distance field",&vsize);
- cmd.opt("border",'B',"distance field border size relative to the object's BBox size (or negative for exact size)",&border);
- bool error=false;
- if (!cmd.parse(argc,argv,&error))
- return error?1:0;
-
- if (cmd.args.size()<1 || cmd.args.size()>2)
- {
- std::cerr << cmd.help() << std::endl;
- return 1;
- }
-
- std::string file_in = cmd.args[0];
- std::string file_out;
- if (cmd.args.size()>=2) file_out = cmd.args[1];
-
- Mesh obj;
-
- if (!obj.load(file_in.c_str()))
- {
- std::cerr << "Failed to read "<size) size=d;
- }
- Vec3f center = (bb.a+bb.b)*0.5;
- float sc = 2/size;
- Mat4x4f m; m.identity();
- m(0,0)=sc; m(0,3) = -center[0]*sc;
- m(1,1)=sc; m(1,3) = -center[1]*sc;
- m(2,2)=sc; m(2,3) = -center[2]*sc;
- for (int i=0; imat<<" = ";
- obj.distmap->mat = xform * obj.distmap->mat;
- std::cout << obj.distmap->mat<= 0.0f)
- {
- std::cout << "Merging vertices closer than " << mergedist < 0 || sphere)
- {
- std::cout << "Tesselating mesh..."< 0)
- {
- rx = (int)ceilf(bbsize[0]/vsize);
- ry = (int)ceilf(bbsize[1]/vsize);
- rz = (int)ceilf(bbsize[2]/vsize);
- }
- std::cout << "Computing "<. *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include
-#include
-#include
-
-#include
-
-using namespace flowvr::render;
-using ftl::Vec3f;
-
-void projectOnSphere(Vec3f& p, Vec3f center, float radius)
-{
- p -= center;
- p = p * (radius / p.norm());
- p += center;
-}
-
-void projectOnSphere(Vec3f& p,Vec3f& n, Vec3f center, float radius)
-{
- p -= center;
- n = p;
- n.normalize();
- p = p * (radius / p.norm());
- p += center;
-}
-
-void tesselateMesh(Mesh& obj, int rec=1, bool onSphere=false)
-{
- BBox bb = obj.calcBBox();
- std::cout << "Mesh bbox="<::iterator it = obj.edges[e1].begin(), itend = obj.edges[e1].end(); it != itend; ++it)
- {
- int g2 = obj.getPG(it->first);
-
- int f1p1 = -1, f1p2 = -1;
- int f2p1 = -1, f2p2 = -1;
- int i1 = -1;
- if (it->second.f1 >= 0)
- {
- Vec3i fp = obj.getFP(it->second.f1);
- f1p1 = fp[0], f1p2 = fp[1];
- if (obj.getPG(f1p1) != g1)
- {
- f1p1 = fp[1]; f1p2 = fp[2];
- if (obj.getPG(f1p1) != g1)
- {
- f1p1 = fp[2]; f1p2 = fp[0]; // this is the last possible edge
- }
- }
- if (obj.getPG(f1p1) != g1 || (obj.getPG(f1p2) != g2))
- {
- std::cerr << "ERROR: Edge "<second.f1<<" = "<second.f1 = -1;
- continue;
- }
- Mesh::Vertex v;
- v = obj.getP(f1p1);
- v += obj.getP(f1p2);
- v.mean(2);
- //if (onSphere) v.p = projectOnSphere(v.p, center, radius[0]);
- i1 = obj.addP(v);
- // replace the face index with the new point index
- it->second.f1 = i1;
- }
-
- if (it->second.f2 >= 0)
- {
- Vec3i fp = obj.getFP(it->second.f2);
- f2p1 = fp[0], f2p2 = fp[1];
- if (obj.getPG(f2p1) != g2)
- {
- f2p1 = fp[1]; f2p2 = fp[2];
- if (obj.getPG(f2p1) != g2)
- {
- f2p1 = fp[2]; f2p2 = fp[0]; // this is the last possible edge
- }
- }
- if (obj.getPG(f2p1) != g2 || (obj.getPG(f2p2) != g1))
- {
- std::cerr << "ERROR: Edge "<second.f2<<" = "<second.f2 = -1;
- continue;
- }
- if (f1p1 == f2p2 && f1p2 == f2p1)
- {
- // same point as other face
- it->second.f2 = i1;
- }
- else
- {
- Mesh::Vertex v;
- v = obj.getP(f2p1);
- v += obj.getP(f2p2);
- v.mean(2);
- //if (onSphere) v.p = projectOnSphere(v.p, center, radius[0]);
- int i2;
- if (i1 == -1) // no other edge, create a new group
- i2 = obj.addP(v);
- else
- {
- // see of all points are on the same subgroup
- int gf1p1 = g1;
- while (obj.getGP0(gf1p1+1) <= -f1p1) ++gf1p1;
- int gf1p2 = g2;
- while (obj.getGP0(gf1p2+1) <= -f1p2) ++gf1p2;
- int gf2p1 = g2;
- while (obj.getGP0(gf2p1+1) <= -f2p1) ++gf2p1;
- int gf2p2 = g1;
- while (obj.getGP0(gf2p2+1) <= -f2p2) ++gf2p2;
- int g = obj.getPG(i1);
- i2 = obj.addP(v,g);
- if (gf1p1 != gf2p2 || gf1p2 != gf2p1)
- {
- // create a subgroup
- obj.GP0(obj.nbg()) = -i2;
- }
- }
- // replace the face index with the new point index
- it->second.f2 = i2;
- }
- }
- }
- }
-
- // then create new faces
- std::cout << "Creating new faces..."< faces_p;
- int nbf0 = obj.nbf();
- faces_p.reserve(nbf0*4);
- for(int i=0; i= (unsigned)obj.nbp() || (unsigned)edges[1] >= (unsigned)obj.nbp() || (unsigned)edges[2] >= (unsigned)obj.nbp())
- {
- std::cerr << "ERROR: invalid edge points " << edges << " in face " <
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/applications/plugins/SofaDistanceGrid/examples/FFDDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn b/applications/plugins/SofaDistanceGrid/examples/FFDDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn
deleted file mode 100644
index 7fa3491bd6c..00000000000
--- a/applications/plugins/SofaDistanceGrid/examples/FFDDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn
+++ /dev/null
@@ -1,92 +0,0 @@
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/applications/plugins/SofaDistanceGrid/examples/FFDDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn b/applications/plugins/SofaDistanceGrid/examples/FFDDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn
deleted file mode 100644
index 55564bf78c7..00000000000
--- a/applications/plugins/SofaDistanceGrid/examples/FFDDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn
+++ /dev/null
@@ -1,97 +0,0 @@
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/applications/plugins/SofaDistanceGrid/examples/README.md b/applications/plugins/SofaDistanceGrid/examples/README.md
deleted file mode 100644
index f55d3e021dc..00000000000
--- a/applications/plugins/SofaDistanceGrid/examples/README.md
+++ /dev/null
@@ -1,2 +0,0 @@
-= Examples
-- ...
diff --git a/applications/plugins/SofaDistanceGrid/examples/RigidDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn b/applications/plugins/SofaDistanceGrid/examples/RigidDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn
deleted file mode 100644
index 345637a7303..00000000000
--- a/applications/plugins/SofaDistanceGrid/examples/RigidDistanceGridCollisionModel_liver_DefaultAnimationLoop.scn
+++ /dev/null
@@ -1,80 +0,0 @@
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/applications/plugins/SofaDistanceGrid/examples/RigidDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn b/applications/plugins/SofaDistanceGrid/examples/RigidDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn
deleted file mode 100644
index ea0d05a2e8b..00000000000
--- a/applications/plugins/SofaDistanceGrid/examples/RigidDistanceGridCollisionModel_liver_FreeMotionAnimationLoop.scn
+++ /dev/null
@@ -1,83 +0,0 @@
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/extensions/CUDA/CMakeLists.txt
deleted file mode 100644
index 2cc359986e0..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/CMakeLists.txt
+++ /dev/null
@@ -1,43 +0,0 @@
-cmake_minimum_required(VERSION 3.22)
-project(SofaDistanceGrid.CUDA LANGUAGES CUDA CXX)
-
-set(HEADER_FILES
- src/SofaDistanceGrid/CUDA/init.h
- src/SofaDistanceGrid/CUDA/config.h.in
-
- src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h
- src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h
- src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h
-)
-
-set(SOURCE_FILES
- src/SofaDistanceGrid/CUDA/init.cpp
-
- src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.cpp
- src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp
- src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cpp
-)
-
-set(CUDA_SOURCES
- src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu
- src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu
-)
-
-sofa_find_package(SofaDistanceGrid REQUIRED)
-sofa_find_package(SofaCUDA REQUIRED)
-
-add_library(${PROJECT_NAME} SHARED ${HEADER_FILES} ${SOURCE_FILES} ${CUDA_SOURCES})
-
-target_link_libraries(${PROJECT_NAME} SofaDistanceGrid)
-target_link_libraries(${PROJECT_NAME} SofaCUDA)
-get_property(SOFACUDA_CUDA_ARCHITECTURES TARGET SofaCUDA PROPERTY CUDA_ARCHITECTURES)
-set_target_properties(${PROJECT_NAME} PROPERTIES CUDA_ARCHITECTURES "${SOFACUDA_CUDA_ARCHITECTURES}" )
-
-sofa_create_package_with_targets(
- PACKAGE_NAME ${PROJECT_NAME}
- PACKAGE_VERSION ${Sofa_VERSION}
- TARGETS ${PROJECT_NAME} AUTO_SET_TARGET_PROPERTIES
- INCLUDE_SOURCE_DIR "src"
- INCLUDE_INSTALL_DIR "${PROJECT_NAME}"
- RELOCATABLE "plugins"
-)
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/SofaDistanceGrid.CUDAConfig.cmake.in b/applications/plugins/SofaDistanceGrid/extensions/CUDA/SofaDistanceGrid.CUDAConfig.cmake.in
deleted file mode 100644
index 62faad7f3c6..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/SofaDistanceGrid.CUDAConfig.cmake.in
+++ /dev/null
@@ -1,9 +0,0 @@
-# CMake package configuration file for the VolumetricRendering.CUDA library
-
-@PACKAGE_GUARD@
-@PACKAGE_INIT@
-
-find_package(SofaDistanceGrid QUIET REQUIRED)
-find_package(SofaCUDA QUIET REQUIRED)
-
-check_required_components(SofaDistanceGrid.CUDA)
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cpp
deleted file mode 100644
index 88718d82630..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cpp
+++ /dev/null
@@ -1,381 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include "CudaCollisionDetection.h"
-#include
-
-namespace sofa
-{
-
-namespace gpu
-{
-
-namespace cuda
-{
-
-using namespace sofa::core::collision;
-using namespace sofa::component::collision::geometry;
-
-extern "C"
-{
- void CudaCollisionDetection_runTests(unsigned int nbTests, unsigned int maxPoints, const void* tests, void* nresults);
-}
-
-void registerCudaCollisionDetection(sofa::core::ObjectFactory* factory)
-{
- factory->registerObjects(sofa::core::ObjectRegistrationData("GPU-based collision detection using CUDA.")
- .add< CudaCollisionDetection >());
-}
-
-void CudaCollisionDetection::beginNarrowPhase()
-{
- Inherit2::beginNarrowPhase();
- for (TestMap::iterator it = tests.begin(), itend = tests.end(); it != itend; ++it)
- if (it->second.index >= 0)
- it->second.index = -1;
- else
- --it->second.index;
-}
-
-void CudaCollisionDetection::addCollisionPair( const std::pair& cmPair )
-{
- core::CollisionModel *finalcm1 = cmPair.first->getLast();
- core::CollisionModel *finalcm2 = cmPair.second->getLast();
-
- Entry& entry = tests[std::make_pair(finalcm1, finalcm2)];
- if (entry.test == NULL)
- {
- entry.test = createTest(finalcm1, finalcm2);
- }
- entry.index = 0; // active pair
-}
-
-void CudaCollisionDetection::endNarrowPhase()
-{
- // first clean up old pairs
- for (TestMap::iterator it = tests.begin(), itend = tests.end(); it != itend;)
- if (it->second.index < -100)
- {
- TestMap::iterator it2 = it;
- ++it;
- tests.erase(it2);
- }
- else ++it;
-
- // compute number of tests
- int ntests = 0;
- for (TestMap::iterator it = tests.begin(), itend = tests.end(); it != itend; ++it)
- if (it->second.index >= 0)
- {
- Test* test = it->second.test;
- if (test->useGPU())
- {
- it->second.index = ntests;
- ntests += test->init();
- }
- }
- if (ntests > 0)
- {
- // resize vectors
- gputests.fastResize(ntests);
- gpuresults.fastResize(ntests);
-
- // init tests
- for (TestMap::iterator it = tests.begin(), itend = tests.end(); it != itend; ++it)
- if (it->second.index >= 0)
- {
- Test* test = it->second.test;
- if (test->useGPU())
- {
- test->fillInfo(&gputests[it->second.index]);
- }
- }
-
- // Launch GPU test
- int maxPoints = 0;
- for (int t=0; t maxPoints)
- maxPoints = gputests[t].nbPoints;
- CudaCollisionDetection_runTests(ntests, maxPoints, gputests.deviceRead(), gpuresults.deviceWrite());
-
- }
-
- // compute CPU-side collisions
- for (TestMap::iterator it = tests.begin(), itend = tests.end(); it != itend; ++it)
- if (it->second.index >= 0)
- {
- Test* test = it->second.test;
- if (!test->useGPU())
- {
- core::CollisionModel *cm1 = it->first.first->getFirst(); //->getNext();
- core::CollisionModel *cm2 = it->first.second->getFirst(); //->getNext();
- addCollisionPair( std::make_pair(cm1, cm2) );
- }
- }
-
- if (ntests > 0)
- {
- // gather gpu results
- const int* results = (const int*)gpuresults.hostRead();
- int total = 0;
- for (int i=0; isecond.index >= 0)
- {
- Test* test = it->second.test;
- if (test->useGPU())
- {
- GPUOutputVector* tresults = &it->second.test->results;
- this->getDetectionOutputs(it->first.first, it->first.second) = tresults;
- int newIndex = 0;
- for (unsigned int t=0; tnbTests(); t++)
- {
- tresults->wtest(t).curSize = results[it->second.index + t];
- tresults->wtest(t).newIndex = newIndex;
- newIndex += tresults->rtest(t).curSize;
- }
- //test->fillContacts(this->outputsMap[it->first], results+it->second.index);
- }
- }
- }
-
- Inherit2::endNarrowPhase();
-}
-
-
-CudaCollisionDetection::Test* CudaCollisionDetection::createTest(core::CollisionModel* model1, core::CollisionModel* model2)
-{
- if (CudaRigidDistanceGridCollisionModel* rigid1 = dynamic_cast(model1))
- {
- if (CudaRigidDistanceGridCollisionModel* rigid2 = dynamic_cast(model2))
- return new RigidRigidTest(rigid1, rigid2);
- else if (CudaSphereCollisionModel* sphere2 = dynamic_cast(model2))
- return new SphereRigidTest(sphere2, rigid1);
- else if (CudaPointCollisionModel* point2 = dynamic_cast(model2))
- return new PointRigidTest(point2, rigid1);
- }
- else if (CudaSphereCollisionModel* sphere1 = dynamic_cast(model1))
- {
- if (CudaRigidDistanceGridCollisionModel* rigid2 = dynamic_cast(model2))
- return new SphereRigidTest(sphere1, rigid2);
- }
- else if (CudaPointCollisionModel* point1 = dynamic_cast(model1))
- {
- if (CudaRigidDistanceGridCollisionModel* rigid2 = dynamic_cast(model2))
- return new PointRigidTest(point1, rigid2);
- }
- std::cout << "CudaCollisionDetection::CPUTest "<getClassName()<<" - "<getClassName()<getClassName()<<" - "<getClassName()<isActive() || !model2->isActive()) return 0;
- bool useP1 = model1->usePoints.getValue();
- bool useP2 = model2->usePoints.getValue();
- if (!useP1 && !useP2) return 0;
- int i0 = model1->getSize();
- for (CudaRigidDistanceGridCollisionElement e1 = CudaRigidDistanceGridCollisionElement(model1->begin()); e1!=model1->end(); ++e1)
- for (CudaRigidDistanceGridCollisionElement e2 = CudaRigidDistanceGridCollisionElement(model2->begin()); e2!=model2->end(); ++e2)
- {
- CudaDistanceGrid* g1 = e1.getGrid();
- CudaDistanceGrid* g2 = e2.getGrid();
- if (g1 && g2)
- {
- if (useP1 && !g1->meshPts.empty()) results.addTest(std::make_pair(e1.getIndex(), i0+e2.getIndex()), g1->meshPts.size());
- if (useP2 && !g2->meshPts.empty()) results.addTest(std::make_pair(i0+e2.getIndex(), e1.getIndex()), g2->meshPts.size());
- }
- }
- return results.nbTests();
-}
-
-/// Fill the info to send to the graphics card
-void CudaCollisionDetection::RigidRigidTest::fillInfo(GPUTest* tests)
-{
- if (results.nbTests()==0) return;
- GPUContact* gresults = (GPUContact*)results.results.deviceWrite();
- GPUContactPoint* gresults1 = (GPUContactPoint*)results.results1.deviceWrite();
- GPUContactPoint* gresults2 = (GPUContactPoint*)results.results2.deviceWrite();
- int i0 = model1->getSize();
- for (unsigned int i=0; i& p1 = elem1.getGrid()->meshPts;
- CudaDistanceGrid& g2 = *elem2.getGrid();
- test.nbPoints = p1.size();
- test.result = gresults + e.firstIndex;
- test.result1 = gresults1 + e.firstIndex;
- test.result2 = gresults2 + e.firstIndex;
- test.points = p1.deviceRead();
- test.radius = NULL;
- test.gridnx = g2.getNx();
- test.gridny = g2.getNy();
- test.gridnz = g2.getNz();
- test.gridbbmin = g2.getBBMin();
- test.gridbbmax = g2.getBBMax();
- test.gridp0 = g2.getPMin();
- test.gridinvdp = g2.getInvCellWidth();
- test.grid = g2.getDists().deviceRead();
- test.margin = 0;
- test.rotation = elem2.getRotation().multTranspose(elem1.getRotation());
- test.translation = elem2.getRotation().multTranspose(elem1.getTranslation()-elem2.getTranslation());
- }
-}
-
-CudaCollisionDetection::SphereRigidTest::SphereRigidTest(CudaSphereCollisionModel*model1, CudaRigidDistanceGridCollisionModel* model2 )
- : model1(model1), model2(model2)
-{
- std::cout << "CudaCollisionDetection::SphereRigidTest "<getClassName()<<" - "<getClassName()<isActive() || !model2->isActive()) return 0;
- const CudaVector& p1 = model1->getMechanicalState()->read(sofa::core::vec_id::read_access::position)->getValue();
- if (p1.empty()) return 0;
-
- for (CudaRigidDistanceGridCollisionElement e2 = CudaRigidDistanceGridCollisionElement(model2->begin()); e2!=model2->end(); ++e2)
- {
- CudaDistanceGrid* g2 = e2.getGrid();
- if (g2)
- results.addTest(std::make_pair(0, e2.getIndex()), p1.size());
- }
- return results.nbTests();
-}
-
-/// Fill the info to send to the graphics card
-void CudaCollisionDetection::SphereRigidTest::fillInfo(GPUTest* tests)
-{
-
- if (results.nbTests()==0) return;
- GPUContact* gresults = (GPUContact*)results.results.deviceWrite();
- GPUContactPoint* gresults1 = (GPUContactPoint*)results.results1.deviceWrite();
- GPUContactPoint* gresults2 = (GPUContactPoint*)results.results2.deviceWrite();
- const CudaVector& p1 = model1->getMechanicalState()->read(sofa::core::vec_id::read_access::position)->getValue();
-
- for (unsigned int i=0; igetR().deviceRead();
- test.grid = g2->getDists().deviceRead();
- test.gridnx = g2->getNx();
- test.gridny = g2->getNy();
- test.gridnz = g2->getNz();
- test.gridbbmin = g2->getBBMin();
- test.gridbbmax = g2->getBBMax();
- test.gridp0 = g2->getPMin();
- test.gridinvdp = g2->getInvCellWidth();
- test.margin = 0; //model1->getRadius(0);
- test.rotation.transpose(Mat3x3f(elem2.getRotation()));
- test.translation = test.rotation*(-elem2.getTranslation());
-
- }
-
-}
-
-CudaCollisionDetection::PointRigidTest::PointRigidTest( CudaPointCollisionModel* model1, CudaRigidDistanceGridCollisionModel* model2 )
- : model1(model1), model2(model2)
-{
- std::cout << "CudaCollisionDetection::PointRigidTest "<getClassName()<<" - "<getClassName()<isActive() || !model2->isActive()) return 0;
- for (CudaPoint e1 = CudaPoint(model1->begin()); e1!=model1->end(); ++e1)
- for (CudaRigidDistanceGridCollisionElement e2 = CudaRigidDistanceGridCollisionElement(model2->begin()); e2!=model2->end(); ++e2)
- {
- CudaDistanceGrid* g2 = e2.getGrid();
- if (g2)
- results.addTest(std::make_pair(e1.getIndex(), e2.getIndex()), e1.getSize());
- }
- return results.nbTests();
-}
-
-/// Fill the info to send to the graphics card
-void CudaCollisionDetection::PointRigidTest::fillInfo(GPUTest* tests)
-{
- if (results.nbTests()==0) return;
- GPUContact* gresults = (GPUContact*)results.results.deviceWrite();
- GPUContactPoint* gresults1 = (GPUContactPoint*)results.results1.deviceWrite();
- GPUContactPoint* gresults2 = (GPUContactPoint*)results.results2.deviceWrite();
- for (unsigned int i=0; i& p1 = model1->getMechanicalState()->read(sofa::core::vec_id::read_access::position)->getValue();
- CudaDistanceGrid& g2 = *elem2.getGrid();
- test.nbPoints = elem1.getSize();
- test.result = gresults + e.firstIndex;
- test.result1 = gresults1 + e.firstIndex;
- test.result2 = gresults2 + e.firstIndex;
- test.points = p1.deviceReadAt(elem1.i0());
- test.radius = NULL;
- test.gridnx = g2.getNx();
- test.gridny = g2.getNy();
- test.gridnz = g2.getNz();
- test.gridbbmin = g2.getBBMin();
- test.gridbbmax = g2.getBBMax();
- test.gridp0 = g2.getPMin();
- test.gridinvdp = g2.getInvCellWidth();
- test.grid = g2.getDists().deviceRead();
- test.margin = 0;
- test.rotation.transpose(elem2.getRotation());
- test.translation = elem2.getRotation().multTranspose(-elem2.getTranslation());
- }
-}
-
-} // namespace cuda
-
-} // namespace gpu
-
-} // namespace sofa
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu
deleted file mode 100644
index e14a17af8b7..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu
+++ /dev/null
@@ -1,235 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include
-#include
-#include
-#include "cuda.h"
-
-extern "C"
-{
- void CudaCollisionDetection_runTests(unsigned int nbTests, unsigned int maxPoints, const void* tests, void* nresults);
-}
-
-struct /*__align__(16)*/ GPUContactPoint
-{
- CudaVec3 p;
- int elem;
-};
-
-struct /*__align__(16)*/ GPUContact
-{
-// int p1;
-// CudaVec3 p2;
- float distance;
- CudaVec3 normal;
-};
-
-struct GPUTest
-{
- GPUContact* result;
- GPUContactPoint* result1;
- GPUContactPoint* result2;
- const CudaVec3* points;
- const float* radius;
- const float* grid;
- //matrix3 rotation;
- CudaVec3 rotation_x,rotation_y,rotation_z;
- CudaVec3 translation;
- float margin;
- int nbPoints;
- int gridnx, gridny, gridnz;
- CudaVec3 gridbbmin, gridbbmax;
- CudaVec3 gridp0, gridinvdp;
-};
-
-struct GPUDeformedCube
-{
- int elem;
- int ix,iy,iz;
- int points0, nbp;
- CudaVec3 initP0, invDP;
-};
-
-struct GPUDeformedCubeState
-{
- CudaVec4 faces[6];
- CudaVec3 C0, Dx, Dy, Dz, Dxy, Dxz, Dyz, Dxyz;
- CudaVec3 center, radius;
-};
-
-struct GPUDeformedCubeBSphere
-{
- CudaVec3 center;
- float radius;
-};
-
-struct GPUTestFFD
-{
- GPUContact* result;
- GPUContactPoint* result1;
- GPUContactPoint* result2;
- const CudaVec3* points;
- const float* radius;
- const float* grid;
- const GPUDeformedCube* ffdCubes;
- GPUDeformedCubeState* ffdState;
- GPUDeformedCubeBSphere* ffdBSphere;
- float margin;
- int nbPoints;
- int nbCubes;
- int gridnx, gridny, gridnz;
- CudaVec3 gridbbmin, gridbbmax;
- CudaVec3 gridp0, gridinvdp;
-};
-
-//////////////////////
-// GPU-side methods //
-//////////////////////
-
-__shared__ GPUTest curTest;
-
-__global__ void CudaCollisionDetection_runTests_kernel(const GPUTest* tests, int* nresults)
-{
- if (threadIdx.x == 0)
- curTest = tests[blockIdx.x];
-
-
- __syncthreads();
-
- //! Dynamically allocated shared memory to compact results
- extern __shared__ int scan[];
-
- CudaVec3 p0,p;
- float distance;
- CudaVec3 grad = CudaVec3::make(0,0,0);
- //CudaVec3 normal;
- int n = 0;
- if (threadIdx.x < curTest.nbPoints)
- {
- p0 = curTest.points[threadIdx.x];
- //p = curTest.rotation * p;
- p = CudaVec3::make(dot(curTest.rotation_x, p0), dot(curTest.rotation_y, p0), dot(curTest.rotation_z, p0));
- p += curTest.translation;
-
- CudaVec3 coefs = mul(p-curTest.gridp0, curTest.gridinvdp);
- int x = __float2int_rd(coefs.x);
- int y = __float2int_rd(coefs.y);
- int z = __float2int_rd(coefs.z);
- if ((unsigned)x < curTest.gridnx-1
- && (unsigned)y < curTest.gridny-1
- && (unsigned)z < curTest.gridnz-1)
- {
- int nx = curTest.gridnx;
- int nxny = nx*curTest.gridny;
- coefs.x -= __int2float_rd(x);
- coefs.y -= __int2float_rd(y);
- coefs.z -= __int2float_rd(z);
- const float* gval = curTest.grid + (x+nx*y+nxny*z);
- float d000 = gval[0 ];
- float d100_d000 = gval[1 ] - d000;
- float d010 = gval[ nx ];
- float d110_d010 = gval[1+nx ] - d010;
- float d001 = gval[ nxny];
- float d101_d001 = gval[1 +nxny] - d001;
- float d011 = gval[ nx+nxny];
- float d111_d011 = gval[1+nx+nxny] - d011;
- float dx00 = d000 + (d100_d000)*coefs.x;
- float dx10_dx00 = d010 + (d110_d010)*coefs.x - dx00;
- float dx01 = d001 + (d101_d001)*coefs.x;
- float dx11_dx01 = d011 + (d111_d011)*coefs.x - dx01;
- float dy0 = dx00 + (dx10_dx00)*coefs.y;
- float dy1_dy0 = dx01 + (dx11_dx01)*coefs.y - dy0;
- distance = dy0 + (dy1_dy0)*coefs.z;
- float r = 0;
- if (curTest.radius)
- {
- r = curTest.radius[threadIdx.x];
- }
- if (distance < curTest.margin+r)
- {
- n = 1;
- grad.z = dy1_dy0;
- grad.y = (dx10_dx00) + ((dx11_dx01)-(dx10_dx00))*coefs.z;
- dy0 = d100_d000 + (d110_d010 - d100_d000)*coefs.y;
- dy1_dy0 = d101_d001 + (d111_d011 - d101_d001)*coefs.y - dy0;
- grad.x = dy0 + (dy1_dy0)*coefs.z;
- grad *= invnorm(grad);
- //normal = grad;
- p -= grad*distance;
- //distance -= r;
- distance = r;
- //grad = CudaVec3::make(0,1,1);
- }
- }
- }
-
- scan[threadIdx.x] = n;
-
- for (int i=1; i=i)
- scan[threadIdx.x] = scan[threadIdx.x] + scan[threadIdx.x - i];
- }
-
- if (n)
- {
- int i = scan[threadIdx.x]-1;
- GPUContact c;
- //c.p1 = threadIdx.x;
- //c.p2 = p;
- c.distance = distance;
- c.normal = -grad;
- //c.normal = normal; //CudaVec3::make(-grad.x,-grad.y,-grad.z); //-grad;
- curTest.result[i] = c;
- //curTest.result[scan[threadIdx.x]-1].p1 = threadIdx.x;
- //curTest.result[scan[threadIdx.x]-1].p2 = p;
- //curTest.result[scan[threadIdx.x]-1].distance = distance;
- //curTest.result[scan[threadIdx.x]-1].normal = normal;
- GPUContactPoint cp1;
- cp1.elem = threadIdx.x;
- cp1.p = p0;
- curTest.result1[i] = cp1;
- GPUContactPoint cp2;
- cp2.elem = 0;
- cp2.p = p;
- curTest.result2[i] = cp2;
- }
- if (threadIdx.x == curTest.nbPoints-1)
- nresults[blockIdx.x] = scan[curTest.nbPoints-1];
-}
-
-//////////////////////
-// CPU-side methods //
-//////////////////////
-
-void CudaCollisionDetection_runTests(unsigned int nbTests, unsigned int maxPoints, const void* tests, void* nresults)
-{
- sofa::gpu::cuda::mycudaPrintf("sizeof(GPUTest)=%d\nsizeof(GPUContact)=%d\nsizeof(matrix3)=%d\n",sizeof(GPUTest),sizeof(GPUContact),sizeof(matrix3));
- const GPUTest* gputests = (const GPUTest*)tests;
- // round up to 16
- //maxPoints = (maxPoints+15)&-16;
- dim3 threads(maxPoints,1);
- dim3 grid(nbTests,1);
- {CudaCollisionDetection_runTests_kernel<<< grid, threads, threads.x*sizeof(int) >>>(gputests, (int*)nresults); mycudaDebugError("CudaCollisionDetection_runTests_kernel");}
-
-}
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h
deleted file mode 100644
index 14a0c9753d9..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h
+++ /dev/null
@@ -1,218 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#ifndef SOFA_GPU_CUDA_CUDACOLLISIONDETECTION_H
-#define SOFA_GPU_CUDA_CUDACOLLISIONDETECTION_H
-
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-
-
-namespace sofa
-{
-
-namespace core
-{
-
-namespace collision
-{
-
-using type::Vec3f;
-using type::Mat3x3f;
-using sofacuda::GPUDetectionOutputVector;
-
-template<>
-class TDetectionOutputVector : public GPUDetectionOutputVector
-{
-};
-
-template<>
-class TDetectionOutputVector : public GPUDetectionOutputVector
-{
-};
-
-template<>
-class TDetectionOutputVector : public GPUDetectionOutputVector
-{
-};
-
-} // namespace collision
-
-} // namespace core
-
-namespace gpu
-{
-
-namespace cuda
-{
-
-
-
-class CudaCollisionDetection
- : public sofa::component::collision::detection::algorithm::BruteForceBroadPhase
- , public sofa::core::collision::NarrowPhaseDetection
-{
-public:
- SOFA_CLASS2(CudaCollisionDetection, sofa::component::collision::detection::algorithm::BruteForceBroadPhase, sofa::core::collision::NarrowPhaseDetection);
- struct GPUTest
- {
- void* result;
- void* result1;
- void* result2;
- const void* points;
- const void* radius;
- const void* grid;
- Mat3x3f rotation;
- Vec3f translation;
- float margin;
- int nbPoints;
- int gridnx, gridny, gridnz;
- Vec3f gridbbmin, gridbbmax;
- Vec3f gridp0, gridinvdp;
- };
-
- /*struct GPUContact
- {
- int p1;
- Vec3f p2;
- float distance;
- Vec3f normal;
- };*/
- //typedef sofa::core::collision::GPUDetectionOutput GPUContact;
- typedef sofa::core::collision::GPUDetectionOutputVector GPUOutputVector;
-
- CudaVector gputests;
- CudaVector gpuresults; ///< number of contact detected on each test
-
- typedef sofa::core::collision::DetectionOutputVector DetectionOutputVector;
-
- class Test
- {
- public:
- GPUOutputVector results;
- Test() {}
- virtual ~Test()
- {
- }
- virtual bool useGPU()=0;
- /// Returns how many tests are required
- virtual int init()=0;
- /// Fill the info to send to the graphics card
- virtual void fillInfo(GPUTest* tests)=0;
- /// Create the list of SOFA contacts from the contacts detected by the GPU
- //virtual void fillContacts(DetectionOutputVector& contacts, const int* nresults)=0;
- };
-
- class CPUTest : public Test
- {
- public:
- CPUTest() {}
- bool useGPU() { return false; }
- int init() { return 0; }
- void fillInfo(GPUTest* /*tests*/) {}
- //void fillContacts(DetectionOutputVector& /*contacts*/, const int* /*nresults*/) {}
- };
-
- class RigidRigidTest : public Test
- {
- public:
- CudaRigidDistanceGridCollisionModel* model1;
- CudaRigidDistanceGridCollisionModel* model2;
- RigidRigidTest(CudaRigidDistanceGridCollisionModel* model1, CudaRigidDistanceGridCollisionModel* model2);
- bool useGPU() { return true; }
- /// Returns how many tests are required
- virtual int init();
- /// Fill the info to send to the graphics card
- virtual void fillInfo(GPUTest* tests);
- /// Create the list of SOFA contacts from the contacts detected by the GPU
- //void fillContacts(DetectionOutputVector& contacts, const int* nresults);
-
- protected:
- //void fillInfo(GPUTest& test, CudaVector& gpucontacts, CudaRigidDistanceGridCollisionElement elem1, CudaRigidDistanceGridCollisionElement elem2);
- //void fillContacts(DetectionOutputVector& contacts, int nresults, CudaVector& gpucontacts, CudaRigidDistanceGridCollisionElement e1, CudaRigidDistanceGridCollisionElement e2, bool invert);
- };
-
- class SphereRigidTest : public Test
- {
- public:
- CudaSphereCollisionModel* model1;
- CudaRigidDistanceGridCollisionModel* model2;
- SphereRigidTest(CudaSphereCollisionModel *model1, CudaRigidDistanceGridCollisionModel* model2);
- bool useGPU() { return true; }
- /// Returns how many tests are required
- virtual int init();
- /// Fill the info to send to the graphics card
- virtual void fillInfo(GPUTest* tests);
- /// Create the list of SOFA contacts from the contacts detected by the GPU
- //void fillContacts(DetectionOutputVector& contacts, const int* nresults);
- };
-
- class PointRigidTest : public Test
- {
- public:
- CudaPointCollisionModel* model1;
- CudaRigidDistanceGridCollisionModel* model2;
- PointRigidTest(CudaPointCollisionModel* model1, CudaRigidDistanceGridCollisionModel* model2);
- bool useGPU() { return true; }
- /// Returns how many tests are required
- virtual int init();
- /// Fill the info to send to the graphics card
- virtual void fillInfo(GPUTest* tests);
- /// Create the list of SOFA contacts from the contacts detected by the GPU
- //void fillContacts(DetectionOutputVector& contacts, const int* nresults);
-
- protected:
- //void fillInfo(GPUTest& test, CudaVector& gpucontacts, CudaRigidDistanceGridCollisionElement elem1, CudaRigidDistanceGridCollisionElement elem2);
- //void fillContacts(DetectionOutputVector& contacts, int nresults, CudaVector& gpucontacts, CudaRigidDistanceGridCollisionElement e1, CudaRigidDistanceGridCollisionElement e2, bool invert);
- };
-
- struct Entry
- {
- int index; // negative if not active
- Test* test;
- Entry() : index(-1), test(NULL) {}
- ~Entry() { if (test!=NULL) delete test; }
- };
-
- typedef std::map< std::pair, Entry > TestMap;
-
- TestMap tests;
-
- virtual void beginNarrowPhase() override;
- virtual void addCollisionPair (const std::pair& cmPair) override;
- virtual void endNarrowPhase() override;
-
-protected:
- Test* createTest(core::CollisionModel* model1, core::CollisionModel* model2);
-};
-
-} // namespace cuda
-
-} // namespace gpu
-
-} // namespace sofa
-
-#endif
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp
deleted file mode 100644
index 7245bc839ae..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp
+++ /dev/null
@@ -1,162 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include
-#include
-#include
-#include
-#include "CudaDistanceGridCollisionModel.h"
-#include
-#include
-
-#include
-#include
-
-
-namespace sofa::component::collision::response::contact
-{
-
-using namespace sofa::gpu::cuda;
-
-template <>
-void BarycentricPenalityContact::setDetectionOutputs(OutputVector* o)
-{
- GPUDetectionOutputVector& outputs = *dynamic_cast(o);
- //const bool printLog = this->f_printLog.getValue();
- if (ff==NULL)
- {
- MechanicalState1* mstate1 = mapper1.createMapping("contactPointsCUDA");
- MechanicalState2* mstate2 = mapper2.createMapping("contactPointsCUDA");
- ff = sofa::core::objectmodel::New(mstate1,mstate2);
- ff->setName( getName() );
- ff->init();
- }
-
- mapper1.setPoints1(&outputs);
- mapper2.setPoints2(&outputs);
- const double d0 = intersectionMethod->getContactDistance() + model1->getContactDistance() + model2->getContactDistance(); // - 0.001;
-#if 0
- int insize = outputs.size();
- int size = insize;
- ff->clear(size);
- //int i = 0;
- for (int i=0; ielem.first);
- //CollisionElement2 elem2(o->elem.second);
- //int index1 = elem1.getIndex();
- //int index2 = elem2.getIndex();
- int index1 = index;
- int index2 = index;
- //// Create mapping for first point
- //index1 = mapper1.addPoint(o->point[0], index1);
- //// Create mapping for second point
- //index2 = mapper2.addPoint(o->point[1], index2);
- double distance = d0 + outputs.get(i)->distance; // + mapper1.radius(elem1) + mapper2.radius(elem2);
- double stiffness = (model1->getContactStiffness(0) * model1->getContactStiffness(0))/distance;
- double mu_v = (model1->getContactFriction(0) + model1->getContactFriction(0));
- ff->addContact(index1, index2, outputs.get(i)->normal, (float)distance, (float)stiffness, (float)mu_v, (float)mu_v, index);
- }
-#else
- double distance = d0; // + mapper1.radius(elem1) + mapper2.radius(elem2);
- double stiffness = (model1->getContactStiffness(0) * model1->getContactStiffness(0)); ///distance;
- ff->setContacts((float)distance, (float)stiffness, &outputs, true);
-#endif
- // Update mappings
- mapper1.update();
- mapper2.update();
-}
-
-template <>
-void BarycentricPenalityContact::setDetectionOutputs(OutputVector* o)
-{
-
- GPUDetectionOutputVector& outputs = *dynamic_cast(o);
- //const bool printLog = this->f_printLog.getValue();
- if (ff==NULL)
- {
- MechanicalState1* mstate1 = mapper1.createMapping("contactPointsCUDA");
- MechanicalState2* mstate2 = mapper2.createMapping("contactPointsCUDA");
- ff = sofa::core::objectmodel::New(mstate1,mstate2);
- ff->setName( getName() );
- ff->init();
- }
-
- mapper1.setPoints1(&outputs);
- mapper2.setPoints2(&outputs);
- const double d0 = intersectionMethod->getContactDistance() + model1->getContactDistance() + model2->getContactDistance(); // - 0.001;
-#if 0
- int insize = outputs.size();
- int size = insize;
- ff->clear(size);
- //int i = 0;
- for (int i=0; ielem.first);
- //CollisionElement2 elem2(o->elem.second);
- //int index1 = elem1.getIndex();
- //int index2 = elem2.getIndex();
- int index1 = index;
- int index2 = index;
- //// Create mapping for first point
- //index1 = mapper1.addPoint(o->point[0], index1);
- //// Create mapping for second point
- //index2 = mapper2.addPoint(o->point[1], index2);
- double distance = d0 + outputs.get(i)->distance; // + mapper1.radius(elem1) + mapper2.radius(elem2);
- double stiffness = (model1->getContactStiffness(0) * model1->getContactStiffness(0))/distance;
- double mu_v = (model1->getContactFriction(0) + model1->getContactFriction(0));
- ff->addContact(index1, index2, outputs.get(i)->normal, (float)distance, (float)stiffness, (float)mu_v, (float)mu_v, index);
- }
-#else
- double distance = d0; // + mapper1.radius(elem1) + mapper2.radius(elem2);
- double stiffness = (model1->getContactStiffness(0) * model1->getContactStiffness(0)); ///distance;
- ff->setContacts((float)distance, (float)stiffness, &outputs, true);
-#endif
- // Update mappings
- mapper1.update();
- mapper2.update();
-
-}
-
-} // namespace sofa::component::collision::response::contact
-
-
-namespace sofa::gpu::cuda
-{
-
-using namespace sofa::component::collision;
-using namespace sofa::component::collision::response::contact;
-
-//sofa::helper::Creator > CudaDistanceGridCudaDistanceGridContactClass("PenalityContactForceField", true);
-sofa::helper::Creator > CudaPointCudaDistanceGridContactClass("PenalityContactForceField", true);
-sofa::helper::Creator > CudaSphereCudaDistanceGridContactClass("PenalityContactForceField", true);
-//sofa::helper::Creator > CudaDistanceGridDistanceGridContactClass("PenalityContactForceField", true);
-//sofa::helper::Creator > CudaDistanceGridPointContactClass("PenalityContactForceField", true);
-//sofa::helper::Creator > CudaDistanceGridSphereContactClass("PenalityContactForceField", true);
-//sofa::helper::Creator > CudaDistanceGridTriangleContactClass("PenalityContactForceField", true);
-
-} // namespace sofa::gpu::cuda
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.cpp
deleted file mode 100644
index 084ba36a8eb..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.cpp
+++ /dev/null
@@ -1,810 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-
-#include "CudaDistanceGridCollisionModel.h"
-#include
-#include
-#include
-#include
-#include
-#if SOFADISTANCEGRID_HAVE_MINIFLOWVR
- #include
-#endif // SOFADISTANCEGRID_HAVE_MINIFLOWVR
-#include
-#include
-
-namespace sofa
-{
-
-namespace gpu
-{
-
-namespace cuda
-{
-
-void registerCudaRigidDistanceGridCollisionModel(sofa::core::ObjectFactory* factory)
-{
- factory->registerObjects(sofa::core::ObjectRegistrationData("GPU-based grid distance field using CUDA.")
- .add< CudaRigidDistanceGridCollisionModel >());
-}
-
-using namespace defaulttype;
-
-////////////////////////////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////
-
-CudaDistanceGrid::CudaDistanceGrid(int nx, int ny, int nz, Coord pmin, Coord pmax)
- : nbRef(1), nx(nx), ny(ny), nz(nz), nxny(nx*ny), nxnynz(nx*ny*nz)
- , pmin(pmin), pmax(pmax)
- , cellWidth ((pmax[0]-pmin[0])/(nx-1), (pmax[1]-pmin[1])/(ny-1),(pmax[2]-pmin[2])/(nz-1))
- , invCellWidth((nx-1)/(pmax[0]-pmin[0]), (ny-1)/(pmax[1]-pmin[1]),(nz-1)/(pmax[2]-pmin[2]))
- , cubeDim(0)
-{
- dists.resize(nxnynz);
-}
-
-CudaDistanceGrid::~CudaDistanceGrid()
-{
- std::map& shared = getShared();
- std::map::iterator it = shared.begin();
- while (it != shared.end() && it->second != this) ++it;
- if (it != shared.end())
- shared.erase(it); // remove this grid from the list of already loaded grids
-}
-
-/// Add one reference to this grid. Note that loadShared already does this.
-CudaDistanceGrid* CudaDistanceGrid::addRef()
-{
- ++nbRef;
- return this;
-}
-
-/// Release one reference, deleting this grid if this is the last
-bool CudaDistanceGrid::release()
-{
- if (--nbRef != 0)
- return false;
- delete this;
- return true;
-}
-
-CudaDistanceGrid* CudaDistanceGrid::load(const std::string& filename, double scale, double sampling, int nx, int ny, int nz, Coord pmin, Coord pmax)
-{
- if (filename == "#cube")
- {
- float dim = (float)scale;
- int np = 5;
- Coord bbmin(-dim, -dim, -dim), bbmax(dim,dim,dim);
- std::cout << "bbox = <"<-<"<"< pmax[c]) pmax[c] = bbmax[c];
- }
- }
- std::cout << "Creating cube distance grid in <"<-<"<"<calcCubeDistance(dim, np);
- if (sampling)
- grid->sampleSurface(sampling);
- std::cout << "Distance grid creation DONE."<4 && filename.substr(filename.length()-4) == ".raw")
- {
- CudaDistanceGrid* grid = new CudaDistanceGrid(nx, ny, nz, pmin, pmax);
- std::ifstream in(filename.c_str(), std::ios::in | std::ios::binary);
- in.read((char*)&(grid->dists[0]), grid->nxnynz*sizeof(Real));
- if (scale != 1.0)
- {
- for (int i=0; i< grid->nxnynz; i++)
- grid->dists[i] *= (float)scale;
- }
- grid->computeBBox();
- if (sampling)
- grid->sampleSurface(sampling);
- return grid;
- }
-#if SOFADISTANCEGRID_HAVE_MINIFLOWVR
- else if (filename.length()>6 && filename.substr(filename.length()-6) == ".fmesh")
- {
- flowvr::render::Mesh mesh;
- if (!mesh.load(filename.c_str()))
- {
- std::cerr << "ERROR loading FlowVR mesh file "<nx;
- ny = mesh.distmap->ny;
- nz = mesh.distmap->nz;
- ftl::Vec3f fpmin = ftl::transform(mesh.distmap->mat,ftl::Vec3f(0,0,0))*(float)scale;
- ftl::Vec3f fpmax = ftl::transform(mesh.distmap->mat,ftl::Vec3f((float)(nx-1),(float)(ny-1),(float)(nz-1)))*(float)scale;
- pmin = Coord(fpmin.ptr());
- pmax = Coord(fpmax.ptr());
- std::cout << "Copying "<-<"<"<nxnynz; i++)
- grid->dists[i] = mesh.distmap->data[i]*(float)scale;
-
- if (sampling)
- grid->sampleSurface(sampling);
- else if (mesh.getAttrib(flowvr::render::Mesh::MESH_POINTS_GROUP))
- {
- int nbpos = 0;
- for (int i=0; i= 0)
- ++nbpos;
- }
- std::cout << "Copying "<meshPts.resize(nbpos);
- int p = 0;
- for (int i=0; i= 0)
- grid->meshPts[p++] = Coord(mesh.getPP(p0).ptr())*scale;
- }
- }
- else
- {
- int nbpos = mesh.nbp();
- std::cout << "Copying "<meshPts.resize(nbpos);
- for (int i=0; imeshPts[i] = Coord(mesh.getPP(i).ptr())*scale;
- }
- grid->computeBBox();
- std::cout << "Distance grid creation DONE."<4 && filename.substr(filename.length()-4) == ".obj")
- {
- sofa::helper::io::Mesh* mesh = sofa::helper::io::Mesh::Create(filename);
- const sofa::type::vector & vertices = mesh->getVertices();
-
- std::cout << "Computing bbox."< bbmax[c]) bbmax[c] = (Real)vertices[i][c];
- }
- bbmin *= scale;
- bbmax *= scale;
- }
- std::cout << "bbox = <"<-<"<"< pmax[c]) pmax[c] = bbmax[c];
- }
- }
- std::cout << "Creating distance grid in <"<-<"<"<meshPts.resize(vertices.size());
- for(unsigned int i=0; imeshPts[i] = vertices[i]*scale;
- const auto & facets = mesh->getFacets();
- int nbt = 0;
- int nbq = 0;
- for (unsigned int i=0; i= 3)
- nbt += pts.size()-2;
- }
- grid->meshTriangles.resize(nbt);
- grid->meshQuads.resize(nbq);
- nbt=0;
- nbq=0;
- for (unsigned int i=0; imeshQuads[nbq++] = sofa::core::topology::BaseMeshTopology::Quad(pts[0],pts[1],pts[2],pts[3]);
- else if (pts.size() >= 3)
- for (unsigned int j=2; jmeshTriangles[nbt++] = sofa::core::topology::BaseMeshTopology::Triangle(pts[0],pts[j-1],pts[j]);
- }
- std::cout << "Computing distance field."<calcDistance();
- if (sampling)
- grid->sampleSurface(sampling);
- grid->computeBBox();
- std::cout << "Distance grid creation DONE."<4 && filename.substr(filename.length()-4) == ".raw")
- {
- std::ofstream out(filename.c_str(), std::ios::out | std::ios::binary);
- out.write((const char*)&(dists[0]), nxnynz*sizeof(Real));
- }
- else
- {
- std::cerr << " CudaDistanceGrid::save(): Unsupported extension: "< bbmax[c]) bbmax[c] = (Real)meshPts[i][c];
- }
- }
- else
- {
- bbmin = pmin;
- bbmax = pmax;
- /// \TODO compute the real bbox from the grid content
- }
-}
-
-/// Sample the surface with points approximately separated by the given sampling distance (expressed in voxels if the value is negative)
-void CudaDistanceGrid::sampleSurface(double sampling)
-{
- std::cout << "CudaDistanceGrid: sample surface with sampling distance " << sampling << std::endl;
- int stepX, stepY, stepZ;
- if (sampling < 0)
- {
- stepX = stepY = stepZ = (int)(-sampling);
- }
- else
- {
- stepX = (int)(sampling/cellWidth[0]);
- stepY = (int)(sampling/cellWidth[1]);
- stepZ = (int)(sampling/cellWidth[2]);
- }
- if (stepX < 1) stepX = 1;
- if (stepY < 1) stepY = 1;
- if (stepZ < 1) stepZ = 1;
- std::cout << "CudaDistanceGrid: sampling steps: " << stepX << " " << stepY << " " << stepZ << std::endl;
-
- SReal maxD = (SReal)sqrt((cellWidth[0]*stepX)*(cellWidth[0]*stepX) + (cellWidth[1]*stepY)*(cellWidth[1]*stepY) + (cellWidth[2]*stepZ)*(cellWidth[2]*stepZ));
- std::vector pts;
- for (int z=1; z maxD) continue;
-
- type::Vec3 pos = coord(x,y,z);
- type::Vec3 n = grad(index(x,y,z), Coord()); // note that there are some redundant computations between interp() and grad()
- n.normalize();
- pos -= n * (d * 0.99); // push pos back to the surface
- d = interp(pos);
- int it = 1;
- while (helper::rabs(d) > 0.01f*maxD && it < 10)
- {
- n = grad(pos);
- n.normalize();
- pos -= n * (d * 0.99); // push pos back to the surface
- d = interp(pos);
- ++it;
- }
- if (it == 10 && helper::rabs(d) > 0.1f*maxD)
- {
- std::cout << "Failed to converge at ("< 1)
- {
- int nbp = np*np*np - (np-2)*(np-2)*(np-2);
- //std::cout << "Copying "< dim2) { s[c] = dim2; out = true; }
- }
- Real d;
- if (out)
- d = (p - s).norm();
- else
- d = helper::rmax(helper::rmax(helper::rabs(s[0]),helper::rabs(s[1])),helper::rabs(s[2])) - dim2;
- dists[i] = d - (dim-dim2);
- }
- //computeBBox();
- bbmin = Coord(-dim,-dim,-dim);
- bbmax = Coord( dim, dim, dim);
-}
-
-/// Compute distance field from given mesh
-void CudaDistanceGrid::calcDistance()
-{
-
- if (GLEW_EXT_framebuffer_object && GLEW_ARB_vertex_buffer_object)
- {
- static GLuint fbid = 0;
- if (!fbid)
- {
- glGenFramebuffersEXT(1,&fbid);
- glBindFramebufferEXT(GL_FRAMEBUFFER_EXT, fbid);
- }
-
- /*
- if (!texcolor)
- glGenTextures(1, &texcolor);
-
- glBindTexture(target, texcolor);
-
- glTexParameteri(target, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
- glTexParameteri(target, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
- glTexParameteri(target, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
- glTexParameteri(target, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
- glTexParameteri(target, GL_TEXTURE_WRAP_R, GL_CLAMP_TO_EDGE);
- glTexImage2D(target, 0, GL_RGBA8, nx, ny, 0,
- GL_RGBA, GL_UNSIGNED_BYTE, NULL);
-
- if (!texdepth)
- glGenTextures(1, &texdepth);
- glBindTexture(target, texdepth);
-
- glTexParameteri(target, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
- glTexParameteri(target, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
- glTexParameteri(target, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
- glTexParameteri(target, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
- glTexParameteri(target, GL_TEXTURE_WRAP_R, GL_CLAMP_TO_EDGE);
- glTexImage2D(target, 0, GL_DEPTH_COMPONENT24, nx, ny, 0,
- GL_DEPTH_COMPONENT, GL_FLOAT, NULL);
-
- glFramebufferTexture2DEXT(GL_FRAMEBUFFER_EXT, GL_COLOR_ATTACHMENT0_EXT, target, texcolor, 0);*/
- }
- else
- {
- std::cerr << "ERROR: Unsupported OpenGL extensions EXT_framebuffer_object ARB_vertex_buffer_object" << std::endl;
- }
-
-}
-
-CudaDistanceGrid* CudaDistanceGrid::loadShared(const std::string& filename, double scale, double sampling, int nx, int ny, int nz, Coord pmin, Coord pmax)
-{
- CudaDistanceGridParams params;
- params.filename = filename;
- params.scale = scale;
- params.sampling = sampling;
- params.nx = nx;
- params.ny = ny;
- params.nz = nz;
- params.pmin = pmin;
- params.pmax = pmax;
- std::map& shared = getShared();
- std::map::iterator it = shared.find(params);
- if (it != shared.end())
- return it->second->addRef();
- else
- {
- return shared[params] = load(filename, scale, sampling, nx, ny, nz, pmin, pmax);
- }
-}
-
-std::map& CudaDistanceGrid::getShared()
-{
- static std::map instance;
- return instance;
-}
-
-
-////////////////////////////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////
-
-CudaRigidDistanceGridCollisionModel::CudaRigidDistanceGridCollisionModel()
- : modified(true)
- , fileCudaRigidDistanceGrid( initData( &fileCudaRigidDistanceGrid, "fileCudaRigidDistanceGrid", "load distance grid from specified file"))
- , scale( initData( &scale, 1.0, "scale", "scaling factor for input file"))
- , sampling( initData( &sampling, 0.0, "sampling", "if not zero: sample the surface with points approximately separated by the given sampling distance (expressed in voxels if the value is negative)"))
- , box( initData( &box, "box", "Field bounding box defined by xmin,ymin,zmin, xmax,ymax,zmax") )
- , nx( initData( &nx, 64, "nx", "number of values on X axis") )
- , ny( initData( &ny, 64, "ny", "number of values on Y axis") )
- , nz( initData( &nz, 64, "nz", "number of values on Z axis") )
- , dumpfilename( initData( &dumpfilename, "dumpfilename","write distance grid to specified file"))
- , usePoints( initData( &usePoints, true, "usePoints", "use mesh vertices for collision detection"))
-{
- rigid = NULL;
- addAlias(&fileCudaRigidDistanceGrid,"filename");
-}
-
-CudaRigidDistanceGridCollisionModel::~CudaRigidDistanceGridCollisionModel()
-{
- for (unsigned int i=0; irelease();
- if (elems[i].prevGrid!=NULL) elems[i].prevGrid->release();
- }
-}
-
-void CudaRigidDistanceGridCollisionModel::init()
-{
- std::cout << "> CudaRigidDistanceGridCollisionModel::init()"<core::CollisionModel::init();
- rigid = dynamic_cast< core::behavior::MechanicalState* > (getContext()->getMechanicalState());
-
- CudaDistanceGrid* grid = NULL;
- if (fileCudaRigidDistanceGrid.getValue().empty())
- {
- if (elems.size()==0 || elems[0].grid==NULL)
- std::cerr << "ERROR: CudaRigidDistanceGridCollisionModel requires an input filename.\n";
- // else the grid has already been set
- return;
- }
- std::cout << "CudaRigidDistanceGridCollisionModel: creating "<-<"<";
- std::cout << std::endl;
- grid = CudaDistanceGrid::loadShared(fileCudaRigidDistanceGrid.getFullPath(), scale.getValue(), sampling.getValue(), nx.getValue(),ny.getValue(),nz.getValue(),box.getValue()[0],box.getValue()[1]);
-
- resize(1);
- elems[0].grid = grid;
- if (grid && !dumpfilename.getValue().empty())
- {
- std::cout << "CudaRigidDistanceGridCollisionModel: dump grid to "<save(dumpfilename.getFullPath());
- }
- std::cout << "< CudaRigidDistanceGridCollisionModel::init()"<core::CollisionModel::resize(s);
- elems.resize(s);
-}
-
-void CudaRigidDistanceGridCollisionModel::setGrid(CudaDistanceGrid* surf, Index index)
-{
- if (elems[index].grid == surf) return;
- if (elems[index].grid!=NULL) elems[index].grid->release();
- elems[index].grid = surf->addRef();
- modified = true;
-}
-
-void CudaRigidDistanceGridCollisionModel::setNewState(Index index, double dt, CudaDistanceGrid* grid, const Matrix3& rotation, const type::Vec3& translation)
-{
- grid->addRef();
- if (elems[index].prevGrid!=NULL)
- elems[index].prevGrid->release();
- elems[index].prevGrid = elems[index].grid;
- elems[index].grid = grid;
- elems[index].prevRotation = elems[index].rotation;
- elems[index].rotation = rotation;
- elems[index].prevTranslation = elems[index].translation;
- elems[index].translation = translation;
- if (!elems[index].isTransformed)
- {
- Matrix3 I; I.identity();
- if (!(rotation == I) || !(translation == type::Vec3()))
- elems[index].isTransformed = true;
- }
- elems[index].prevDt = dt;
- modified = true;
-}
-
-using sofa::component::collision::geometry::CubeCollisionModel;
-
-/// Create or update the bounding volume hierarchy.
-void CudaRigidDistanceGridCollisionModel::computeBoundingTree(int maxDepth)
-{
- CubeCollisionModel* cubeModel = this->createPrevious();
-
- if (!modified && !isMoving() && !cubeModel->empty()) return; // No need to recompute BBox if immobile
-
- updateGrid();
-
- cubeModel->resize(size);
- for (int i=0; i(elems[i])->recalcBBox();
- type::Vec3 emin, emax;
- if (rigid)
- {
- const RigidTypes::Coord& xform = rigid->read(sofa::core::vec_id::read_access::position)->getValue()[i];
- elems[i].translation = xform.getCenter();
- xform.getOrientation().toMatrix(elems[i].rotation);
- elems[i].isTransformed = true;
- }
- if (elems[i].isTransformed)
- {
- //std::cout << "Grid "< x + <"<"<getBBCorner(0);
- emin = corner;
- emax = emin;
- for (int j=1; j<8; j++)
- {
- corner = elems[i].translation + elems[i].rotation * elems[i].grid->getBBCorner(j);
- for(int c=0; c<3; c++)
- if (corner[c] < emin[c]) emin[c] = corner[c];
- else if (corner[c] > emax[c]) emax[c] = corner[c];
- }
- }
- else
- {
- emin = elems[i].grid->getBBMin();
- emax = elems[i].grid->getBBMax();
- }
- cubeModel->setParentOf(i, emin, emax); // define the bounding box of the current element
- //std::cout << "Grid "<-<"<"<computeBoundingTree(maxDepth);
- modified = false;
-}
-
-void CudaRigidDistanceGridCollisionModel::updateGrid()
-{
-}
-
-void CudaRigidDistanceGridCollisionModel::drawCollisionModel(const core::visual::VisualParams* vparams)
-{
- if (vparams->displayFlags().getShowWireFrame())
- glPolygonMode(GL_FRONT_AND_BACK, GL_LINE);
- glDisable(GL_LIGHTING);
- glColor4fv(getColor4f());
- glPointSize(3);
- for (unsigned int i = 0; i < elems.size(); i++)
- {
- draw(vparams, i);
- }
- glPointSize(1);
- if (vparams->displayFlags().getShowWireFrame())
- glPolygonMode(GL_FRONT_AND_BACK, GL_FILL);
-}
-
-void CudaRigidDistanceGridCollisionModel::draw(const core::visual::VisualParams* , Index index)
-{
-#if SOFACUDA_HAVE_SOFA_GL == 1
- if (elems[index].isTransformed)
- {
- glPushMatrix();
- // float m[16];
- // (*rigid->getX())[index].writeOpenGlMatrix( m );
- // glMultMatrixf(m);
- Matrix4 m;
- m.identity();
- m = elems[index].rotation;
- m.transpose();
- m[3] = Vec4(elems[index].translation,1.0);
- sofa::gl::glMultMatrix(m.ptr());
- }
-
- CudaDistanceGrid* grid = getGrid(index);
- CudaDistanceGrid::Coord corners[8];
- for(unsigned int i=0; i<8; i++)
- corners[i] = grid->getCorner(i);
- //glEnable(GL_BLEND);
- //glDepthMask(0);
- if (!isMoving())
- glColor4f(0.25f, 0.25f, 0.25f, 0.1f);
- else
- glColor4f(0.5f, 0.5f, 0.5f, 0.1f);
- glBegin(GL_LINES);
- {
- glVertex3fv(corners[0].ptr()); glVertex3fv(corners[4].ptr());
- glVertex3fv(corners[1].ptr()); glVertex3fv(corners[5].ptr());
- glVertex3fv(corners[2].ptr()); glVertex3fv(corners[6].ptr());
- glVertex3fv(corners[3].ptr()); glVertex3fv(corners[7].ptr());
- glVertex3fv(corners[0].ptr()); glVertex3fv(corners[2].ptr());
- glVertex3fv(corners[1].ptr()); glVertex3fv(corners[3].ptr());
- glVertex3fv(corners[4].ptr()); glVertex3fv(corners[6].ptr());
- glVertex3fv(corners[5].ptr()); glVertex3fv(corners[7].ptr());
- glVertex3fv(corners[0].ptr()); glVertex3fv(corners[1].ptr());
- glVertex3fv(corners[2].ptr()); glVertex3fv(corners[3].ptr());
- glVertex3fv(corners[4].ptr()); glVertex3fv(corners[5].ptr());
- glVertex3fv(corners[6].ptr()); glVertex3fv(corners[7].ptr());
- }
- glEnd();
- glDisable(GL_BLEND);
- glDepthMask(1);
- for(unsigned int i=0; i<8; i++)
- corners[i] = grid->getBBCorner(i);
- //glEnable(GL_BLEND);
- //glDepthMask(0);
-
- if (!isMoving())
- glColor4f(0.5f, 0.5f, 0.5f, 1.0f);
- else
- glColor4f(1.0f, 1.0f, 1.0f, 1.0f);
- glBegin(GL_LINES);
- {
- glVertex3fv(corners[0].ptr()); glVertex3fv(corners[4].ptr());
- glVertex3fv(corners[1].ptr()); glVertex3fv(corners[5].ptr());
- glVertex3fv(corners[2].ptr()); glVertex3fv(corners[6].ptr());
- glVertex3fv(corners[3].ptr()); glVertex3fv(corners[7].ptr());
- glVertex3fv(corners[0].ptr()); glVertex3fv(corners[2].ptr());
- glVertex3fv(corners[1].ptr()); glVertex3fv(corners[3].ptr());
- glVertex3fv(corners[4].ptr()); glVertex3fv(corners[6].ptr());
- glVertex3fv(corners[5].ptr()); glVertex3fv(corners[7].ptr());
- glVertex3fv(corners[0].ptr()); glVertex3fv(corners[1].ptr());
- glVertex3fv(corners[2].ptr()); glVertex3fv(corners[3].ptr());
- glVertex3fv(corners[4].ptr()); glVertex3fv(corners[5].ptr());
- glVertex3fv(corners[6].ptr()); glVertex3fv(corners[7].ptr());
- }
- glEnd();
-
- const float mindist = -(grid->getPMax()-grid->getPMin()).norm()*0.1f;
- const float maxdist = (grid->getPMax()-grid->getPMin()).norm()*0.025f;
-
- if (grid->meshPts.empty())
- {
- glBegin(GL_POINTS);
- {
- for (int z=0, ind=0; zgetNz(); z++)
- for (int y=0; ygetNy(); y++)
- for (int x=0; xgetNx(); x++, ind++)
- {
- CudaDistanceGrid::Coord p = grid->coord(x,y,z);
- CudaDistanceGrid::Real d = (*grid)[ind];
- if (d < mindist || d > maxdist) continue;
- d /= maxdist;
- if (d<0)
- glColor3d(1+d*0.25, 0, 1+d);
- else
- glColor3d(0, 1-d*0.25, 1-d);
- glVertex3fv(p.ptr());
- }
- }
- glEnd();
- }
- else
- {
- glColor3d(1, 1 ,1);
- glBegin(GL_POINTS);
- for (unsigned int i=0; imeshPts.size(); i++)
- {
- glVertex3fv(grid->meshPts[i].ptr());
- }
- glEnd();
- glBegin(GL_LINES);
- for (unsigned int i=0; imeshPts.size(); i++)
- {
- CudaDistanceGrid::Coord p ( grid->meshPts[i].ptr() );
- glColor3d(1, 1 ,1);
- CudaDistanceGrid::Coord grad = grid->grad(p);
- grad.normalize();
- for (int j = -2; j <= 2; j++)
- {
- CudaDistanceGrid::Coord p2 = p + grad * (j*maxdist/2);
- CudaDistanceGrid::Real d = grid->eval(p2);
- //if (rabs(d) > maxdist) continue;
- d /= maxdist;
- if (d<0)
- glColor3d(1+d*0.25, 0, 1+d);
- else
- glColor3d(0, 1-d*0.25, 1-d);
- glVertex3fv(p2.ptr());
- if (j>-2 && j < 2)
- glVertex3fv(p2.ptr());
- }
- }
- glEnd();
- }
- if (elems[index].isTransformed)
- {
- glPopMatrix();
- }
-#endif // SOFACUDA_HAVE_SOFA_GL == 1
-}
-
-
-
-
-} // namespace cuda
-
-} // namespace gpu
-
-} // namespace sofa
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h
deleted file mode 100644
index 8da7b24b5d6..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h
+++ /dev/null
@@ -1,591 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#ifndef SOFA_GPU_CUDA_CUDADISTANCEGRIDCOLLISIONMODEL_H
-#define SOFA_GPU_CUDA_CUDADISTANCEGRIDCOLLISIONMODEL_H
-
-#include
-
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-
-namespace sofa
-{
-
-namespace gpu
-{
-
-namespace cuda
-{
-
-using namespace sofa::type;
-using namespace sofa::defaulttype;
-//using namespace sofa::component::collision;
-
-class CudaDistanceGrid
-{
-public:
- typedef float Real;
- static Real maxDist() { return (Real)1e10; }
- typedef Vec3f Coord;
- //typedef Vec4f SamplingPoint; ///< 3D coordinates + radius
- typedef CudaVector VecReal;
- typedef CudaVector VecCoord;
-
- CudaDistanceGrid(int nx, int ny, int nz, Coord pmin, Coord pmax);
-
-protected:
- ~CudaDistanceGrid();
-
-public:
-
- /// Load a distance grid
- static CudaDistanceGrid* load(const std::string& filename, double scale=1.0, double sampling=0.0, int nx=64, int ny=64, int nz=64, Coord pmin = Coord(), Coord pmax = Coord());
-
- /// Load or reuse a distance grid
- static CudaDistanceGrid* loadShared(const std::string& filename, double scale=1.0, double sampling=0.0, int nx=64, int ny=64, int nz=64, Coord pmin = Coord(), Coord pmax = Coord());
-
- /// Add one reference to this grid. Note that loadShared already does this.
- CudaDistanceGrid* addRef();
-
- /// Release one reference, deleting this grid if this is the last
- bool release();
-
- /// Save current grid
- bool save(const std::string& filename);
-
- /// Compute distance field from stored mesh
- void calcDistance();
-
- /// Compute distance field for a cube of the given half-size.
- /// Also create a mesh of points using np points per axis
- void calcCubeDistance(Real dim=1, int np=5);
-
- /// Sample the surface with points approximately separated by the given sampling distance (expressed in voxels if the value is negative)
- void sampleSurface(double sampling=-1.0);
-
- /// Update bbox
- void computeBBox();
-
- int getNx() const { return nx; }
- int getNy() const { return ny; }
- int getNz() const { return nz; }
- const Coord& getCellWidth() const { return cellWidth; }
- const Coord& getInvCellWidth() const { return invCellWidth; }
- const VecReal& getDists() const { return dists; }
- VecReal& getDists() { return dists; }
-
- int size() const { return nxnynz; }
-
- const Coord& getBBMin() const { return bbmin; }
- const Coord& getBBMax() const { return bbmax; }
- void setBBMin(const Coord& val) { bbmin = val; }
- void setBBMax(const Coord& val) { bbmax = val; }
- Coord getBBCorner(int i) const { return Coord((i&1)?bbmax[0]:bbmin[0],(i&2)?bbmax[1]:bbmin[1],(i&4)?bbmax[2]:bbmin[2]); }
- bool inBBox(const Coord& p, Real margin=0.0f) const
- {
- for (int c=0; c<3; ++c)
- if (p[c] < bbmin[c]-margin || p[c] > bbmax[c]+margin) return false;
- return true;
- }
-
- const Coord& getPMin() const { return pmin; }
- const Coord& getPMax() const { return pmax; }
- Coord getCorner(int i) const { return Coord((i&1)?pmax[0]:pmin[0],(i&2)?pmax[1]:pmin[1],(i&4)?pmax[2]:pmin[2]); }
-
- bool isCube() const { return cubeDim != 0; }
- Real getCubeDim() const { return cubeDim; }
-
- bool inGrid(const Coord& p) const
- {
- Coord epsilon = cellWidth*0.1;
- for (int c=0; c<3; ++c)
- if (p[c] < pmin[c]+epsilon[c] || p[c] > pmax[c]-epsilon[c]) return false;
- return true;
- }
-
- Coord clamp(Coord p) const
- {
- for (int c=0; c<3; ++c)
- if (p[c] < pmin[c]) p[c] = pmin[c];
- else if (p[c] > pmax[c]) p[c] = pmax[c];
- return p;
- }
-
- Index ix(const Coord& p) const
- {
- return helper::rfloor((p[0]-pmin[0])*invCellWidth[0]);
- }
-
- Index iy(const Coord& p) const
- {
- return helper::rfloor((p[1]-pmin[1])*invCellWidth[1]);
- }
-
- Index iz(const Coord& p) const
- {
- return helper::rfloor((p[2]-pmin[2])*invCellWidth[2]);
- }
-
- Index index(const Coord& p, Coord& coefs) const
- {
- coefs[0] = (p[0]-pmin[0])*invCellWidth[0];
- coefs[1] = (p[1]-pmin[1])*invCellWidth[1];
- coefs[2] = (p[2]-pmin[2])*invCellWidth[2];
- int x = helper::rfloor(coefs[0]);
- if (x==-1) x=0; else if (x==nx-1) x=nx-2;
- coefs[0] -= x;
- int y = helper::rfloor(coefs[1]);
- if (y==-1) y=0; else if (y==ny-1) y=ny-2;
- coefs[1] -= y;
- int z = helper::rfloor(coefs[2]);
- if (z==-1) z=0; else if (z==nz-1) z=nz-2;
- coefs[2] -= z;
- return x+nx*(y+ny*(z));
- }
-
- Index index(const Coord& p) const
- {
- Coord coefs;
- return index(p, coefs);
- }
-
- Index index(int x, int y, int z)
- {
- return x+nx*(y+ny*(z));
- }
-
- Coord coord(int x, int y, int z)
- {
- return pmin+Coord(x*cellWidth[0], y*cellWidth[1], z*cellWidth[2]);
- }
-
- Real operator[](Index index) const { return dists[index]; }
- Real& operator[](Index index) { return dists[index]; }
-
- static Real interp(Real coef, Real a, Real b)
- {
- return a+coef*(b-a);
- }
-
- Real interp(Index index, const Coord& coefs) const
- {
- return interp(coefs[2],interp(coefs[1],interp(coefs[0],dists[index ],dists[index+1 ]),
- interp(coefs[0],dists[index +nx ],dists[index+1+nx ])),
- interp(coefs[1],interp(coefs[0],dists[index +nxny],dists[index+1 +nxny]),
- interp(coefs[0],dists[index +nx+nxny],dists[index+1+nx+nxny])));
- }
-
- Real interp(const Coord& p) const
- {
- Coord coefs;
- int i = index(p, coefs);
- return interp(i, coefs);
- }
-
- Coord grad(Index index, const Coord& coefs) const
- {
- // val = dist[0][0][0] * (1-x) * (1-y) * (1-z)
- // + dist[1][0][0] * ( x) * (1-y) * (1-z)
- // + dist[0][1][0] * (1-x) * ( y) * (1-z)
- // + dist[1][1][0] * ( x) * ( y) * (1-z)
- // + dist[0][0][1] * (1-x) * (1-y) * ( z)
- // + dist[1][0][1] * ( x) * (1-y) * ( z)
- // + dist[0][1][1] * (1-x) * ( y) * ( z)
- // + dist[1][1][1] * ( x) * ( y) * ( z)
- // dval / dx = (dist[1][0][0]-dist[0][0][0]) * (1-y) * (1-z)
- // + (dist[1][1][0]-dist[0][1][0]) * ( y) * (1-z)
- // + (dist[1][0][1]-dist[0][0][1]) * (1-y) * ( z)
- // + (dist[1][1][1]-dist[0][1][1]) * ( y) * ( z)
- const Real dist000 = dists[index ];
- const Real dist100 = dists[index+1 ];
- const Real dist010 = dists[index +nx ];
- const Real dist110 = dists[index+1+nx ];
- const Real dist001 = dists[index +nxny];
- const Real dist101 = dists[index+1 +nxny];
- const Real dist011 = dists[index +nx+nxny];
- const Real dist111 = dists[index+1+nx+nxny];
- return Coord(
- interp(coefs[2],interp(coefs[1],dist100-dist000,dist110-dist010),interp(coefs[1],dist101-dist001,dist111-dist011)), //*invCellWidth[0],
- interp(coefs[2],interp(coefs[0],dist010-dist000,dist110-dist100),interp(coefs[0],dist011-dist001,dist111-dist101)), //*invCellWidth[1],
- interp(coefs[1],interp(coefs[0],dist001-dist000,dist101-dist100),interp(coefs[0],dist011-dist010,dist111-dist110))); //*invCellWidth[2]);
- }
-
- Coord grad(const Coord& p) const
- {
- Coord coefs;
- Index i = index(p, coefs);
- return grad(i, coefs);
- }
-
- Real eval(const Coord& x) const
- {
- Real d;
- if (inGrid(x))
- {
- d = interp(x);
- }
- else
- {
- Coord xclamp = clamp(x);
- d = interp(xclamp);
- d = helper::rsqrt((x-xclamp).norm2() + d*d); // we underestimate the distance
- }
- return d;
- }
-
- Real quickeval(const Coord& x) const
- {
- Real d;
- if (inGrid(x))
- {
- d = dists[index(x)] - cellWidth[0]; // we underestimate the distance
- }
- else
- {
- Coord xclamp = clamp(x);
- d = dists[index(xclamp)] - cellWidth[0]; // we underestimate the distance
- d = helper::rsqrt((x-xclamp).norm2() + d*d);
- }
- return d;
- }
-
- Real eval2(const Coord& x) const
- {
- Real d2;
- if (inGrid(x))
- {
- Real d = interp(x);
- d2 = d*d;
- }
- else
- {
- Coord xclamp = clamp(x);
- Real d = interp(xclamp);
- d2 = ((x-xclamp).norm2() + d*d); // we underestimate the distance
- }
- return d2;
- }
-
- Real quickeval2(const Coord& x) const
- {
- Real d2;
- if (inGrid(x))
- {
- Real d = dists[index(x)] - cellWidth[0]; // we underestimate the distance
- d2 = d*d;
- }
- else
- {
- Coord xclamp = clamp(x);
- Real d = dists[index(xclamp)] - cellWidth[0]; // we underestimate the distance
- d2 = ((x-xclamp).norm2() + d*d);
- }
- return d2;
- }
-
- //CudaVector meshPts;
- CudaVector meshPts;
- sofa::core::topology::BaseMeshTopology::SeqTriangles meshTriangles;
- sofa::core::topology::BaseMeshTopology::SeqQuads meshQuads;
-
-protected:
- int nbRef;
- VecReal dists;
- const int nx,ny,nz, nxny, nxnynz;
- const Coord pmin, pmax;
- const Coord cellWidth, invCellWidth;
- Coord bbmin, bbmax; ///< bounding box of the object, smaller than the grid
-
- Real cubeDim; ///< Cube dimension (!=0 if this is actually a cube
- //bool updated;
-
- // Grid shared resources
-
- struct CudaDistanceGridParams
- {
- std::string filename;
- double scale;
- double sampling;
- int nx,ny,nz;
- Coord pmin,pmax;
- bool operator==(const CudaDistanceGridParams& v) const
- {
- if (!(filename == v.filename)) return false;
- if (!(scale == v.scale )) return false;
- if (!(sampling == v.sampling)) return false;
- if (!(nx == v.nx )) return false;
- if (!(ny == v.ny )) return false;
- if (!(nz == v.nz )) return false;
- if (!(pmin[0] == v.pmin[0] )) return false;
- if (!(pmin[1] == v.pmin[1] )) return false;
- if (!(pmin[2] == v.pmin[2] )) return false;
- if (!(pmax[0] == v.pmax[0] )) return false;
- if (!(pmax[1] == v.pmax[1] )) return false;
- if (!(pmax[2] == v.pmax[2] )) return false;
- return true;
- }
- bool operator<(const CudaDistanceGridParams& v) const
- {
- if (filename < v.filename) return false;
- if (filename > v.filename) return true;
- if (scale < v.scale ) return false;
- if (scale > v.scale ) return true;
- if (sampling < v.sampling) return false;
- if (sampling > v.sampling) return true;
- if (nx < v.nx ) return false;
- if (nx > v.nx ) return true;
- if (ny < v.ny ) return false;
- if (ny > v.ny ) return true;
- if (nz < v.nz ) return false;
- if (nz > v.nz ) return true;
- if (pmin[0] < v.pmin[0] ) return false;
- if (pmin[0] > v.pmin[0] ) return true;
- if (pmin[1] < v.pmin[1] ) return false;
- if (pmin[1] > v.pmin[1] ) return true;
- if (pmin[2] < v.pmin[2] ) return false;
- if (pmin[2] > v.pmin[2] ) return true;
- if (pmax[0] < v.pmax[0] ) return false;
- if (pmax[0] > v.pmax[0] ) return true;
- if (pmax[1] < v.pmax[1] ) return false;
- if (pmax[1] > v.pmax[1] ) return true;
- if (pmax[2] < v.pmax[2] ) return false;
- if (pmax[2] > v.pmax[2] ) return true;
- return false;
- }
- bool operator>(const CudaDistanceGridParams& v) const
- {
- if (filename > v.filename) return false;
- if (filename < v.filename) return true;
- if (scale > v.scale ) return false;
- if (scale < v.scale ) return true;
- if (sampling < v.sampling) return false;
- if (sampling > v.sampling) return true;
- if (nx > v.nx ) return false;
- if (nx < v.nx ) return true;
- if (ny > v.ny ) return false;
- if (ny < v.ny ) return true;
- if (nz > v.nz ) return false;
- if (nz < v.nz ) return true;
- if (pmin[0] > v.pmin[0] ) return false;
- if (pmin[0] < v.pmin[0] ) return true;
- if (pmin[1] > v.pmin[1] ) return false;
- if (pmin[1] < v.pmin[1] ) return true;
- if (pmin[2] > v.pmin[2] ) return false;
- if (pmin[2] < v.pmin[2] ) return true;
- if (pmax[0] > v.pmax[0] ) return false;
- if (pmax[0] < v.pmax[0] ) return true;
- if (pmax[1] > v.pmax[1] ) return false;
- if (pmax[1] < v.pmax[1] ) return true;
- if (pmax[2] > v.pmax[2] ) return false;
- if (pmax[2] < v.pmax[2] ) return true;
- return false;
- }
- };
- static std::map& getShared();
-
-};
-////////////////////////////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////
-////////////////////////////////////////////////////////////////////////////////
-
-class CudaRigidDistanceGridCollisionModel;
-
-class CudaRigidDistanceGridCollisionElement : public core::TCollisionElementIterator
-{
-public:
- CudaRigidDistanceGridCollisionElement(CudaRigidDistanceGridCollisionModel* model, Index index);
-
- explicit CudaRigidDistanceGridCollisionElement(const core::CollisionElementIterator& i);
-
- CudaDistanceGrid* getGrid();
-
- bool isTransformed();
- const Matrix3& getRotation();
- const Vec3& getTranslation();
-
- void setGrid(CudaDistanceGrid* surf);
-
- /// @name Previous state data
- /// Used to estimate velocity in case the distance grid itself is dynamic
- /// @{
- CudaDistanceGrid* getPrevGrid();
- const Matrix3& getPrevRotation();
- const Vec3& getPrevTranslation();
- double getPrevDt();
- /// @}
-
- /// Set new grid and transform, keeping the old state to estimate velocity
- void setNewState(double dt, CudaDistanceGrid* grid, const Matrix3& rotation, const Vec3& translation);
-};
-
-class CudaRigidDistanceGridCollisionModel : public core::CollisionModel
-{
-public:
- SOFA_CLASS(CudaRigidDistanceGridCollisionModel,core::CollisionModel);
-protected:
-
- class ElementData
- {
- public:
- Matrix3 rotation;
- Vec3 translation;
- CudaDistanceGrid* grid;
-
- /// @name Previous state data
- /// Used to estimate velocity in case the distance grid itself is dynamic
- /// @{
- CudaDistanceGrid* prevGrid; ///< Previous grid
- Matrix3 prevRotation; ///< Previous rotation
- Vec3 prevTranslation; ///< Previous translation
- double prevDt; ///< Time difference between previous and current state
- /// @}
-
- bool isTransformed; ///< True if translation/rotation was set
- ElementData() : grid(NULL), prevGrid(NULL), prevDt(0.0), isTransformed(false) { rotation.identity(); prevRotation.identity(); }
- };
-
- std::vector elems;
- bool modified;
- core::behavior::MechanicalState* rigid;
-
- void updateGrid();
-
- void drawCollisionModel(const core::visual::VisualParams*) override;
-
-public:
- // Input data parameters
- sofa::core::objectmodel::DataFileName fileCudaRigidDistanceGrid;
- Data< double > scale; ///< scaling factor for input file
- Data< double > sampling; ///< if not zero: sample the surface with points approximately separated by the given sampling distance (expressed in voxels if the value is negative)
- Data< type::fixed_array > box; ///< Field bounding box defined by xmin,ymin,zmin, xmax,ymax,zmax
- Data< int > nx; ///< number of values on X axis
- Data< int > ny; ///< number of values on Y axis
- Data< int > nz; ///< number of values on Z axis
- sofa::core::objectmodel::DataFileName dumpfilename;
-
- typedef Rigid3Types InDataTypes;
- typedef Vec3Types DataTypes;
- typedef CudaRigidDistanceGridCollisionElement Element;
-
- Data< bool > usePoints; ///< use mesh vertices for collision detection
-
- CudaRigidDistanceGridCollisionModel();
-
- ~CudaRigidDistanceGridCollisionModel();
-
- core::behavior::MechanicalState* getRigidModel() { return rigid; }
- core::behavior::MechanicalState* getMechanicalState() { return rigid; }
-
- void init() override;
-
- CudaDistanceGrid* getGrid(Index index=0)
- {
- return elems[index].grid;
- }
- bool isTransformed(Index index=0)
- {
- return elems[index].isTransformed;
- }
- const Matrix3& getRotation(Index index=0)
- {
- return elems[index].rotation;
- }
- const Vec3& getTranslation(Index index=0)
- {
- return elems[index].translation;
- }
-
- void setGrid(CudaDistanceGrid* surf, Index index=0);
-
- CudaDistanceGrid* getPrevGrid(Index index=0)
- {
- return elems[index].prevGrid;
- }
- const Matrix3& getPrevRotation(Index index=0)
- {
- return elems[index].prevRotation;
- }
- const Vec3& getPrevTranslation(Index index=0)
- {
- return elems[index].prevTranslation;
- }
- double getPrevDt(Index index=0)
- {
- return elems[index].prevDt;
- }
-
- /// Set new grid and transform, keeping the old state to estimate velocity
- void setNewState(Index index, double dt, CudaDistanceGrid* grid, const Matrix3& rotation, const Vec3& translation);
-
- /// @}
-
- /// Set new grid and transform, keeping the old state to estimate velocity
- void setNewState(double dt, CudaDistanceGrid* grid, const Matrix3& rotation, const Vec3& translation);
-
- // -- CollisionModel interface
-
- void resize(Size size) override;
-
- /// Create or update the bounding volume hierarchy.
- void computeBoundingTree(int maxDepth=0) override;
-
- void draw(const core::visual::VisualParams*, Index index) override;
-};
-
-inline CudaRigidDistanceGridCollisionElement::CudaRigidDistanceGridCollisionElement(CudaRigidDistanceGridCollisionModel* model, Index index)
- : core::TCollisionElementIterator(model, index)
-{}
-
-inline CudaRigidDistanceGridCollisionElement::CudaRigidDistanceGridCollisionElement(const core::CollisionElementIterator& i)
- : core::TCollisionElementIterator(static_cast(i.getCollisionModel()), i.getIndex())
-{
-}
-
-inline CudaDistanceGrid* CudaRigidDistanceGridCollisionElement::getGrid() { return model->getGrid(index); }
-inline void CudaRigidDistanceGridCollisionElement::setGrid(CudaDistanceGrid* surf) { return model->setGrid(surf, index); }
-
-inline bool CudaRigidDistanceGridCollisionElement::isTransformed() { return model->isTransformed(index); }
-inline const Matrix3& CudaRigidDistanceGridCollisionElement::getRotation() { return model->getRotation(index); }
-inline const Vec3& CudaRigidDistanceGridCollisionElement::getTranslation() { return model->getTranslation(index); }
-
-inline CudaDistanceGrid* CudaRigidDistanceGridCollisionElement::getPrevGrid() { return model->getPrevGrid(index); }
-inline const Matrix3& CudaRigidDistanceGridCollisionElement::getPrevRotation() { return model->getPrevRotation(index); }
-inline const Vec3& CudaRigidDistanceGridCollisionElement::getPrevTranslation() { return model->getPrevTranslation(index); }
-inline double CudaRigidDistanceGridCollisionElement::getPrevDt() { return model->getPrevDt(index); }
-
-inline void CudaRigidDistanceGridCollisionElement::setNewState(double dt, CudaDistanceGrid* grid, const Matrix3& rotation, const Vec3& translation)
-{
- return model->setNewState(dt, grid, rotation, translation);
-}
-
-} // namespace cuda
-
-} // namespace gpu
-
-} // namespace sofa
-
-#endif
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu
deleted file mode 100644
index 28965400760..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu
+++ /dev/null
@@ -1,71 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include
-#include
-#include "cuda.h"
-
-extern "C"
-{
- void RigidContactMapperCuda3f_setPoints2(unsigned int size, unsigned int nbTests, unsigned int maxPoints, const void* tests, const void* contacts, void* map);
-}
-
-struct /*__align__(16)*/ GPUContact
-{
- int p1;
- float3 p2;
- float distance;
- float3 normal;
-};
-
-struct /*__align__(8)*/ GPUTestEntry
-{
- int firstIndex;
- int curSize;
- int maxSize;
- int newIndex;
- int elem1,elem2;
-};
-
-__shared__ GPUTestEntry curTestEntry;
-
-__global__ void RigidContactMapperCuda3f_setPoints2_kernel(const GPUTestEntry* tests, const GPUContact* contacts, float3* map)
-{
- if (threadIdx.x == 0)
- curTestEntry = tests[blockIdx.x];
-
- __syncthreads();
-
- GPUContact c = contacts[curTestEntry.firstIndex + threadIdx.x];
- if (threadIdx.x < curTestEntry.curSize)
- {
- map[curTestEntry.newIndex + threadIdx.x] = c.p2;
- }
-}
-
-void RigidContactMapperCuda3f_setPoints2(unsigned int size, unsigned int nbTests, unsigned int maxPoints, const void* tests, const void* contacts, void* map)
-{
- // round up to 16
- //maxPoints = (maxPoints+15)&-16;
- dim3 threads(maxPoints,1);
- dim3 grid(nbTests,1);
- {RigidContactMapperCuda3f_setPoints2_kernel<<< grid, threads >>>((const GPUTestEntry*)tests, (GPUContact*)contacts, (float3*)map); mycudaDebugError("RigidContactMapperCuda3f_setPoints2_kernel");}
-}
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h
deleted file mode 100644
index e98718ceee4..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h
+++ /dev/null
@@ -1,102 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#pragma once
-#include
-#include
-#include
-
-namespace sofa::gpu::cuda
-{
-extern "C"
-{
- void RigidContactMapperCuda3f_setPoints2(unsigned int size, unsigned int nbTests, unsigned int maxPoints, const void* tests, const void* contacts, void* map);
-}
-}
-
-/// Mapper for CudaRigidDistanceGridCollisionModel
-template
-class sofa::component::collision::response::mapper::ContactMapper :
- public sofa::component::collision::response::mapper::RigidContactMapper
-{
-public:
- typedef typename DataTypes::Real Real;
- typedef typename DataTypes::Coord Coord;
- typedef typename DataTypes::VecCoord VecCoord;
- typedef typename DataTypes::VecDeriv VecDeriv;
- typedef RigidContactMapper Inherit;
- typedef typename Inherit::MMechanicalState MMechanicalState;
- typedef typename Inherit::MCollisionModel MCollisionModel;
-
- int addPoint(const Coord& P, int index, Real& r)
- {
- int i = this->Inherit::addPoint(P, index, r);
- if (!this->mapping)
- {
- MCollisionModel* model = this->model;
- MMechanicalState* outmodel = this->outmodel.get();
- Data* d_x = outmodel->write(sofa::core::vec_id::write_access::position);
- VecDeriv& vx = *d_x->beginEdit();
- Data* d_v = outmodel->write(sofa::core::vec_id::write_access::velocity);
- VecCoord& vv = *d_v->beginEdit();
-
- typename DataTypes::Coord& x = vx[i];
- typename DataTypes::Deriv& v = vv[i];
- if (model->isTransformed(index))
- {
- x = model->getTranslation(index) + model->getRotation(index) * P;
- }
- else
- {
- x = P;
- }
- v = typename DataTypes::Deriv();
-
- d_x->endEdit();
- d_v->endEdit();
- }
- return i;
- }
-
- void setPoints2(sofacuda::GPUDetectionOutputVector* outputs)
- {
- int n = outputs->size();
- int nt = outputs->nbTests();
- int maxp = 0;
- for (int i=0; irtest(i).curSize > maxp) maxp = outputs->rtest(i).curSize;
- if (this->outmodel)
- this->outmodel->resize(n);
- if (this->mapping)
- {
- this->mapping->d_points.beginEdit()->fastResize(n);
- this->mapping->m_rotatedPoints.fastResize(n);
- gpu::cuda::RigidContactMapperCuda3f_setPoints2(n, nt, maxp, outputs->tests.deviceRead(), outputs->results.deviceRead(), this->mapping->d_points.beginEdit()->deviceWrite());
- }
- else
- {
- Data* d_x = this->outmodel->write(sofa::core::vec_id::write_access::position);
- VecCoord& vx = *d_x->beginEdit();
- gpu::cuda::RigidContactMapperCuda3f_setPoints2(n, nt, maxp, outputs->tests.deviceRead(), outputs->results.deviceRead(), vx.deviceWrite());
- d_x->endEdit();
- }
- }
-};
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/config.h.in b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/config.h.in
deleted file mode 100644
index ed81edd83ac..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/config.h.in
+++ /dev/null
@@ -1,37 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#pragma once
-
-#include
-
-#ifdef SOFA_BUILD_SOFADISTANCEGRID_CUDA
-# define SOFA_TARGET @PROJECT_NAME@
-# define SOFA_SOFADISTANCEGRID_CUDA_API SOFA_EXPORT_DYNAMIC_LIBRARY
-#else
-# define SOFA_SOFADISTANCEGRID_CUDA_API SOFA_IMPORT_DYNAMIC_LIBRARY
-#endif
-
-namespace sofadistancegrid::cuda
-{
- constexpr const char* MODULE_NAME = "@PROJECT_NAME@";
- constexpr const char* MODULE_VERSION = "@PROJECT_VERSION@";
-} // namespace sofadistancegrid::cuda
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.cpp
deleted file mode 100644
index fd24b4bdf2a..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.cpp
+++ /dev/null
@@ -1,80 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#include
-#include
-#include
-
-#include
-#include
-
-namespace sofa::gpu::cuda
-{
- extern void registerCudaCollisionDetection(sofa::core::ObjectFactory* factory);
- extern void registerCudaRigidDistanceGridCollisionModel(sofa::core::ObjectFactory* factory);
-}
-
-namespace sofadistancegrid::cuda
-{
-
-extern "C" {
- SOFA_EXPORT_DYNAMIC_LIBRARY void initExternalModule();
- SOFA_EXPORT_DYNAMIC_LIBRARY const char* getModuleName();
- SOFA_EXPORT_DYNAMIC_LIBRARY const char* getModuleVersion();
- SOFA_SOFADISTANCEGRID_API void registerObjects(sofa::core::ObjectFactory* factory);
-}
-
-void initExternalModule()
-{
- init();
-}
-
-const char* getModuleName()
-{
- return MODULE_NAME;
-}
-
-const char* getModuleVersion()
-{
- return MODULE_VERSION;
-}
-
-void init()
-{
- static bool first = true;
- if (first)
- {
- // make sure that this plugin is registered into the PluginManager
- sofa::helper::system::PluginManager::getInstance().registerPlugin(MODULE_NAME);
-
- sofadistancegrid::initSofaDistanceGrid();
- sofa::gpu::cuda::init();
- first = false;
- }
-}
-
-void registerObjects(sofa::core::ObjectFactory* factory)
-{
- sofa::gpu::cuda::registerCudaCollisionDetection(factory);
- sofa::gpu::cuda::registerCudaRigidDistanceGridCollisionModel(factory);
-}
-
-} // namespace volumetricrendering::cuda
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.h
deleted file mode 100644
index df76bdbe8e5..00000000000
--- a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.h
+++ /dev/null
@@ -1,29 +0,0 @@
-/******************************************************************************
-* SOFA, Simulation Open-Framework Architecture *
-* (c) 2006 INRIA, USTL, UJF, CNRS, MGH *
-* *
-* This program is free software; you can redistribute it and/or modify it *
-* under the terms of the GNU Lesser General Public License as published by *
-* the Free Software Foundation; either version 2.1 of the License, or (at *
-* your option) any later version. *
-* *
-* This program is distributed in the hope that it will be useful, but WITHOUT *
-* ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or *
-* FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
-* for more details. *
-* *
-* You should have received a copy of the GNU Lesser General Public License *
-* along with this program. If not, see . *
-*******************************************************************************
-* Authors: The SOFA Team and external contributors (see Authors.txt) *
-* *
-* Contact information: contact@sofa-framework.org *
-******************************************************************************/
-#pragma once
-
-#include
-
-namespace sofadistancegrid::cuda
-{
-SOFA_SOFADISTANCEGRID_CUDA_API void init();
-} // namespace sofadistancegrid::cuda
diff --git a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/CMakeLists.txt
deleted file mode 100644
index d0f6f280175..00000000000
--- a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/CMakeLists.txt
+++ /dev/null
@@ -1,59 +0,0 @@
-cmake_minimum_required(VERSION 3.22)
-project(miniFlowVR)
-
-set(HEADER_FILES
- include/ftl/cmdline.h
- include/ftl/crc.h
- include/ftl/fixed_array.h
- include/ftl/mat.h
- include/ftl/quat.h
- include/ftl/rmath.h
- include/ftl/type.h
- include/ftl/vec.h
- include/flowvr/render/bbox.h
- include/flowvr/render/mesh.h
- include/flowvr/render/mesh.inl
- include/flowvr/render/noise.h
-)
-
-set(SOURCE_FILES
- src/ftlm/cmdline.cpp
- src/ftlm/crc.cpp
- src/ftlm/mat.cpp
- src/ftlm/quat.cpp
- src/ftlm/type.cpp
- src/librender/bbox.cpp
- src/librender/mesh.cpp
- src/librender/mesh_dist.cpp
- src/librender/mesh_io_mesh.cpp
- src/librender/mesh_io_obj.cpp
- src/librender/mesh_io_off.cpp
- src/librender/mesh_io_stl.cpp
- src/librender/mesh_io_vtk.cpp
- src/librender/mesh_io_lwo.cpp
-)
-
-add_library(${PROJECT_NAME} STATIC ${HEADER_FILES} ${SOURCE_FILES})
-target_include_directories(${PROJECT_NAME} PUBLIC "$")
-target_include_directories(${PROJECT_NAME} PUBLIC "$")
-
-target_compile_features(${PROJECT_NAME} PUBLIC cxx_std_17) # static_assert needs at least C++11 (not by default with clang5)
-
-# The code must be relocatable if we want to link a shared library against it
-if("x${CMAKE_CXX_COMPILER_ID}" STREQUAL "xGNU" OR "x${CMAKE_CXX_COMPILER_ID}" STREQUAL "xClang")
- target_compile_options(${PROJECT_NAME} PRIVATE "-fPIC")
-endif()
-
-if(WIN32)
- # remove warnings about deprecation (CRT,etc)
- target_compile_options(${PROJECT_NAME} PRIVATE "/wd4996")
-endif()
-
-include(SofaMacros)
-sofa_create_package_with_targets(
- PACKAGE_NAME MiniFlowVR
- PACKAGE_VERSION 0.1
- TARGETS ${PROJECT_NAME}
- INCLUDE_SOURCE_DIR "include"
- INCLUDE_INSTALL_DIR "extlibs/MiniFlowVR"
- )
diff --git a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/MiniFlowVRConfig.cmake.in b/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/MiniFlowVRConfig.cmake.in
deleted file mode 100644
index 7f5ef1db2ef..00000000000
--- a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/MiniFlowVRConfig.cmake.in
+++ /dev/null
@@ -1,9 +0,0 @@
-# CMake package configuration file for the miniFlowVR library
-
-@PACKAGE_INIT@
-
-if(NOT TARGET miniFlowVR)
- include("${CMAKE_CURRENT_LIST_DIR}/MiniFlowVRTargets.cmake")
-endif()
-
-check_required_components(miniFlowVR)
diff --git a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/README.txt b/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/README.txt
deleted file mode 100644
index fe0d26848a1..00000000000
--- a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/README.txt
+++ /dev/null
@@ -1,2 +0,0 @@
-This folder contains a small subset of tools, mainly mesh manipulations,
-available in FlowVR. See http://flowvr.sf.net/ for the complete package.
diff --git a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/include/flowvr/render/bbox.h b/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/include/flowvr/render/bbox.h
deleted file mode 100644
index 1886ab3fa4e..00000000000
--- a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/include/flowvr/render/bbox.h
+++ /dev/null
@@ -1,112 +0,0 @@
-/******* COPYRIGHT ************************************************
-* *
-* FlowVR Render *
-* Parallel Rendering Library *
-* *
-*-----------------------------------------------------------------*
-* COPYRIGHT (C) 2005 by *
-* Laboratoire Informatique et Distribution (UMR5132) and *
-* INRIA Project MOVI. ALL RIGHTS RESERVED. *
-* *
-* This source is covered by the GNU LGPL, please refer to the *
-* COPYING-LIB file for further information. *
-* *
-*-----------------------------------------------------------------*
-* *
-* Original Contributors: *
-* Jeremie Allard, *
-* Clement Menier. *
-* *
-*******************************************************************
-* *
-* File: ./include/flowvr/render/bbox.h *
-* *
-* Contacts: *
-* *
-******************************************************************/
-#ifndef FLOWVR_RENDER_BBOX_H
-#define FLOWVR_RENDER_BBOX_H
-
-#include
-#include
-
-#include
-
-namespace flowvr
-{
-
-namespace render
-{
-
-class BBox
-{
- public:
- ftl::Vec3f a,b;
-
- BBox();
- BBox(const ftl::Vec3f& _a,const ftl::Vec3f& _b);
-
- bool isEmpty() const;
-
- void clear();
-
- BBox& operator+=(const BBox& p);
-
- BBox& operator+=(const ftl::Vec3f& p);
-
- BBox operator+(const BBox& p);
-
- void apply(const BBox& p,const ftl::Mat4x4f& m);
-
- BBox operator*(const ftl::Mat4x4f& m);
-
- float size() const
- {
- if (isEmpty()) return 0;
- ftl::Vec3f v=b-a;
- return v.norm();
- }
-
- bool in(const ftl::Vec3f& p) const
- {
- return (p.x()>=a.x() && p.x()<=b.x() && p.y()>=a.y() && p.y()<=b.y() && p.z()>=a.z() && p.z()<=b.z());
- }
-
- bool in(const BBox& bb) const
- {
- return (bb.b.x()>=a.x() && bb.a.x()<=b.x() && bb.b.y()>=a.y() && bb.a.y()<=b.y() && bb.b.z()>=a.z() && bb.a.z()<=b.z());
- }
-
- /// Test if bbox intersect with half-space defined as eq.P >= 0
- bool inHalfspace(const ftl::Vec4f& eq) const;
-
- /// Test if bbox intersect with frustum defined by the given matrix
- /// (interpreted as an OpenGL projection matrix)
- /// if exact is false then the result is overestimated
- /// if exact is true then a 3x3 matrix inversion is required
- bool inFrustum(const ftl::Mat4x4f& frustum, bool exact=false) const;
-
- bool operator==(const BBox& p) const
- {
- return a==p.a && b==p.b;
- }
-
- bool operator!=(const BBox& p) const
- {
- return a!=p.a || b!=p.b;
- }
-
- std::string toString() const;
- void fromString(const std::string& text);
-
-};
-
-} // namespace render
-
-} // namespace flowvr
-
-// iostream
-std::ostream& operator<<(std::ostream& o, const flowvr::render::BBox& b);
-std::istream& operator>>(std::istream& in, flowvr::render::BBox& b);
-
-#endif
diff --git a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/include/flowvr/render/mesh.h b/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/include/flowvr/render/mesh.h
deleted file mode 100644
index 4e1a6f9463d..00000000000
--- a/applications/plugins/SofaDistanceGrid/extlibs/miniFlowVR/include/flowvr/render/mesh.h
+++ /dev/null
@@ -1,426 +0,0 @@
-/******* COPYRIGHT ************************************************
-* *
-* FlowVR Render *
-* Parallel Rendering Library *
-* *
-*-----------------------------------------------------------------*
-* COPYRIGHT (C) 2005 by *
-* Laboratoire Informatique et Distribution (UMR5132) and *
-* INRIA Project MOVI. ALL RIGHTS RESERVED. *
-* *
-* This source is covered by the GNU LGPL, please refer to the *
-* COPYING-LIB file for further information. *
-* *
-*-----------------------------------------------------------------*
-* *
-* Original Contributors: *
-* Jeremie Allard, *
-* Clement Menier. *
-* *
-*******************************************************************
-* *
-* File: ./include/flowvr/render/mesh.h *
-* *
-* Contacts: *
-* *
-******************************************************************/
-#ifndef FLOWVR_RENDER_MESH_H
-#define FLOWVR_RENDER_MESH_H
-
-#define MINI_FLOWVR
-
-#include
-#include
-#include "bbox.h"
-#ifndef MINI_FLOWVR
-#include "chunkwriter.h"
-#endif
-
-#include
-#include