コード例 #1
0
void PotentialFieldSolver::evaluateScalarPotential()
{
	double h = m_SpatialHasher_mass.getCellSize().x;
	m_particle_dphidn->memset(0);



	PotentialInterpolateFarFieldScalar(m_evalPos->getDevicePtr(),
		m_grid_phi->getDevicePtr(),m_particle_dphidn->getDevicePtr(),
		m_SpatialHasher_mass.getCellSize().x,
		m_gridx,m_gridy,m_gridz,m_M_eval,m_origin);

	Potential_PPCorrMNScalar(m_SpatialHasher_mass.getStartTable(),
		m_SpatialHasher_mass.getEndTable(),
		m_evalPos->getDevicePtr(),
		m_p_massPos_Reorder->getDevicePtr(),
		m_particle_mass_Reorder->getDevicePtr(),
		m_particle_dphidn->getDevicePtr(),
		1.0/h,
		h,
		1.0,
		1.0,
		make_uint3(m_gridx,m_gridy,m_gridz),
		make_uint3(m_gridx,m_gridy,m_gridz),
		make_uint2(m_gridx*m_gridy,m_gridx),
		make_uint2(m_gridx*m_gridy,m_gridx),
		m_K,
		m_M_eval,
		m_N_mass,
		m_origin);

}
コード例 #2
0
ファイル: Particles.cpp プロジェクト: cdf1989/FluidSim
void Particles::InitParams()
{
	mparams.radius = 0.5f;
	mparams.gridSize = make_uint3(64,64,256);
	mparams.cellSize = mparams.radius * 2.0f;
	mparams.cellNum = make_uint3(mparams.gridSize.x / mparams.cellSize, mparams.gridSize.y / mparams.cellSize, mparams.gridSize.z / mparams.cellSize);
	
	mparams.wholeNumCells = mparams.cellNum.x * mparams.cellNum.y*mparams.cellNum.z;
	mparams.worldPos = make_float3(0,0,0);
	mparams.colliderRadius = 0.9f * mparams.radius;
	mparams.gravity = make_float3(0.0f, -9.81f, 0.0f);
	mparams.timeStep = 0.001f;
	mparams.boundaryDamping = -0.5f;
	mparams.globalDamping = 1.0f;
	numParticles = 300000;

	//fluid coeffecient
	mparams.cutoffdist = 1.0f;
	mparams.stiffness = 500.0f;
	mparams.scale = 1.0f;
	mparams.restRHO = 1.7f;
	mparams.visalocityScale = 0.1f;
	mparams.tensionScale = 0.001f;
	mparams.spring = 10.0f;
	mparams.shear = 0.1f;
	mparams.damping = 0.0f;
	mparams.attraction = 0.0f;
}
コード例 #3
0
void
PotentialFieldSolver::computeFarFieldBuffer()
{
	m_p_massPos->copy(m_p_massPos->DEVICE_TO_HOST);
	setDomain(m_origin,m_p_massPos->getHostPtr(),m_N_mass,m_L);
	printf("%lf,%lf,%lf,%lf\n",m_L,m_origin.x,m_origin.y,m_origin.z);
	

	m_ParticleToMesh();
	m_SolvePoisson();
	m_ComputeGradient();
	
	PotentialComputeFarField(m_SpatialHasher_mass.getStartTable(),
							 m_SpatialHasher_mass.getEndTable(),
							 m_p_massPos->getDevicePtr(),
							 m_particle_mass_Reorder->getDevicePtr(),
							 m_grid_density->getDevicePtr(),
							 m_grid_phi->getDevicePtr(),
							 m_grid_gradPhi->getDevicePtr(),
							 m_far_gradPhi->getDevicePtr(),
							 m_SpatialHasher_mass.getCellSize().x,
							 1.0,1.0,
							 make_uint3(m_gridx,m_gridy,m_gridz),
							 make_uint3(m_gridx,m_gridy,m_gridz),
							 m_K,
							 m_M_eval,
							 m_origin);
							 
}
コード例 #4
0
void PotentialFieldSolver::computeFarFieldPotentialBuffer()
{
	m_p_massPos->copy(m_p_massPos->DEVICE_TO_HOST);
	setDomain(m_origin,m_p_massPos->getHostPtr(),m_N_mass,m_L);

	m_ParticleToMesh();
	m_SolvePoisson();
	PotentialComputeScalarFarField(m_particle_mass->getDeviceWritePtr(),m_grid_density->getDeviceWritePtr(),m_grid_phi->getDeviceWritePtr(),m_SpatialHasher_mass.getCellSize().x, 1.0,1.0,
		make_uint3(m_gridx,m_gridy,m_gridz),
		make_uint3(m_gridx,m_gridy,m_gridz),
		m_K);
}
コード例 #5
0
ファイル: volumegpu.hpp プロジェクト: guo2004131/freesurfer
      //! Hidden copy constructor
      VolumeGPU( const VolumeGPU& src ) : dims(make_uint3(0,0,0)),
					  d_data(make_cudaPitchedPtr(NULL,0,0,0)) {
	std::cerr << __FUNCTION__
		  << ": Please don't use copy constructor"
		  << std::endl;
	exit( EXIT_FAILURE );
      }
