TAPsCUDA.cu File Reference

#include "../Support/TAPsListOfAdvSimConsts.hpp"
#include <cstdlib>
#include <cstdio>
#include <cutil.h>
#include "TAPsCUDA_Def.cu"

Include dependency graph for TAPsCUDA.cu:

Go to the source code of this file.

Classes
class	TAPs_CUDA_CLASS_DATA_Pool
class	TAPs_CUDA_CLASS_DATA_Vertex_List
Defines
#define	TAPs_CUDA_DATA_VERTEX_LIST_TEXTURE_FROM_LIN_MEM
Functions
bool	AddToGlobal_Pool_Of_TAPs_CUDA_CLASS_DATA_Vertex_List (TAPs_CUDA_CLASS_DATA_Vertex_List *dataObj)
	Add Vertex List Data Object to the CUDA Global Data Pool.
template<typename T>
__device__ void	TAPsCUDA__device__F1 ()
template<typename T>
__global__ void	TAPsCUDA__global__F1 ()
void	TAPsCUDA__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)
	For SimPropForMultiPartMeshModel_HalfEdge AdvSim Function.
void	TAPsCUDA__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)
	For Suture Model (ModelStrand & ModelSuture).
void	TAPsCUDA__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)
void	TAPsCUDA__global__SimPropForMultiPartMeshModel_HalfEdge_AdvSim (int numOfVertices, int numOfThreads, float tCurrent, float tNext, float *host_data)
	For Suture Model (ModelStrand & ModelSuture).
template<typename T>
__host__ __device__ void	TAPsCUDA__host__device__F1 ()
void	TAPsCUDA__host__Init (int argc, char **argv)
	Initialize CUDA Device.
bool	TAPsCUDA__InitailizeDataForSutureModel (unsigned int &cudaID, unsigned int numOfVertices, float *vertexList, float *prevVertexList)
	Initialize CUDA Data for a Suture Model (ModelStrand & ModelSuture).
bool	TAPsCUDA__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).
bool	TAPsCUDA__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.
Variables
texture< float4, 1, cudaReadModeElementType >	CudaTexHomeVertexList
	texture for list of home vertices
texture< float4, 1, cudaReadModeElementType >	CudaTexPosConstraintList
	texture for list of position constraints (xyz) plus force ratio (w)
texture< float4, 1, cudaReadModeElementType >	CudaTexPrevVertexList
	texture for list of previous vertices
texture< uint1, 1, cudaReadModeElementType >	CudaTexSimFlagsList
	texture for list of simulation flags
texture< int1, 1, cudaReadModeElementType >	CudaTexVertexConnectionList
	texture for list of vertex connections (x:=vertex#, y:=rest length, z:=stiffness, w:=damper)
texture< float4, 1, cudaReadModeElementType >	CudaTexVertexList
	For HETriMeshOneModelMultiParts AdvSim Function.
TAPs_CUDA_CLASS_DATA_Pool	TAPs_CUDA_DATA_GlobalPool
	CUDA Global Data Pool.

---

Define Documentation

#define TAPs_CUDA_DATA_VERTEX_LIST_TEXTURE_FROM_LIN_MEM

Definition at line 243 of file TAPsCUDA.cu.

---

Function Documentation

bool AddToGlobal_Pool_Of_TAPs_CUDA_CLASS_DATA_Vertex_List

(

TAPs_CUDA_CLASS_DATA_Vertex_List *

dataObj

)

Add Vertex List Data Object to the CUDA Global Data Pool.

Definition at line 150 of file TAPsCUDA_Def.cu.

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

    unsigned int current_size = TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool;
    if ( current_size == 0 ) {
        // SHOULD NEVER REACHED THIS!!!
        //TAPs_CUDA_DATA_GlobalPool.DataForVertexList = (TAPs_CUDA_CLASS_DATA_Vertex_List *)malloc( ++TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool );
        TAPs_CUDA_DATA_GlobalPool.DataForVertexList = new TAPs_CUDA_CLASS_DATA_Vertex_List * [ ++TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool ];
    }
    else {
        TAPs_CUDA_CLASS_DATA_Vertex_List ** old_list = TAPs_CUDA_DATA_GlobalPool.DataForVertexList;
        //TAPs_CUDA_DATA_GlobalPool.DataForVertexList = (TAPs_CUDA_CLASS_DATA_Vertex_List *)malloc( ++TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool );
        TAPs_CUDA_DATA_GlobalPool.DataForVertexList = new TAPs_CUDA_CLASS_DATA_Vertex_List * [ ++TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool ];
        for ( int i = 0; i < current_size; ++i ) {
            TAPs_CUDA_DATA_GlobalPool.DataForVertexList[ i ] = old_list[ i ];
        }
        //free ( old_list );
        delete [] old_list;
    }
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[ current_size ] = dataObj;
    return true;
}

