Skip to content

Commit

Permalink
collision checking works for a triangle in opencl and cpu
Browse files Browse the repository at this point in the history
  • Loading branch information
@enjalot
enjalot committed Jul 29, 2010
1 parent d72ef74 commit 28f0fa7
Show file tree
Hide file tree
Showing 6 changed files with 264 additions and 233 deletions.
10 changes: 9 additions & 1 deletion enjacl/enja.cpp
View file
Expand Up @@ -229,6 +229,7 @@ EnjaParticles::~EnjaParticles()
if(ckKernel)clReleaseKernel(ckKernel);
if(cpProgram)clReleaseProgram(cpProgram);
if(cqCommandQueue)clReleaseCommandQueue(cqCommandQueue);

if(v_vbo)
{
glBindBuffer(1, v_vbo);
Expand All @@ -248,18 +249,25 @@ EnjaParticles::~EnjaParticles()
i_vbo = 0;
}

//if(vbo_cl)clReleaseMemObject(vbo_cl);
printf("seg fault 1?\n");
if(cl_vbos[0])clReleaseMemObject(cl_vbos[0]);
if(cl_vbos[1])clReleaseMemObject(cl_vbos[1]);
if(cl_vbos[2])clReleaseMemObject(cl_vbos[2]);
printf("seg fault 1.5?\n");
//why are these arrays segfaulting when released?
if(cl_vert_gen)clReleaseMemObject(cl_vert_gen);
if(cl_velo_gen)clReleaseMemObject(cl_velo_gen);
if(cl_velocities)clReleaseMemObject(cl_velocities);
printf("seg fault 1.6?\n");
if(cl_indices)clReleaseMemObject(cl_indices);
printf("seg fault 1.7?\n");
//if(cl_life)clReleaseMemObject(cl_life);
if(cxGPUContext)clReleaseContext(cxGPUContext);

printf("seg fault 1.8?\n");
if(cdDevices)delete(cdDevices);

printf("seg fault 2?\n");
}


Expand Down
1 change: 1 addition & 0 deletions enjacl/main.cpp
View file
Expand Up @@ -184,6 +184,7 @@ void appDestroy()

delete enjas;
if(glutWindowHandle)glutDestroyWindow(glutWindowHandle);
printf("about to exit!\n");

exit(0);
}
Expand Down
8 changes: 4 additions & 4 deletions enjacl/opencl.cpp
View file
Expand Up @@ -119,17 +119,17 @@ void EnjaParticles::popCorn()
//support arrays for the particle system
cl_vert_gen = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, vbo_size, NULL, &ciErrNum);
cl_velo_gen = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, vbo_size, NULL, &ciErrNum);
cl_velo_gen= clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, vbo_size, NULL, &ciErrNum);
//cl_life = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(float) * num, NULL, &ciErrNum);
cl_velo_gen = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, vbo_size, NULL, &ciErrNum);
cl_indices = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(int) * num, NULL, &ciErrNum);

ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cl_vert_gen, CL_TRUE, 0, vbo_size, &vert_gen[0], 0, NULL, &evt);
clReleaseEvent(evt);
ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cl_velo_gen, CL_TRUE, 0, vbo_size, &velo_gen[0], 0, NULL, &evt);
clReleaseEvent(evt);
ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cl_velo_gen, CL_TRUE, 0, vbo_size, &velo_gen[0], 0, NULL, &evt);
clReleaseEvent(evt);
//ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cl_life, CL_TRUE, 0, sizeof(float) * num, life, 0, NULL, &evt);
//clReleaseEvent(evt);
ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cl_indices, CL_TRUE, 0, sizeof(int) * num, &indices[0], 0, NULL, &evt);
clReleaseEvent(evt);
clFinish(cqCommandQueue);


Expand Down
197 changes: 89 additions & 108 deletions enjacl/physics/collision.cl
View file
Expand Up @@ -4,6 +4,47 @@ float4 cross_product(float4 a, float4 b)
return (float4)(a.y*b.z - a.z*b.y, a.z*b.x - a.x*b.z, a.x*b.y - a.y*b.x, 0);
}

