diff --git a/Sofa/framework/Config/CMakeLists.txt b/Sofa/framework/Config/CMakeLists.txt
index fffe07ef454..91202e1a826 100644
--- a/Sofa/framework/Config/CMakeLists.txt
+++ b/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
 
diff --git a/applications/plugins/CMakeLists.txt b/applications/plugins/CMakeLists.txt
index 5b0ef1c4bbb..cfd74c7f771 100644
--- a/applications/plugins/CMakeLists.txt
+++ b/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
diff --git a/applications/plugins/SofaCUDA/CMakeLists.txt b/applications/plugins/SofaCUDA/CMakeLists.txt
index b752d64562b..343aa18e187 100644
--- a/applications/plugins/SofaCUDA/CMakeLists.txt
+++ b/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()
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cpp b/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cpp
new file mode 100644
index 00000000000..d5cfa67d104
--- /dev/null
+++ b/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: contact@sofa-framework.org                             *
+******************************************************************************/
+#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>;
+
+}
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cu b/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cu
index a8e6be42025..c7651fd5bd8 100644
--- a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cu
+++ b/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
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h b/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h
index 9f4ffedf604..33196df2843 100644
--- a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h
+++ b/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h
@@ -24,13 +24,12 @@
 #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
@@ -38,7 +37,6 @@ 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
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/GPUDetectionOutputVector.h b/applications/plugins/SofaCUDA/sofa/gpu/cuda/GPUDetectionOutputVector.h
new file mode 100644
index 00000000000..2a32d83bbb5
--- /dev/null
+++ b/applications/plugins/SofaCUDA/sofa/gpu/cuda/GPUDetectionOutputVector.h
@@ -0,0 +1,182 @@
+/******************************************************************************
+*                 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: contact@sofa-framework.org                             *
+******************************************************************************/
+#pragma once
+#include <sofa/core/collision/DetectionOutput.h>
+#include <sofa/gpu/cuda/CudaTypes.h>
+
+
+namespace sofacuda
+{
+
+using sofa::type::Vec3f;
+
+/**
+ *  \brief Generic description of a contact point using GPU.
+ *
+ *  Each contact point is described by :
+ *
+ *  \item p1: index of the collision element
+ *  \item P2: position of the contact point
+ *  \item distance: estimated penetration distance
+ *  \item normal: contact normal in global space
+ */
+struct GPUContactPoint
+{
+    Vec3f p;
+    int elem;
+};
+
+/**
+ *  \brief Generic description of a contact using GPU.
+ *
+ *  Each contact is described by :
+ *
+ *  \item distance: estimated penetration distance
+ *  \item normal: contact normal in global space
+ */
+struct GPUContact
+{
+    Vec3f normal;
+    float distance;
+};
+
+
+/**
+ *  \brief Abstract description of a set of contact point using GPU.
+ */
+class GPUDetectionOutputVector : public sofa::core::collision::DetectionOutputVector
+{
+public:
+    ~GPUDetectionOutputVector() override
+    {
+    }
+
+    sofa::gpu::cuda::CudaVector<GPUContactPoint> results1, results2;
+    sofa::gpu::cuda::CudaVector<GPUContact> results;
+    struct TestEntry
+    {
+        int firstIndex; ///< Index of the first result in the results array
+        int maxSize; ///< Maximum number of contacts resulting from this test
+        int curSize; ///< Current number of detected contacts
+        int newIndex; ///< Index of the first result in a new compacted array
+        std::pair<int,int> elems; ///< Elements
+        TestEntry() : firstIndex(0), maxSize(0), curSize(0), newIndex(0), elems(std::make_pair(0,0)) {}
+    };
+    sofa::gpu::cuda::CudaVector< TestEntry > tests;
+
+    unsigned int size() const
+    {
+        if (results.empty()) return 0;
+        int s = 0;
+        for (unsigned int i=0; i<tests.size(); i++)
+            s += tests[i].curSize;
+        return s;
+    }
+
+    void clear()
+    {
+        results1.clear();
+        results2.clear();
+        results.clear();
+        tests.clear();
+    }
+
+    void release()
+    {
+        // GPU vectors are stored in other data structures, they should not be deleted by the pipeline
+    }
+
+    unsigned int nbTests() { return tests.size(); }
+
+    const TestEntry& rtest(int i) { return tests[i]; }
+    TestEntry& wtest(int i) { return tests[i]; }
+
+    int addTest(std::pair<int,int> elems, int maxSize)
+    {
+        int t = tests.size();
+        TestEntry e;
+        e.elems = elems;
+        e.firstIndex = results.size();
+        e.maxSize = maxSize;
+        e.curSize = 0;
+        e.newIndex = 0;
+        results.fastResize(e.firstIndex+maxSize);
+        results1.fastResize(e.firstIndex+maxSize);
+        results2.fastResize(e.firstIndex+maxSize);
+        tests.push_back(e);
+        return t;
+    }
+
+    const GPUContact* get(int i)
+    {
+        unsigned int t=0;
+        while(t<nbTests() && rtest(t).newIndex > i) ++t;
+        if (t<nbTests())
+            return &(results[rtest(t).firstIndex + (i-rtest(t).newIndex)]);
+        else
+            return NULL;
+    }
+
+    const GPUContactPoint* getP1(int i)
+    {
+        unsigned int t=0;
+        while(t<nbTests() && rtest(t).newIndex > i) ++t;
+        if (t<nbTests())
+            return &(results1[rtest(t).firstIndex + (i-rtest(t).newIndex)]);
+        else
+            return NULL;
+    }
+
+    const GPUContactPoint* getP2(int i)
+    {
+        unsigned int t=0;
+        while(t<nbTests() && rtest(t).newIndex > i) ++t;
+        if (t<nbTests())
+            return &(results2[rtest(t).firstIndex + (i-rtest(t).newIndex)]);
+        else
+            return nullptr;
+    }
+
+    /// Const iterator to iterate the detection pairs
+    virtual sofa::type::Vec3 getFirstPosition(unsigned idx) override
+    {
+        return sofa::type::Vec3(getP1(idx)->p[0],getP1(idx)->p[1],getP1(idx)->p[2]);
+    }
+
+    /// Const iterator end to iterate the detection pairs
+    virtual sofa::type::Vec3 getSecondPosition(unsigned idx) override
+    {
+        return sofa::type::Vec3(getP2(idx)->p[0],getP2(idx)->p[1],getP2(idx)->p[2]);
+    }
+
+};
+
+}
+
+namespace sofa::core::collision
+{
+
+using GPUContactPoint = sofacuda::GPUContactPoint;
+using GPUContact = sofacuda::GPUContact;
+using GPUDetectionOutputVector = sofacuda::GPUDetectionOutputVector;
+
+}
diff --git a/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/geometry/CudaPointModel.h b/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/geometry/CudaPointModel.h
index 1be2ccbcbbc..eb9327ac8d6 100644
--- a/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/geometry/CudaPointModel.h
+++ b/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/geometry/CudaPointModel.h
@@ -36,7 +36,7 @@ using namespace sofa::defaulttype;
 
 class CudaPointCollisionModel;
 
