#include "TAPsCUDA_GL_ModelElasticRod.cu"

Include dependency graph for TAPsCUDA_GL_ModelElasticRod_Def.cu:

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

Functions
BEGIN_NAMESPACE_TAPs__CUDA __global__ void	GL__GenCylinderForElasticRodModel_CU (unsigned int numOfNodes, unsigned int numOfThreads, float4 *pVertexData, unsigned int numOfCrossSectionVertices, float radius)
__global__ void	GL__GenCylinderForElasticRodModel_PLHM_CU (unsigned int numOfNodes, unsigned int numOfThreads, float pCenterlines, float pOrientations, float4 *pVertexData, unsigned int numOfCrossSectionVertices, float radius)
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 (with CUDA's PLHM)

Function Documentation

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().

Referenced by ElasticRod_CompByCUDA< T >::GenCylinderVertexDataForDrawing().

{
    // 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:

Here is the caller graph for this function:

BEGIN_NAMESPACE_TAPs__CUDA __global__ void GL__GenCylinderForElasticRodModel_CU	(	unsigned int	numOfNodes,
		unsigned int	numOfThreads,
		float4 *	pVertexData,
		unsigned int	numOfCrossSectionVertices,
		float	radius
	)

Parameters:

numOfNodes	number of nodes
numOfThreads	number of threads
pVertexData	OpenGL vertex buffer object
numOfCrossSectionVertices	number of cross section's vertices
radius	radius

Definition at line 19 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CudaTexGLElasticRod, CudaTexOrientationList, CudaTexVertexList, and QuatenionToRotationMatrix4x4().

{
    // Block index
    int bx = blockIdx.x;
    // Thread index
    int tx = threadIdx.x;
    // Vertex number
    int vertexNo = (numOfThreads * bx + tx);
    if ( vertexNo >= numOfNodes-1 ) return;
    
    // Index for output
    int idx = vertexNo * (numOfCrossSectionVertices+1) * 4 * 3;
    
    float4 curr_ori = tex1Dfetch( CudaTexOrientationList, vertexNo );
    matrix4x4 rotMat0 = QuatenionToRotationMatrix4x4( curr_ori );
    float4 next_ori = tex1Dfetch( CudaTexOrientationList, vertexNo+1 );
    matrix4x4 rotMat1 = QuatenionToRotationMatrix4x4( next_ori );
    float4 curr_pos = tex1Dfetch( CudaTexVertexList, vertexNo );
    matrix4x4 rotMatAvg = (rotMat1 + rotMat0) / 2.0f;
    float4 next_pos = tex1Dfetch( CudaTexVertexList, vertexNo+1 );
    float4 middlePt0 = (next_pos + curr_pos) / 2.0f;
    float4 next2_pos = tex1Dfetch( CudaTexVertexList, vertexNo+2 );
    float4 middlePt1 = (next2_pos + next_pos) / 2.0f;
    
    //__shared__ float3 cs_vertices[numOfCrossSectionVertices];
    float4 vertices[3][20];
    float4 normals[3][20];
    float4 texCoords[3][20];
    
    // Compute the cross section vertices, normals, and texture coordinates for middle point from pos0 to pos1
    for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
        vertices[0][i] = rotMat0 * tex2D( CudaTexGLElasticRod, i, 0.0f );
        vertices[0][i].x *= radius;
        vertices[0][i].y *= radius;
        vertices[0][i].z *= radius;
        vertices[0][i].x += middlePt0.x;
        vertices[0][i].y += middlePt0.y;
        vertices[0][i].z += middlePt0.z;
        normals[0][i]  = rotMat0 * tex2D( CudaTexGLElasticRod, i, 1.0f );
        texCoords[0][i] = tex2D( CudaTexGLElasticRod, i, 2.0f );
    }
    // Compute the cross section vertices, normals, and texture coordinates for pos1
    for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
        vertices[1][i] = rotMatAvg * tex2D( CudaTexGLElasticRod, i, 0.0f );
        vertices[1][i].x *= radius;
        vertices[1][i].y *= radius;
        vertices[1][i].z *= radius;
        vertices[1][i].x += next_pos.x;
        vertices[1][i].y += next_pos.y;
        vertices[1][i].z += next_pos.z;
        normals[1][i]  = rotMatAvg * tex2D( CudaTexGLElasticRod, i, 1.0f );
        texCoords[1][i] = tex2D( CudaTexGLElasticRod, i, 2.0f );
        texCoords[1][i].x = 0.5f;
    }
    // Compute the cross section vertices, normals, and texture coordinates for middle point from pos1 to pos2
    for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
        vertices[2][i] = rotMat1 * tex2D( CudaTexGLElasticRod, i, 0.0f );
        vertices[2][i].x *= radius;
        vertices[2][i].y *= radius;
        vertices[2][i].z *= radius;
        vertices[2][i].x += middlePt1.x;
        vertices[2][i].y += middlePt1.y;
        vertices[2][i].z += middlePt1.z;
        normals[2][i]  = rotMat1 * tex2D( CudaTexGLElasticRod, i, 1.0f );
        texCoords[2][i] = tex2D( CudaTexGLElasticRod, i, 2.0f );
        texCoords[2][i].x = 1.0f;
    }

    {
        // Two triangle strips from middlePt0 to pos1 and from pos1 to middlePt1
        // Draw triangle strip from middlePt0 to pos1
        for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
            pVertexData[idx++] = vertices[1][i];
            pVertexData[idx++] = normals[1][i];
            pVertexData[idx++] = texCoords[1][i];

            pVertexData[idx++] = vertices[0][i];
            pVertexData[idx++] = normals[0][i];
            pVertexData[idx++] = texCoords[0][i];
        }
        pVertexData[idx++] = vertices[1][0];
        pVertexData[idx++] = normals[1][0];
        pVertexData[idx++] = texCoords[1][0];

        pVertexData[idx++] = vertices[0][0];
        pVertexData[idx++] = normals[0][0];
        pVertexData[idx++] = texCoords[0][0];

        // Draw triangle strip from pos1 to middlePt1
        for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
            pVertexData[idx++] = vertices[2][i];
            pVertexData[idx++] = normals[2][i];
            pVertexData[idx++] = texCoords[2][i];

            pVertexData[idx++] = vertices[1][i];
            pVertexData[idx++] = normals[1][i];
            pVertexData[idx++] = texCoords[1][i];

        }
        pVertexData[idx++] = vertices[2][0];
        pVertexData[idx++] = normals[2][0];
        pVertexData[idx++] = texCoords[2][0];

        pVertexData[idx++] = vertices[1][0];
        pVertexData[idx++] = normals[1][0];
        pVertexData[idx++] = texCoords[1][0];
    }
}

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 (with CUDA's PLHM)

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().

Referenced by ElasticRod_CompByCUDA< T >::GenCylinderVertexDataForDrawing().

{
    // 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:

Here is the caller graph for this function:

__global__ void GL__GenCylinderForElasticRodModel_PLHM_CU	(	unsigned int	numOfNodes,
		unsigned int	numOfThreads,
		float *	pCenterlines,
		float *	pOrientations,
		float4 *	pVertexData,
		unsigned int	numOfCrossSectionVertices,
		float	radius
	)

Parameters:

numOfNodes	number of nodes
numOfThreads	number of threads
pCenterlines	centerlines
pOrientations	orientations
pVertexData	OpenGL vertex buffer object
numOfCrossSectionVertices	number of cross section's vertices
radius	radius

Definition at line 136 of file TAPsCUDA_GL_ModelElasticRod_Def.cu.

References CudaTexGLElasticRod, and QuatenionToRotationMatrix4x4().

{
    // Block index
    int bx = blockIdx.x;
    // Thread index
    int tx = threadIdx.x;
    // Vertex number
    int vertexNo = (numOfThreads * bx + tx);
    if ( vertexNo >= numOfNodes-1 ) return;
    
    // Index for output
    int idx = vertexNo * (numOfCrossSectionVertices+1) * 4 * 3;
    
    int idx0 = vertexNo * 4;
    int idx1 = idx0+1;
    int idx2 = idx1+1;
    int idx3 = idx2+1;
    
    //float4 curr_ori = tex1Dfetch( CudaTexOrientationList, vertexNo );
    float4 curr_ori = make_float4( pOrientations[idx0], pOrientations[idx1], pOrientations[idx2], pOrientations[idx3] );
    matrix4x4 rotMat0 = QuatenionToRotationMatrix4x4( curr_ori );
    //float4 next_ori = tex1Dfetch( CudaTexOrientationList, vertexNo+1 );
    float4 next_ori = make_float4( pOrientations[idx0+4], pOrientations[idx1+4], pOrientations[idx2+4], pOrientations[idx3+4] );
    matrix4x4 rotMat1 = QuatenionToRotationMatrix4x4( next_ori );
    //float4 curr_pos = tex1Dfetch( CudaTexVertexList, vertexNo );
    float4 curr_pos = make_float4( pCenterlines[idx0], pCenterlines[idx1], pCenterlines[idx2], pCenterlines[idx3] );
    matrix4x4 rotMatAvg = (rotMat1 + rotMat0) / 2.0f;
    //float4 next_pos = tex1Dfetch( CudaTexVertexList, vertexNo+1 );
    float4 next_pos = make_float4( pCenterlines[idx0+4], pCenterlines[idx1+4], pCenterlines[idx2+4], pCenterlines[idx3+4] );
    float4 middlePt0 = (next_pos + curr_pos) / 2.0f;
    //float4 next2_pos = tex1Dfetch( CudaTexVertexList, vertexNo+2 );
    float4 next2_pos = make_float4( pCenterlines[idx0+8], pCenterlines[idx1+8], pCenterlines[idx2+8], pCenterlines[idx3+8] );
    float4 middlePt1 = (next2_pos + next_pos) / 2.0f;
    
    //__shared__ float3 cs_vertices[numOfCrossSectionVertices];
    float4 vertices[3][20];
    float4 normals[3][20];
    float4 texCoords[3][20];
    
    // Compute the cross section vertices, normals, and texture coordinates for middle point from pos0 to pos1
    for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
        vertices[0][i] = rotMat0 * tex2D( CudaTexGLElasticRod, i, 0.0f );
        vertices[0][i].x *= radius;
        vertices[0][i].y *= radius;
        vertices[0][i].z *= radius;
        vertices[0][i].x += middlePt0.x;
        vertices[0][i].y += middlePt0.y;
        vertices[0][i].z += middlePt0.z;
        normals[0][i]  = rotMat0 * tex2D( CudaTexGLElasticRod, i, 1.0f );
        texCoords[0][i] = tex2D( CudaTexGLElasticRod, i, 2.0f );
    }
    // Compute the cross section vertices, normals, and texture coordinates for pos1
    for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
        vertices[1][i] = rotMatAvg * tex2D( CudaTexGLElasticRod, i, 0.0f );
        vertices[1][i].x *= radius;
        vertices[1][i].y *= radius;
        vertices[1][i].z *= radius;
        vertices[1][i].x += next_pos.x;
        vertices[1][i].y += next_pos.y;
        vertices[1][i].z += next_pos.z;
        normals[1][i]  = rotMatAvg * tex2D( CudaTexGLElasticRod, i, 1.0f );
        texCoords[1][i] = tex2D( CudaTexGLElasticRod, i, 2.0f );
        texCoords[1][i].x = 0.5f;
    }
    // Compute the cross section vertices, normals, and texture coordinates for middle point from pos1 to pos2
    for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
        vertices[2][i] = rotMat1 * tex2D( CudaTexGLElasticRod, i, 0.0f );
        vertices[2][i].x *= radius;
        vertices[2][i].y *= radius;
        vertices[2][i].z *= radius;
        vertices[2][i].x += middlePt1.x;
        vertices[2][i].y += middlePt1.y;
        vertices[2][i].z += middlePt1.z;
        normals[2][i]  = rotMat1 * tex2D( CudaTexGLElasticRod, i, 1.0f );
        texCoords[2][i] = tex2D( CudaTexGLElasticRod, i, 2.0f );
        texCoords[2][i].x = 1.0f;
    }

    {
        // Two triangle strips from middlePt0 to pos1 and from pos1 to middlePt1
        // Draw triangle strip from middlePt0 to pos1
        for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
            pVertexData[idx++] = vertices[1][i];
            pVertexData[idx++] = normals[1][i];
            pVertexData[idx++] = texCoords[1][i];

            pVertexData[idx++] = vertices[0][i];
            pVertexData[idx++] = normals[0][i];
            pVertexData[idx++] = texCoords[0][i];
        }
        pVertexData[idx++] = vertices[1][0];
        pVertexData[idx++] = normals[1][0];
        pVertexData[idx++] = texCoords[1][0];

        pVertexData[idx++] = vertices[0][0];
        pVertexData[idx++] = normals[0][0];
        pVertexData[idx++] = texCoords[0][0];

        // Draw triangle strip from pos1 to middlePt1
        for ( int i = 0; i < numOfCrossSectionVertices; ++i ) {
            pVertexData[idx++] = vertices[2][i];
            pVertexData[idx++] = normals[2][i];
            pVertexData[idx++] = texCoords[2][i];

            pVertexData[idx++] = vertices[1][i];
            pVertexData[idx++] = normals[1][i];
            pVertexData[idx++] = texCoords[1][i];

        }
        pVertexData[idx++] = vertices[2][0];
        pVertexData[idx++] = normals[2][0];
        pVertexData[idx++] = texCoords[2][0];

        pVertexData[idx++] = vertices[1][0];
        pVertexData[idx++] = normals[1][0];
        pVertexData[idx++] = texCoords[1][0];
    }
}

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:

Functions

Function Documentation