A collection of available functions using CUDA in TAPs. More...

#include "../Core/TAPsListOfNamespaces.hpp"

Include dependency graph for TAPsCUDAListOfFunctions.hpp:

This graph shows which files directly or indirectly include this file:

Functions
BEGIN_NAMESPACE_TAPs__CUDA void	Host__Init (int argc, char **argv)
	Initialize CUDA.
void	Host__Init_withPLHM (int argc, char **argv)
	Initialize CUDA (with Page-Locked Host Memory)
int	DeviceCount ()
	Get number of devices.
bool	GetDevice (int &device)
	Get the device on which the active host thread executes the device code.
bool	SetDevice (int device)
bool	GetDeviceProperties (struct cudaDeviceProp *prop, int device)
	Get Device's Properties.
bool	CanDeviceMapHostMemory (int device)
	Can Map Host Memory (i.e. can utilize CUDA's Page-Locked Host Memory)
void	Global__ModelStrand_AdvSim (unsigned int cudaID, unsigned int numOfVertices, unsigned int numOfThreads, float currentTime, float timeStep, int numOfSubSteps, float ptMass, float Ks, float Kd, float Lrest, float host_vertex_data, float host_prev_vertex_data)
	CUDA Wrapper Function -- For ModelStrand AdvanceSimulation Function.
void	Global__ModelStrand_AdvSim_ADVSIM (unsigned int cudaID, unsigned int numOfVertices, unsigned int numOfThreads, float currentTime, float timeStep, int numOfSubSteps, float ptMass, float Ks, float Kd, float Lrest, float host_vertex_data, float host_prev_vertex_data, unsigned int host_sim_flags_data, float host_pos_constraint_data)
bool	InitailizeDataForSutureModel (unsigned int &cudaID, unsigned int numOfVertices, float vertexList, float prevVertexList)
	Initialize CUDA Data for a Suture Model (ModelStrand & ModelSuture)
bool	InitailizeDataForSutureModel_ADVSIM (unsigned int &cudaID, unsigned int numOfVertices, float vertexList, float prevVertexList, unsigned int simFlagsList, float posConstraintList)
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)
	CUDA Wrapper Function -- For ModelElasticRod AdvanceSimulation Function.
bool	InitailizeDataForPLHMElasticRodModel (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	ClearDataForPLHMElasticRodModel (unsigned int &cudaID)
void	Global__PLHMModelElasticRod_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))
void	GL__UnregisterBufferObject (GLuint &vbo_GL)
	Clear the CUDA/GL for elastic rod's drawing.
bool	GL__IsInitailizeDataForElasticRodModel_Drawing ()
	Is the texture for cross section's vertex data initialized?
bool	GL__InitailizeDataForElasticRodModel_Drawing (unsigned int numberOfCrossSectionVertices, float *crossSectionVertexData, GLuint vbo_GL)
	Initialize the CUDA/GL for elastic rod's drawing.
void	GL__ClearDataForElasticRodModel_Drawing ()
	Clear the CUDA/GL for elastic rod's drawing.
void	GL__GenCylinderForElasticRodModel (unsigned int cudaID, unsigned int numOfNodes, unsigned int numOfThreads, GLuint vbo_GL, unsigned int numOfCrossSectionVertices, float radius)
	Compute the generalized cylinder for the elastic rod.
void	GL__GenCylinderForElasticRodModel_PLHM (unsigned int cudaID, unsigned int numOfNodes, unsigned int numOfThreads, GLuint vbo_GL, unsigned int numOfCrossSectionVertices, float radius)
	Compute the generalized cylinder for the elastic rod.
void	Global__SimPropForMultiPartMeshModel_HalfEdge_AdvSim (int numOfVertices, int numOfThreads, float tCurrent, float tNext, float *data)
	CUDA Wrapper Function -- For SimPropForMultiPartMeshModel_HalfEdge AdvSim Function.
void	Global__HETriMeshOneModelMultiParts_AdvSim (unsigned int cudaID, unsigned int numOfVertices, unsigned int numOfThreads, float currentTime, float timeStep, int numOfSubSteps, float ptMass, float Ks, float Kd, float HKs, float HKd, float host_vertex_data, float host_prev_vertex_data, float host_home_vertex_data, int host_connection_list, unsigned int max_connection_size)
	CUDA Wrapper Function -- For HETriMeshOneModelMultiParts AdvSim Function.
bool	InitailizeDataForVertexList (unsigned int &cudaID, unsigned int numOfVertices, unsigned int max_connection_size, float vertexList, float prevVertexList, float homeVertexList, int vertexConnectionList)
	Initialize CUDA Data for a Mesh Model by Using Vertex List.

Detailed Description

A collection of available functions using CUDA in TAPs.

"CUDA/TAPsCUDAListOfFunctions.hpp"

Definition in file TAPsCUDAListOfFunctions.hpp.

Function Documentation

bool CanDeviceMapHostMemory ( int device )

