I think, as an absolute CUDA novice, I may have taken on a bit much with this one - as such, please forgive the unbelievable number of good practice violations you’re almost certain to find staring at you (but please point them out anyway!) - however, if anyone has any idea why my simple iterative raycaster is doing the following, I’d be forever beholden. Essentially, when the ray skips a non-existing child and then impacts on one of its siblings, something’s going wrong with the intersection test somewhere along the line, so the ray is going straight through the nearest plane and (generally) impacting on the far plane. The top of my simple octree is complete, so there are no graphical glitches when viewing it from above - however, from the sides or bottom, there are planes missing left, right and centre. Pictures are attached.
EDIT: Got rid of previous rendering method, just using kd-restart for simplicity. Same problem.
Ray generation has been borrowed from elsewhere, as I was having an enormous mental block on how to actually do the damn thing.
// method to handle intersection with voxel, starts by intersecting with root
__device__ int intersectOctree(Ray traceRay, float3 octreeMin, float3 octreeMax, Voxel* cudaOctree, float3 translation)
{
Voxel* currentVoxel = cudaOctree + 1; //0 left empty for simplistic root - null pointer
// compute intersection of ray with all six voxel planes
float3 invRay = make_float3(1.0f) / traceRay.direction;
float3 bottomIntersection = invRay * (octreeMin - traceRay.origin);
float3 topIntersection = invRay * (octreeMax - traceRay.origin);
// re-order intersections to find smallest and largest on each axis
float3 minIntersection = fminf(topIntersection, bottomIntersection);
float3 maxIntersection = fmaxf(topIntersection, bottomIntersection);
// find the largest min and the smallest max
float largest_min = fmaxf(fmaxf(minIntersection.x, minIntersection.y), minIntersection.z);
float smallest_max = fminf(fminf(maxIntersection.x, maxIntersection.y), maxIntersection.z);
// if the ray doesn't hit then return
if (smallest_max < largest_min)
{
return 0;
}
float3 pos = traceRay.origin + traceRay.direction*largest_min;
float3 voxelMin = octreeMin;
float3 voxelMax = octreeMax;
int childIndex = 0;
int childID = 0;
float3 childMin;
float3 childMax;
float3 splitPlane;
float3 range;
while (!cudaIsLeaf(currentVoxel))
{
range = (voxelMax - voxelMin)/2;
splitPlane = voxelMin + range;
childID = 0;
childIndex = 0;
while (childIndex == 0)
{
childID = 0;
if (pos.x >= splitPlane.x)
{
childMin.x = splitPlane.x;
childMax.x = voxelMax.x;
childID += 4;
}
else
{
childMin.x = voxelMin.x;
childMax.x = splitPlane.x;
}
if (pos.y >= splitPlane.y)
{
childMin.y = splitPlane.y;
childMax.y = voxelMax.y;
childID += 2;
}
else
{
childMin.y = voxelMin.y;
childMax.y = splitPlane.y;
}
if (pos.z >= splitPlane.z)
{
childMin.z = splitPlane.z;
childMax.z = voxelMax.z;
childID += 1;
}
else
{
childMin.z = voxelMin.z;
childMax.z = splitPlane.z;
}
// find the next child to check based on the ray intersection
childIndex = cudaGetChildIndexByID(currentVoxel, childID);
if (childIndex==0) //no collision! step ray.
{
pos += traceRay.direction*range;
if (pos.x >= voxelMax.x || pos.y >= voxelMax.y || pos.z >= voxelMax.z ||
pos.x <= voxelMin.x || pos.y <= voxelMin.y || pos.z <= voxelMin.z) { break; } //If step takes you out of the voxel
}
else
{
currentVoxel = cudaOctree + childIndex;
if (cudaIsLeaf(currentVoxel))
{
return cudaGetColour(currentVoxel);
}
voxelMax = childMax;
voxelMin = childMin;
}
}
if (childIndex==0) //if ray does not intersect any children
{
if (currentVoxel->parentIndex == 0) //if there is no parent to move to
{
return 0;
}
else //if there are potential intersections with siblings, restart
{
voxelMin = octreeMin;
voxelMax = octreeMax;
currentVoxel = &cudaOctree[1];
}
}
}
return cudaGetColour(currentVoxel);
}
Apologies for the verbosity, I’m just a tad nervous as this is my first post and I’m aware that my code’s not the best, so I’m trying to not sound like an idiot as well as looking like one! I don’t want to overload you with code, but if there’s anything you need to know, just say and I’ll post it up ASAP!
Ed
EDIT: Forgot system specs. D’oh.
Basics:
Core 2 Quad 2.4GHz/core
3GB RAM
GTX 460
If you need to know any more, do say :)
Win7 Pro x64