-class CudaPoint : public core::TCollisionElementIterator<CudaPointCollisionModel>
+class SOFA_GPU_CUDA_API CudaPoint : public core::TCollisionElementIterator<CudaPointCollisionModel>
 {
 public:
     CudaPoint(CudaPointCollisionModel* model, Index index);
@@ -47,7 +47,7 @@ class CudaPoint : public core::TCollisionElementIterator<CudaPointCollisionModel
     explicit CudaPoint(const core::CollisionElementIterator& i);
 };
 
-class CudaPointCollisionModel : public core::CollisionModel
+class SOFA_GPU_CUDA_API CudaPointCollisionModel : public core::CollisionModel
 {
 public:
     SOFA_CLASS(CudaPointCollisionModel,core::CollisionModel);
diff --git a/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.h b/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.h
index ec0a4fe4ee1..17385757d47 100644
--- a/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.h
+++ b/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.h
@@ -23,7 +23,7 @@
 
 #include <sofa/gpu/cuda/CudaTypes.h>
 #include <sofa/component/collision/response/contact/PenalityContactForceField.h>
-#include <sofa/gpu/cuda/CudaCollisionDetection.h>
+#include <sofa/gpu/cuda/GPUDetectionOutputVector.h>
 
 namespace sofa::component::collision::response::contact
 {
@@ -81,7 +81,7 @@ class PenalityContactForceField<CudaVec3fTypes> : public core::behavior::PairInt
 
     void addContact(int m1, int m2, const Deriv& norm, Real dist, Real ks, Real mu_s = 0.0f, Real mu_v = 0.0f, int oldIndex = 0);
 
-    void setContacts(Real distance, Real ks, sofa::core::collision::GPUDetectionOutputVector* outputs, bool useDistance, type::Mat3x3f* normXForm = NULL);
+    void setContacts(Real distance, Real ks, sofacuda::GPUDetectionOutputVector* outputs, bool useDistance, type::Mat3x3f* normXForm = nullptr);
 
     virtual void addForce(const core::MechanicalParams* mparams, DataVecDeriv& d_f1, DataVecDeriv& d_f2, const DataVecCoord& d_x1, const DataVecCoord& d_x2, const DataVecDeriv& d_v1, const DataVecDeriv& d_v2) override;
 
diff --git a/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.inl b/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.inl
index 5c68526fb1d..6042beabb6e 100644
--- a/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.inl
+++ b/applications/plugins/SofaCUDA/src/SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.inl
@@ -86,7 +86,7 @@ void PenalityContactForceField<CudaVec3fTypes>::addContact(int /*m1*/, int /*m2*
     pen.push_back(0);
 }
 
-void PenalityContactForceField<CudaVec3fTypes>::setContacts(Real d0, Real stiffness, sofa::core::collision::GPUDetectionOutputVector* outputs, bool useDistance, type::Mat3x3f* normXForm)
+void PenalityContactForceField<CudaVec3fTypes>::setContacts(Real d0, Real stiffness, sofacuda::GPUDetectionOutputVector* outputs, bool useDistance, type::Mat3x3f* normXForm)
 {
 #if 1
     int n = outputs->size();
diff --git a/applications/plugins/SofaCUDA/src/SofaCUDA/component/mapping/linear/CudaSubsetMapping.h b/applications/plugins/SofaCUDA/src/SofaCUDA/component/mapping/linear/CudaSubsetMapping.h
index 0ecf1f53bbe..3c17ecbc25a 100644
--- a/applications/plugins/SofaCUDA/src/SofaCUDA/component/mapping/linear/CudaSubsetMapping.h
+++ b/applications/plugins/SofaCUDA/src/SofaCUDA/component/mapping/linear/CudaSubsetMapping.h
@@ -146,10 +146,10 @@ void SubsetMapping<gpu::cuda::CudaVec3fTypes, gpu::cuda::CudaVec3f1Types>::apply
 
 
 #ifndef SOFA_GPU_CUDA_CUDASUBSETMAPPING_CPP
-extern template class SOFA_GPU_CUDA_API SubsetMapping< CudaVec3fTypes, CudaVec3fTypes >;
-extern template class SOFA_GPU_CUDA_API SubsetMapping< CudaVec3f1Types, CudaVec3f1Types >;
-extern template class SOFA_GPU_CUDA_API SubsetMapping< CudaVec3f1Types, CudaVec3fTypes >;
-extern template class SOFA_GPU_CUDA_API SubsetMapping< CudaVec3fTypes, CudaVec3f1Types >;
+extern template class SOFA_GPU_CUDA_API SubsetMapping< sofa::gpu::cuda::CudaVec3fTypes, sofa::gpu::cuda::CudaVec3fTypes >;
+extern template class SOFA_GPU_CUDA_API SubsetMapping< sofa::gpu::cuda::CudaVec3f1Types, sofa::gpu::cuda::CudaVec3f1Types >;
+extern template class SOFA_GPU_CUDA_API SubsetMapping< sofa::gpu::cuda::CudaVec3f1Types, sofa::gpu::cuda::CudaVec3fTypes >;
+extern template class SOFA_GPU_CUDA_API SubsetMapping< sofa::gpu::cuda::CudaVec3fTypes, sofa::gpu::cuda::CudaVec3f1Types >;
 #endif
 
 } // namespace sofa::component::mapping::linear
diff --git a/applications/plugins/SofaDistanceGrid/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/CMakeLists.txt
index d6bc2b621ed..0d67279ce5d 100644
--- a/applications/plugins/SofaDistanceGrid/CMakeLists.txt
+++ b/applications/plugins/SofaDistanceGrid/CMakeLists.txt
@@ -22,6 +22,7 @@ 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
@@ -86,3 +87,6 @@ sofa_create_package_with_targets(
     INCLUDE_INSTALL_DIR "${PROJECT_NAME}"
     RELOCATABLE "plugins"
     )
+
+set(PLUGIN_SOFADISTANCEGRID_CUDA ON BOOL STRING "Build the SofaDistanceGrid.CUDA plugin")
+sofa_add_subdirectory(plugin extensions/CUDA SofaDistanceGrid.CUDA)
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/CMakeLists.txt b/applications/plugins/SofaDistanceGrid/extensions/CUDA/CMakeLists.txt
new file mode 100644
index 00000000000..b5f9f033313
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/CMakeLists.txt
@@ -0,0 +1,41 @@
+cmake_minimum_required(VERSION 3.22)
+project(SofaDistanceGrid.CUDA 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)
+
+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
new file mode 100644
index 00000000000..62faad7f3c6
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/SofaDistanceGrid.CUDAConfig.cmake.in
@@ -0,0 +1,9 @@
+# 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/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cpp
similarity index 100%
rename from applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.cpp
rename to applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cpp
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.cu b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu
similarity index 99%
rename from applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.cu
rename to applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu
index ce2c2fcb108..06aa701ed14 100644
--- a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.cu
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.cu
@@ -21,7 +21,7 @@
 ******************************************************************************/
 #include <sofa/gpu/cuda/CudaCommon.h>
 #include <sofa/gpu/cuda/CudaMath.h>
-#include "mycuda.h"
+#include <sofa/gpu/cuda/mycuda.h>
 #include "cuda.h"
 
 #if defined(__cplusplus) && CUDA_VERSION < 2000
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h
similarity index 67%
rename from applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.h
rename to applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h
index 856286391c5..14a0c9753d9 100644
--- a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDetection.h
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDetection.h
@@ -23,11 +23,12 @@
 #define SOFA_GPU_CUDA_CUDACOLLISIONDETECTION_H
 
 #include <sofa/core/collision/DetectionOutput.h>
-#include <sofa/gpu/cuda/CudaDistanceGridCollisionModel.h>
+#include <SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h>
 #include <SofaCUDA/component/collision/geometry/CudaSphereModel.h>
 #include <SofaCUDA/component/collision/geometry/CudaPointModel.h>
 #include <sofa/component/collision/detection/algorithm/BruteForceBroadPhase.h>
 #include <sofa/core/collision/NarrowPhaseDetection.h>
+#include <sofa/gpu/cuda/GPUDetectionOutputVector.h>
 
 
 namespace sofa
@@ -41,161 +42,7 @@ namespace collision
 
 using type::Vec3f;
 using type::Mat3x3f;
-
-/**
- *  \brief Generic description of a contact point using GPU.
- *
- *  Each contact point is described by :
- *
- *  \item p1: index of the collision element
- *  \item P2: position of the contact point
- *  \item distance: estimated penetration distance
- *  \item normal: contact normal in global space
- */
-
-struct GPUContactPoint
-{
-    Vec3f p;
-    int elem;
-};
-
-
-/**
- *  \brief Generic description of a contact using GPU.
- *
- *  Each contact is described by :
- *
- *  \item distance: estimated penetration distance
- *  \item normal: contact normal in global space
- */
-struct GPUContact
-{
-    Vec3f normal;
-    float distance;
-};
-
-/*
-struct GPUDetectionOutput
-{
-    int p1;
-    Vec3f p2;
-    //union {
-        float distance;
-    //    int i2;
-    //};
-    Vec3f normal;
-};
-*/
-
-/**
- *  \brief Abstract description of a set of contact point using GPU.
- */
-class GPUDetectionOutputVector : public DetectionOutputVector
-{
-public:
-    ~GPUDetectionOutputVector()
-    {
-    }
-
-    sofa::gpu::cuda::CudaVector<GPUContactPoint> results1, results2;
-    sofa::gpu::cuda::CudaVector<GPUContact> results;
-    struct TestEntry
-    {
-        int firstIndex; ///< Index of the first result in the results array
-        int maxSize; ///< Maximum number of contacts resulting from this test
-        int curSize; ///< Current number of detected contacts
-        int newIndex; ///< Index of the first result in a new compacted array
-        std::pair<int,int> elems; ///< Elements
-        TestEntry() : firstIndex(0), maxSize(0), curSize(0), newIndex(0), elems(std::make_pair(0,0)) {}
-    };
-    sofa::gpu::cuda::CudaVector< TestEntry > tests;
-
-    unsigned int size() const
-    {
-        if (results.empty()) return 0;
-        int s = 0;
-        for (unsigned int i=0; i<tests.size(); i++)
-            s += tests[i].curSize;
-        return s;
-    }
-
-    void clear()
-    {
-        results1.clear();
-        results2.clear();
-        results.clear();
-        tests.clear();
-    }
-
-    void release()
-    {
-        // GPU vectors are stored in other data structures, they should not be deleted by the pipeline
-    }
-
-    unsigned int nbTests() { return tests.size(); }
-
-    const TestEntry& rtest(int i) { return tests[i]; }
-    TestEntry& wtest(int i) { return tests[i]; }
-
-    int addTest(std::pair<int,int> elems, int maxSize)
-    {
-        int t = tests.size();
-        TestEntry e;
-        e.elems = elems;
-        e.firstIndex = results.size();
-        e.maxSize = maxSize;
-        e.curSize = 0;
-        e.newIndex = 0;
-        results.fastResize(e.firstIndex+maxSize);
-        results1.fastResize(e.firstIndex+maxSize);
-        results2.fastResize(e.firstIndex+maxSize);
-        tests.push_back(e);
-        return t;
-    }
-
-    const GPUContact* get(int i)
-    {
-        unsigned int t=0;
-        while(t<nbTests() && rtest(t).newIndex > i) ++t;
-        if (t<nbTests())
-            return &(results[rtest(t).firstIndex + (i-rtest(t).newIndex)]);
-        else
-            return NULL;
-    }
-
-    const GPUContactPoint* getP1(int i)
-    {
-        unsigned int t=0;
-        while(t<nbTests() && rtest(t).newIndex > i) ++t;
-        if (t<nbTests())
-            return &(results1[rtest(t).firstIndex + (i-rtest(t).newIndex)]);
-        else
-            return NULL;
-    }
-
-    const GPUContactPoint* getP2(int i)
-    {
-        unsigned int t=0;
-        while(t<nbTests() && rtest(t).newIndex > i) ++t;
-        if (t<nbTests())
-            return &(results2[rtest(t).firstIndex + (i-rtest(t).newIndex)]);
-        else
-            return NULL;
-    }
-
-    /// Const iterator to iterate the detection pairs
-    virtual type::Vec3 getFirstPosition(unsigned idx) override
-    {
-        return type::Vec3(getP1(idx)->p[0],getP1(idx)->p[1],getP1(idx)->p[2]);
-    }
-
-    /// Const iterator end to iterate the detection pairs
-    virtual type::Vec3 getSecondPosition(unsigned idx) override
-    {
-        return type::Vec3(getP2(idx)->p[0],getP2(idx)->p[1],getP2(idx)->p[2]);
-    }
-
-};
+using sofacuda::GPUDetectionOutputVector;
 
 template<>
 class TDetectionOutputVector<sofa::gpu::cuda::CudaRigidDistanceGridCollisionModel,sofa::gpu::cuda::CudaRigidDistanceGridCollisionModel> : public GPUDetectionOutputVector
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDistanceGrid.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp
similarity index 97%
rename from applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDistanceGrid.cpp
rename to applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp
index fe0f57ea6fa..566d72c8305 100644
--- a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaCollisionDistanceGrid.cpp
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaCollisionDistanceGrid.cpp
@@ -23,7 +23,8 @@
 #include <SofaCUDA/component/solidmechanics/spring/CudaSpringForceField.inl>
 #include <SofaCUDA/component/statecontainer/CudaMechanicalObject.inl>
 #include <SofaCUDA/component/mapping/linear/CudaIdentityMapping.inl>
-#include "CudaContactMapper.h"
+#include <SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h>
+#include <SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h>
 #include <SofaCUDA/component/collision/response/contact/CudaPenalityContactForceField.h>
 #include "CudaDistanceGridCollisionModel.h"
 #include <SofaCUDA/component/collision/geometry/CudaSphereModel.h>
@@ -56,7 +57,7 @@ using namespace sofa::gpu::cuda;
 template <>
 void BarycentricPenalityContact<CudaPointCollisionModel,CudaRigidDistanceGridCollisionModel,CudaVec3fTypes>::setDetectionOutputs(OutputVector* o)
 {
-    TOutputVector& outputs = *static_cast<TOutputVector*>(o);
+    GPUDetectionOutputVector& outputs = *dynamic_cast<GPUDetectionOutputVector*>(o);
     //const bool printLog = this->f_printLog.getValue();
     if (ff==NULL)
     {
@@ -109,7 +110,7 @@ template <>
 void BarycentricPenalityContact<CudaSphereCollisionModel,CudaRigidDistanceGridCollisionModel,CudaVec3fTypes>::setDetectionOutputs(OutputVector* o)
 {
 
-    TOutputVector& outputs = *static_cast<TOutputVector*>(o);
+    GPUDetectionOutputVector& outputs = *dynamic_cast<GPUDetectionOutputVector*>(o);
     //const bool printLog = this->f_printLog.getValue();
     if (ff==NULL)
     {
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaDistanceGridCollisionModel.cpp b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.cpp
similarity index 100%
rename from applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaDistanceGridCollisionModel.cpp
rename to applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.cpp
diff --git a/applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaDistanceGridCollisionModel.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h
similarity index 100%
rename from applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaDistanceGridCollisionModel.h
rename to applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.h
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu
new file mode 100644
index 00000000000..f72bf5d03c7
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.cu
@@ -0,0 +1,86 @@
+/******************************************************************************
+*                 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: contact@sofa-framework.org                             *
+******************************************************************************/
+#include <vector_types.h>
+#include <sofa/gpu/cuda/mycuda.h>
+#include "cuda.h"
+
+#if defined(__cplusplus) && CUDA_VERSION < 2000
+namespace sofa
+{
+namespace gpu
+{
+namespace cuda
+{
+#endif
+
+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");}
+}
+
+#if defined(__cplusplus) && CUDA_VERSION < 2000
+}
+}
+}
+#endif
diff --git a/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h
new file mode 100644
index 00000000000..c14d6e47f61
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/CudaDistanceGridContactMapper.h
@@ -0,0 +1,102 @@
+/******************************************************************************
+*                 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: contact@sofa-framework.org                             *
+******************************************************************************/
+#pragma once
+#include <sofa/component/collision/response/mapper/RigidContactMapper.h>
+#include <sofa/gpu/cuda/GPUDetectionOutputVector.h>
+#include <SofaDistanceGrid/CUDA/CudaDistanceGridCollisionModel.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);
+}
+}
+
+/// Mapper for CudaRigidDistanceGridCollisionModel
+template <class DataTypes>
+class sofa::component::collision::response::mapper::ContactMapper<sofa::gpu::cuda::CudaRigidDistanceGridCollisionModel,DataTypes> :
+    public sofa::component::collision::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(sofacuda::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();
+        }
+    }
+};
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
new file mode 100644
index 00000000000..ed81edd83ac
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/config.h.in
@@ -0,0 +1,37 @@
+/******************************************************************************
+*                 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: contact@sofa-framework.org                             *
+******************************************************************************/
+#pragma once
+
+#include <sofa/config.h>
+
+#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
new file mode 100644
index 00000000000..4d6783616c7
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.cpp
@@ -0,0 +1,62 @@
+/******************************************************************************
+*                 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: contact@sofa-framework.org                             *
+******************************************************************************/
+#include <SofaDistanceGrid/CUDA/init.h>
+#include <SofaDistanceGrid/initSofaDistanceGrid.h>
+#include <SofaCUDA/init.h>
+
+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_EXPORT_DYNAMIC_LIBRARY const char* getModuleComponentList();
+}
+
+void initExternalModule()
+{
+    init();
+}
+
+const char* getModuleName()
+{
+    return MODULE_NAME;
+}
+
+const char* getModuleVersion()
+{
+    return MODULE_VERSION;
+}
+
+void init()
+{
+    static bool first = true;
+    if (first)
+    {
+        sofadistancegrid::initSofaDistanceGrid();
+        sofa::gpu::cuda::init();
+        first = false;
+    }
+}
+
+} // 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
new file mode 100644
index 00000000000..df76bdbe8e5
--- /dev/null
+++ b/applications/plugins/SofaDistanceGrid/extensions/CUDA/src/SofaDistanceGrid/CUDA/init.h
@@ -0,0 +1,29 @@
+/******************************************************************************
+*                 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: contact@sofa-framework.org                             *
+******************************************************************************/
+#pragma once
+
+#include <SofaDistanceGrid/CUDA/config.h>
+
+namespace sofadistancegrid::cuda
+{
+SOFA_SOFADISTANCEGRID_CUDA_API void init();
+} // namespace sofadistancegrid::cuda
diff --git a/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/config.h.in b/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/config.h.in
index 5f0cc3c2d77..512419c34c3 100644
--- a/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/config.h.in
+++ b/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/config.h.in
@@ -28,10 +28,16 @@
 #cmakedefine01 SOFADISTANCEGRID_HAVE_SOFA_GL
 
 #ifdef SOFA_BUILD_SOFADISTANCEGRID
