Hello!
I've been toying with OpenGL and OpenCL trying to simulate Hair on the GPU.
However I've run into issues that are not OpenGL/OpenCL specific - I seem to be incapable of writing a correct solving algorithm for distance constraints.
I'm not sure where the issue is - maybe I'm generating the constraints incorrectly, maybe it's the solving.
That's how I'm generating the constraints:
[source lang="cpp"]// create constraints from original positions
__kernel void buildConstraints(__global const float4* in, __global DistanceConstraint* out, const int verticesPerHair)
{
int id = get_global_id(0);
if((id % verticesPerHair) != 0)
{
out[id].id1 = id-1;
out[id].id2 = id;
out[id].d = length(in[id-1]-in[id]);
}
else
{
out[id].id1 = id;
out[id].id2 = id;
out[id].d = 0;
}
}[/source]
(id % verticesPerHair) != 0 means that the hair is a root. All but the root is constrained to the preceding vertex in the hair. The root is constrained to itself so it won't move
Prior to constraining the vertices the verlet integration is computed on every vertex except the root:
[source lang="cpp"]
// verlet
if(idl != 0)
{
float4 difference = localpos[idl] - oldpos[idx];
oldpos[idx] = localpos[idl];
localpos[idl] += (1-friction)*difference + force*pdt*pdt;
}[/source]
Then to satisfy the constraints I run fixed number of iterations of this code:
[source lang="cpp"]barrier(CLK_LOCAL_MEM_FENCE);
DistanceConstraint constraint = constraints[idx];
for(int i = 0; i < dcIter; i++)
{
if((idl%2) == 1)
{
distanceConstraint(localpos, constraint.id1 % verticesPerHair, constraint.id2 % verticesPerHair, constraint.d);
}
barrier(CLK_LOCAL_MEM_FENCE);
if((idl%2) == 0)
{
distanceConstraint(localpos, constraint.id1 % verticesPerHair, constraint.id2 % verticesPerHair, constraint.d);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
[/source]
constraint is what was generated for this vertex in buildConstraints(). The barrier() function is there because the algorithm runs in parallel, so it solves the constraint for odd vertices then waits for all threads to complete and then solves the constraints for even vertices.
the function distanceConstraint() looks like this:
[source lang="cpp"]
// satisfy distance constraints
void distanceConstraint(__local HairVertex* pos, int idx1, int idx2, const float conDist)
{
// vector from vertex 1 to vertex 2
float4 difference = pos[idx2].position-pos[idx1].position;
// By what amount move which vertex
float2 cFactors = constraintFactors(idx1, idx2);
// length of the vector between the vertices
float len = length(difference);
// distance the vertices must move towards each other
float distance = len-conDist;
// no corruption through div by zero
len = max(len,1e-7f);
// move
pos[idx1].position += cFactors.x*distance/len*difference;
pos[idx2].position += cFactors.y*distance/len*difference;
}[/source]
constraintFactors() returns factors that determine how to handle vertices depending on whether or not they are roots:
[source lang="cpp"]
// for given indices id1 and id2 return constraint factors
float2 constraintFactors(int id1, int id2)
{
if(id1 != 0)
{
// first vertex is not a root
if(id2 != 0)
{
// both vertices not a root
return (float2)(0.5f, -0.5f);
}
else
{
// first vertex is not root, second is root
return (float2)(0.0f, 0.0f);
}
}
else
{
// first vertex is root
if(id2 != 0)
{
// first vertex is root, second is not
return (float2)(0.0f, -1.0f);
}
else
{
// both vertices are roots (shouldn't happen)
return (float2)(0.0f, 0.0f);
}
}
}[/source]
I have looked through NVIDIAs DirectX implementation of the simulation (NVIDIAs SDK 11) and the only deviation from the constraint code was that when the first vertex is not root and the second is, the factors were (1.0f,0.0f). I am however not certain how the constraints were generated, so I predict that my mistake is somewhere there.
The actual problem is that after a while the simulation goes full spaghetti mode... To be a bit more precise - the control vertices seem to move from the root towards the tip and the tip seems to move away in the direction the force is applied... Here's what it looks like in point mode:
http://sw-ores.de/pr...s/spaghetti.mkv
I have no clue what I'm doing wrong
