aaron/ispc
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

fixed cuda kernel

Browse Source
This commit is contained in:
Evghenii
2013-11-18 13:28:12 +01:00
parent 915dc4be7f
commit 4bc8c79bd3
1 changed files with 41 additions and 50 deletions
Download Patch File Download Diff File Expand all files Collapse all files

83
examples_cuda/rt/rt.cu
View File

@@ -58,17 +58,12 @@ struct Float3
} }
}; };
#if 0 #define int8 char
#define DIRISNEG #define int16 short
#endif
struct Ray { struct Ray {
float3 origin, dir, invDir; float3 origin, dir, invDir;
#ifdef DIRISNEG /* this fails to compile with nvvm */
unsigned int dirIsNeg[3];
#else
unsigned int dirIsNeg0, dirIsNeg1, dirIsNeg2; unsigned int dirIsNeg0, dirIsNeg1, dirIsNeg2;
#endif
float mint, maxt; float mint, maxt;
int hitId; int hitId;
}; };
@@ -79,8 +74,6 @@ struct Triangle {
int pad[3]; int pad[3];
}; };
#define int8 char
#define int16 short
struct LinearBVHNode { struct LinearBVHNode {
float bounds[2][3]; float bounds[2][3];
unsigned int offset; // num primitives for leaf, second child for interior unsigned int offset; // num primitives for leaf, second child for interior
@@ -106,7 +99,8 @@ static inline float Dot(const float3 a, const float3 b) {
} }
__device__ __device__
static inline void generateRay( const float raster2camera[4][4], inline
static void generateRay( const float raster2camera[4][4],
const float camera2world[4][4], const float camera2world[4][4],
float x, float y, Ray &ray) { float x, float y, Ray &ray) {
ray.mint = 0.f; ray.mint = 0.f;
@@ -136,7 +130,7 @@ static inline void generateRay( const float raster2camera[4][4],
ray.invDir = 1.f / ray.dir; ray.invDir = 1.f / ray.dir;
#ifdef DIRISNEG #if 0
ray.dirIsNeg[0] = any(ray.invDir.x < 0) ? 1 : 0; ray.dirIsNeg[0] = any(ray.invDir.x < 0) ? 1 : 0;
ray.dirIsNeg[1] = any(ray.invDir.y < 0) ? 1 : 0; ray.dirIsNeg[1] = any(ray.invDir.y < 0) ? 1 : 0;
ray.dirIsNeg[2] = any(ray.invDir.z < 0) ? 1 : 0; ray.dirIsNeg[2] = any(ray.invDir.z < 0) ? 1 : 0;
@@ -147,9 +141,9 @@ static inline void generateRay( const float raster2camera[4][4],
#endif #endif
} }
__device__ __device__
static inline bool BBoxIntersect(const float bounds[2][3], inline
static bool BBoxIntersect(const float bounds[2][3],
const Ray &ray) { const Ray &ray) {
float3 bounds0 = { bounds[0][0], bounds[0][1], bounds[0][2] }; float3 bounds0 = { bounds[0][0], bounds[0][1], bounds[0][2] };
float3 bounds1 = { bounds[1][0], bounds[1][1], bounds[1][2] }; float3 bounds1 = { bounds[1][0], bounds[1][1], bounds[1][2] };
@@ -187,9 +181,9 @@ static inline bool BBoxIntersect(const float bounds[2][3],
} }
__device__ __device__
static inline bool TriIntersect(const Triangle &tri, Ray &ray) { inline
static bool TriIntersect(const Triangle &tri, Ray &ray) {
float3 p0 = { tri.p[0][0], tri.p[0][1], tri.p[0][2] }; float3 p0 = { tri.p[0][0], tri.p[0][1], tri.p[0][2] };
float3 p1 = { tri.p[1][0], tri.p[1][1], tri.p[1][2] }; float3 p1 = { tri.p[1][0], tri.p[1][1], tri.p[1][2] };
float3 p2 = { tri.p[2][0], tri.p[2][1], tri.p[2][2] }; float3 p2 = { tri.p[2][0], tri.p[2][1], tri.p[2][2] };
@@ -228,20 +222,15 @@ static inline bool TriIntersect(const Triangle &tri, Ray &ray) {
return hit; return hit;
} }
__device__ __device__
static inline bool BVHIntersect(const LinearBVHNode nodes[], inline
const Triangle tris[], Ray &r) { bool BVHIntersect(const LinearBVHNode nodes[],
const Triangle tris[], Ray &r,
int todo[]) {
Ray ray = r; Ray ray = r;
bool hit = false; bool hit = false;
// Follow ray through BVH nodes to find primitive intersections // Follow ray through BVH nodes to find primitive intersections
int todoOffset = 0, nodeNum = 0; int todoOffset = 0, nodeNum = 0;
#if 0
__shared__ int todoX[64*4];
volatile int * todo = &todoX[warpIdx * 64];
#else
int todo[64];
#endif
while (true) { while (true) {
// Check ray against BVH node // Check ray against BVH node
@@ -255,20 +244,17 @@ static inline bool BVHIntersect(const LinearBVHNode nodes[],
if (TriIntersect(tris[primitivesOffset+i], ray)) if (TriIntersect(tris[primitivesOffset+i], ray))
hit = true; hit = true;
} }
if (todoOffset == 0)
break;
nodeNum = todo[--todoOffset]; nodeNum = todo[--todoOffset];
} }
else { else {
// Put far BVH node on _todo_ stack, advance to near node // Put far BVH node on _todo_ stack, advance to near node
#ifdef DIRISNEG
const int dirIsNeg = r.dirIsNeg[node.splitAxis];
#else
int dirIsNeg; int dirIsNeg;
if (node.splitAxis == 0) dirIsNeg = r.dirIsNeg0; if (node.splitAxis == 0) dirIsNeg = r.dirIsNeg0;
if (node.splitAxis == 1) dirIsNeg = r.dirIsNeg1; if (node.splitAxis == 1) dirIsNeg = r.dirIsNeg1;
if (node.splitAxis == 2) dirIsNeg = r.dirIsNeg2; if (node.splitAxis == 2) dirIsNeg = r.dirIsNeg2;
#endif if (dirIsNeg) {
if (dirIsNeg)
{
todo[todoOffset++] = nodeNum + 1; todo[todoOffset++] = nodeNum + 1;
nodeNum = node.offset; nodeNum = node.offset;
} }
@@ -279,10 +265,10 @@ static inline bool BVHIntersect(const LinearBVHNode nodes[],
} }
} }
else { else {
nodeNum = todo[--todoOffset];
}
if (todoOffset == 0) if (todoOffset == 0)
break; break;
nodeNum = todo[--todoOffset];
}
} }
r.maxt = ray.maxt; r.maxt = ray.maxt;
r.hitId = ray.hitId; r.hitId = ray.hitId;
@@ -290,9 +276,9 @@ static inline bool BVHIntersect(const LinearBVHNode nodes[],
return hit; return hit;
} }
__device__ __device__
static inline void raytrace_tile( int x0, int x1, inline
static void raytrace_tile( int x0, int x1,
int y0, int y1, int y0, int y1,
int width, int height, int width, int height,
int baseWidth, int baseHeight, int baseWidth, int baseHeight,
@@ -304,23 +290,30 @@ static inline void raytrace_tile( int x0, int x1,
float widthScale = (float)(baseWidth) / (float)(width); float widthScale = (float)(baseWidth) / (float)(width);
float heightScale = (float)(baseHeight) / (float)(height); float heightScale = (float)(baseHeight) / (float)(height);
// foreach_tiled (y = y0 ... y1, x = x0 ... x1) #if 0
for ( int y = y0; y < y1; y++) int * todo = new int[64];
for ( int xb = x0; xb < x1; xb += programCount) #define ALLOC
#else
int todo[64];
#endif
for (int y = y0 ;y < y1; y++)
for (int x = x0 + programIndex; x < x1; x += programCount)
if (x < x1)
{ {
const int x = xb + programIndex;
Ray ray; Ray ray;
generateRay(raster2camera, camera2world, x*widthScale, generateRay(raster2camera, camera2world, x*widthScale,
y*heightScale, ray); y*heightScale, ray);
BVHIntersect(nodes, triangles, ray); BVHIntersect(nodes, triangles, ray, todo);
int offset = y * width + x; int offset = y * width + x;
if (x < x1)
{
image[offset] = ray.maxt; image[offset] = ray.maxt;
id[offset] = ray.hitId; id[offset] = ray.hitId;
} }
}
#ifdef ALLOC
delete todo;
#endif
} }
@@ -333,8 +326,6 @@ void raytrace_tile_task( int width, int height,
float image[], int id[], float image[], int id[],
const LinearBVHNode nodes[], const LinearBVHNode nodes[],
const Triangle triangles[]) { const Triangle triangles[]) {
if (taskIndex >= taskCount) return;
int dx = 64, dy = 8; // must match dx, dy below int dx = 64, dy = 8; // must match dx, dy below
int xBuckets = (width + (dx-1)) / dx; int xBuckets = (width + (dx-1)) / dx;
int x0 = (taskIndex % xBuckets) * dx; int x0 = (taskIndex % xBuckets) * dx;
@@ -345,7 +336,6 @@ void raytrace_tile_task( int width, int height,
raytrace_tile(x0, x1, y0, y1, width, height, baseWidth, baseHeight, raytrace_tile(x0, x1, y0, y1, width, height, baseWidth, baseHeight,
raster2camera, camera2world, image, raster2camera, camera2world, image,
id, nodes, triangles); id, nodes, triangles);
cudaDeviceSynchronize();
} }
@@ -361,8 +351,9 @@ extern "C" __global__ void raytrace_ispc_tasks( int width, int height,
int yBuckets = (height + (dy-1)) / dy; int yBuckets = (height + (dy-1)) / dy;
int nTasks = xBuckets * yBuckets; int nTasks = xBuckets * yBuckets;
if (programIndex == 0) if (programIndex == 0)
raytrace_tile_task<<<(nTasks-1+4)/4,128>>>(width, height, baseWidth, baseHeight, raytrace_tile_task<<<(nTasks+4-1)/4,128>>>(width, height, baseWidth, baseHeight,
raster2camera, camera2world, raster2camera, camera2world,
image, id, nodes, triangles); image, id, nodes, triangles);
cudaDeviceSynchronize(); cudaDeviceSynchronize();
} }