-#  define SOFA_TARGET SofaDistanceGrid
+#  define SOFA_TARGET @PROJECT_NAME@
 #  define SOFA_SOFADISTANCEGRID_API SOFA_EXPORT_DYNAMIC_LIBRARY
 #else
 #  define SOFA_SOFADISTANCEGRID_API SOFA_IMPORT_DYNAMIC_LIBRARY
 #endif
 
+namespace sofadistancegrid
+{
+    constexpr const char* MODULE_NAME = "@PROJECT_NAME@";
+    constexpr const char* MODULE_VERSION = "@PROJECT_VERSION@";
+}
+
 #endif
diff --git a/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.cpp b/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.cpp
index 80ff252f960..473b0ccdec0 100644
--- a/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.cpp
+++ b/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.cpp
@@ -19,15 +19,12 @@
 *                                                                             *
 * Contact information: contact@sofa-framework.org                             *
 ******************************************************************************/
-#include "initSofaDistanceGrid.h"
+#include <SofaDistanceGrid/initSofaDistanceGrid.h>
 #include "components/collision/DistanceGridCollisionModel.h"
 #include "components/forcefield/DistanceGridForceField.h"
 #include "RegisterModelToCollisionFactory.h"
 
-namespace sofa
-{
-
-namespace component
+namespace sofadistancegrid
 {
 extern "C" {
 SOFA_SOFADISTANCEGRID_API void initExternalModule();
@@ -39,6 +36,11 @@ SOFA_SOFADISTANCEGRID_API const char* getModuleComponentList();
 }
 
 void initExternalModule()
+{
+    initSofaDistanceGrid();
+}
+
+void initSofaDistanceGrid()
 {
     static bool first = true;
     if (first)
@@ -50,12 +52,12 @@ void initExternalModule()
 
 const char* getModuleName()
 {
-    return "SofaDistanceGrid";
+    return sofadistancegrid::MODULE_NAME;
 }
 
 const char* getModuleVersion()
 {
-    return "1.0";
+    return sofadistancegrid::MODULE_VERSION;
 }
 
 const char* getModuleLicense()
@@ -70,13 +72,4 @@ const char* getModuleDescription()
            "point in space. This is why it is often used to implement collisions.     ";
 }
 
-const char* getModuleComponentList()
-{
-    return "DistanceGridCollisionModel FFDDistanceGridDiscreteIntersection RayDistanceGridContact "
-           "RigidDistanceGridDiscreteIntersection DistanceGridForceField";
-}
-
-} /// component
-
-} /// sofa
-
+} /// namespace sofadistancegrid
diff --git a/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.h b/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.h
index 1c6c362ca61..7d3fcf19084 100644
--- a/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.h
+++ b/applications/plugins/SofaDistanceGrid/src/SofaDistanceGrid/initSofaDistanceGrid.h
@@ -19,22 +19,13 @@
 *                                                                             *
 * Contact information: contact@sofa-framework.org                             *
 ******************************************************************************/
-#ifndef SOFA_INIT_SOFADISTANCEGRID_H
-#define SOFA_INIT_SOFADISTANCEGRID_H
+#pragma once
 #include <SofaDistanceGrid/config.h>
 
-namespace sofa
+namespace sofadistancegrid
 {
 
-namespace component
-{
-
-
 void SOFA_SOFADISTANCEGRID_API initSofaDistanceGrid();
 
-} // namespace component
-
-} // namespace sofa
-
-#endif
+} // namespace sofadistancegrid