Can Map Host Memory (i.e. can utilize CUDA's Page-Locked Host Memory)

Definition at line 86 of file TAPsCUDA_GlobalFns_Def.cu.

References GetDeviceProperties().

{
    //if ( !SetDevice( device ) )   return false;
    cudaDeviceProp prop;
    if ( !GetDeviceProperties( &prop, device ) )    return false;
    return prop.canMapHostMemory == 1;
}

Here is the call graph for this function:

void ClearDataForElasticRodModel ( unsigned int & cudaID )

Clear CUDA Data for an Elastic Rod Model (ModelElasticRod)

Parameters:

cudaID CUDA ID for the object

Definition at line 72 of file TAPsCUDA_VertexListModelElasticRod_Def.cu.

References DATA_GlobalPool, and DATA_Pool::DataForVertexList.

{
    if ( cudaID != 0 ) {
        delete DATA_GlobalPool.DataForVertexList[cudaID];
        cudaID = 0;
    }
}

void ClearDataForPLHMElasticRodModel ( unsigned int & cudaID )

Clear CUDA Data for an Elastic Rod Model (ModelElasticRod) Utilizing CUDA's Page-Locked Host Memory

Clear CUDA Data for an Elastic Rod Model (ModelElasticRod) Clear CUDA Data for an Elastic Rod Model (ModelElasticRod)

Parameters:

cudaID CUDA ID for the object

Definition at line 83 of file TAPsCUDA_VertexListPLHMModelElasticRod_Def.cu.

References DATA_GlobalPool, and DATA_Pool::DataForVertexListPLHM.

{
    if ( cudaID != 0 ) {
        delete DATA_GlobalPool.DataForVertexListPLHM[cudaID];
        cudaID = 0;
    }
}

int DeviceCount ( )

Get number of devices.

Definition at line 48 of file TAPsCUDA_GlobalFns_Def.cu.

{
    int count;
    cudaGetDeviceCount( &count );
    return count;
}

bool GetDevice ( int & device )

Get the device on which the active host thread executes the device code.

Definition at line 56 of file TAPsCUDA_GlobalFns_Def.cu.

References GetErrorString().

{
    cudaError_t err = cudaGetDevice( & device );
    if ( err == cudaSuccess ) return true;
    printf( "FAILED: GetDevice(%i) -- %s\n", device, GetErrorString( err ) );
    return false;
}

Here is the call graph for this function:

bool GetDeviceProperties	(	struct cudaDeviceProp *	prop,
		int	device
	)

Get Device's Properties.

Definition at line 77 of file TAPsCUDA_GlobalFns_Def.cu.

References GetErrorString().

{
    cudaError_t err = cudaGetDeviceProperties( prop, device );
    if ( err == cudaSuccess ) return true;
    printf( "FAILED: GetDeviceProperties( ..., %i) -- %s\n", device, GetErrorString( err ) );
    return false;
}

Here is the call graph for this function:

void GL__ClearDataForElasticRodModel_Drawing ( )

Clear the CUDA/GL for elastic rod's drawing.

Definition at line 279 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CUDA_GL_ElasticRod::Clear(), and Cuda_GL_ElasticRod.

{
    Cuda_GL_ElasticRod.Clear();
}

Here is the call graph for this function:

void GL__GenCylinderForElasticRodModel	(	unsigned int	cudaID,
		unsigned int	numOfNodes,
		unsigned int	numOfThreads,
		GLuint	vbo_GL,
		unsigned int	numOfCrossSectionVertices,
		float	radius
	)

Compute the generalized cylinder for the elastic rod.

Parameters:

cudaID	CUDA ID
numOfNodes	number of nodes
numOfThreads	number of threads per CUDA thread block
vbo_GL	OpenGL vertex buffer object
numOfCrossSectionVertices	number of cross section's vertices
radius	radius

Definition at line 285 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CUDA_GL_ElasticRod::BindCrossSectionVertices(), DATA_Vertex_List::BindOrientationList(), DATA_Vertex_List::BindVertexList(), Cuda_GL_ElasticRod, DATA_GlobalPool, DATA_Pool::DataForVertexList, CUDA_GL_ElasticRod::UnbindCrossSectionVertices(), DATA_Vertex_List::UnbindOrientationList(), and DATA_Vertex_List::UnbindVertexList().

{
    // 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]->BindOrientationList( numOfNodes );
    Cuda_GL_ElasticRod.BindCrossSectionVertices();
    // Map to GL
    float4 * pVertexData;
    CUDA_SAFE_CALL( cudaGLMapBufferObject( (void**)&pVertexData, vbo_GL ) );
    CUT_CHECK_ERROR( "Mapping an OpenGL buffer object FAILED!" );

    // Call the CUDA kernel
    GL__GenCylinderForElasticRodModel_CU<<< Dg, Db, Ns, S >>>( 
        numOfNodes, numOfThreads, pVertexData, numOfCrossSectionVertices, radius
    );
    CUT_CHECK_ERROR( "GenCylinderForElasticRodModel Kernel execution failed!" );

    // Unbind Textures
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindOrientationList();
    Cuda_GL_ElasticRod.UnbindCrossSectionVertices();
    // Unmap from GL
    CUDA_SAFE_CALL( cudaGLUnmapBufferObject( vbo_GL ) );
}

Here is the call graph for this function:

void GL__GenCylinderForElasticRodModel_PLHM	(	unsigned int	cudaID,
		unsigned int	numOfNodes,
		unsigned int	numOfThreads,
		GLuint	vbo_GL,
		unsigned int	numOfCrossSectionVertices,
		float	radius
	)

Compute the generalized cylinder for the elastic rod.

Parameters:

cudaID	CUDA ID
numOfNodes	number of nodes
numOfThreads	number of threads per CUDA thread block
vbo_GL	OpenGL vertex buffer object
numOfCrossSectionVertices	number of cross section's vertices
radius	radius

Definition at line 326 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CUDA_GL_ElasticRod::BindCrossSectionVertices(), Cuda_GL_ElasticRod, DATA_GlobalPool, DATA_Pool::DataForVertexListPLHM, DATA_Vertex_ListPLHM::GetDevOrientationList(), DATA_Vertex_ListPLHM::GetDevVertexList(), and CUDA_GL_ElasticRod::UnbindCrossSectionVertices().

{
    // 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 Texture
    Cuda_GL_ElasticRod.BindCrossSectionVertices();
    // Map to GL
    float4 * pVertexData;
    CUDA_SAFE_CALL( cudaGLMapBufferObject( (void**)&pVertexData, vbo_GL ) );
    CUT_CHECK_ERROR( "Mapping an OpenGL buffer object FAILED!" );

    // Call the CUDA kernel
    GL__GenCylinderForElasticRodModel_PLHM_CU<<< Dg, Db, Ns, S >>>( 
        numOfNodes, numOfThreads,
        DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevVertexList(),
        DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevOrientationList(),
        pVertexData, numOfCrossSectionVertices, radius
    );
    CUT_CHECK_ERROR( "GenCylinderForElasticRodModel Kernel execution failed!" );

    // Unbind Texture
    Cuda_GL_ElasticRod.UnbindCrossSectionVertices();
    // Unmap from GL
    CUDA_SAFE_CALL( cudaGLUnmapBufferObject( vbo_GL ) );
}

Here is the call graph for this function:

bool GL__InitailizeDataForElasticRodModel_Drawing	(	unsigned int	numberOfCrossSectionVertices,
		float *	crossSectionVertexData,
		GLuint	vbo_GL
	)

Initialize the CUDA/GL for elastic rod's drawing.

Parameters:

numberOfCrossSectionVertices	number of cross section's vertices
crossSectionVertexData	cross section's vertex data
vbo_GL	the OpenGL buffer object for vertex data

Definition at line 269 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CUDA_GL_ElasticRod::Create(), and Cuda_GL_ElasticRod.

Referenced by ModelElasticRod< T >::InitGL().

{
    return Cuda_GL_ElasticRod.Create( numberOfCrossSectionVertices, crossSectionVertexData, vbo_GL );
}

Here is the call graph for this function:

Here is the caller graph for this function:

bool GL__IsInitailizeDataForElasticRodModel_Drawing ( )

Is the texture for cross section's vertex data initialized?

Definition at line 265 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CUDA_GL_ElasticRod::IsInitialized().

{ return CUDA_GL_ElasticRod::IsInitialized(); }

Here is the call graph for this function:

void GL__UnregisterBufferObject ( GLuint & vbo_GL )

Clear the CUDA/GL for elastic rod's drawing.

Definition at line 44 of file TAPsCUDA_GL_ModelElasticRod.cu.

Referenced by ModelElasticRod< T >::ClearGL().

{
    CUDA_SAFE_CALL( cudaGLUnregisterBufferObject( vbo_GL ) );
}

Here is the caller graph for this function:

void Global__HETriMeshOneModelMultiParts_AdvSim	(	unsigned int	cudaID,
		unsigned int	numOfVertices,
		unsigned int	numOfThreads,
		float	currentTime,
		float	timeStep,
		int	numOfSubSteps,
		float	ptMass,
		float	Ks,
		float	Kd,
		float	HKs,
		float	HKd,
		float *	host_vertex_data,
		float *	host_prev_vertex_data,
		float *	host_home_vertex_data,
		int *	host_connection_list,
		unsigned int	max_connection_size
	)

CUDA Wrapper Function -- For HETriMeshOneModelMultiParts AdvSim Function.

For HETriMeshOneModelMultiParts AdvSim Function CUDA Wrapper Function -- For HETriMeshOneModelMultiParts AdvSim Function

Parameters:

cudaID	CUDA ID
numOfVertices	number of vertices
numOfThreads	number of threads per CUDA thread block
currentTime	current time
timeStep	time step
numOfSubSteps	number of sub-steps
ptMass	point mass
Ks	spring stiffness
Kd	spring damper
HKs	home spring stiffness
HKd	home spring damper
host_vertex_data	host's vertex data
host_prev_vertex_data	host's previous vertex data
host_home_vertex_data	host's home vertex data
host_connection_list	host's connection index list
max_connection_size	maximum connection size

Definition at line 1191 of file TAPsCUDA_VertexListMSM_Def.cu.

References DATA_Vertex_List::BindHomeVertexList(), DATA_Vertex_List::BindPrevVertexList(), DATA_Vertex_List::BindVertexConnectionList(), DATA_Vertex_List::BindVertexList(), DATA_GlobalPool, DATA_Pool::DataForVertexList, DATA_Vertex_List::GetVertexList(), DATA_Vertex_List::SwapVertexList(), DATA_Vertex_List::UnbindHomeVertexList(), DATA_Vertex_List::UnbindPrevVertexList(), DATA_Vertex_List::UnbindVertexConnectionList(), and DATA_Vertex_List::UnbindVertexList().

{
    //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 );

    // Allocate device memory for out data
    unsigned int size = sizeof(float) * numOfVertices * 4;
    float * out_data;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data, 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_vertex_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_vertex_data, size, cudaMemcpyHostToDevice ) );

    // Kernel invocation
    unsigned int numOfThreadBlocks = ( numOfVertices + 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( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindHomeVertexList( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindVertexConnectionList( numOfVertices, max_connection_size );

    int CHOICE = 2;

    // HOME VERTEX
    if ( CHOICE == 1 ) {
        // Call the CUDA kernel (Choice 1 -- Home Vertex Position Only)
        Global__HETriMeshOneModelMultiParts_AdvSim_CU<1><<< Dg, Db, Ns, S >>>( 
            numOfVertices, numOfThreads, 
            currentTime, timeStep, 
            ptMass, Ks, Kd, HKs, HKd, 
            max_connection_size, 
            out_data 
        );
        // 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();
    }

    // MASS SPRING SYSTEM
    else if ( CHOICE == 2 ) {
        float subTimeStep = timeStep / numOfSubSteps;
        for ( int i = 0; i < numOfSubSteps; ++i ) {
            //printf( "delta time: %g\n", timeStep/numOfSubSteps );

            // Call the CUDA kernel (Choice 2 -- Mass Spring System)
            Global__HETriMeshOneModelMultiParts_AdvSim_CU<2><<< Dg, Db, Ns, S >>>( 
                numOfVertices, numOfThreads, 
                currentTime, subTimeStep, 
                ptMass, Ks, Kd, HKs, HKd, 
                max_connection_size, 
                out_data 
            );
            // 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, size, cudaMemcpyDeviceToDevice ) );
            currentTime += subTimeStep;
        }
    }

    // Unbind Textures
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindHomeVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexConnectionList();

    // 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)
    if ( CHOICE > 0 ) {
        CUDA_SAFE_CALL( cudaMemcpy( host_prev_vertex_data, out_data, size, cudaMemcpyDeviceToHost ) );
    }

    //DATA_GlobalPool.DataForVertexList[cudaID]->PrintDebug( numOfVertices, host_vertex_data, host_prev_vertex_data, host_home_vertex_data );
    //DATA_GlobalPool.DataForVertexList[cudaID]->PrintDebug_ConnectionList( numOfVertices, max_connection_size, host_connection_list );

    // Free device memory
    CUDA_SAFE_CALL( cudaFree( out_data ) );
}

Here is the call 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
	)

CUDA Wrapper Function -- For ModelElasticRod AdvanceSimulation Function.

Parameters:

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.

{
    //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 call graph for this function:

void Global__ModelStrand_AdvSim	(	unsigned int	cudaID,
		unsigned int	numOfVertices,
		unsigned int	numOfThreads,
		float	currentTime,
		float	timeStep,
		int	numOfSubSteps,
		float	ptMass,
		float	Ks,
		float	Kd,
		float	Lrest,
		float *	host_vertex_data,
		float *	host_prev_vertex_data
	)

CUDA Wrapper Function -- For ModelStrand AdvanceSimulation Function.

CUDA Wrapper Function -- For ModelStrand AdvanceSimulation Function

Parameters:

cudaID	CUDA ID
numOfVertices	number of vertices
numOfThreads	number of threads per CUDA thread block
currentTime	current time
timeStep	time step
numOfSubSteps	number of sub-steps
ptMass	point mass
Ks	spring stiffness
Kd	spring damper
Lrest	spring rest length
host_vertex_data	host's vertex data
host_prev_vertex_data	host's previous vertex data

Definition at line 630 of file TAPsCUDA_VertexListMSM_Def.cu.

References DATA_Vertex_List::BindPrevVertexList(), DATA_Vertex_List::BindVertexList(), DATA_GlobalPool, DATA_Pool::DataForVertexList, DATA_Vertex_List::GetVertexList(), DATA_Vertex_List::SwapVertexList(), DATA_Vertex_List::UnbindPrevVertexList(), and DATA_Vertex_List::UnbindVertexList().

{
    //printf( "FILE:%s LINE:%d Global__ModelStrand_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 memory for out data
    unsigned int size = sizeof(float) * numOfVertices * 4;
    float * out_data;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data, 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_vertex_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_vertex_data, size, cudaMemcpyHostToDevice ) );

    // Kernel invocation
    unsigned int numOfThreadBlocks = ( numOfVertices + 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( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfVertices );
    //DATA_GlobalPool.DataForVertexList[cudaID]->BindHomeVertexList( numOfVertices );

    //*
    int CHOICE = 1;

    // MASS SPRING SYSTEM
    if ( CHOICE == 1 ) {
        float subTimeStep = timeStep / numOfSubSteps;
        for ( int i = 0; i < numOfSubSteps; ++i ) {
            // Call the CUDA kernel (Choice 1 -- Mass Spring System)
            Global__ModelStrand_AdvSim_CU<1><<< Dg, Db, Ns, S >>>( 
                numOfVertices, numOfThreads, 
                currentTime, subTimeStep, 
                ptMass, Ks, Kd, Lrest, 
                out_data 
            );
            // 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, size, cudaMemcpyDeviceToDevice ) );
            currentTime += subTimeStep;
        }
    }
    //*/

    // Unbind Textures
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
    //DATA_GlobalPool.DataForVertexList[cudaID]->UnbindHomeVertexList();

    // 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)
    if ( CHOICE > 0 ) {
        CUDA_SAFE_CALL( cudaMemcpy( host_prev_vertex_data, out_data, size, cudaMemcpyDeviceToHost ) );
    }

    //DATA_GlobalPool.DataForVertexList[cudaID]->PrintDebug( numOfVertices, host_vertex_data );

    // Free device memory
    CUDA_SAFE_CALL( cudaFree( out_data ) );
}

Here is the call graph for this function:

void Global__ModelStrand_AdvSim_ADVSIM	(	unsigned int	cudaID,
		unsigned int	numOfVertices,
		unsigned int	numOfThreads,
		float	currentTime,
		float	timeStep,
		int	numOfSubSteps,
		float	ptMass,
		float	Ks,
		float	Kd,
		float	Lrest,
		float *	host_vertex_data,
		float *	host_prev_vertex_data,
		unsigned int *	host_sim_flags_data,
		float *	host_pos_constraint_data
	)

CUDA Wrapper Function -- For ModelStrand AdvanceSimulation Function with TAPs_ADVANCED_SIMULATION

Parameters:

cudaID	CUDA ID
numOfVertices	number of vertices
numOfThreads	number of threads per CUDA thread block
currentTime	current time
timeStep	time step
numOfSubSteps	number of sub-steps
ptMass	point mass
Ks	spring stiffness
Kd	spring damper
Lrest	spring rest length
host_vertex_data	host's vertex data
host_prev_vertex_data	host's previous vertex data
host_sim_flags_data	host's simulation flags data
host_pos_constraint_data	host's position constraint data

Definition at line 726 of file TAPsCUDA_VertexListMSM_Def.cu.

{
    printf( "FILE:%s LINE:%d Global__ModelStrand_AdvSim_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 memory for out data
    unsigned int size = sizeof(float) * numOfVertices * 4;
    float * out_data;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data, size ) );

    // Allocate device memory for out data for changing suture's simulation flags list
    unsigned int size_SimFlagsList = sizeof(unsigned int) * numOfVertices;
    unsigned int * out_data_SimFlagsList;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data_SimFlagsList, size_SimFlagsList ) );

    // Allocate device memory for out data for changing suture's position constraint list
    float * out_data_PosConstraintList;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data_PosConstraintList, 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_vertex_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_vertex_data, size, cudaMemcpyHostToDevice ) );

    CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetSimFlagsList(), host_sim_flags_data, sizeof(unsigned int)*numOfVertices, cudaMemcpyHostToDevice ) );
    CUDA_SAFE_CALL( cudaMemcpy( DATA_GlobalPool.DataForVertexList[cudaID]->GetPosConstraintList(), host_pos_constraint_data, size, cudaMemcpyHostToDevice ) );

    // Kernel invocation
    unsigned int numOfThreadBlocks = ( numOfVertices + 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( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindSimFlagsList( numOfVertices );
    DATA_GlobalPool.DataForVertexList[cudaID]->BindPosConstraintList( numOfVertices );

    //*
    // Add 10 for TAPs_ADVANCED_SIMULATION
    int CHOICE = 100 + 1;

    // MASS SPRING SYSTEM
    if ( CHOICE == 101 ) {
        float subTimeStep = timeStep / numOfSubSteps;
        for ( int i = 0; i < numOfSubSteps; ++i ) {
            //printf( "delta time: %g\n", timeStep/numOfSubSteps );

            // Call the CUDA kernel (Choice 100+1 -- Mass Spring System)
            // with TAPs_ADVANCED_SIMULATION
            Global__ModelStrand_AdvSim_CU<101><<< Dg, Db, Ns, S >>>( 
                numOfVertices, numOfThreads, 
                currentTime, subTimeStep, 
                ptMass, Ks, Kd, Lrest, 
                out_data 
            );
            // 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, size, cudaMemcpyDeviceToDevice ) );
            currentTime += subTimeStep;
        }

        // Call the CUDA kernel to enforce constraints
        // with TAPs_ADVANCED_SIMULATION
        Global__ModelStrand_AdvSim_Enforce_Constraint_CU<101><<< Dg, Db, Ns, S >>>( 
            numOfVertices, numOfThreads, 
            currentTime, subTimeStep, 
            ptMass, Ks, Kd, Lrest, 
            out_data_SimFlagsList, 
            out_data_PosConstraintList 
        );
    }
    //*/

    // Unbind Textures
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindSimFlagsList();
    DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPosConstraintList();

    // 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)
    if ( CHOICE > 0 ) {
        CUDA_SAFE_CALL( cudaMemcpy( host_prev_vertex_data, out_data, size, cudaMemcpyDeviceToHost ) );
        CUDA_SAFE_CALL( cudaMemcpy( host_sim_flags_data, out_data_SimFlagsList, size_SimFlagsList, cudaMemcpyDeviceToHost ) );
        CUDA_SAFE_CALL( cudaMemcpy( host_pos_constraint_data, out_data_PosConstraintList, size, cudaMemcpyDeviceToHost ) );
    }

    /*
    // DEBUG
    {
        float * data = (float *)malloc( size );
        CUDA_SAFE_CALL( cudaMemcpy( data, out_data, size, cudaMemcpyDeviceToHost ) );
        for ( int i = 0, n = 0; i < 2; ++i, n+=4 ) {
            printf( "Out Vertex# %d: %g %g %g %g \n", i, data[n], data[n+1], data[n+2], data[n+3] );
        }
        for ( int i = numOfVertices-2, n = i*4; i < numOfVertices; ++i, n+=4 ) {
            printf( "Out Vertex# %d: %g %g %g %g \n", i, data[n], data[n+1], data[n+2], data[n+3] );
        }
        free( data );
    }
    */

    DATA_GlobalPool.DataForVertexList[cudaID]->PrintDebug( numOfVertices, host_vertex_data, NULL, NULL, host_sim_flags_data, host_pos_constraint_data );

    // Free device memory
    CUDA_SAFE_CALL( cudaFree( out_data ) );
    CUDA_SAFE_CALL( cudaFree( out_data_SimFlagsList ) );
    CUDA_SAFE_CALL( cudaFree( out_data_PosConstraintList ) );
}

Here is the call graph for this function:

void Global__PLHMModelElasticRod_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
	)