template<typename T>

__device__ void TAPsCUDA__device__F1

(

)

[inline]

__device__F1 __device__ means a function that is executed on the device and callable from the device only.

Definition at line 116 of file TAPsCUDA.cu.

{}

template<typename T>

__global__ void TAPsCUDA__global__F1

(

)

[inline]

__global__F1 __global__ means a function as being a kernel, which is executed on the device and callable from the host only.

Definition at line 110 of file TAPsCUDA.cu.

{}

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

For SimPropForMultiPartMeshModel_HalfEdge 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 1217 of file TAPsCUDA_Def.cu.

{
    //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( TAPs_CUDA_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( TAPs_CUDA_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
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindVertexList( numOfVertices );
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfVertices );
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindHomeVertexList( numOfVertices );
    TAPs_CUDA_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)
        TAPsCUDA__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.
        TAPs_CUDA_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)
            TAPsCUDA__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.
            TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->SwapVertexList();

            // Copy data to the device's memory for (current) Vertex List
            CUDA_SAFE_CALL( cudaMemcpy( TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->GetVertexList(), out_data, size, cudaMemcpyDeviceToDevice ) );
            currentTime += subTimeStep;
        }
    }

    // Unbind Textures
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindHomeVertexList();
    TAPs_CUDA_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 ) );
    }

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

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

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

For Suture Model (ModelStrand & ModelSuture).

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 652 of file TAPsCUDA_Def.cu.

{
    //printf( "FILE:%s LINE:%d TAPsCUDA__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( TAPs_CUDA_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( TAPs_CUDA_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
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindVertexList( numOfVertices );
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfVertices );
    //TAPs_CUDA_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)
            TAPsCUDA__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.
            TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->SwapVertexList();

            // Copy data to the device's memory for (current) Vertex List
            CUDA_SAFE_CALL( cudaMemcpy( TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->GetVertexList(), out_data, size, cudaMemcpyDeviceToDevice ) );
            currentTime += subTimeStep;
        }
    }
    //*/

    // Unbind Textures
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
    //TAPs_CUDA_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 ) );
    }

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

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

void TAPsCUDA__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

CUDA Wrapper Function -- For ModelStrand AdvSim 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 748 of file TAPsCUDA_Def.cu.

