|
|
@ -25,35 +25,24 @@ __kernel void min_kern(
|
|
|
|
__write_only image2d_t image
|
|
|
|
__write_only image2d_t image
|
|
|
|
){
|
|
|
|
){
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Get the pixel position of this worker
|
|
|
|
size_t id = get_global_id(0);
|
|
|
|
size_t id = get_global_id(0);
|
|
|
|
int2 pixel = {id % resolution->x, id / resolution->x};
|
|
|
|
int2 pixel = {id % resolution->x, id / resolution->x};
|
|
|
|
//int2 pixel = {1, 1};
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Slew the ray into it's correct position based on the view matrix's starting position
|
|
|
|
|
|
|
|
// and the camera's current direction
|
|
|
|
|
|
|
|
|
|
|
|
float3 ray_dir = projection_matrix[pixel.x + resolution->x * pixel.y];
|
|
|
|
float3 ray_dir = projection_matrix[pixel.x + resolution->x * pixel.y];
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Yaw
|
|
|
|
ray_dir = (float3)(
|
|
|
|
ray_dir = (float3)(
|
|
|
|
ray_dir.z * sin(cam_dir->y) + ray_dir.x * cos(cam_dir->y),
|
|
|
|
ray_dir.z * sin(cam_dir->y) + ray_dir.x * cos(cam_dir->y),
|
|
|
|
ray_dir.y,
|
|
|
|
ray_dir.y,
|
|
|
|
ray_dir.z * cos(cam_dir->y) - ray_dir.x * sin(cam_dir->y)
|
|
|
|
ray_dir.z * cos(cam_dir->y) - ray_dir.x * sin(cam_dir->y)
|
|
|
|
);
|
|
|
|
);
|
|
|
|
//
|
|
|
|
|
|
|
|
// float a = cam_dir->x;
|
|
|
|
|
|
|
|
// float b = cam_dir->y;
|
|
|
|
|
|
|
|
// float c = cam_dir->z;
|
|
|
|
|
|
|
|
//
|
|
|
|
|
|
|
|
// ray_dir.x = ray_dir.z * sin(b) + ray_dir.x * cos(b);
|
|
|
|
|
|
|
|
// ray_dir.y = ray_dir.y;
|
|
|
|
|
|
|
|
// ray_dir.z = ray_dir.z * cos(b) - ray_dir.x * sin(b);
|
|
|
|
|
|
|
|
//
|
|
|
|
|
|
|
|
//
|
|
|
|
|
|
|
|
// float3 ray_dir2 = (float3)(
|
|
|
|
|
|
|
|
// ray_dir.x * cos(c) - ray_dir.y * sin(c),
|
|
|
|
|
|
|
|
// ray_dir.x * sin(c) + ray_dir.y * cos(c),
|
|
|
|
|
|
|
|
// ray_dir.z);
|
|
|
|
|
|
|
|
//
|
|
|
|
|
|
|
|
// printf("%f, %f, %f", ray_dir2.x, ray_dir2.y, ray_dir2.z);
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Pitch
|
|
|
|
ray_dir = (float3)(
|
|
|
|
ray_dir = (float3)(
|
|
|
|
ray_dir.x * cos(cam_dir->z) - ray_dir.y * sin(cam_dir->z),
|
|
|
|
ray_dir.x * cos(cam_dir->z) - ray_dir.y * sin(cam_dir->z),
|
|
|
|
ray_dir.x * sin(cam_dir->z) + ray_dir.y * cos(cam_dir->z),
|
|
|
|
ray_dir.x * sin(cam_dir->z) + ray_dir.y * cos(cam_dir->z),
|
|
|
@ -64,10 +53,6 @@ __kernel void min_kern(
|
|
|
|
int3 voxel_step = {1, 1, 1};
|
|
|
|
int3 voxel_step = {1, 1, 1};
|
|
|
|
voxel_step *= (ray_dir > 0) - (ray_dir < 0);
|
|
|
|
voxel_step *= (ray_dir > 0) - (ray_dir < 0);
|
|
|
|
|
|
|
|
|
|
|
|
/*voxel_step.x *= (ray_dir.x > 0) - (ray_dir.x < 0);
|
|
|
|
|
|
|
|
voxel_step.y *= (ray_dir.y > 0) - (ray_dir.y < 0);
|
|
|
|
|
|
|
|
voxel_step.z *= (ray_dir.z > 0) - (ray_dir.z < 0);*/
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Setup the voxel coords from the camera origin
|
|
|
|
// Setup the voxel coords from the camera origin
|
|
|
|
int3 voxel = convert_int3(*cam_pos);
|
|
|
|
int3 voxel = convert_int3(*cam_pos);
|
|
|
|
|
|
|
|
|
|
|
@ -79,21 +64,26 @@ __kernel void min_kern(
|
|
|
|
// for all 3 axis XYZ.
|
|
|
|
// for all 3 axis XYZ.
|
|
|
|
float3 intersection_t = delta_t;
|
|
|
|
float3 intersection_t = delta_t;
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Create a psuedo random number for view fog
|
|
|
|
int2 randoms = { 3, 14 };
|
|
|
|
int2 randoms = { 3, 14 };
|
|
|
|
uint seed = randoms.x + id;
|
|
|
|
uint seed = randoms.x + id;
|
|
|
|
uint t = seed ^ (seed << 11);
|
|
|
|
uint t = seed ^ (seed << 11);
|
|
|
|
uint result = randoms.y ^ (randoms.y >> 19) ^ (t ^ (t >> 8));
|
|
|
|
uint result = randoms.y ^ (randoms.y >> 19) ^ (t ^ (t >> 8));
|
|
|
|
|
|
|
|
|
|
|
|
int max_dist = 500 + result % 50;
|
|
|
|
// Distance a ray can travel before it terminates
|
|
|
|
|
|
|
|
int max_dist = 200 + result % 50;
|
|
|
|
int dist = 0;
|
|
|
|
int dist = 0;
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Bitmask to keep track of which axis was tripped
|
|
|
|
int3 mask = { 0, 0, 0 };
|
|
|
|
int3 mask = { 0, 0, 0 };
|
|
|
|
|
|
|
|
|
|
|
|
// Andrew Woo's raycasting algo
|
|
|
|
// Andrew Woo's raycasting algo
|
|
|
|
do {
|
|
|
|
do {
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// Non-branching test of the lowest delta_t value
|
|
|
|
mask = intersection_t.xyz <= min(intersection_t.yzx, intersection_t.zxy);
|
|
|
|
mask = intersection_t.xyz <= min(intersection_t.yzx, intersection_t.zxy);
|
|
|
|
float3 thing = delta_t * fabs(convert_float3(mask.xyz));
|
|
|
|
|
|
|
|
|
|
|
|
// Based on the result increment the voxel and intersection
|
|
|
|
intersection_t += delta_t * fabs(convert_float3(mask.xyz));
|
|
|
|
intersection_t += delta_t * fabs(convert_float3(mask.xyz));
|
|
|
|
voxel.xyz += voxel_step.xyz * mask.xyz;
|
|
|
|
voxel.xyz += voxel_step.xyz * mask.xyz;
|
|
|
|
|
|
|
|
|
|
|
@ -101,12 +91,15 @@ __kernel void min_kern(
|
|
|
|
int3 overshoot = voxel <= *map_dim;
|
|
|
|
int3 overshoot = voxel <= *map_dim;
|
|
|
|
int3 undershoot = voxel > 0;
|
|
|
|
int3 undershoot = voxel > 0;
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// "Sky"
|
|
|
|
if (overshoot.x == 0 || overshoot.y == 0 || overshoot.z == 0 || undershoot.x == 0 || undershoot.y == 0){
|
|
|
|
if (overshoot.x == 0 || overshoot.y == 0 || overshoot.z == 0 || undershoot.x == 0 || undershoot.y == 0){
|
|
|
|
write_imageui(image, pixel, (uint4)(50, 50, 50, 255));
|
|
|
|
write_imageui(image, pixel, (uint4)(135, 206, 235, 255));
|
|
|
|
return;
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
// "Water"
|
|
|
|
if (undershoot.z == 0) {
|
|
|
|
if (undershoot.z == 0) {
|
|
|
|
write_imageui(image, pixel, (uint4)(14, 30, 50, 255));
|
|
|
|
write_imageui(image, pixel, (uint4)(64, 164, 223, 255));
|
|
|
|
return;
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
@ -130,7 +123,7 @@ __kernel void min_kern(
|
|
|
|
return;
|
|
|
|
return;
|
|
|
|
case 5:
|
|
|
|
case 5:
|
|
|
|
//write_imageui(image, pixel, (uint4)(200, 200, 200, 255));
|
|
|
|
//write_imageui(image, pixel, (uint4)(200, 200, 200, 255));
|
|
|
|
write_imageui(image, pixel, white_light((uint4)(225, 232, 214, 100), (float3)(lights[7], lights[8], lights[9]), mask));
|
|
|
|
write_imageui(image, pixel, white_light((uint4)(44, 176, 55, 100), (float3)(lights[7], lights[8], lights[9]), mask));
|
|
|
|
return;
|
|
|
|
return;
|
|
|
|
case 6:
|
|
|
|
case 6:
|
|
|
|
write_imageui(image, pixel, (uint4)(30, 80, 10, 255));
|
|
|
|
write_imageui(image, pixel, (uint4)(30, 80, 10, 255));
|
|
|
@ -141,6 +134,6 @@ __kernel void min_kern(
|
|
|
|
dist++;
|
|
|
|
dist++;
|
|
|
|
} while (dist < max_dist);
|
|
|
|
} while (dist < max_dist);
|
|
|
|
|
|
|
|
|
|
|
|
write_imageui(image, pixel, (uint4)(73, 81, 89, 255));
|
|
|
|
write_imageui(image, pixel, (uint4)(135, 206, 235, 255));
|
|
|
|
return;
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
}
|