コード例 #6
0
bool
BiotSavartSolver::m_Intepolate()
{
	for (int c=0;c<3;c++)
	{
		m_particle_U[c]->memset(0);
		MeshToParticle(m_evalPos->getDevicePtr(),
			m_grid_U[c]->getDevicePtr(),
			m_particle_U[c]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x,
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint3(m_gridx,m_gridy,m_gridz),
			m_M_eval,
			m_origin);
	}
	return true;
}
コード例 #7
0
ファイル: main.cpp プロジェクト: hoopoe/cmag
void SystemInit()
{

	psystem = new PoiseuilleFlowSystem(
		make_uint3(16, 64 -  2 * 3, 1),
		3,
		make_uint3(16, 64, 4), 1.0f / (2 * (64 - 6) * 1000),
		true);				

	psystem->reset();

	renderer = new ParticleRenderer;
	renderer->setradius(psystem->getParticleRadius());
	renderer->setColorBuffer(psystem->getColorBuffer());		

    //cutilCheckError(cutCreateTimer(&timer));
    sdkCreateTimer(&timer);
}
コード例 #8
0
ファイル: MarchingCubes.cpp プロジェクト: DengzhiLiu/CG
void CMarchingCubes::InitMC(int _width, int _height, int _depth, ElemType* _pVolume)
{
	// Data Array A[:, :, 1], A[:, :, 2], ... , A[:, : , n]
	//m_GridSize = make_uint3(_depth, _width, _height);
	m_GridSize = make_uint3(_width, _height, _depth);
	m_NumVoxels = m_GridSize.x * m_GridSize.y * m_GridSize.z;
	m_MaxVerts = m_GridSize.x * m_GridSize.y * 30;					// Num of MaxVerts need change
#ifdef _DEBUG
	printf("grids: %d * %d * %d = %d voxels\n", m_GridSize.x, m_GridSize.y, m_GridSize.z, m_NumVoxels);
#endif // _DEBUG

	// needed change
	int size = m_GridSize.x * m_GridSize.y * m_GridSize.z * sizeof(float);
	//////////////////////////////////////////////////////////////////////////
	int len = m_GridSize.x * m_GridSize.y * m_GridSize.z * 3;
	float *pVolTemp = new float[len];
	for (int i = 0; i < len; i++)
		pVolTemp[i] = _pVolume[i];
	//////////////////////////////////////////////////////////////////////////
	cutilSafeCall(cudaMalloc((void**) &m_pdVolume, size * 3));
	cutilSafeCall(cudaMemcpy(m_pdVolume, pVolTemp, size * 3, cudaMemcpyHostToDevice) );
	bindVolumeTexture(m_pdVolume);				// map the coordinates to the texture directly
	delete []pVolTemp;

	// allocate textures
	allocateTextures( &m_pdEdgeTable, &m_pdTriTable, &m_pdNumVertsTable );

	// allocate device memory
	unsigned int memSize = sizeof(uint) * m_NumVoxels;
	cutilSafeCall(cudaMalloc((void**) &m_pdVoxelVerts, memSize));
	cutilSafeCall(cudaMalloc((void**) &m_pdVoxelVertsScan, memSize));
	cutilSafeCall(cudaMalloc((void**) &m_pdVoxelOccupied, memSize));
	cutilSafeCall(cudaMalloc((void**) &m_pdVoxelOccupiedScan, memSize));
	cutilSafeCall(cudaMalloc((void**) &m_pdCompactedVoxelArray, memSize));

	// initialize CUDPP scan
	CUDPPConfiguration config;
	config.algorithm = CUDPP_SCAN;
	config.datatype = CUDPP_UINT;
	config.op = CUDPP_ADD;
	config.options = CUDPP_OPTION_FORWARD | CUDPP_OPTION_EXCLUSIVE;
	cudppPlan(&m_Scanplan, config, m_NumVoxels, 1, 0);
}
コード例 #9
0
VR_Window::VR_Window()
{
    set_orientation( Gtk::ORIENTATION_HORIZONTAL );

    pc_file_open = false;

    vrender = new VRender;
    volume_origin = make_float3( 0.5, 0, 0 );

    cloud = new Cloud;

    cloud->world.resolution = make_float3( RENDER_RESOLUTION, RENDER_RESOLUTION, RENDER_RESOLUTION );
    printf("\n World Resolution: %f x %f x %f", cloud->world.resolution.x, cloud->world.resolution.y, cloud->world.resolution.z );

    cloud->world.size = make_uint3( MAX_VOLUME_SIDE, MAX_VOLUME_SIDE, MAX_VOLUME_SIDE );
    printf("\n World Size: %d x %d x %d", cloud->world.size.x, cloud->world.size.y, cloud->world.size.z );

    cloud->world.count = cloud->world.size.x * cloud->world.size.y * cloud->world.size.z;
    printf("\n World Count: %d",cloud->world.count);

    cloud->world.dimension = make_float3( cloud->world.resolution.x * (float)cloud->world.size.x,
                                          cloud->world.resolution.y * (float)cloud->world.size.y,
                                          cloud->world.resolution.z * (float)cloud->world.size.z );
    printf("\n World Dimension: %f x %f x %f", cloud->world.dimension.x, cloud->world.dimension.y, cloud->world.dimension.z );

    cloud->world.min.x = 1000.f * volume_origin.x - 0.5 * cloud->world.dimension.x;
    cloud->world.min.y = 1000.f * volume_origin.y - 0.5 * cloud->world.dimension.y;
    cloud->world.min.z = 1000.f * volume_origin.z - 0.5 * cloud->world.dimension.z;
    printf("\n World Minimum: %f %f %f", cloud->world.min.x, cloud->world.min.y, cloud->world.min.z );

    cloud->world.max.x = cloud->world.min.x + cloud->world.dimension.x;
    cloud->world.max.y = cloud->world.min.y + cloud->world.dimension.y;
    cloud->world.max.z = cloud->world.min.z + cloud->world.dimension.z;
    printf("\n World Maximum: %f %f %f\n", cloud->world.max.x, cloud->world.max.y, cloud->world.max.z );

    adaptive_world_sizing = false;
    socket_timer_idx = 0;
    memset( socket_timer, 0, TIMER_SIZE * sizeof(double) );

    numKinects = 1;
}
コード例 #10
0
 DINLINE float3_X getPosition( UNIRNG& rng,
                                const uint32_t totalNumParsPerCell,
                                const uint32_t curParticle )
 {
     // spacing between particles in each direction in the cell
     const float3_X spacing = float3_X( float_X(1.0) / float_X(numParsPerCell_X),
                                         float_X(1.0) / float_X(numParsPerCell_Y),
                                         float_X(1.0) / float_X(numParsPerCell_Z) );
     // length of the x lattice, number of particles in the xy plane
     const uint32_t lineX   = numParsPerCell_X;
     const uint32_t planeXY = numParsPerCell_X * numParsPerCell_Y;
     
     // coordinate in the local in-cell lattice
     //   x = [0, numParsPerCell_X-1]
     //   y = [0, numParsPerCell_Y-1]
     //   z = [0, numParsPerCell_Z-1]
     const uint3 inCellCoordinate = make_uint3( curParticle % lineX,
                                                (curParticle % planeXY) / lineX,
                                                curParticle / planeXY );
     
     return float3_X( float_X(inCellCoordinate.x) * spacing.x() + spacing.x() * 0.5,
                         float_X(inCellCoordinate.y) * spacing.y() + spacing.y() * 0.5,
                         float_X(inCellCoordinate.z) * spacing.z() + spacing.z() * 0.5 );
 }
