[SofaCUDA,SofaDistanceGrid] Move DistanceGrid CUDA files in a SofaDistanceGrid extension

Sofa/framework/Config/CMakeLists.txt

@@ -296,13 +296,22 @@ set_target_properties(${PROJECT_NAME} PROPERTIES
   MAP_IMPORTED_CONFIG_MINSIZEREL Release
   MAP_IMPORTED_CONFIG_RELWITHDEBINFO Release
   )
-target_compile_options(${PROJECT_NAME} PUBLIC "$<$<CONFIG:RELEASE>:${SOFACONFIG_COMPILE_OPTIONS_RELEASE}>")
-target_compile_options(${PROJECT_NAME} PUBLIC "$<$<CONFIG:DEBUG>:${SOFACONFIG_COMPILE_OPTIONS_DEBUG}>")
-target_compile_options(${PROJECT_NAME} PUBLIC "${SOFACONFIG_COMPILE_OPTIONS}")
 
-target_link_options(${PROJECT_NAME} PUBLIC "$<$<CONFIG:RELEASE>:${SOFACONFIG_LINK_OPTIONS_RELEASE}>")
-target_link_options(${PROJECT_NAME} PUBLIC "$<$<CONFIG:DEBUG>:${SOFACONFIG_LINK_OPTIONS_DEBUG}>")
-target_link_options(${PROJECT_NAME} PUBLIC "${SOFACONFIG_LINK_OPTIONS}")
+set(is_cxx "$<COMPILE_LANGUAGE:CXX>")
+set(is_c "$<COMPILE_LANGUAGE:C>")
+set(is_c_cxx "$<OR:${is_cxx},${is_c}>")
+set(is_release "$<CONFIG:RELEASE>")
+set(is_debug "$<CONFIG:DEBUG>")
+set(is_c_cxx_release "$<AND:${is_c_cxx},${is_release}>")
+set(is_c_cxx_debug "$<AND:${is_c_cxx},${is_debug}>")
+
+target_compile_options(${PROJECT_NAME} PUBLIC "$<${is_c_cxx_release}:${SOFACONFIG_COMPILE_OPTIONS_RELEASE}>")
+target_compile_options(${PROJECT_NAME} PUBLIC "$<${is_c_cxx_debug}:${SOFACONFIG_COMPILE_OPTIONS_DEBUG}>")
+target_compile_options(${PROJECT_NAME} PUBLIC "$<${is_c_cxx}:${SOFACONFIG_COMPILE_OPTIONS}>")
+
+target_link_options(${PROJECT_NAME} PUBLIC "$<${is_c_cxx_release}:${SOFACONFIG_LINK_OPTIONS_RELEASE}>")
+target_link_options(${PROJECT_NAME} PUBLIC "$<${is_c_cxx_debug}:${SOFACONFIG_LINK_OPTIONS_DEBUG}>")
+target_link_options(${PROJECT_NAME} PUBLIC "$<${is_c_cxx}:${SOFACONFIG_LINK_OPTIONS}>")
 
 set_target_properties(${PROJECT_NAME} PROPERTIES FOLDER Sofa.Framework) # IDE folder
 

applications/plugins/CMakeLists.txt

@@ -16,8 +16,6 @@ sofa_add_subdirectory(plugin CImgPlugin CImgPlugin) # ON by default and first as
 sofa_add_subdirectory(plugin ArticulatedSystemPlugin ArticulatedSystemPlugin ON)
 sofa_add_subdirectory(plugin SofaEulerianFluid SofaEulerianFluid)
 sofa_add_subdirectory(plugin SofaSphFluid SofaSphFluid EXTERNAL GIT_REF master)
-sofa_add_subdirectory(plugin SofaDistanceGrid SofaDistanceGrid) # Depends on SofaMiscCollision
-sofa_add_subdirectory(plugin SofaImplicitField SofaImplicitField)
 sofa_add_subdirectory(plugin MultiThreading MultiThreading ON)
 sofa_add_subdirectory(plugin DiffusionSolver DiffusionSolver) # Depends on CImgPlugin
 sofa_add_subdirectory(plugin image image) # Depends on CImgPlugin, DiffusionSolver, MultiThreading (soft)
@@ -73,3 +71,6 @@ if(Sofa.GL_FOUND)
 else()
     message("Sofa.GL not found; disabling SofaSimpleGUI and VolumetricRendering plugins")
 endif()
+
+sofa_add_subdirectory(plugin SofaDistanceGrid SofaDistanceGrid) # Also defines SofaDistanceGrid.CUDA
+sofa_add_subdirectory(plugin SofaImplicitField SofaImplicitField) # Depends on SofaDistanceGrid

applications/plugins/SofaCUDA/CMakeLists.txt