{
    printf( "FILE:%s LINE:%d TAPsCUDA__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( TAPs_CUDA_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( TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->GetPrevVertexList(), host_prev_vertex_data, size, cudaMemcpyHostToDevice ) );

    CUDA_SAFE_CALL( cudaMemcpy( TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->GetSimFlagsList(), host_sim_flags_data, sizeof(unsigned int)*numOfVertices, cudaMemcpyHostToDevice ) );
    CUDA_SAFE_CALL( cudaMemcpy( TAPs_CUDA_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
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindVertexList( numOfVertices );
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindPrevVertexList( numOfVertices );
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->BindSimFlagsList( numOfVertices );
    TAPs_CUDA_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
            TAPsCUDA__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.
            TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->SwapVertexList();

            // Copy data to the device's memory for (current) Vertex List
            CUDA_SAFE_CALL( cudaMemcpy( TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->GetVertexList(), out_data, size, cudaMemcpyDeviceToDevice ) );
            currentTime += subTimeStep;
        }

        // Call the CUDA kernel to enforce constraints
        // with TAPs_ADVANCED_SIMULATION
        TAPsCUDA__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
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindVertexList();
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindPrevVertexList();
    TAPs_CUDA_DATA_GlobalPool.DataForVertexList[cudaID]->UnbindSimFlagsList();
    TAPs_CUDA_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 );
    }
    //*/

    TAPs_CUDA_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 ) );
}

void TAPsCUDA__global__SimPropForMultiPartMeshModel_HalfEdge_AdvSim	(	int	numOfVertices,
		int	numOfThreads,
		float	currentTime,
		float	timeStep,
		float *	host_data
	)

For Suture Model (ModelStrand & ModelSuture).

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 963 of file TAPsCUDA_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
    TAPsCUDA__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 ) );
}

template<typename T>

__host__ __device__ void TAPsCUDA__host__device__F1

(

)

[inline]

__host__device__F1 __host__ and __device__ means a function that is compiled for both the host and the device.

Definition at line 122 of file TAPsCUDA.cu.

{}

void TAPsCUDA__host__Init	(	int	argc,
		char **	argv
	)

Initialize CUDA Device.

To use CUDA (Compute Unified Device Architecture by NVIDIA)

Add G:G: Files Corporation CUDA SDK/common/inc; to the "(C/C++) Additional Include Directories".

Add G:;G: Files Corporation CUDA SDK/common/lib; to the "(Linker) Additional Library Directories".

Add this file to the project solution.

Configure the "Configuration Properties -- General -- Tool" of this file to "Custom Build Tool".

In "Custom Build Step -- General", set the "Command Line" to [Debug Version -- In the program directory] "$(CUDA_BIN_PATH)/nvcc.exe" -ccbin "$(VC80_InstallDir)/bin" -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I"G:G: Files Corporation CUDA SDK/common/inc" -o $(InputName).obj [Debug Version -- In the TAPs/Cuda directory] "$(CUDA_BIN_PATH)/nvcc.exe" -ccbin "$(VC80_InstallDir)/bin" -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I"G:G: Files Corporation CUDA SDK/common/inc" -o $(InputName).obj G:/TAPs/Cuda/ [Release Version -- In the TAPs/Cuda directory] "$(CUDA_BIN_PATH)/nvcc.exe" -ccbin "$(VC80_InstallDir)/bin" -c -Xcompiler /EHsc,/nologo,/MD -I"G:G: Files Corporation CUDA SDK/common/inc" -o $(InputName).obj G:/TAPs/Cuda/

In "Custom Build Step -- General", set the "Description" to Performing Custom Build Step

In "Custom Build Step -- General", set the "Outputs" to $(InputName).obj

In "Custom Build Step -- General", set the "Additional Dependencies" to _Def.cu

From CUDA Programming Guide

Any call to a __global__ function must specify the execution configuration for that call.

The execution configuration defines the dimension of the grid and blocks that will be used to execute the function on the device, as well as the associated stream (see Section 4.5.1.5 for a description of streams). It is specified by inserting an expression of the form <<< Dg, Db, Ns, S >>> between the function name and the parenthesized argument list, where:

• Dg is of type dim3 (see Section 4.3.1.2) and specifies the dimension and size of the grid, such that Dg.x * Dg.y equals the number of blocks being launched; Dg.z is unused;

• Db is of type dim3 (see Section 4.3.1.2) and specifies the dimension and size of each block, such that Db.x * Db.y * Db.z equals the number of threads per block;