CUDA Wrapper Function -- For ModelElasticRod AdvanceSimulation Function Utilizing CUDA's Page-Locked Host Memory

Parameters:

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 310 of file TAPsCUDA_VertexListPLHMModelElasticRod_Def.cu.

{
    //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 );
    
    //printf( "090921 -- PLHM\n" );

    // Allocate device memories for out data
    unsigned int size = sizeof(float) * numOfNodes * 4;

    // // Bind Textures
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindVertexList( numOfNodes );
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindPrevVertexList( numOfNodes );
    // //DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindHomeVertexList( numOfNodes );
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindOrientationList( numOfNodes );
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindPrevOrientationList( numOfNodes );
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindForceList_1( numOfNodes );
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->BindForceList_2( numOfNodes );
    
    // 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;
    
    #if ER_TIMING_ADV_SIM_PLHM != 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 (w/ PLHM): numOfSubSteps %i\n", i );

            #if ER_TIMING_ADV_SIM_PLHM != 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_next_orientation
            );
            */

            /*
            // 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_next_centerline
            );
            */
            
            // Call the CUDA kernel (Choice 1) -- combine centerline and orientation
            Global__ModelElasticRod_AdvSimPLHM_CU<<< Dg, Db, Ns, S >>>( 
                numOfNodes, numOfThreads, currentTime, subTimeStep, 
                radius, length, length_ori, mass, k_mdr,
                Kconstraint_3rdDirAlignTangent, Kvdamping,
                Ps, Pb,
                DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevVertexList(),
                DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevOrientationList(),
                DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevForceList_1(),
                DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevForceList_2(),
                DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevPrevVertexList(),
                DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetDevPrevOrientationList()
            );
            
            // Swap pointers for Vertex List and Previous Vertex List
            // So that the Previous Vertex List is remain unchanged in the CUDA memory.
            DATA_GlobalPool.DataForVertexListPLHM[cudaID]->SwapVertexList();

            currentTime += subTimeStep;

            #if ER_TIMING_ADV_SIM_PLHM != 0
                cudaEventRecord( stop, 0 );
                cudaEventSynchronize( stop ); 
                cudaEventElapsedTime( &time, start, stop );
                if ( !isFirstRun ) {
                    totalTime += time;
                    ++counts;
                }
            #endif
        }
    }
    
    #if ER_TIMING_ADV_SIM_PLHM != 0
        if ( isFirstRun )   isFirstRun = false;
        cudaEventDestroy( start );
        cudaEventDestroy( stop );
        printf( "CUDA Kernal (w/ PLHM) TotalTime %f; Counts %i\n", totalTime, counts );
    #endif

    // Need this to wait for all threads to finish
    //cudaThreadSynchronize();
    cudaStreamSynchronize(S);
    
    //cudaSetDeviceFlags( cudaDeviceBlockingSync );
    
    // // Unbind Textures
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindVertexList();
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindPrevVertexList();
    // //DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindHomeVertexList();
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindOrientationList();
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindPrevOrientationList();
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindForceList_1();
    // DATA_GlobalPool.DataForVertexListPLHM[cudaID]->UnbindForceList_2();

    // Check if kernel execution generated and error
    CUT_CHECK_ERROR( "Kernel execution failed!" );

    // Reset the pointers to data
    host_pos_data           = DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetVertexList();
    host_prev_pos_data      = DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetPrevVertexList();
    host_ori_data           = DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetOrientationList();
    host_prev_ori_data      = DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetPrevOrientationList();
    //host_int_force_data       = DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetForceList_1();
    //host_ext_force_data       = DATA_GlobalPool.DataForVertexListPLHM[cudaID]->GetForceList_2();
}

