PotentialFieldSolver::PotentialFieldSolver()
{
m_M_eval=0;
m_N_mass=0;
m_cell_h=0;
m_gridx=m_gridy=m_gridz=0;
m_hashx=m_hashy=m_hashz=0;
m_initialized=false;
m_origin=make_double4(0,0,0,0);
m_L=0;
m_K=0;
m_center=make_double3(0,0,0);
m_total_mass=0;
m_evalPos = new GpuArrayf4;
m_evalNormal = new GpuArrayd3;
m_evalPos_Reorder = new GpuArrayf4;
m_p_massPos = new GpuArrayf4;
m_p_massPos_Reorder = new GpuArrayf4;
m_particle_mass = new GpuArrayd;
m_particle_mass_Reorder = new GpuArrayd;
m_grid_density = new GpuArrayd;
m_grid_Rhs = new GpuArrayd;
m_grid_phi = new GpuArrayd;
m_particle_gradPhi = new GpuArrayd3;
m_particle_dphidn = new GpuArrayd;
m_particle_gradPhi_deorder = new GpuArrayd3;
m_grid_gradPhi = new GpuArrayd3;
m_far_gradPhi = new GpuArrayd3;
m_SLP_area = new GpuArrayd;
}
BiotSavartSolver::BiotSavartSolver()
{
m_N_vort = 0;
m_M_eval = 0;
m_gridx = m_gridy = m_gridz = 0;
m_hashx = m_hashy = m_hashz = 0;
m_cell_h = 0.0;
m_origin = make_double4(0,0,0,0);
m_center = make_double3(0,0,0);
m_L = 0.0;
m_initialized = false;
m_K = 0;
m_evalPos = new gf_GpuArray<float4>;
m_evalPos_Reorder = new gf_GpuArray<float4>;
m_p_vortPos = new gf_GpuArray<float4>;
m_p_vortPos_Reorder = new GpuArrayf4;
m_isVIC = false;
for (int i=0;i<NUM_COMPONENTS; i++)
{
m_total_vort[i]=0;
m_grid_Rhs[i] = new GpuArrayd;
m_particle_vort[i] = new GpuArrayd;
m_particle_vort_Reorder[i] = new GpuArrayd;
m_grid_vort[i] = new GpuArrayd;
m_grid_Psi[i] = new GpuArrayd;
m_particle_U[i] = new GpuArrayd;
m_particle_U_deorder[i] = new GpuArrayd;
m_grid_U[i] = new GpuArrayd;
m_far_U[i] = new GpuArrayd;
}
}
bool
PotentialFieldSolver::m_ComputeGradient()
{
m_grid_gradPhi->memset(make_double3(0,0,0));
ComputeGradient(m_grid_phi->getDevicePtr(),m_grid_gradPhi->getDevicePtr(),
m_SpatialHasher_mass.getCellSize().x,1.0,m_gridx,m_gridy,m_gridz);
return true;
}
bool
PotentialFieldSolver::setEvalParameter(uint m, GpuArrayf4 * pos)
{
if(m!=m_M_eval)
{
m_M_eval = m;
m_SpatialHasher_eval.endSpatialHash();
m_SpatialHasher_eval.initSpatialHash(m_M_eval, m_gridx,m_gridy,m_gridz);
m_evalPos->free();
m_evalPos->alloc(m_M_eval);
m_evalPos->memset(make_float4(0,0,0,0));
m_evalPos_Reorder->free();
m_evalPos_Reorder->alloc(m_M_eval);
m_evalPos_Reorder->memset(make_float4(0,0,0,0));
m_particle_gradPhi->free();
m_particle_gradPhi->alloc(m_M_eval);
m_particle_gradPhi->memset(make_double3(0,0,0));
m_particle_gradPhi_deorder->free();
m_particle_gradPhi_deorder->alloc(m_M_eval);
m_particle_gradPhi_deorder->memset(make_double3(0,0,0));
}
if( setEvalPos(pos))
{
return true;
}
else
{
return false;
}
}
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__ double3 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z) {return make_double3(x,y,z);}
template<> inline __host__ __device__ double3 _pixMakeZero<double3>() {return make_double3(0.,0.,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;
}
bool
PotentialFieldSolver::initializeSolver(uint gdx, uint gdy, uint gdz, uint K, uint M, uint N)
{
m_K = K;
if(m_initialized)
{
if(gdx==m_gridx && gdy==m_gridy && gdz==m_gridz && M==m_M_eval && N==m_N_mass)
{
//zerofy memory
m_evalPos->memset(make_float4(0,0,0,0));
m_evalNormal->memset(make_double3(0,0,0));
m_evalPos_Reorder->memset(make_float4(0,0,0,0));
m_particle_gradPhi->memset(make_double3(0,0,0));
m_particle_dphidn->memset(0);
m_particle_gradPhi_deorder->memset(make_double3(0,0,0));
m_p_massPos->memset(make_float4(0,0,0,0));
m_p_massPos_Reorder->memset(make_float4(0,0,0,0));
m_particle_mass->memset(0);
m_particle_mass_Reorder->memset(0);
m_grid_density->memset(0);
m_grid_Rhs->memset(0);
m_grid_phi->memset(0);
m_grid_gradPhi->memset(make_double3(0,0,0));
m_far_gradPhi->memset(make_double3(0,0,0));
m_SLP_area->memset(0);
}
else //just reinitialize everything
{
m_gridx = gdx;
m_gridy = gdy;
m_gridz = gdz;
m_M_eval = M;
m_N_mass = N;
m_PoissonSolver.m_InitialSystem(m_gridx, m_gridy, m_gridz);
m_SpatialHasher_eval.endSpatialHash();
m_SpatialHasher_eval.initSpatialHash(m_M_eval, m_gridx,m_gridy,m_gridz);
m_SpatialHasher_mass.endSpatialHash();
m_SpatialHasher_mass.initSpatialHash(m_N_mass, m_gridx,m_gridy,m_gridz);
m_evalPos->free();
m_evalPos->alloc(m_M_eval);
m_evalPos->memset(make_float4(0,0,0,0));
m_evalPos_Reorder->free();
m_evalPos_Reorder->alloc(m_M_eval);
m_evalPos_Reorder->memset(make_float4(0,0,0,0));
m_evalNormal->free();
m_evalNormal->alloc(m_M_eval);
m_evalNormal->memset(make_double3(0,0,0));
m_particle_gradPhi->free();
m_particle_gradPhi->alloc(m_M_eval);
m_particle_gradPhi->memset(make_double3(0,0,0));
m_particle_gradPhi_deorder->free();
m_particle_gradPhi_deorder->alloc(m_M_eval);
m_particle_gradPhi_deorder->memset(make_double3(0,0,0));
m_particle_dphidn->free();
m_particle_dphidn->alloc(m_M_eval);
m_particle_dphidn->memset(0);
m_p_massPos->free();
m_p_massPos->alloc(m_N_mass);
m_p_massPos->memset(make_float4(0,0,0,0));
m_p_massPos_Reorder->free();
m_p_massPos_Reorder->alloc(m_N_mass);
m_p_massPos_Reorder->memset(make_float4(0,0,0,0));
m_particle_mass->free();
m_particle_mass->alloc(m_N_mass);
m_particle_mass->memset(0);
m_particle_mass_Reorder->free();
m_particle_mass_Reorder->alloc(m_N_mass);
m_particle_mass_Reorder->memset(0);
m_grid_density->free();
m_grid_density->alloc(m_gridx*m_gridy*m_gridz);
m_grid_density->memset(0);
m_grid_Rhs->free();
m_grid_Rhs->alloc(m_gridx*m_gridy*m_gridz);
m_grid_Rhs->memset(0);
m_grid_phi->free();
m_grid_phi->alloc(m_gridx*m_gridy*m_gridz);
m_grid_phi->memset(0);
m_grid_gradPhi->free();
m_grid_gradPhi->alloc(m_gridx*m_gridy*m_gridz);
m_grid_gradPhi->memset(make_double3(0,0,0));
m_far_gradPhi->free();
m_far_gradPhi->alloc(m_gridx*m_gridy*m_gridz);
m_far_gradPhi->memset(make_double3(0,0,0));
m_SLP_area->free();
m_SLP_area->alloc(m_M_eval);
m_SLP_area->memset(0);
}
}
else
{
m_gridx = gdx;
m_gridy = gdy;
m_gridz = gdz;
m_M_eval = M;
m_N_mass = N;
m_PoissonSolver.m_InitialSystem(m_gridx, m_gridy, m_gridz);
m_SpatialHasher_eval.initSpatialHash(m_M_eval, m_gridx,m_gridy,m_gridz);
m_SpatialHasher_mass.initSpatialHash(m_N_mass, m_gridx,m_gridy,m_gridz);
m_evalPos->alloc(m_M_eval);
m_evalPos->memset(make_float4(0,0,0,0));
m_evalPos_Reorder->alloc(m_M_eval);
m_evalPos_Reorder->memset(make_float4(0,0,0,0));
m_evalNormal->alloc(m_M_eval);
m_evalNormal->memset(make_double3(0,0,0));
m_particle_gradPhi->alloc(m_M_eval);
m_particle_gradPhi->memset(make_double3(0,0,0));
m_particle_gradPhi_deorder->alloc(m_M_eval);
m_particle_gradPhi_deorder->memset(make_double3(0,0,0));
m_particle_dphidn->alloc(m_M_eval);
m_particle_dphidn->memset(0);
m_p_massPos->alloc(m_N_mass);
m_p_massPos->memset(make_float4(0,0,0,0));
m_p_massPos_Reorder->alloc(m_N_mass);
m_p_massPos_Reorder->memset(make_float4(0,0,0,0));
m_particle_mass->alloc(m_N_mass);
m_particle_mass->memset(0);
m_particle_mass_Reorder->alloc(m_N_mass);
m_particle_mass_Reorder->memset(0);
m_grid_density->alloc(m_gridx*m_gridy*m_gridz);
m_grid_density->memset(0);
m_grid_Rhs->alloc(m_gridx*m_gridy*m_gridz);
m_grid_Rhs->memset(0);
m_grid_phi->alloc(m_gridx*m_gridy*m_gridz);
m_grid_phi->memset(0);
m_grid_gradPhi->alloc(m_gridx*m_gridy*m_gridz);
m_grid_gradPhi->memset(make_double3(0,0,0));
m_far_gradPhi->alloc(m_gridx*m_gridy*m_gridz);
m_far_gradPhi->memset(make_double3(0,0,0));
m_SLP_area->alloc(m_M_eval);
m_SLP_area->memset(0);
m_initialized = true;
}
return true;
}
