Move CUDA code related to SPH from SofaCUDA to a SofaSphFluid extension

CMakeLists.txt

@@ -11,6 +11,7 @@ sofa_find_package(Sofa.GL QUIET)
 set(PLUGIN_SPH_SRC_DIR src/SofaSphFluid)
 set(HEADER_FILES
     ${PLUGIN_SPH_SRC_DIR}/config.h.in
+    ${PLUGIN_SPH_SRC_DIR}/initSPHFluid.h
     ${PLUGIN_SPH_SRC_DIR}/ParticleSink.h
     ${PLUGIN_SPH_SRC_DIR}/ParticleSink.inl
     ${PLUGIN_SPH_SRC_DIR}/ParticleSource.h
@@ -74,3 +75,5 @@ sofa_create_package_with_targets(
     INCLUDE_INSTALL_DIR ${PROJECT_NAME}
     RELOCATABLE "plugins"
     )
+
+sofa_add_subdirectory(plugin extensions/CUDA SofaSphFluid.CUDA)

extensions/CUDA/CMakeLists.txt

@@ -0,0 +1,48 @@
+cmake_minimum_required(VERSION 3.12)
+project(SofaSphFluid.CUDA CUDA CXX)
+
+set(HEADER_FILES
+    src/SofaSphFluid/CUDA/init.h
+    src/SofaSphFluid/CUDA/config.h.in
+
+    src/SofaSphFluid/CUDA/CudaParticleSource.h
+    src/SofaSphFluid/CUDA/CudaParticleSource.inl
+    src/SofaSphFluid/CUDA/CudaSPHFluidForceField.h
+    src/SofaSphFluid/CUDA/CudaSPHFluidForceField.inl
+    src/SofaSphFluid/CUDA/CudaParticlesRepulsionForceField.h
+    src/SofaSphFluid/CUDA/CudaParticlesRepulsionForceField.inl
+    src/SofaSphFluid/CUDA/CudaSpatialGridContainer.h
+    src/SofaSphFluid/CUDA/CudaSpatialGridContainer.inl
+)
+
+set(SOURCE_FILES
+    src/SofaSphFluid/CUDA/init.cpp
+
+    src/SofaSphFluid/CUDA/CudaParticleSource.cpp
+    src/SofaSphFluid/CUDA/CudaSPHFluidForceField.cpp
+    src/SofaSphFluid/CUDA/CudaParticlesRepulsionForceField.cpp
+    src/SofaSphFluid/CUDA/CudaSpatialGridContainer.cpp
+)
+
+set(CUDA_SOURCES
+    src/SofaSphFluid/CUDA/CudaParticleSource.cu
+    src/SofaSphFluid/CUDA/CudaSPHFluidForceField.cu
+    src/SofaSphFluid/CUDA/CudaParticlesRepulsionForceField.cu
+    src/SofaSphFluid/CUDA/CudaSpatialGridContainer.cu
+)
+
+sofa_find_package(SofaSphFluid REQUIRED)
+sofa_find_package(SofaCUDA REQUIRED)
+
+add_library(${PROJECT_NAME} SHARED ${HEADER_FILES} ${SOURCE_FILES} ${CUDA_SOURCES})
+target_link_libraries(${PROJECT_NAME} SofaSphFluid)
+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"
+)

extensions/CUDA/SofaSphFluid.CUDAConfig.cmake.in

@@ -0,0 +1,9 @@
+# CMake package configuration file for the SofaSphFluid.CUDA library
+
+@PACKAGE_GUARD@
+@PACKAGE_INIT@
+
+find_package(SofaSphFluid QUIET REQUIRED)
+find_package(SofaCUDA QUIET REQUIRED)
+
+check_required_components(SofaSphFluid.CUDA)

extensions/CUDA/src/SofaSphFluid/CUDA/CudaParticleSource.cpp