• Ns is of type size_t and specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory; this dynamically allocated memory is used by any of the variables declared as an external array as mentioned in Section 4.2.2.3; Ns is an optional argument which defaults to 0;

• S is of type cudaStream_t and specifies the associated stream; S is an optional argument which defaults to 0.

Initialize CUDA Device.

Initialize CUDA Device

Definition at line 21 of file TAPsCUDA_Def.cu.

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

bool TAPsCUDA__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 68 of file TAPsCUDA_Def.cu.

{
    // 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 = TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool;

    // Allocate CUDA data for the object.
    TAPs_CUDA_CLASS_DATA_Vertex_List * newData = new TAPs_CUDA_CLASS_DATA_Vertex_List( numOfVertices, false, 0, true );
    AddToGlobal_Pool_Of_TAPs_CUDA_CLASS_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;
}

bool TAPsCUDA__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).

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

Definition at line 106 of file TAPsCUDA_Def.cu.

{
    // 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 = TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool;

    // Allocate CUDA data for the object.
    TAPs_CUDA_CLASS_DATA_Vertex_List * newData = new TAPs_CUDA_CLASS_DATA_Vertex_List( numOfVertices, false, 0, true );
    AddToGlobal_Pool_Of_TAPs_CUDA_CLASS_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;
}

bool TAPsCUDA__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.

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 30 of file TAPsCUDA_Def.cu.

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

    // Assign CUDA ID to the object
    // The object has to use cudaID to communicate with the CUDA.
    cudaID = TAPs_CUDA_DATA_GlobalPool.SizeOfGlobal_Pool;

    // Allocate CUDA data for the object.
    TAPs_CUDA_CLASS_DATA_Vertex_List * newData = new TAPs_CUDA_CLASS_DATA_Vertex_List( numOfVertices, true, max_connection_size, false );
    AddToGlobal_Pool_Of_TAPs_CUDA_CLASS_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;
}

---

Variable Documentation

texture<float4, 1, cudaReadModeElementType> CudaTexHomeVertexList

texture for list of home vertices

Definition at line 236 of file TAPsCUDA.cu.

texture<float4, 1, cudaReadModeElementType> CudaTexPosConstraintList

texture for list of position constraints (xyz) plus force ratio (w)

Definition at line 241 of file TAPsCUDA.cu.

texture<float4, 1, cudaReadModeElementType> CudaTexPrevVertexList

texture for list of previous vertices

Definition at line 235 of file TAPsCUDA.cu.

texture<uint1, 1, cudaReadModeElementType> CudaTexSimFlagsList

texture for list of simulation flags

Definition at line 240 of file TAPsCUDA.cu.

texture<int1, 1, cudaReadModeElementType> CudaTexVertexConnectionList

texture for list of vertex connections (x:=vertex#, y:=rest length, z:=stiffness, w:=damper)

Definition at line 237 of file TAPsCUDA.cu.

texture<float4, 1, cudaReadModeElementType> CudaTexVertexList

For HETriMeshOneModelMultiParts AdvSim Function.

CUDA texture dimension supports only 1-, 2-, and 4-component vector types, but not 3.

Use texture from CUDA Linear Memory by calling cudaBindTexture(...). CUDA Linear Memory is allocated by calling cudaMalloc(...). To enter or change the data in the Linear Memory, must use cudaMemcpy(...). and free by calling cudaFree(...).

Use texture from CUDA Arrays by calling cudaBindTextureToArray(...). CUDA Array is allocated by calling cudaMallocArray(...). To enter or change the data in the Array, must use cudaMemcpyToArray(...). and free by calling cudaFreeArray(...).

The texture reference has to be global, otherwise NVCC will generate error! Texture References texture for list of vertices

Definition at line 234 of file TAPsCUDA.cu.

TAPs_CUDA_CLASS_DATA_Pool TAPs_CUDA_DATA_GlobalPool

CUDA Global Data Pool.

Definition at line 523 of file TAPsCUDA.cu.