Here is the call graph for this function:

void Global__SimPropForMultiPartMeshModel_HalfEdge_AdvSim	(	int	numOfVertices,
		int	numOfThreads,
		float	tCurrent,
		float	tNext,
		float *	data
	)

CUDA Wrapper Function -- For SimPropForMultiPartMeshModel_HalfEdge AdvSim Function.

For SimPropForMultiPartMeshModel_HalfEdge AdvSim Function CUDA Wrapper Function -- For SimPropForMultiPartMeshModel_HalfEdge AdvSim Function

Parameters:

numOfVertices	number of vertices
numOfThreads	number of threads per CUDA thread block
currentTime	current time
timeStep	time step
host_data	host data

Definition at line 939 of file TAPsCUDA_VertexListMSM_Def.cu.

{
    /*
    // CUDA device properties
    int dev_no;
    CUDA_SAFE_CALL( cudaGetDevice( &dev_no ) );
    cudaDeviceProp * prop;
    CUDA_SAFE_CALL( cudaGetDeviceProperties( prop, dev_no ) );
    //printf( "maxThreadsPerBlock: %d\n", prop->maxThreadsPerBlock );
    */

    // Kernel invocation
    unsigned int numOfThreadBlocks = ( numOfVertices + numOfThreads - 1 ) / numOfThreads;
    dim3 Dg( numOfThreadBlocks, 1, 1 );
    dim3 Db( numOfThreads, 1, 1 );
    size_t Ns = 0;
    cudaStream_t S = 0;

    printf( "Number of vertices: %d\n", numOfVertices );
    printf( "Number of threads: %d\n", numOfThreads );
    printf( "Number of grids: %d\n", numOfVertices/numOfThreads );

    // Allocate device memory
    unsigned int size = sizeof(float) * numOfVertices * 3;
    float * device_data;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&device_data, size ) );
    float * out_data;
    CUDA_SAFE_CALL( cudaMalloc( (void **)&out_data, size ) );

    // Copy host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( device_data, host_data, size, cudaMemcpyHostToDevice ) );

    // Call the CUDA kernel
    Global__SimPropForMultiPartMeshModel_HalfEdge_AdvSim_CU<1><<< Dg, Db, Ns, S >>>
    ( numOfVertices, numOfThreads, currentTime, timeStep, device_data, out_data );

    // Check if kernel execution generated and error
    CUT_CHECK_ERROR( "Kernel execution failed" );

    // Copy output data to host data
    CUDA_SAFE_CALL( cudaMemcpy( host_data, out_data, size, cudaMemcpyDeviceToHost ) );


    // Free device memory
    CUDA_SAFE_CALL( cudaFree( device_data ) );
    CUDA_SAFE_CALL( cudaFree( out_data ) );
}