@@ -0,0 +1,56 @@
+/******************************************************************************
+*                 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 <SofaCUDA/sofa/gpu/cuda/CudaTypes.h>
+#include "CudaParticleSource.inl"
+#include <sofa/core/ObjectFactory.h>
+#include <sofa/core/behavior/ProjectiveConstraintSet.h>
+#include <SofaSphFluid/ParticleSource.inl>
+
+namespace sofa
+{
+
+
+namespace core::behavior
+{
+template class ProjectiveConstraintSet<gpu::cuda::CudaVec3fTypes>;
+#ifdef SOFA_GPU_CUDA_DOUBLE
+template class ProjectiveConstraintSet<gpu::cuda::CudaVec3dTypes>;
+#endif
+} // namespace core::behavior
+
+
+namespace gpu::cuda
+{
+
+
+
+int ParticleSourceCudaClass = core::RegisterObject("Supports GPU-side computations using CUDA")
+        .add< component::misc::ParticleSource<CudaVec3fTypes> >()
+#ifdef SOFA_GPU_CUDA_DOUBLE
+        .add< component::misc::ParticleSource<CudaVec3dTypes> >()
+#endif // SOFA_GPU_CUDA_DOUBLE
+        ;
+
+} // namespace gpu::cuda
+
+
+} // namespace sofa

extensions/CUDA/src/SofaSphFluid/CUDA/CudaParticleSource.cu

@@ -0,0 +1,132 @@
+/******************************************************************************
+*                 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 <SofaCUDA/sofa/gpu/cuda/CudaCommon.h>
+#include <SofaCUDA/sofa/gpu/cuda/CudaMath.h>
+#include "cuda.h"
+
+#if defined(__cplusplus) && CUDA_VERSION < 2000
+namespace sofa
+{
+namespace gpu
+{
+namespace cuda
+{
+#endif
+
+extern "C"
+{
+
+    void ParticleSourceCuda3f_fillValues(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, float fx, float fy, float fz);
+    void ParticleSourceCuda3f_copyValuesWithOffset(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, const void* src, float fx, float fy, float fz);
+
+#ifdef SOFA_GPU_CUDA_DOUBLE
+
+    void ParticleSourceCuda3d_fillValues(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, double fx, double fy, double fz);
+    void ParticleSourceCuda3d_copyValuesWithOffset(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, const void* src, double fx, double fy, double fz);
+
+#endif // SOFA_GPU_CUDA_DOUBLE
+}
+
+//////////////////////
+// GPU-side methods //
+//////////////////////
+
+
+template<class real>
+__global__ void ParticleSourceCuda3t_fillValues_kernel(unsigned int totalsize, unsigned int subsetsize, real* dest, const unsigned int* indices, real fx, real fy, real fz)
+{
+    unsigned int index0 = blockIdx.x * BSIZE;
+    unsigned int index = index0+threadIdx.x;
+    if (index < subsetsize)
+    {
+        unsigned int dindex = indices[index];
+        if (dindex < totalsize)
+        {
+            unsigned int dindex3 = dindex * 3;
+            dest[dindex3+0] = fx;
+            dest[dindex3+1] = fy;
+            dest[dindex3+2] = fz;
+        }
+    }
+}
+
+template<class real>
+__global__ void ParticleSourceCuda3t_copyValuesWithOffset_kernel(unsigned int totalsize, unsigned int subsetsize, real* dest, const unsigned int* indices, const real* src, real fx, real fy, real fz)
+{
+    unsigned int index0 = blockIdx.x * BSIZE;
+    unsigned int index = index0+threadIdx.x;
+    if (index < subsetsize)
+    {
+        unsigned int dindex = indices[index];
+        if (dindex < totalsize)
+        {
+            unsigned int dindex3 = dindex * 3;
+            unsigned int index3 = index * 3;
+            dest[dindex3+0] = src[index3+0]+fx;
+            dest[dindex3+1] = src[index3+1]+fy;
+            dest[dindex3+2] = src[index3+2]+fz;
+        }
+    }
+}
+
+//////////////////////
+// CPU-side methods //
+//////////////////////
+
+
+void ParticleSourceCuda3f_fillValues(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, float fx, float fy, float fz)
+{
+    dim3 threads(BSIZE,1);
+    dim3 grid((subsetsize+BSIZE-1)/BSIZE,1);
+    {ParticleSourceCuda3t_fillValues_kernel<float><<< grid, threads >>>(totalsize, subsetsize, (float*)dest, (const unsigned int*)indices, fx, fy, fz); mycudaDebugError("ParticleSourceCuda3t_fillValues_kernel<float>");}
+}
+
+void ParticleSourceCuda3f_copyValuesWithOffset(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, const void* src, float fx, float fy, float fz)
+{
+    dim3 threads(BSIZE,1);
+    dim3 grid((subsetsize+BSIZE-1)/BSIZE,1);
+    {ParticleSourceCuda3t_copyValuesWithOffset_kernel<float><<< grid, threads >>>(totalsize, subsetsize, (float*)dest, (const unsigned int*)indices, (const float*)src, fx, fy, fz); mycudaDebugError("ParticleSourceCuda3t_copyValuesWithOffset_kernel<float>");}
+}
+
+#ifdef SOFA_GPU_CUDA_DOUBLE
+
+void ParticleSourceCuda3d_fillValues(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, double fx, double fy, double fz)
+{
+    dim3 threads(BSIZE,1);
+    dim3 grid((subsetsize+BSIZE-1)/BSIZE,1);
+    {ParticleSourceCuda3t_fillValues_kernel<double><<< grid, threads >>>(totalsize, subsetsize, (double*)dest, (const unsigned int*)indices, fx, fy, fz); mycudaDebugError("ParticleSourceCuda3t_fillValues_kernel<double>");}
+}
+
+void ParticleSourceCuda3d_copyValuesWithOffset(unsigned int totalsize, unsigned int subsetsize, void* dest, const void* indices, const void* src, double fx, double fy, double fz)
+{
+    dim3 threads(BSIZE,1);
+    dim3 grid((subsetsize+BSIZE-1)/BSIZE,1);
+    {ParticleSourceCuda3t_copyValuesWithOffset_kernel<double><<< grid, threads >>>(totalsize, subsetsize, (double*)dest, (const unsigned int*)indices, (const double*)src, fx, fy, fz); mycudaDebugError("ParticleSourceCuda3t_copyValuesWithOffset_kernel<double>");}
+}
+
+#endif // SOFA_GPU_CUDA_DOUBLE
+
+#if defined(__cplusplus) && CUDA_VERSION < 2000
+} // namespace cuda
+} // namespace gpu
+} // namespace sofa
+#endif

extensions/CUDA/src/SofaSphFluid/CUDA/CudaParticleSource.h

@@ -0,0 +1,27 @@
+/******************************************************************************
+*                 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                             *
+******************************************************************************/
+#ifndef SOFA_GPU_CUDA_CUDAPARTICLESOURCE_H
+#define SOFA_GPU_CUDA_CUDAPARTICLESOURCE_H
+
+#include <SofaCUDA/sofa/gpu/cuda/CudaTypes.h>
+
+#endif