@@ -19,6 +19,7 @@ set(HEADER_FILES
     ### Common
     sofa/gpu/cuda/CudaBaseVector.h
     sofa/gpu/cuda/CudaCommon.h
+    sofa/gpu/cuda/CudaContactMapper.h
     sofa/gpu/cuda/CudaMath.h
     sofa/gpu/cuda/CudaMath.inl
     sofa/gpu/cuda/CudaMathRigid.h
@@ -105,6 +106,7 @@ set(SOURCE_FILES
     ### Common
     init.cpp
     sofa/gpu/cuda/CudaBaseVector.cpp
+    sofa/gpu/cuda/CudaContactMapper.cpp
     sofa/gpu/cuda/mycuda.cpp
 
     ### Mechanical
@@ -182,6 +184,7 @@ set(CUDA_SOURCES
     ### Common
     sofa/gpu/cuda/mycuda.cu
     sofa/gpu/cuda/CudaBaseVector.cu
+    sofa/gpu/cuda/CudaContactMapper.cu
     sofa/gpu/cuda/CudaScan.cu
     sofa/gpu/cuda/CudaSort.cu
 
@@ -269,33 +272,6 @@ if(Sofa.GUI.Qt_FOUND)
     list(APPEND SOURCE_FILES sofa/gpu/gui/CudaDataWidget.cpp)
 endif()
 
-
-find_package(SofaDistanceGrid QUIET)
-if(SofaDistanceGrid_FOUND)
-    sofa_find_package(MiniFlowVR QUIET)
-    if(MiniFlowVR_FOUND)
-        message(STATUS "SofaCUDA: MiniFlowVR enabled.")
-    else()
-        message(STATUS "SofaCUDA: MiniFlowVR was not enabled, therefore some feature of CudaDistanceGridCollisionModel will not be compiled.")
-    endif()
-
-    list(APPEND HEADER_FILES
-        sofa/gpu/cuda/CudaDistanceGridCollisionModel.h
-        sofa/gpu/cuda/CudaContactMapper.h
-        sofa/gpu/cuda/CudaCollisionDetection.h)
-
-    list(APPEND SOURCE_FILES
-        sofa/gpu/cuda/CudaDistanceGridCollisionModel.cpp
-        sofa/gpu/cuda/CudaCollisionDistanceGrid.cpp
-        sofa/gpu/cuda/CudaCollisionDetection.cpp)
-
-    list(APPEND CUDA_SOURCES
-        sofa/gpu/cuda/CudaContactMapper.cu
-        sofa/gpu/cuda/CudaCollisionDetection.cu)
-else()
-    message(STATUS "SofaCUDA: SofaDistanceGrid was not enabled, therefore CudaDistanceGridCollisionModel will not be compiled.")
-endif()
-
 find_package(SofaSphFluid QUIET)
 if(SofaSphFluid_FOUND)
     list(APPEND HEADER_FILES
@@ -426,14 +402,6 @@ if(Sofa.GL_FOUND)
     target_link_libraries(${PROJECT_NAME} Sofa.GL)
 endif()
 
-if(SofaDistanceGrid_FOUND)
-    target_link_libraries(${PROJECT_NAME} SofaDistanceGrid)
-    if(MiniFlowVR_FOUND)
-        target_link_libraries(${PROJECT_NAME} miniFlowVR)
-    endif()
-endif()
-
-
 if(Sofa.GUI_FOUND)
     target_link_libraries(${PROJECT_NAME} Sofa.GUI)
 endif()

applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cpp

@@ -0,0 +1,32 @@
+/******************************************************************************
+*                 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 <http://www.gnu.org/licenses/>.        *
+*******************************************************************************
+* Authors: The SOFA Team and external contributors (see Authors.txt)          *
+*                                                                             *
+* Contact information: [email protected]                             *
+******************************************************************************/
+#define SOFACUDA_CUDACONTACTMAPPER_CPP
+
+#include <SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h>
+
+namespace sofa::component::collision
+{
+
+template class SOFA_GPU_CUDA_API response::mapper::ContactMapper<sofa::gpu::cuda::CudaPointCollisionModel, CudaVec3fTypes>;
+template class SOFA_GPU_CUDA_API response::mapper::ContactMapper<sofa::gpu::cuda::CudaSphereCollisionModel, CudaVec3fTypes>;
+
+}

applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cu

@@ -34,7 +34,6 @@ namespace cuda
 
 extern "C"
 {
-    void RigidContactMapperCuda3f_setPoints2(unsigned int size, unsigned int nbTests, unsigned int maxPoints, const void* tests, const void* contacts, void* map);
     void SubsetContactMapperCuda3f_setPoints1(unsigned int size, unsigned int nbTests, unsigned int maxPoints, unsigned int nbPointsPerElem, const void* tests, const void* contacts, void* map);
 }
 
@@ -61,20 +60,6 @@ struct /*__align__(8)*/ GPUTestEntry
 
 __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;
-    }
-}
-
 __global__ void SubsetContactMapperCuda3f_setPoints1_kernel(unsigned int nbPointsPerElem, const GPUTestEntry* tests, const GPUContact* contacts, int* map)
 {
     if (threadIdx.x == 0)
@@ -94,15 +79,6 @@ __global__ void SubsetContactMapperCuda3f_setPoints1_kernel(unsigned int nbPoint
 // CPU-side methods //
 //////////////////////
 
-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");}
-}
-
 void SubsetContactMapperCuda3f_setPoints1(unsigned int size, unsigned int nbTests, unsigned int maxPoints, unsigned int nbPointsPerElem, const void* tests, const void* contacts, void* map)
 {
     // round up to 16

applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h

@@ -24,21 +24,19 @@
 #include <sofa/component/collision/response/mapper/BarycentricContactMapper.h>
 #include <sofa/component/collision/response/mapper/RigidContactMapper.inl>
 #include <sofa/component/collision/response/mapper/SubsetContactMapper.inl>
-#include <sofa/gpu/cuda/CudaDistanceGridCollisionModel.h>
 #include <SofaCUDA/component/collision/geometry/CudaSphereModel.h>
 #include <SofaCUDA/component/collision/geometry/CudaPointModel.h>
-#include <sofa/gpu/cuda/CudaCollisionDetection.h>
+#include <sofa/gpu/cuda/GPUDetectionOutputVector.h>
 #include <SofaCUDA/component/mapping/nonlinear/CudaRigidMapping.h>
 #include <SofaCUDA/component/mapping/linear/CudaSubsetMapping.h>
-
+#include <sofa/gpu/cuda/CudaTypes.h>
 
 
 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);
     void SubsetContactMapperCuda3f_setPoints1(unsigned int size, unsigned int nbTests, unsigned int maxPoints, unsigned int nbPointsPerElem, const void* tests, const void* contacts, void* map);
 }
 
@@ -52,76 +50,6 @@ using namespace sofa::defaulttype;
 using namespace sofa::gpu::cuda;
 using sofa::core::collision::GPUDetectionOutputVector;
 
-
-/// Mapper for CudaRigidDistanceGridCollisionModel
-template <class DataTypes>
-class response::mapper::ContactMapper<sofa::gpu::cuda::CudaRigidDistanceGridCollisionModel,DataTypes> : public response::mapper::RigidContactMapper<sofa::gpu::cuda::CudaRigidDistanceGridCollisionModel,DataTypes>
-{
-public:
-    typedef typename DataTypes::Real Real;
-    typedef typename DataTypes::Coord Coord;
-    typedef typename DataTypes::VecCoord VecCoord;
-    typedef typename DataTypes::VecDeriv VecDeriv;
-    typedef RigidContactMapper<sofa::gpu::cuda::CudaRigidDistanceGridCollisionModel,DataTypes> 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<VecCoord>* d_x = outmodel->write(core::VecCoordId::position());
-            VecDeriv& vx = *d_x->beginEdit();
-            Data<VecDeriv>* d_v = outmodel->write(core::VecDerivId::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(GPUDetectionOutputVector* outputs)
-    {
-        int n = outputs->size();
-        int nt = outputs->nbTests();
-        int maxp = 0;
-        for (int i=0; i<nt; i++)
-            if (outputs->rtest(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<VecCoord>* d_x = this->outmodel->write(core::VecCoordId::position());
-            VecCoord& vx = *d_x->beginEdit();
-            gpu::cuda::RigidContactMapperCuda3f_setPoints2(n, nt, maxp, outputs->tests.deviceRead(), outputs->results.deviceRead(), vx.deviceWrite());
-            d_x->endEdit();
-        }
-    }
-};
-
-
 /// Mapper for CudaPointDistanceGridCollisionModel
 template <class DataTypes>
 class response::mapper::ContactMapper<sofa::gpu::cuda::CudaPointCollisionModel,DataTypes> : public response::mapper::SubsetContactMapper<sofa::gpu::cuda::CudaPointCollisionModel,DataTypes>
@@ -186,4 +114,9 @@ class response::mapper::ContactMapper<CudaSphereCollisionModel, DataTypes> : pub
     }
 };
 
+#if !defined(SOFACUDA_CUDACONTACTMAPPER_CPP)
+extern template class SOFA_GPU_CUDA_API response::mapper::ContactMapper<sofa::gpu::cuda::CudaPointCollisionModel, sofa::gpu::cuda::CudaVec3fTypes>;
+extern template class SOFA_GPU_CUDA_API response::mapper::ContactMapper<sofa::gpu::cuda::CudaSphereCollisionModel, sofa::gpu::cuda::CudaVec3fTypes>;
+#endif
+
 } // namespace sofa::component::collision