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

+1

Browse Source
This commit is contained in:
Evghenii
2013-11-18 13:11:57 +01:00
parent ee61a265f4
commit 4f9b8ebc73
1 changed files with 18 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
examples_cuda/rt/rt.cu
View File

@@ -319,7 +319,6 @@ static inline void raytrace_tile( int x0, int x1,
extern "C"
__global__ __global__
void raytrace_tile_task( int width, int height, void raytrace_tile_task( int width, int height,
int baseWidth, int baseHeight, int baseWidth, int baseHeight,
@@ -330,7 +329,7 @@ void raytrace_tile_task( int width, int height,
const Triangle triangles[]) { const Triangle triangles[]) {
if (taskIndex >= taskCount) return; if (taskIndex >= taskCount) return;
int dx = 32, dy = 16; // 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;
int x1 = min(x0 + dx, width); int x1 = min(x0 + dx, width);
@@ -343,3 +342,20 @@ void raytrace_tile_task( int width, int height,
} }
extern "C" __global__ void raytrace_ispc_tasks( int width, int height,
int baseWidth, int baseHeight,
const float raster2camera[4][4],
const float camera2world[4][4],
float image[], int id[],
const LinearBVHNode nodes[],
const Triangle triangles[]) {
int dx = 64, dy = 8;
int xBuckets = (width + (dx-1)) / dx;
int yBuckets = (height + (dy-1)) / dy;
int nTasks = xBuckets * yBuckets;
if (programIndex == 0)
raytrace_tile_task<<<(nTasks-1+4)/4,128>>>(width, height, baseWidth, baseHeight,
raster2camera, camera2world,
image, id, nodes, triangles);
cudaDeviceSynchronize();
}