 |
TAPs 0.7.7.3
|
|
TAPs 0.7.7.3
|
A CUDA Vertex List for the Elastic Rod model. More...
Include dependency graph for TAPsCUDA_VertexListModelElasticRod.cu: This graph shows which files directly or indirectly include this file:Go to the source code of this file.
Defines | |
| #define | TAPs_CUDA_VERTEX_LIST_MODEL_ELASTIC_ROD_HPP |
| #define | ER_TIMING_ADV_SIM 0 |
Functions | |
| bool | InitailizeDataForElasticRodModel (unsigned int &cudaID, unsigned int numOfNodes, float *posList, float *prevPosList, float *oriList, float *prevOriList, float *intForceList, float *extForceList) |
| Initialize CUDA Data for an Elastic Rod Model (ModelElasticRod) | |
| void | ClearDataForElasticRodModel (unsigned int &cudaID) |
| Clear CUDA Data for an Elastic Rod Model (ModelElasticRod) | |
| void | Global__ModelElasticRod_AdvSim (unsigned int cudaID, unsigned int numOfNodes, unsigned int numOfThreads, float currentTime, float timeStep, int numOfSubSteps, float radius, float length, float length_ori, float mass, float material_density, float Kconstraint_3rdDirAlignTangent, float Kvdamping, float Ps, float Pb_x, float Pb_y, float Pb_z, float *host_pos_data, float *host_prev_pos_data, float *host_ori_data, float *host_prev_ori_data, float *host_int_force_data, float *host_ext_force_data) |
A CUDA Vertex List for the Elastic Rod model.
"CUDA/TAPsCUDA_VertexListModelElasticRod.cu"
Definition in file TAPsCUDA_VertexListModelElasticRod.cu.
| #define ER_TIMING_ADV_SIM 0 |
Definition at line 19 of file TAPsCUDA_VertexListModelElasticRod.cu.
| #define TAPs_CUDA_VERTEX_LIST_MODEL_ELASTIC_ROD_HPP |
Definition at line 13 of file TAPsCUDA_VertexListModelElasticRod.cu.
| void ClearDataForElasticRodModel | ( | unsigned int & | cudaID | ) |
Clear CUDA Data for an Elastic Rod Model (ModelElasticRod)
Definition at line 72 of file TAPsCUDA_VertexListModelElasticRod_Def.cu.
Referenced by ElasticRod_CompByCUDA< T >::Clean().
{
if ( cudaID != 0 ) {
delete DATA_GlobalPool.DataForVertexList[cudaID];
cudaID = 0;
}
}
Here is the caller graph for this function:| void Global__ModelElasticRod_AdvSim | ( | unsigned int | cudaID, |
| unsigned int | numOfNodes, | ||
| unsigned int | numOfThreads, | ||
| float | currentTime, | ||
| float | timeStep, | ||
| int | numOfSubSteps, | ||
| float | radius, | ||
| float | length, | ||
| float | length_ori, | ||
| float | mass, | ||
| float | material_density, | ||
| float | Kconstraint_3rdDirAlignTangent, | ||
| float | Kvdamping, | ||
| float | Ps, | ||
| float | Pb_x, | ||
| float | Pb_y, | ||
| float | Pb_z, | ||
| float * | host_pos_data, | ||
| float * | host_prev_pos_data, | ||
| float * | host_ori_data, | ||
| float * | host_prev_ori_data, | ||
| float * | host_int_force_data, | ||
| float * | host_ext_force_data | ||
| ) |
| cudaID | CUDA ID |
| numOfNodes | number of nodes |
| numOfThreads | number of threads per CUDA thread block |
| currentTime | current time |
| timeStep | time step |
| numOfSubSteps | number of sub-steps |
| radius | radius |
| length | rest length |
| length_ori | orientation's rest length |
| mass | point mass |
| material_density | material density |
| Kconstraint_3rdDirAlignTangent | Kconstraint for aligning centerline's tangent with orientation's 3rd direction |
| Kvdamping | centerline's velocity damper |
| Ps | potential stretch constant |
| Pb_x | potential bend constant -- x |
| Pb_y | potential bend constant -- y |
| Pb_z | potential bend constant -- z |
| host_pos_data | host's position data |
| host_prev_pos_data | host's previous position data |
| host_ori_data | host's orientation data |
| host_prev_ori_data | host's previous orientation data |
| host_int_force_data | host's (internal) force data -- xyzw |
| host_ext_force_data | host's (external) force data -- xyzw |
Definition at line 478 of file TAPsCUDA_VertexListModelElasticRod_Def.cu.
Referenced by ElasticRod_CompByCUDA< T >::AdvSim().
{
//printf( "FILE:%s LINE:%d Global__ModelElasticRod_AdvSim\n", __FILE__, __LINE__ );
//printf( "currentTime: %g, timeStep: %g,\n", currentTime, timeStep );
//printf( "ptMass: %g, Ks: %g, Kd: %g,\n", ptMass, Ks, Kd );
//printf( "Number of vertices: %d\n", numOfVertices );
//printf( "Number of threads: %d\n", numOfThreads );
//printf( "Number of grids: %d\n", ( numOfVertices + numOfThreads - 1 ) / numOfThreads );
// Allocate device memories for out data
unsigned int size = sizeof(float) * numOfNodes * 4;
float * out_data_1; // for centerlines' position
CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data_1, size ) );
float * out_data_2; // for orientations
CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data_2, size ) );
// Copy Vertex List (and Previous Vertex List) from host data to device data,
// since vertices are dynamically changed.
// While home vertices and connection list are unchanged.
CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetVertexList(), host_pos_data, size, cudaMemcpyHostToDevice ) );
// The Previous Vertex List can be saved in CUDA memory.
// Hence, the Previous Vertex List does not have to be updated.
//CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetPrevVertexList(), host_prev_pos_data, size, cudaMemcpyHostToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetOrientationList(), host_ori_data, size, cudaMemcpyHostToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetForceList_1(), host_int_force_data, size, cudaMemcpyHostToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetForceList_2(), host_ext_force_data, size, cudaMemcpyHostToDevice ) );
// Kernel invocation
unsigned int numOfThreadBlocks = ( numOfNodes + numOfThreads - 1 ) / numOfThreads;
dim3 Dg( numOfThreadBlocks, 1, 1 );
dim3 Db( numOfThreads, 1, 1 );
size_t Ns = 0;
cudaStream_t S = 0;
// Bind Textures
DATA_GlobalPool.DataForVertexList[cudaID]->BindVertexList( numOfNodes );
DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfNodes );
//DATA_GlobalPool.DataForVertexList[cudaID]->BindHomeVertexList( numOfNodes );
DATA_GlobalPool.DataForVertexList[cudaID]->BindOrientationList( numOfNodes );
DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevOrientationList( numOfNodes );
DATA_GlobalPool.DataForVertexList[cudaID]->BindForceList_1( numOfNodes );
DATA_GlobalPool.DataForVertexList[cudaID]->BindForceList_2( numOfNodes );
#if ER_TIMING_ADV_SIM != 0
static bool isFirstRun = true;
cudaEvent_t start, stop;
float time;
static unsigned int counts = 0;
static float totalTime = 0.0f;
cudaEventCreate( &start );
cudaEventCreate( &stop );
#endif
{
//float k_mdr = material_density * radius * radius * K_PI;
float k_mdr = material_density * radius * radius * 3.1415926535897932384626433832795f;
//printf( "%f = %f * %f * %f * %f\n", k_mdr, material_density, radius, radius, K_PI );
float subTimeStep = timeStep / numOfSubSteps;
float3 Pb = make_float3( Pb_x, Pb_y, Pb_z );
for ( int i = 1; i <= numOfSubSteps; ++i ) {
//printf( "CUDA: numOfSubSteps %i\n", i );
#if ER_TIMING_ADV_SIM != 0
cudaEventRecord( start, 0 );
#endif
/*
// Call the CUDA kernel
Global__ModelElasticRod_AdvSim_CU_Orientation<<< Dg, Db, Ns, S >>>(
numOfNodes, numOfThreads, currentTime, subTimeStep,
radius, length, length_ori, mass, k_mdr,
Kconstraint_3rdDirAlignTangent, Kvdamping,
Ps, Pb,
out_data_2
);
*/
/*
// Call the CUDA kernel
Global__ModelElasticRod_AdvSim_CU_Centerline<<< Dg, Db, Ns, S >>>(
numOfNodes, numOfThreads, currentTime, subTimeStep,
radius, length, length_ori, mass, k_mdr,
Kconstraint_3rdDirAlignTangent, Kvdamping,
Ps, Pb,
out_data_1
);
*/
// Call the CUDA kernel (Choice 1) -- combine centerline and orientation
Global__ModelElasticRod_AdvSim_CU<<< Dg, Db, Ns, S >>>(
numOfNodes, numOfThreads, currentTime, subTimeStep,
radius, length, length_ori, mass, k_mdr,
Kconstraint_3rdDirAlignTangent, Kvdamping,
Ps, Pb,
out_data_1, out_data_2
);
//cudaThreadSynchronize();
// Swap pointers for Vertex List and Previous Vertex List
// So that the Previous Vertex List is remain unchanged in the CUDA memory.
DATA_GlobalPool.DataForVertexList[cudaID]->SwapVertexList();
// Copy data to the device's memory for (current) Vertex List
CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetVertexList(), out_data_1, size, cudaMemcpyDeviceToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetOrientationList(), out_data_2, size, cudaMemcpyDeviceToDevice ) );
currentTime += subTimeStep;
#if ER_TIMING_ADV_SIM != 0
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
if ( !isFirstRun ) {
totalTime += time;
++counts;
}
#endif
}
}
#if ER_TIMING_ADV_SIM != 0
if ( isFirstRun ) isFirstRun = false;
cudaEventDestroy( start );
cudaEventDestroy( stop );
printf( "CUDA Kernal TotalTime %f; Counts %i\n", totalTime, counts );
#endif
// Unbind Textures
DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
//DATA_GlobalPool.DataForVertexList[cudaID]->UnbindHomeVertexList();
DATA_GlobalPool.DataForVertexList[cudaID]->UnbindOrientationList();
DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevOrientationList();
DATA_GlobalPool.DataForVertexList[cudaID]->UnbindForceList_1();
DATA_GlobalPool.DataForVertexList[cudaID]->UnbindForceList_2();
// Check if kernel execution generated and error
CUT_CHECK_ERROR( "Kernel execution failed!" );
// Copy output data to host data Previous Vertex List
// The host will have to swap its pointers (for current and previous vertex list)
CUDA_SAFE_CALL( cudaMemcpy( host_prev_pos_data, out_data_1, size, cudaMemcpyDeviceToHost ) );
CUDA_SAFE_CALL( cudaMemcpy( host_prev_ori_data, out_data_2, size, cudaMemcpyDeviceToHost ) );
//DATA_GlobalPool.DataForVertexList[cudaID]->PrintDebug( numOfVertices, host_vertex_data );
// Free device memories
CUDA_SAFE_CALL( cudaFree( out_data_1 ) );
CUDA_SAFE_CALL( cudaFree( out_data_2 ) );
}
Here is the caller graph for this function:| bool InitailizeDataForElasticRodModel | ( | unsigned int & | cudaID, |
| unsigned int | numOfNodes, | ||
| float * | posList, | ||
| float * | prevPosList, | ||
| float * | oriList, | ||
| float * | prevOriList, | ||
| float * | intForceList, | ||
| float * | extForceList | ||
| ) |
Initialize CUDA Data for an Elastic Rod Model (ModelElasticRod)
| cudaID | CUDA ID for the object |
| numOfNodes | number of nodes |
| posList | list of centerlines' position -- xyzw |
| prevPosList | list of centerlines' previous position -- xyzw |
| oriList | list of orientations -- xyzw |
| prevOriList | list of previous orientations -- xyzw |
| intForceList | list of (internal) forces -- xyzw |
| extForceList | list of (external) forces -- xyzw |
Definition at line 17 of file TAPsCUDA_VertexListModelElasticRod_Def.cu.
Referenced by ElasticRod_CompByCUDA< T >::Init().
{
// Here the Elastic Rod Model uses Vertex List Data
// However, it has implicit connection (each vertex connect to the previous and next vertices).
// So the connection list is set to zero.
// Home vertex is unused for now. Plan is to use home vertex for sticking suture to a surface.
// Assign CUDA ID to the object
// The object has to use cudaID to communicate with the CUDA.
cudaID = DATA_GlobalPool.SizeOfGlobal_Pool;
// Allocate CUDA data for the object.
DATA_Vertex_List * newData = new DATA_Vertex_List(
numOfNodes, // total number of nodes
false, // include home positions into the data
0, // maximum vertex connection
false, // include simulation flags constraint into the data
true, // include orientations
true, // include previous orientations
true, // include forces (set 1) as for internal force
true // include forces (set 1) as for external force
);
AddToGlobal_Pool_Of_DATA_Vertex_List( newData );
// Size of memory for (xyzw) vertices
int size = numOfNodes * 4 * sizeof(float);
// Copy Vertex List from host data to device data
CUDA_SAFE_CALL( cudaMemcpy( newData->GetVertexList(), posList, size, cudaMemcpyHostToDevice ) );
// Copy Previous Vertex List from host data to device data
CUDA_SAFE_CALL( cudaMemcpy( newData->GetPrevVertexList(), prevPosList, size, cudaMemcpyHostToDevice ) );
// Copy Home Vertex List from host data to device data
//CUDA_SAFE_CALL( cudaMemcpy( newData->GetHomeVertexList(), homeVertexList, size, cudaMemcpyHostToDevice ) );
// Copy Vertex Connection List from host data to device data
CUDA_SAFE_CALL( cudaMemcpy( newData->GetOrientationList(), oriList, size, cudaMemcpyHostToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( newData->GetPrevOrientationList(), prevOriList, size, cudaMemcpyHostToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( newData->GetForceList_1(), intForceList, size, cudaMemcpyHostToDevice ) );
CUDA_SAFE_CALL( cudaMemcpy( newData->GetForceList_2(), extForceList, size, cudaMemcpyHostToDevice ) );
return true;
}
Here is the caller graph for this function:
1.7.4 