コード例 #11
0
ファイル: ParticleSystem.cpp プロジェクト: dougct/ocelot-ufmg
void
ParticleSystem::reset(ParticleConfig config)
{
    switch (config)
    {
        default:
        case CONFIG_RANDOM:
            initCubeRandom(vec3f(0.0, 1.0, 0.0), vec3f(1.0, 1.0, 1.0), vec3f(0.0f), 100.0);
            break;

        case CONFIG_GRID:
            {
                float jitter = m_particleRadius*0.01f;
                uint s = (int) ceilf(powf((float) m_numParticles, 1.0f / 3.0f));
                uint gridSize[3];
                gridSize[0] = gridSize[1] = gridSize[2] = s;
                initGrid(vec3f(-1.0, 0.0, -1.0), make_uint3(s, s, s), vec3f(m_particleRadius*2.0f), jitter, vec3f(0.0), m_numParticles, 100.0);
            }
            break;
    }

    m_pos.copy(GpuArray<float4>::HOST_TO_DEVICE);
    m_vel.copy(GpuArray<float4>::HOST_TO_DEVICE);
}
コード例 #12
0
bool PotentialFieldSolver::evaluateGradient( bool large_eval )
{
	double h = m_SpatialHasher_mass.getCellSize().x;
	m_particle_gradPhi->memset(make_double3(0,0,0));


	if(large_eval)
	{
		m_SpatialHasher_eval.endSpatialHash();
		m_SpatialHasher_eval.initSpatialHash(m_M_eval,m_gridx,m_gridy,m_gridz);


		m_SpatialHasher_eval.setSpatialHashGrid(m_gridx, h,
			m_SpatialHasher_mass.getWorldOrigin(),
			m_M_eval);
		m_SpatialHasher_eval.setHashParam();
		m_SpatialHasher_eval.doSpatialHash(m_evalPos->getDevicePtr(),m_M_eval);
		m_SpatialHasher_eval.reorderData(m_M_eval, m_evalPos->getDevicePtr(),m_evalPos_Reorder->getDevicePtr(),4,1);

		m_particle_gradPhi_deorder->memset(make_double3(0,0,0));

		PotentialInterpolateFarField(m_evalPos_Reorder->getDevicePtr(),
			m_far_gradPhi->getDevicePtr(),m_particle_gradPhi_deorder->getDevicePtr(),
			m_SpatialHasher_mass.getCellSize().x,
			m_gridx,m_gridy,m_gridz,m_M_eval,m_origin);

		Potential_PPCorrMN(m_SpatialHasher_mass.getStartTable(),
			m_SpatialHasher_mass.getEndTable(),
			m_evalPos_Reorder->getDevicePtr(),
			m_p_massPos_Reorder->getDevicePtr(),
			m_particle_mass_Reorder->getDevicePtr(),
			m_particle_gradPhi_deorder->getDevicePtr(),
			1.0/h,
			h,
			1.0,
			1.0,
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint2(m_gridx*m_gridy,m_gridx),
			make_uint2(m_gridx*m_gridy,m_gridx),
			m_K,
			m_M_eval,
			m_N_mass,
			m_origin);
		m_SpatialHasher_eval.deorderData(m_M_eval,m_particle_gradPhi_deorder->getDevicePtr(),m_particle_gradPhi->getDevicePtr(),3,2);
	}
	else
	{
		PotentialInterpolateFarField(m_evalPos->getDevicePtr(),
			m_far_gradPhi->getDevicePtr(),m_particle_gradPhi->getDevicePtr(),
			m_SpatialHasher_mass.getCellSize().x,
			m_gridx,m_gridy,m_gridz,m_M_eval,m_origin);

		Potential_PPCorrMN(m_SpatialHasher_mass.getStartTable(),
			m_SpatialHasher_mass.getEndTable(),
			m_evalPos->getDevicePtr(),
			m_p_massPos_Reorder->getDevicePtr(),
			m_particle_mass_Reorder->getDevicePtr(),
			m_particle_gradPhi->getDevicePtr(),
			1.0/h,
			h,
			1.0,
			1.0,
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint2(m_gridx*m_gridy,m_gridx),
			make_uint2(m_gridx*m_gridy,m_gridx),
			m_K,
			m_M_eval,
			m_N_mass,
			m_origin);
	}


	

	PotentialComputeGradForOutParticle(m_evalPos->getDevicePtr(),m_total_mass, m_center,
		m_SpatialHasher_mass.getWorldOrigin(), 
		make_float3(m_SpatialHasher_mass.getWorldOrigin().x+m_L,
		m_SpatialHasher_mass.getWorldOrigin().y+m_L,
		m_SpatialHasher_mass.getWorldOrigin().z+m_L),
		1.0,1.0,m_particle_gradPhi->getDevicePtr(),m_M_eval);


	return true;
}
コード例 #13
0
ファイル: voxelizer_utils.cpp プロジェクト: gramez/Voxelizer
void testOverlap(vtkPolyData *data, uint y, uint z, bool returnAllResults, uint res)
{
    double3* vertices;
    uint3* indices;
    double3* normals;
    double3 minVertex;
    double3 maxVertex;
    double voxelDistance;
    uint3 resolution;
    bool testState;

    // Get triangles with indices (polys) and the vertices (points).
    vtkCellArray* polys = data->GetPolys();
    vtkPoints* points = data->GetPoints();

    uint numberOfVertices = (uint)points->GetNumberOfPoints();
    cout << numberOfVertices << " vertices found.\n";
    uint numberOfPolygons = (uint)polys->GetNumberOfCells();
    cout << numberOfPolygons << " polygons found.\n";

    // Allocate vertex and index arrays in main memory.
    vertices = new double3[numberOfVertices];
    indices = new uint3[numberOfPolygons];
    normals = new double3[numberOfPolygons];

    // Calculate the bounding box of the input model.
    double vert[3];
    points->GetPoint(0, vert);

    minVertex.x = vert[0];
    maxVertex.x = vert[0];
    minVertex.y = vert[1];
    maxVertex.y = vert[1];
    minVertex.z = vert[2];
    maxVertex.z = vert[2];

    for (vtkIdType i = 1; i < points->GetNumberOfPoints(); i++)
    {
        points->GetPoint(i, vert);
        // Determine minimum and maximum coordinates.

        if (vert[0] > maxVertex.x)
            maxVertex.x = vert[0];
        if (vert[1] > maxVertex.y)
            maxVertex.y = vert[1];
        if (vert[2] > maxVertex.z)
            maxVertex.z = vert[2];

        if (vert[0] < minVertex.x)
            minVertex.x = vert[0];
        if (vert[1] < minVertex.y)
            minVertex.y = vert[1];
        if (vert[2] < minVertex.z)
            minVertex.z = vert[2];
    }

    cout << "Minimum corner: " << minVertex.x << ", " << minVertex.y << ", " << minVertex.z << "\n";
    cout << "Maximum corner: " << maxVertex.x << ", " << maxVertex.y << ", " << maxVertex.z << "\n";

    // Copy vertex data to the device.
    for (vtkIdType i = 0; i < points->GetNumberOfPoints(); i++)
    {
        points->GetPoint(i, vert);

        vertices[i] = make_double3(vert[0], vert[1], vert[2]); // - minVertex;
    }

    // Copy index data to the array.
    vtkIdList* idlist = vtkIdList::New();
    uint currentPoly = 0;
    polys->InitTraversal();
    while(polys->GetNextCell(idlist))
    {
        uint index[3] = { 0, 0, 0 };
        for (vtkIdType i = 0; i < idlist->GetNumberOfIds(); i++)
        {
            if (i > 2)
            {
                cout << "Too many indices! Only triangle-meshes are supported.";
                return;
            }

            index[i] = idlist->GetId(i);
        }

        indices[currentPoly] = make_uint3(index[0], index[1], index[2]);

        // Calculate the surface normal of the current polygon.
        double3 U = make_double3(vertices[indices[currentPoly].x].x - vertices[indices[currentPoly].z].x,
                               vertices[indices[currentPoly].x].y - vertices[indices[currentPoly].z].y, 
                               vertices[indices[currentPoly].x].z - vertices[indices[currentPoly].z].z);
        double3 V = make_double3(vertices[indices[currentPoly].y].x - vertices[indices[currentPoly].x].x,
                               vertices[indices[currentPoly].y].y - vertices[indices[currentPoly].x].y, 
                               vertices[indices[currentPoly].y].z - vertices[indices[currentPoly].x].z);

        normals[currentPoly] = normalize(cross(U, V));

        currentPoly++;
    }

    double3 diffVertex = maxVertex - minVertex;

    if (diffVertex.x > diffVertex.y)
    {
        if (diffVertex.x > diffVertex.z)
        {
            voxelDistance = diffVertex.x / double(res - 1);
            resolution.x = res;
            resolution.y = uint(diffVertex.y / voxelDistance) + 1;
            resolution.z = uint(diffVertex.z / voxelDistance) + 1;
        }
        else
        {
            voxelDistance = diffVertex.z / double(res - 1);
            resolution.x = uint(diffVertex.x / voxelDistance) + 1;
            resolution.y = uint(diffVertex.y / voxelDistance) + 1;
            resolution.z = res;
        }
    }
    else
    {
        if (diffVertex.y > diffVertex.z)
        {
            voxelDistance = diffVertex.y / double(res - 1);
            resolution.x = uint(diffVertex.x / voxelDistance) + 1;
            resolution.y = res;
            resolution.z = uint(diffVertex.z / voxelDistance) + 1;
        }
        else
        {
            voxelDistance = diffVertex.z / double(res - 1);
            resolution.x = uint(diffVertex.x / voxelDistance) + 1;
            resolution.y = uint(diffVertex.y / voxelDistance) + 1;
            resolution.z = res;
        }
    }

    if (resolution.x % 32 != 0)
        resolution.x += 32 - (resolution.x % 32);
    if (resolution.y % 16 != 0)
        resolution.y += 16 - (resolution.y % 16);
    if (resolution.z % 16 != 0)
        resolution.z += 16 - (resolution.z % 16);

    cout << "Voxel width: " << voxelDistance << "\n";
    cout << "Resolution: " << resolution.x << " X " << resolution.y << " X " << resolution.z << "\n";
    cout << "YZ-coordinates: (" << y << ", " << z << ")\n\n";

    uint intersectionCount = 0;
    std::vector<uint> intersections = std::vector<uint>();

    // Shoot a ray from the voxel coordinates, and test for intersection against every triangle.
    for (uint i = 0; i < numberOfPolygons; i++)
    {
        double2 vertexProjections[3];
        double2 edgeNormals[3];
        double distanceToEdge[3];
        double testResult[3];

        if (normals[i].x == 0.0)
            continue;

        vertexProjections[0] = make_double2(vertices[indices[i].x].y, vertices[indices[i].x].z);
        vertexProjections[1] = make_double2(vertices[indices[i].y].y, vertices[indices[i].y].z);
        vertexProjections[2] = make_double2(vertices[indices[i].z].y, vertices[indices[i].z].z);

        double2 p = make_double2(minVertex.y + double(y) * voxelDistance, minVertex.z + double(z) * voxelDistance);

        double2 vMin = vertexProjections[0];
        double2 vMax = vertexProjections[0];

        if (vertexProjections[1].x < vMin.x)
            vMin.x = vertexProjections[1].x;
        if (vertexProjections[1].y < vMin.y)
            vMin.y = vertexProjections[1].y;
        if (vertexProjections[2].x < vMin.x)
            vMin.x = vertexProjections[2].x;
        if (vertexProjections[2].y < vMin.y)
            vMin.y = vertexProjections[2].y;

        if (vertexProjections[1].x > vMax.x)
            vMax.x = vertexProjections[1].x;
        if (vertexProjections[1].y > vMax.y)
            vMax.y = vertexProjections[1].y;
        if (vertexProjections[2].x > vMax.x)
            vMax.x = vertexProjections[2].x;
        if (vertexProjections[2].y > vMax.y)
            vMax.y = vertexProjections[2].y;

        if ((p.x < vMin.x) || (p.y < vMin.y) || (p.x > vMax.x) || (p.y > vMax.y))
            continue;

        for (uint j = 0; j < 3; j++)
        {
            double2 edge = make_double2(vertexProjections[(j+1)%3].x - vertexProjections[j].x, 
                                        vertexProjections[(j+1)%3].y - vertexProjections[j].y);
            if (normals[i].x >= 0.0)
                edgeNormals[j] = normalize(make_double2(-edge.y, edge.x));
            else
                edgeNormals[j] = normalize(make_double2(edge.y, -edge.x));

            distanceToEdge[j] = edgeNormals[j].x * vertexProjections[j].x + edgeNormals[j].y * vertexProjections[j].y;
            distanceToEdge[j] *= -1.0;

            testResult[j] = distanceToEdge[j] + edgeNormals[j].x * p.x + edgeNormals[j].y * p.y;

            if ((edgeNormals[j].x > 0.0) || ((edgeNormals[j].x == 0) && (edgeNormals[j].y < 0.0)))
                testResult[j] += DBL_MIN;
            
        }

        if ((testResult[0] > 0.0) && (testResult[1] > 0.0) && (testResult[2] > 0.0))
            testState = true;
        else
            testState = false;

        if (testState || returnAllResults)
        {
            if (testState)
            {
                intersectionCount++;
                cout << "Intersection nr. " << intersectionCount << " found with triangle " << i << "\n";
            }
            else
                cout << "No intersection found with triangle " << i << "\n";
            cout << "Vertex 1      : (" << vertexProjections[0].x << ", " << vertexProjections[0].y << ")\n";
            cout << "Vertex 2      : (" << vertexProjections[1].x << ", " << vertexProjections[1].y << ")\n";
            cout << "Vertex 3      : (" << vertexProjections[2].x << ", " << vertexProjections[2].y << ")\n";
            cout << "Normal        : (" << normals[i].x << ", " << normals[i].y << ", " << normals[i].z << ")\n";
            cout << "Edge normal 1 : (" << edgeNormals[0].x << ", " << edgeNormals[0].y << ")\n";
            cout << "Edge normal 2 : (" << edgeNormals[1].x << ", " << edgeNormals[1].y << ")\n";
            cout << "Edge normal 3 : (" << edgeNormals[2].x << ", " << edgeNormals[2].y << ")\n";
            cout << "Distances     : " << distanceToEdge[0] << ", " << distanceToEdge[1] << ", " << distanceToEdge[2] << "\n";
            cout << "Test results  : " << testResult[0] << ", " << testResult[1] << ", " << testResult[2] << "\n";
            cout << "P             : " << p.x << ", " << p.y << "\n";
            cout << (testState ? "Intersection found" : "No intersection") << " at X: ";

            double3 v = vertices[indices[i].x];
            double3 n = normals[i];

            double3 A = make_double3(v.x - minVertex.x, v.y - p.x, v.z - p.y);
            double B = dot(A, n);
            double px = B / n.x;

            uint voxelX = ceilf(px / voxelDistance);

            cout << voxelX << "\n\n";

            if (testState)
                intersections.push_back(voxelX);
        }
    }

    uint voxels[UTIL_RES / 32];

    for (uint i = 0; i < UTIL_RES / 32; i++)
        voxels[i] = 0;

    if (intersectionCount > 0)
    {
        for (std::vector<unsigned int>::iterator it = intersections.begin(); it < intersections.end(); it++)
        {
            uint intersection = *it;

            uint relevantInteger = intersection / 32;
            uint relevantBit = intersection % 32;

            uint bitmask = UINT_MAX >> relevantBit;
            voxels[relevantInteger] ^= bitmask;

            for (uint i = relevantInteger + 1; i < UTIL_RES / 32; i++)
                voxels[i] ^= UINT_MAX;
        }

        cout << "The voxelization for Y: " << y << ", Z: " << z << "\n\n";
        char hexRep[12];
        for (uint i = 0; i < UTIL_RES / 32; i++)
        {
            sprintf( hexRep, "%X", voxels[i] );
            //sprintf_s(hexRep, 12*sizeof(char), "%X", voxels[i]);
            cout << hexRep << " ";
        }

        cout << "\n\n";
    }
    else if (returnAllResults)
コード例 #14
0
ファイル: NCVPixelOperations.hpp プロジェクト: marketto89/pcl
static inline __host__ __device__ uint3 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z) {return make_uint3(x,y,z);}
コード例 #15
0
ファイル: vecvolgpu.hpp プロジェクト: guo2004131/freesurfer
 //! Default constructor
 VecVolGPU( void ) : dims(make_uint3(0,0,0)),
                     d_x(make_cudaPitchedPtr(NULL,0,0,0)),
                     d_y(make_cudaPitchedPtr(NULL,0,0,0)),
                     d_z(make_cudaPitchedPtr(NULL,0,0,0)) {};
