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);
}
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;
}
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);
}
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);
}
//! 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 );
}
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;
}
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);
}
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);
}
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;
}
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 );
}
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);
}
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;
}
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)
static inline __host__ __device__ uint3 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z) {return make_uint3(x,y,z);}
//! 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)) {};
//! Default constructor
VolumeArgGPU( void ) : dims(make_uint3(0,0,0)),
pitchedPtr(NULL),
dataPitch(0) {};
template<> inline __host__ __device__ uint3 _pixMakeZero<uint3>() {return make_uint3(0,0,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;
}
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;
}
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);
}
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;
}
//! Default constructor
VolumeGPU( void ) : dims(make_uint3(0,0,0)),
d_data(make_cudaPitchedPtr(NULL,0,0,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;
////////////////////////////////////////////////////////////////////////////////
// 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);
}