BEGIN_NAMESPACE_TAPs__CUDA void Host__Init	(	int	argc,
		char **	argv
	)

Initialize CUDA.

Definition at line 17 of file TAPsCUDA_GlobalFns_Def.cu.

{
    printf( "START: CUT_DEVICE_INIT\n" );
    CUT_DEVICE_INIT(argc, argv);
    printf( "END:   CUT_DEVICE_INIT\n" );
}

void Host__Init_withPLHM	(	int	argc,
		char **	argv
	)

Initialize CUDA (with Page-Locked Host Memory)

Definition at line 25 of file TAPsCUDA_GlobalFns_Def.cu.

References CanDeviceMapHostMemory(), GetDevice(), and SetDevice().

{
    if ( !SetDevice( 0 ) ) {
        printf( "FAILED: setting the device 0! -- File: %s Line: %i\n", __FILE__, __LINE__ );
        exit(-501);
    }
    int device;
    if ( !GetDevice( device ) ) {
        printf( "FAILED: getting the current device! -- File: %s Line: %i\n", __FILE__, __LINE__ );
        exit(-501);
    }
    if ( !CanDeviceMapHostMemory( device ) ) {
        printf( "FAILED: the current device cannot map host memory! -- File: %s Line: %i\n", __FILE__, __LINE__ );
        exit(-502);
    }
    cudaError_t err = cudaSetDeviceFlags( cudaDeviceMapHost );
    if ( err != cudaSuccess ) {
        printf( "FAILED: cannot set device's map host memory flag! (%s) -- File: %s Line: %i\n", cudaGetErrorString( err ), __FILE__, __LINE__ );
        exit(-503);
    }
}