コード例 #16
0
ファイル: volumegpu.hpp プロジェクト: ewong718/freesurfer
 //! Default constructor
 VolumeArgGPU( void ) : dims(make_uint3(0,0,0)),
   pitchedPtr(NULL),
   dataPitch(0) {};
コード例 #17
0
ファイル: NCVPixelOperations.hpp プロジェクト: marketto89/pcl
template<> inline __host__ __device__ uint3 _pixMakeZero<uint3>() {return make_uint3(0,0,0);}
コード例 #18
0
bool
BiotSavartSolver::m_ParticleToMesh()
{
	m_SpatialHasher_vort.setSpatialHashGrid(m_gridx, m_L/(double)m_gridx,
		make_float3(m_origin.x,m_origin.y,m_origin.z),
		m_N_vort);
	m_SpatialHasher_vort.setHashParam();
	m_SpatialHasher_vort.doSpatialHash(m_p_vortPos->getDevicePtr(),m_N_vort);

	m_p_vortPos_Reorder->memset(make_float4(0,0,0,0));
	m_SpatialHasher_vort.reorderData(m_N_vort, (void*)(m_p_vortPos->getDevicePtr()),
		(void*)(m_p_vortPos_Reorder->getDevicePtr()), 4, 1);


	for(int i=0;i<NUM_COMPONENTS;i++)
	{
		m_particle_vort_Reorder[i]->memset(0);
		m_SpatialHasher_vort.reorderData(m_N_vort, (void*)(m_particle_vort[i]->getDevicePtr()),
			(void*)(m_particle_vort_Reorder[i]->getDevicePtr()), 1, 2);

	}

	for (int c=0;c<NUM_COMPONENTS;c++)
	{
		m_grid_vort[c]->memset(0);
		ParticleToMesh(m_SpatialHasher_vort.getStartTable(),
			m_SpatialHasher_vort.getEndTable(),
			m_p_vortPos_Reorder->getDevicePtr(),
			m_particle_vort_Reorder[c]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x,
			m_grid_vort[c]->getDevicePtr(),
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint3(m_gridx,m_gridy,m_gridz),
			m_N_vort,
			m_origin);
		cudaMemcpy(m_grid_Rhs[c]->getDevicePtr(),
			m_grid_vort[c]->getDevicePtr(),
			m_grid_Rhs[c]->getSize()*m_grid_Rhs[c]->typeSize(),
			cudaMemcpyDeviceToDevice);
		ComputeRHS(m_grid_Rhs[c]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x*m_SpatialHasher_vort.getCellSize().x,
			-1.0,
			m_gridx*m_gridy*m_gridz);
		//m_p_vortPos_Reorder->copy(GpuArrayf4::DEVICE_TO_HOST);
		//m_particle_vort_Reorder[c]->copy(GpuArrayd::DEVICE_TO_HOST);
		//double total_weight = 0;
		//double total_mass = 0;
		//for(int i=0; i<m_N_vort; i++)
		//{
		//	double *host = m_particle_vort_Reorder[c]->getHostPtr();
		//	total_weight += fabs(host[i]);
		//	total_mass += host[i];
		//}
		//double cx=0, cy=0, cz=0;
		//for(int i=0; i<m_N_vort; i++)
		//{
		//	float4 *hpos = m_p_vortPos_Reorder->getHostPtr();
		//	double *hmass = m_particle_vort_Reorder[c]->getHostPtr();
		//	cx+=hpos[i].x*fabs(hmass[i]);
		//	cy+=hpos[i].y*fabs(hmass[i]);
		//	cz+=hpos[i].z*fabs(hmass[i]);
		//	//printf("%f,%f,%f\n",cx,cy,cz);
		//}
		//cx=cx/total_weight;
		//cy=cy/total_weight;
		//cz=cz/total_weight;

		//m_center.x = cx;
		//m_center.y = cy;
		//m_center.z = cz;
		//m_total_vort[c] = total_mass;
		////printf("%f,%f,%f,%f\n",cx,cy,cz,total_mass);
		//applyDirichlet(m_grid_Rhs[c]->getDevicePtr(), 
		//	make_double4(cx,cy,cz,0),
		//	total_mass,
		//	m_origin,
		//	m_SpatialHasher_vort.getCellSize().x,
		//	m_gridx,
		//	m_gridy,
		//	m_gridz);
	}


	return true;

}
コード例 #19
0
bool BiotSavartSolver::evaluateVelocity( GpuArrayf4 *another_end, uint is_segment )
{

	//m_p_vortPos->copy(gf_GpuArray<float4>::DEVICE_TO_HOST);
	//setDomain(m_origin, m_p_vortPos->getHostPtr(),m_N_vort,m_L);
	////printf("%f,%f,%f,%f,\n",m_origin.x,m_origin.y,m_origin.z,m_L);
	//m_SpatialHasher_eval.setSpatialHashGrid(m_gridx, m_L/(double)m_gridx,
	//	make_float3(m_origin.x,m_origin.y,m_origin.z),
	//	m_M_eval);
	//m_SpatialHasher_eval.setHashParam();
	//m_SpatialHasher_eval.doSpatialHash(m_evalPos->getDevicePtr(),m_M_eval);
	//m_SpatialHasher_eval.reorderData(m_M_eval,m_evalPos->getDevicePtr(),m_evalPos_Reorder->getDevicePtr(),4,1);


	//m_ParticleToMesh();
	//m_SolvePoisson();
	//m_ComputeCurl();
	//m_Intepolate();
	//m_LocalCorrection(another_end);
	//m_unsortResult();
	GpuArrayf4 *temp_pos=new GpuArrayf4;
	temp_pos->alloc(m_M_eval);
	temp_pos->memset(make_float4(0,0,0,0));

	for(int i=0;i<NUM_COMPONENTS;i++)
	{
		m_particle_U[i]->memset(0);
		m_particle_U_deorder[i]->memset(0);
	}
	if(!m_isVIC){
		
		m_SpatialHasher_eval.setSpatialHashGrid(m_gridx, m_L/(double)m_gridx,
			make_float3(m_origin.x,m_origin.y,m_origin.z),
			m_M_eval);
		m_SpatialHasher_eval.setHashParam();
		m_SpatialHasher_eval.doSpatialHash(m_evalPos->getDevicePtr(),m_M_eval);
		m_SpatialHasher_eval.reorderData(m_M_eval, m_evalPos->getDevicePtr(),m_evalPos_Reorder->getDevicePtr(),4,1);
		if(is_segment==1)
		{
			m_SpatialHasher_eval.reorderData(m_M_eval,another_end->getDevicePtr(),temp_pos->getDevicePtr(),4,1);
		}

		BiotSavartInterpolateFarField(m_evalPos_Reorder->getDevicePtr(),
			m_far_U[0]->getDevicePtr(),m_far_U[1]->getDevicePtr(), m_far_U[2]->getDevicePtr(),
			m_particle_U_deorder[0]->getDevicePtr(), m_particle_U_deorder[1]->getDevicePtr(),m_particle_U_deorder[2]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x,
			m_gridx,m_gridy,m_gridz,
			m_M_eval,
			m_origin);
		BiotSavartPPCorrScaleMN(m_SpatialHasher_vort.getStartTable(),
			m_SpatialHasher_vort.getEndTable(),
			m_evalPos_Reorder->getDevicePtr(),
			temp_pos->getDevicePtr(),
			is_segment,
			m_p_vortPos_Reorder->getDevicePtr(),
			m_particle_vort_Reorder[0]->getDevicePtr(),
			m_particle_vort_Reorder[1]->getDevicePtr(),
			m_particle_vort_Reorder[2]->getDevicePtr(),
			m_particle_U_deorder[0]->getDevicePtr(),
			m_particle_U_deorder[1]->getDevicePtr(),
			m_particle_U_deorder[2]->getDevicePtr(),
			m_grid_U[0]->getDevicePtr(),
			m_grid_U[1]->getDevicePtr(),
			m_grid_U[2]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x,
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint3(m_gridx,m_gridy,m_gridz),
			m_K,
			m_M_eval,
			m_N_vort,
			m_origin);
		for (int c=0;c<3;c++)
		{
			m_SpatialHasher_eval.deorderData(m_M_eval,m_particle_U_deorder[c]->getDevicePtr(),m_particle_U[c]->getDevicePtr(),1,2);
		}
		
	}
	else
	{
		
		BiotSavartInterpolateFarField(m_evalPos->getDevicePtr(),
			m_grid_U[0]->getDevicePtr(),m_grid_U[1]->getDevicePtr(), m_grid_U[2]->getDevicePtr(),
			m_particle_U[0]->getDevicePtr(), m_particle_U[1]->getDevicePtr(),m_particle_U[2]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x,
			m_gridx,m_gridy,m_gridz,
			m_M_eval,
			m_origin);
		
	}

	//BiotSavartComputeVelocityForOutParticle(m_evalPos->getDevicePtr(),
	//	make_double3(m_total_vort[0],m_total_vort[1],m_total_vort[2]), 
	//	m_center,
	//	m_SpatialHasher_vort.getWorldOrigin(), 
	//	make_float3(m_SpatialHasher_vort.getWorldOrigin().x+m_L,
	//				 m_SpatialHasher_vort.getWorldOrigin().y+m_L,
	//				 m_SpatialHasher_vort.getWorldOrigin().z+m_L),
	//    m_particle_U[0]->getDevicePtr(),
	//	m_particle_U[1]->getDevicePtr(),
	//	m_particle_U[2]->getDevicePtr(),
	//	m_M_eval);

	temp_pos->free();

	return true;

}
コード例 #20
0
void
BiotSavartSolver::computeFarFieldBuffer()
{
	//in order to do this, we need 
	//grid based velocity,
	//sorted vortex and their start end table
	//a double3 5x5x5xDXxDYxDZ local velocity field 
	m_p_vortPos->copy(gf_GpuArray<float4>::DEVICE_TO_HOST);
	setDomain(m_origin, m_p_vortPos->getHostPtr(),m_N_vort,m_L);
	//printf("%f,%f,%f,%f,\n",m_origin.x,m_origin.y,m_origin.z,m_L);
	if(!m_isVIC){

		m_SpatialHasher_eval.setSpatialHashGrid(m_gridx, m_L/(double)m_gridx,
			make_float3(m_origin.x,m_origin.y,m_origin.z),
			m_M_eval);
		m_SpatialHasher_eval.setHashParam();
		m_SpatialHasher_eval.doSpatialHash(m_evalPos->getDevicePtr(),m_M_eval);
		m_SpatialHasher_eval.reorderData(m_M_eval,m_evalPos->getDevicePtr(),m_evalPos_Reorder->getDevicePtr(),4,1);
	}


	m_ParticleToMesh();
	m_SolvePoisson();
	m_ComputeCurl();
	m_grid_U[0]->copy(m_grid_U[0]->DEVICE_TO_HOST);
	m_grid_U[1]->copy(m_grid_U[1]->DEVICE_TO_HOST);
	m_grid_U[2]->copy(m_grid_U[2]->DEVICE_TO_HOST);
	
	if(!m_isVIC){
		for(int i=0; i<3; i++)
		{
			cudaMemcpy(m_far_U[i]->getDevicePtr(),
				m_grid_U[i]->getDevicePtr(),
				sizeof(double)*m_grid_U[i]->getSize(),
				cudaMemcpyDeviceToDevice);
		}
		BiotSavartComputeFarField(m_SpatialHasher_eval.getStartTable(),m_SpatialHasher_eval.getEndTable(),
			m_evalPos->getDevicePtr(),
			m_grid_vort[0]->getDevicePtr(),
			m_grid_vort[1]->getDevicePtr(),
			m_grid_vort[2]->getDevicePtr(),
			m_grid_Psi[0]->getDevicePtr(),
			m_grid_Psi[1]->getDevicePtr(),
			m_grid_Psi[2]->getDevicePtr(),
			m_grid_U[0]->getDevicePtr(),
			m_grid_U[1]->getDevicePtr(),
			m_grid_U[2]->getDevicePtr(),
			m_far_U[0]->getDevicePtr(),m_far_U[1]->getDevicePtr(),m_far_U[2]->getDevicePtr(),
			m_SpatialHasher_vort.getCellSize().x,
			make_uint3(m_gridx,m_gridy,m_gridz),
			make_uint3(m_gridx,m_gridy,m_gridz),
			m_K,
			m_M_eval,
			m_origin);

	}
	m_grid_U[0]->copy(m_grid_U[0]->HOST_TO_DEVICE);
	m_grid_U[1]->copy(m_grid_U[1]->HOST_TO_DEVICE);
	m_grid_U[2]->copy(m_grid_U[2]->HOST_TO_DEVICE);


}
コード例 #21
0
bool
PotentialFieldSolver::m_ParticleToMesh()
{
	m_SpatialHasher_mass.setSpatialHashGrid(m_gridx, m_L/(double)m_gridx,
		make_float3(m_origin.x,m_origin.y,m_origin.z),
		m_N_mass);
	m_SpatialHasher_mass.setHashParam();
	m_SpatialHasher_mass.doSpatialHash(m_p_massPos->getDevicePtr(),m_N_mass);

	m_p_massPos_Reorder->memset(make_float4(0,0,0,0));
	m_SpatialHasher_mass.reorderData(m_N_mass, (void*)(m_p_massPos->getDevicePtr()),
		(void*)(m_p_massPos_Reorder->getDevicePtr()), 4, 1);



	m_particle_mass_Reorder->memset(0);
	m_SpatialHasher_mass.reorderData(m_N_mass, (void*)(m_particle_mass->getDevicePtr()),
			(void*)(m_particle_mass_Reorder->getDevicePtr()), 1, 2);


	m_grid_density->memset(0);
	ParticleToMesh(m_SpatialHasher_mass.getStartTable(),
		m_SpatialHasher_mass.getEndTable(),
		m_p_massPos_Reorder->getDevicePtr(),
		m_particle_mass_Reorder->getDevicePtr(),
		m_SpatialHasher_mass.getCellSize().x,
		m_grid_density->getDevicePtr(),
		make_uint3(m_gridx,m_gridy,m_gridz),
		make_uint3(m_gridx,m_gridy,m_gridz),
		m_N_mass,
		m_origin);
	cudaMemcpy(m_grid_Rhs->getDevicePtr(),
		m_grid_density->getDevicePtr(),
		m_grid_Rhs->getSize()*m_grid_Rhs->typeSize(),
		cudaMemcpyDeviceToDevice);
	ComputeRHS(m_grid_Rhs->getDevicePtr(),
		m_SpatialHasher_mass.getCellSize().x*m_SpatialHasher_mass.getCellSize().x,
		-1.0,
		m_gridx*m_gridy*m_gridz);

	m_p_massPos_Reorder->copy(gf_GpuArray<float4>::DEVICE_TO_HOST);
	m_particle_mass_Reorder->copy(gf_GpuArray<double>::DEVICE_TO_HOST);

	double total_weight = 0;
	double total_mass = 0;
	for(int i=0; i<m_N_mass; i++)
	{
		double *host = m_particle_mass_Reorder->getHostPtr();
		total_weight += fabs(host[i]);
		total_mass += host[i];
	}
	double cx=0, cy=0, cz=0;
	for(int i=0; i<m_N_mass; i++)
	{
		float4 *hpos = m_p_massPos_Reorder->getHostPtr();
		double *hmass = m_particle_mass_Reorder->getHostPtr();
		cx+=hpos[i].x*fabs(hmass[i]);
		cy+=hpos[i].y*fabs(hmass[i]);
		cz+=hpos[i].z*fabs(hmass[i]);
		//printf("%f,%f,%f\n",cx,cy,cz);
	}
	cx=cx/total_weight;
	cy=cy/total_weight;
	cz=cz/total_weight;

	m_center.x = cx;
	m_center.y = cy;
	m_center.z = cz;
	m_total_mass = total_mass;

	applyDirichlet(m_grid_Rhs->getDevicePtr(), 
		make_double4(cx,cy,cz,0),
		m_total_mass,
		make_double4(m_origin.x,m_origin.y,m_origin.z,0),
		m_SpatialHasher_mass.getCellSize().x,
		m_gridx,
		m_gridy,
		m_gridz);

	return true;

}
コード例 #22
0
ファイル: volumegpu.hpp プロジェクト: ewong718/freesurfer
 //! Default constructor
 VolumeGPU( void ) : dims(make_uint3(0,0,0)),
   d_data(make_cudaPitchedPtr(NULL,0,0,0)) {};
コード例 #23
0
						uint3 gridSize, uint3 gridSizeShift, uint3 gridSizeMask,
						float3 voxelSize, float isoValue, uint activeVoxels, uint maxVerts);

extern "C" void allocateTextures(	uint **d_edgeTable, uint **d_triTable,  uint **d_numVertsTable );
extern "C" void bindVolumeTexture(uchar *d_volume);


#include "cudpp/cudpp.h"

// constants
const unsigned int window_width = 512;
const unsigned int window_height = 512;

const char *volumeFilename = "Bucky.raw";

uint3 gridSizeLog2 = make_uint3(5, 5, 5);
uint3 gridSizeShift;
uint3 gridSize;
uint3 gridSizeMask;

float3 voxelSize;
uint numVoxels    = 0;
uint maxVerts     = 0;
uint activeVoxels = 0;
uint totalVerts   = 0;

float isoValue		= 0.2f;
float dIsoValue		= 0.005f;

CUDPPHandle scanplan;
コード例 #24
0
////////////////////////////////////////////////////////////////////////////////
// initialize marching cubes
////////////////////////////////////////////////////////////////////////////////
void
initMC(int argc, char** argv)
{
    // parse command line arguments
    int n;
    if (cutGetCmdLineArgumenti( argc, (const char**) argv, "grid", &n)) {
        gridSizeLog2.x = gridSizeLog2.y = gridSizeLog2.z = n;
    }
    if (cutGetCmdLineArgumenti( argc, (const char**) argv, "gridx", &n)) {
        gridSizeLog2.x = n;
    }
    if (cutGetCmdLineArgumenti( argc, (const char**) argv, "gridy", &n)) {
        gridSizeLog2.y = n;
    }
    if (cutGetCmdLineArgumenti( argc, (const char**) argv, "gridz", &n)) {
        gridSizeLog2.z = n;
    }

    char *filename;
    if (cutGetCmdLineArgumentstr( argc, (const char**) argv, "file", &filename)) {
        volumeFilename = filename;
    }

    gridSize = make_uint3(1<<gridSizeLog2.x, 1<<gridSizeLog2.y, 1<<gridSizeLog2.z);
    gridSizeMask = make_uint3(gridSize.x-1, gridSize.y-1, gridSize.z-1);
    gridSizeShift = make_uint3(0, gridSizeLog2.x, gridSizeLog2.x+gridSizeLog2.y);

    numVoxels = gridSize.x*gridSize.y*gridSize.z;
    voxelSize = make_float3(2.0f / gridSize.x, 2.0f / gridSize.y, 2.0f / gridSize.z);
    maxVerts = gridSize.x*gridSize.y*100;

    printf("grid: %d x %d x %d = %d voxels\n", gridSize.x, gridSize.y, gridSize.z, numVoxels);
    printf("max verts = %d\n", maxVerts);

#if SAMPLE_VOLUME
    // load volume data
    char* path = cutFindFilePath(volumeFilename, argv[0]);
    if (path == 0) {
        fprintf(stderr, "Error finding file '%s'\n", volumeFilename);
        cudaThreadExit();
        exit(EXIT_FAILURE);
    }

    int size = gridSize.x*gridSize.y*gridSize.z*sizeof(uchar);
    uchar *volume = loadRawFile(path, size);
    cutilSafeCall(cudaMalloc((void**) &d_volume, size));
    cutilSafeCall(cudaMemcpy(d_volume, volume, size, cudaMemcpyHostToDevice) );
    free(volume);

	bindVolumeTexture(d_volume);
#endif

    if (g_bQAReadback) {
        cudaMalloc((void **)&(d_pos),    maxVerts*sizeof(float)*4);
        cudaMalloc((void **)&(d_normal), maxVerts*sizeof(float)*4);
    } else {
        // create VBOs
        createVBO(&posVbo, maxVerts*sizeof(float)*4);
		// DEPRECATED: cutilSafeCall( cudaGLRegisterBufferObject(posVbo) );
		cutilSafeCall(cudaGraphicsGLRegisterBuffer(&cuda_posvbo_resource, posVbo, 
							   cudaGraphicsMapFlagsWriteDiscard));

        createVBO(&normalVbo, maxVerts*sizeof(float)*4);
        // DEPRECATED: cutilSafeCall(cudaGLRegisterBufferObject(normalVbo));
		cutilSafeCall(cudaGraphicsGLRegisterBuffer(&cuda_normalvbo_resource, normalVbo, 
							   cudaGraphicsMapFlagsWriteDiscard));

    }

    // allocate textures
	allocateTextures(	&d_edgeTable, &d_triTable, &d_numVertsTable );

    // allocate device memory
    unsigned int memSize = sizeof(uint) * numVoxels;
    cutilSafeCall(cudaMalloc((void**) &d_voxelVerts,            memSize));
    cutilSafeCall(cudaMalloc((void**) &d_voxelVertsScan,        memSize));
    cutilSafeCall(cudaMalloc((void**) &d_voxelOccupied,         memSize));
    cutilSafeCall(cudaMalloc((void**) &d_voxelOccupiedScan,     memSize));
    cutilSafeCall(cudaMalloc((void**) &d_compVoxelArray,   memSize));

    // initialize CUDPP scan
    CUDPPConfiguration config;
    config.algorithm    = CUDPP_SCAN;
    config.datatype     = CUDPP_UINT;
    config.op           = CUDPP_ADD;
    config.options      = CUDPP_OPTION_FORWARD | CUDPP_OPTION_EXCLUSIVE;
    cudppPlan(&scanplan, config, numVoxels, 1, 0);
}