//Moller and Trumbore
bool intersect_triangle(float4 pos, float4 vel, float4 tri[3], float4 triN, float dist)
{
//take in the particle position and velocity (treated as a Ray)
//also the triangle vertices for the ray intersection
//we take in the precalculated triangle's normal to first test for distance
//dist is the threshold to determine if we are close enough to the triangle
//to even check for distance
float4 edge1, edge2, tvec, pvec, qvec;
float det, inv_det, u, v;
float eps = .000001;

//check distance
tvec = pos - tri[0];
float distance = -dot(tvec, triN) / dot(vel, triN);
if (distance > dist)
return false;


edge1 = tri[1] - tri[0];
edge2 = tri[2] - tri[0];

pvec = cross(vel, edge2);
det = dot(edge1, pvec);
//culling branch
//if(det > -eps && det < eps)
if(det < eps)
return false;

u = dot(tvec, pvec);
if (u < 0.0 || u > det)//1.0)
return false;

qvec = cross(tvec, edge1);
v = dot(vel, qvec);
if (v < 0.0 || u + v > det)//1.0f)
return false;

return true;
}



//update the particle position and color
Expand All @@ -12,8 +53,8 @@ __kernel void enja(__global float4* vertices, __global float4* colors, __global
unsigned int i = get_global_id(0);

float life = velocities[i].w;
life -= h; //should probably depend on time somehow
h = h*10;
life -= h/10; //should probably depend on time somehow
//h = h*10;
if(life <= 0.)
{
//reset this particle
Expand All @@ -26,125 +67,64 @@ __kernel void enja(__global float4* vertices, __global float4* colors, __global
velocities[i].z = velo_gen[i].z;
life = 1.;
}
float xn = vertices[i].x;
float yn = vertices[i].y;
float zn = vertices[i].z;


float vxn = velocities[i].x;
float vyn = velocities[i].y;
float vzn = velocities[i].z;
velocities[i].x = vxn;
velocities[i].y = vyn - h*9.8;
velocities[i].z = vzn;// - h*9.8;

xn += h*velocities[i].x;
yn += h*velocities[i].y;
zn += h*velocities[i].z;
float4 pos = vertices[i];
float4 vel = velocities[i];

float xn = pos.x;
float yn = pos.y;
float zn = pos.z;

float vxn = vel.x;
float vyn = vel.y;
float vzn = vel.z;
vel.x = vxn;
vel.y = vyn - h*9.8;
vel.z = vzn;// - h*9.8;

xn += h*vel.x;
yn += h*vel.y;
zn += h*vel.z;


//set up test plane
float4 plane[4];
plane[0] = (float4)(0,-1,0,0);
plane[1] = (float4)(0,-1,2,0);
plane[0] = (float4)(-2,-1,-2,0);
plane[1] = (float4)(-2,-1,2,0);
plane[2] = (float4)(2,-1,2,0);
plane[3] = (float4)(2,-1,0,0);
plane[3] = (float4)(2,-1,-2,0);

//triangle fan from plane (for handling faces)
float4 tri1[3];
tri1[0] = plane[0];
tri1[1] = plane[1];
tri1[2] = plane[2];

/* do 1 triangle first
float4 tri2[3];
tr2[0] = plane[0];
tr2[1] = plane[2];
tr2[2] = plane[3];
*/
float4 tri[3];
tri[0] = plane[0];
tri[1] = plane[1];
tri[2] = plane[2];

//calculate the normal of the triangle
//might not need to do this if we just have plane's normal
float4 A = tri1[0];
float4 B = tri1[1];
float4 C = tri1[2];
float4 A = tri[0];
float4 B = tri[1];
float4 C = tri[2];

//float4 tri1N = cross_product(B - A, C - A);
float4 tri1N = (float4)(0.0, 1.0, 0.0, 0.0);
//calculate the distnace of pos from triangle
float distance = -dot(vertices[i] - A, tri1N) / dot(velocities[i], tri1N);
float4 P = vertices[i] + distance*velocities[i];
if (distance <= 0.0f)
{
//particle is past the plane so don't do anything
}
else
float4 triN = normalize(cross_product(B - A, C - A));
//float4 tri1N = (float4)(0.0, 1.0, 0.0, 0.0);

if(intersect_triangle(pos, vel, tri, triN, h))
{
int x = 0;
int y = 2;
//these should be projections...
//or at least calculated from the dominant axis of the normal
float2 Ap = (float2)(A.x, A.z);
float2 Bp = (float2)(A.x, A.z);
float2 Cp = (float2)(A.x, A.z);
float2 Pp = (float2)(A.x, A.z);

float2 b = (Bp - Ap);
float2 c = (Cp - Ap);
float2 p = (Pp - Ap);

float u = (p.y*c.x - p.x*c.y)/(b.y*c.x - b.x*c.y);
float v = (p.y*b.x - p.x*b.y)/(c.y*b.x - c.x*b.y);

if(u >= 0 and v >= 0 and u+v <= 1)
{

//particle and triangle collision
//if (yn < -1.0f)

float4 posA = vertices[i];
float radiusA = 5.0f;
float4 posB = vertices[i] + radiusA/2.0f;
float4 relPos = (float4)(posB.x - posA.x, posB.y - posA.y, posB.z - posA.z, 0);
float dist = sqrt(relPos.x * relPos.x + relPos.y * relPos.y + relPos.z * relPos.z);
float collideDist = radiusA + 0.0;//radiusB;


//float4 norm = (-0.707106, 0.707106, 0.0, 0.0); //this is actually a unit vector
float4 norm = (float4)(0.0, 1.0, 0.0, 0.0);
norm = normalize(norm);
float4 velA = velocities[i]; //velocity of particle
float4 velB = (float4)(0,0,0,0); //velocity of object
float4 relVel = (float4)(velB.x - velA.x, velB.y - velA.y, velB.z - velA.z, 0);

float relVelDotNorm = relVel.x * norm.x + relVel.y * norm.y + relVel.z * norm.z;
float4 tanVel = (float4)(relVel.x - relVelDotNorm * norm.x, relVel.y - relVelDotNorm * norm.y, relVel.z - relVelDotNorm * norm.z, 0);
float4 force = (float4)(0,0,0,0);
float springFactor = -.5;//-spring * (collideDist - dist);
float damping = 1.0f;
float shear = 1.0f;
float attraction = 1.0f;
force = (float4)(
springFactor * norm.x + damping * relVel.x + shear * tanVel.x + attraction * relPos.x,
springFactor * norm.y + damping * relVel.y + shear * tanVel.y + attraction * relPos.y,
springFactor * norm.z + damping * relVel.z + shear * tanVel.z + attraction * relPos.z,
0
);

//float mag = sqrt(vel.x*vel.x + vel.y*vel.y + vel.z*vel.z); //store the magnitude of the velocity
//vel /= mag;
//vel = 2.f*(dot(normal, vel))*normal - vel;
////vel *= mag; //we know the direction lets go the right speed

velA += force;

xn = vertices[i].x + h*velA.x;
yn = vertices[i].y + h*velA.y;
zn = vertices[i].z + h*velA.z;
velocities[i] = velA;
}
}
//lets do some specular reflection
float mag = sqrt(vel.x*vel.x + vel.y*vel.y + vel.z*vel.z); //store the magnitude of the velocity
float4 nvel = normalize(vel);
float s = 2.0f*(dot(triN, nvel));
float4 dir = s * triN - nvel; //new direction
float damping = .5f;
mag *= damping;
vel = -mag * dir;

xn = pos.x + h*vel.x;
yn = pos.y + h*vel.y;
zn = pos.z + h*vel.z;

}

vertices[i].x = xn;
vertices[i].y = yn;
vertices[i].z = zn;
Expand All @@ -156,6 +136,7 @@ __kernel void enja(__global float4* vertices, __global float4* colors, __global
//colors[i].w = 1-life;
colors[i].w = 1;

velocities[i] = vel;
//save the life!
velocities[i].w = life;
}
Expand Down

0 comments on commit 28f0fa7

Please sign in to comment.