Skip to content

Commit

Permalink
Move DistanceGrid CUDA
Browse files Browse the repository at this point in the history
  • Loading branch information
@alxbilger
alxbilger committed Aug 2, 2024
1 parent ff3931c commit 17a804c
Show file tree
Hide file tree
Showing 27 changed files with 622 additions and 333 deletions.
3 changes: 2 additions & 1 deletion applications/plugins/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,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
Expand Down Expand Up @@ -69,3 +68,5 @@ if(Sofa.GL_FOUND)
else()
message("Sofa.GL not found; disabling SofaSimpleGUI and VolumetricRendering plugins")
endif()

sofa_add_subdirectory(plugin SofaDistanceGrid SofaDistanceGrid)
38 changes: 3 additions & 35 deletions applications/plugins/SofaCUDA/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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

Expand Down Expand Up @@ -335,33 +338,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
Expand Down Expand Up @@ -492,14 +468,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()
Expand Down
32 changes: 32 additions & 0 deletions applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cpp
View file
Original file line number Diff line number Diff line change
@@ -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>;

}
24 changes: 0 additions & 24 deletions applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}

Expand All @@ -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)
Expand All @@ -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
Expand Down
81 changes: 7 additions & 74 deletions applications/plugins/SofaCUDA/sofa/gpu/cuda/CudaContactMapper.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}

Expand All @@ -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>
Expand Down Expand Up @@ -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
Loading

0 comments on commit 17a804c

Please sign in to comment.