Here is the call 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)

Parameters:

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.

References AddToGlobal_Pool_Of_DATA_Vertex_List(), DATA_GlobalPool, DATA_Vertex_List::GetForceList_1(), DATA_Vertex_List::GetForceList_2(), DATA_Vertex_List::GetOrientationList(), DATA_Vertex_List::GetPrevOrientationList(), DATA_Vertex_List::GetPrevVertexList(), DATA_Vertex_List::GetVertexList(), and DATA_Pool::SizeOfGlobal_Pool.

{
    // 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 call graph for this function:

bool InitailizeDataForPLHMElasticRodModel	(	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)

Initialize CUDA Data for an Elastic Rod Model (ModelElasticRod) Utilizing CUDA's Page-Locked Host Memory

Initialize CUDA Data for an Elastic Rod Model (ModelElasticRod)

Parameters:

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_VertexListPLHMModelElasticRod_Def.cu.

References AddToGlobal_Pool_Of_DATA_Vertex_ListPLHM(), DATA_GlobalPool, DATA_Vertex_ListPLHM::GetForceList_1(), DATA_Vertex_ListPLHM::GetForceList_2(), DATA_Vertex_ListPLHM::GetOrientationList(), DATA_Vertex_ListPLHM::GetPrevOrientationList(), DATA_Vertex_ListPLHM::GetPrevVertexList(), DATA_Vertex_ListPLHM::GetVertexList(), and DATA_Pool::SizeOfGlobal_PoolPLHM.

{
    int size = numOfNodes * 4 * sizeof(float);

    int version;
    cudaRuntimeGetVersion( &version );
    printf( "CUDA runtime version: %i\n", version );

    /*
    // Size of memory for (xyzw) vertices
    float * h_p, d_p;
    cudaError_t err;
    // Pinned memory allocated on the CPU
    // Report and exit the program if error
    // Get the device pointers for the pinned CPU memory mapped into the GPU memory space
    // Report and exit the program if error
    err = cudaHostAlloc( (void **)& h_p, size, cudaHostAllocMapped );
    CheckError( err, "cudaHostAlloc for posList", true );
    err = cudaHostGetDevicePointer( (void **)& d_p, (void *) h_p, 0 );
    CheckError( err, "get the device pointer for posList", true );
    //*/

    // 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_PoolPLHM;

    // Allocate CUDA data for the object.
    DATA_Vertex_ListPLHM * newData = new DATA_Vertex_ListPLHM( 
        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_ListPLHM( newData );
    
    // Set the data pointers
    posList         = newData->GetVertexList();
    prevPosList     = newData->GetPrevVertexList();
    oriList         = newData->GetOrientationList();
    prevOriList     = newData->GetPrevOrientationList();
    intForceList    = newData->GetForceList_1();
    extForceList    = newData->GetForceList_2();
    
    return true;
}

Here is the call graph for this function:

bool InitailizeDataForSutureModel	(	unsigned int &	cudaID,
		unsigned int	numOfVertices,
		float *	vertexList,
		float *	prevVertexList
	)

Initialize CUDA Data for a Suture Model (ModelStrand & ModelSuture)

Parameters:

cudaID	CUDA ID for the object
numOfVertices	number of vertices
vertexList	list of xyzw vertices
prevVertexList	list of previous xyzw vertices

Definition at line 60 of file TAPsCUDA_VertexListMSM_Def.cu.

References AddToGlobal_Pool_Of_DATA_Vertex_List(), DATA_GlobalPool, DATA_Vertex_List::GetPrevVertexList(), DATA_Vertex_List::GetVertexList(), and DATA_Pool::SizeOfGlobal_Pool.

{
    // Here Suture 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( numOfVertices, false, 0, true );
    AddToGlobal_Pool_Of_DATA_Vertex_List( newData );

    // Size of memory for (xyzw) vertices
    int size = numOfVertices * 4 * sizeof(float);

    // Copy Vertex List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetVertexList(), vertexList, size, cudaMemcpyHostToDevice ) );
    // Copy Previous Vertex List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetPrevVertexList(), prevVertexList, 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

    return true;
}

Here is the call graph for this function:

bool InitailizeDataForSutureModel_ADVSIM	(	unsigned int &	cudaID,
		unsigned int	numOfVertices,
		float *	vertexList,
		float *	prevVertexList,
		unsigned int *	simFlagsList,
		float *	posConstraintList
	)

Initialize CUDA Data for a Suture Model (ModelStrand & ModelSuture) with TAPs_ADVANCED_SIMULATION

Parameters:

cudaID	CUDA ID for the object
numOfVertices	number of vertices
vertexList	list of xyzw vertices
prevVertexList	list of previous xyzw vertices
simFlagsList	list of simulation flags
posConstraintList	list of position constraints (xyz) plus force ratio (w)

Definition at line 99 of file TAPsCUDA_VertexListMSM_Def.cu.

References AddToGlobal_Pool_Of_DATA_Vertex_List(), DATA_GlobalPool, DATA_Vertex_List::GetPosConstraintList(), DATA_Vertex_List::GetPrevVertexList(), DATA_Vertex_List::GetSimFlagsList(), DATA_Vertex_List::GetVertexList(), and DATA_Pool::SizeOfGlobal_Pool.

{
    // Here Suture 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( numOfVertices, false, 0, true );
    AddToGlobal_Pool_Of_DATA_Vertex_List( newData );

    // Size of memory for (xyzw) vertices
    int size = numOfVertices * 4 * sizeof(float);

    // Copy Vertex List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetVertexList(), vertexList, size, cudaMemcpyHostToDevice ) );
    // Copy Previous Vertex List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetPrevVertexList(), prevVertexList, 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->GetVertexConnectionList(), vertexConnectionList, numOfVertices*max_connection_size*sizeof(int), cudaMemcpyHostToDevice ) );
    // Copy Simulation Flags List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetSimFlagsList(), simFlagsList, numOfVertices*sizeof(unsigned int), cudaMemcpyHostToDevice ) );
    // Copy Position Constraint List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetPosConstraintList(), posConstraintList, size, cudaMemcpyHostToDevice ) );

    return true;
}

Here is the call graph for this function:

bool InitailizeDataForVertexList	(	unsigned int &	cudaID,
		unsigned int	numOfVertices,
		unsigned int	max_connection_size,
		float *	vertexList,
		float *	prevVertexList,
		float *	homeVertexList,
		int *	vertexConnectionList
	)

Initialize CUDA Data for a Mesh Model by Using Vertex List.

Initialize CUDA Data for a Mesh Model by Using Vertex List

Parameters:

cudaID	CUDA ID for the object
numOfVertices	number of vertices
max_connection_size	maximum connection size
vertexList	list of xyz vertices
prevVertexList	list of previous xyzw vertices
homeVertexList	list of xyz home vertices
vertexConnectionList	vertex connection (1D array of 2D array -- vertex id and vertex connections)

Definition at line 21 of file TAPsCUDA_VertexListMSM_Def.cu.

References AddToGlobal_Pool_Of_DATA_Vertex_List(), DATA_GlobalPool, DATA_Vertex_List::GetHomeVertexList(), DATA_Vertex_List::GetPrevVertexList(), DATA_Vertex_List::GetVertexConnectionList(), DATA_Vertex_List::GetVertexList(), and DATA_Pool::SizeOfGlobal_Pool.

{
    //printf( "InitailizeDataForVertexList\n" );
    //fflush( stdout );

    // 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( numOfVertices, true, max_connection_size, false );
    AddToGlobal_Pool_Of_DATA_Vertex_List( newData );

    // Size of memory for (xyzw) vertices
    int size = numOfVertices * 4 * sizeof(float);

    // Copy Vertex List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetVertexList(), vertexList, size, cudaMemcpyHostToDevice ) );
    // Copy Previous Vertex List from host data to device data
    CUDA_SAFE_CALL( cudaMemcpy( newData->GetPrevVertexList(), prevVertexList, 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->GetVertexConnectionList(), vertexConnectionList, numOfVertices*max_connection_size*sizeof(int), cudaMemcpyHostToDevice ) );

    return true;
}

Here is the call graph for this function:

bool SetDevice ( int device )

Set the device on which the active host thread executes the device code If the host thread has already initialized the CUDA runtime by calling non-device management runtime functions, this call results in cudaErrorSetOnActiveProcess.

Definition at line 68 of file TAPsCUDA_GlobalFns_Def.cu.

References GetErrorString().

{
    cudaError_t err = cudaSetDevice( device );
    if ( err == cudaSuccess ) return true;
    printf( "FAILED: SetDevice(%i) -- %s\n", device, GetErrorString( err ) );
    return false;
}

Here is the call graph for this function: