From cf607382a96fb07c988ff564883b7248c9ea8e33 Mon Sep 17 00:00:00 2001 From: mitchellhansen Date: Thu, 1 Sep 2016 22:45:48 -0700 Subject: [PATCH] WORKING! Awesome! It now casts fully inside the gpu, context is then switched to gl and then rendered via sfml. It has no loop, no controls, and the aspect ratio is off, but holy hell it works! --- include/CL_Wrapper.h | 2 +- include/Map.h | 6 ++ kernels/minimal_kernel.c | 156 +++++++++++++++++++++++++++++++++++++-- src/CL_Wrapper.cpp | 4 +- src/main.cpp | 21 ++++-- 5 files changed, 172 insertions(+), 17 deletions(-) diff --git a/include/CL_Wrapper.h b/include/CL_Wrapper.h index 8348536..7069b7a 100644 --- a/include/CL_Wrapper.h +++ b/include/CL_Wrapper.h @@ -39,7 +39,7 @@ public: int compile_kernel(std::string kernel_source, bool is_path, std::string kernel_name); int set_kernel_arg(std::string kernel_name, int index, std::string buffer_name); int store_buffer(cl_mem, std::string buffer_name); - int run_kernel(std::string kernel_name); + int run_kernel(std::string kernel_name, const int work_size); bool assert(int error_code, std::string function_name); diff --git a/include/Map.h b/include/Map.h index df3b246..8949512 100644 --- a/include/Map.h +++ b/include/Map.h @@ -34,6 +34,12 @@ public: char *list; sf::Vector3i dimensions; + void setVoxel(sf::Vector3i position, int val){ + + list[position.x + dimensions.x * (position.y + dimensions.z * position.z)] = val; + + }; + void moveLight(sf::Vector2f in); sf::Vector3f global_light; diff --git a/kernels/minimal_kernel.c b/kernels/minimal_kernel.c index 0e21f34..670e02c 100644 --- a/kernels/minimal_kernel.c +++ b/kernels/minimal_kernel.c @@ -10,12 +10,159 @@ __kernel void min_kern( size_t id = get_global_id(0); - float4 black = (float4)(.49, .68, .81, 1); + int2 pixel = {id % resolution->x, id / resolution->x}; - int2 pixelcoord = (int2) (id, id); - write_imagef(image, pixelcoord, black); + float3 ray_dir = projection_matrix[pixel.x + resolution->x * pixel.y]; + //printf("%i === %f, %f, %f\n", id, ray_dir.x, ray_dir.y, ray_dir.z); + // Y axis, pitch + //ray_dir.x = ray_dir.z * sin(cam_dir->y) + ray_dir.x * cos(cam_dir->y); + //ray_dir.y = ray_dir.y; + //ray_dir.z = ray_dir.z * cos(cam_dir->y) - ray_dir.x * sin(cam_dir->y); + + + ray_dir = (float3)( + ray_dir.z * sin(cam_dir->y) + ray_dir.x * cos(cam_dir->y), + ray_dir.y, + ray_dir.z * cos(cam_dir->y) - ray_dir.x * sin(cam_dir->y) + ); + + // Z axis, yaw + //ray_dir.x = ray_dir.x * cos(cam_dir->z) - ray_dir.y * sin(cam_dir->z); + //ray_dir.y = ray_dir.x * sin(cam_dir->z) + ray_dir.y * cos(cam_dir->z); + //ray_dir.z = ray_dir.z; + + ray_dir = (float3)( + 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.z + ); + + + // Setup the voxel step based on what direction the ray is pointing + int3 voxel_step = {1, 1, 1}; + 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 + int3 voxel = { + floorf(cam_pos->x), + floorf(cam_pos->y), + floorf(cam_pos->z) + }; + + // Delta T is the units a ray must travel along an axis in order to + // traverse an integer split + float3 delta_t = { + fabsf(1.0f / ray_dir.x), + fabsf(1.0f / ray_dir.y), + fabsf(1.0f / ray_dir.z) + }; + + // Intersection T is the collection of the next intersection points + // for all 3 axis XYZ. + float3 intersection_t = { + delta_t.x, + delta_t.y, + delta_t.z + }; + + + int dist = 0; + int face = -1; + // X:0, Y:1, Z:2 + + // Andrew Woo's raycasting algo + do { + if ((intersection_t.x) < (intersection_t.y)) { + if ((intersection_t.x) < (intersection_t.z)) { + + face = 0; + voxel.x += voxel_step.x; + intersection_t.x = intersection_t.x + delta_t.x; + } else { + + face = 2; + voxel.z += voxel_step.z; + intersection_t.z = intersection_t.z + delta_t.z; + } + } else { + if ((intersection_t.y) < (intersection_t.z)) { + + face = 1; + voxel.y += voxel_step.y; + intersection_t.y = intersection_t.y + delta_t.y; + } else { + + face = 2; + voxel.z += voxel_step.z; + intersection_t.z = intersection_t.z + delta_t.z; + } + } + + // If the ray went out of bounds + if (voxel.z >= map_dim->z) { + write_imagef(image, pixel, (float4)(.5, .50, .00, 1)); + return; + } + if (voxel.x >= map_dim->x) { + write_imagef(image, pixel, (float4)(.00, .00, .99, 1)); + return; + } + if (voxel.y >= map_dim->x) { + write_imagef(image, pixel, (float4)(.00, .99, .00, 1)); + return; + } + + if (voxel.x < 0) { + write_imagef(image, pixel, (float4)(.99, .00, .99, 1)); + return; + } + if (voxel.y < 0) { + write_imagef(image, pixel, (float4)(.99, .99, .00, 1)); + return; + } + if (voxel.z < 0) { + write_imagef(image, pixel, (float4)(.00, .99, .99, 1)); + return; + } + + // If we hit a voxel + int index = voxel.x + map_dim->x * (voxel.y + map_dim->z * voxel.z); + int voxel_data = map[index]; + + if (id == 100) + printf("%i, %i, %i\n", voxel.x, voxel.y, voxel.z); + + switch (voxel_data) { + case 1: + write_imagef(image, pixel, (float4)(.50, .00, .00, 1)); + return; + case 2: + write_imagef(image, pixel, (float4)(.00, .50, .00, 1.00)); + return; + case 3: + write_imagef(image, pixel, (float4)(.00, .00, .50, 1.00)); + return; + case 4: + write_imagef(image, pixel, (float4)(.25, .00, .25, 1.00)); + return; + case 5: + write_imagef(image, pixel, (float4)(.10, .30, .80, 1.00)); + return; + case 6: + write_imagef(image, pixel, (float4)(.30, .80, .10, 1.00)); + return; + } + + dist++; + } while (dist < 200); + + + write_imagef(image, pixel, (float4)(.00, .00, .00, .00)); + return; //printf("%i %i -- ", id, map[id]); //printf("%i, %i, %i\n", map_dim->x, map_dim->y, map_dim->z); @@ -23,7 +170,4 @@ __kernel void min_kern( //printf("%f, %f, %f\n", cam_dir->x, cam_dir->y, cam_dir->z); //printf("%f, %f, %f\n", cam_pos->x, cam_pos->y, cam_pos->z); - - - } \ No newline at end of file diff --git a/src/CL_Wrapper.cpp b/src/CL_Wrapper.cpp index 2427d85..4df21e5 100644 --- a/src/CL_Wrapper.cpp +++ b/src/CL_Wrapper.cpp @@ -230,10 +230,10 @@ int CL_Wrapper::store_buffer(cl_mem buffer, std::string buffer_name){ buffer_map.emplace(std::make_pair(buffer_name, buffer)); } -int CL_Wrapper::run_kernel(std::string kernel_name){ +int CL_Wrapper::run_kernel(std::string kernel_name, const int work_size){ const int WORKER_SIZE = 10; - size_t global_work_size[1] = { WORKER_SIZE }; + size_t global_work_size[1] = { static_cast(work_size) }; cl_kernel kernel = kernel_map.at(kernel_name); diff --git a/src/main.cpp b/src/main.cpp index 718cb4e..ca8f0eb 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -29,8 +29,8 @@ #include "RayCaster.h" #include "CL_Wrapper.h" -const int WINDOW_X = 150; -const int WINDOW_Y = 150; +const int WINDOW_X = 100; +const int WINDOW_Y = 100; @@ -76,6 +76,8 @@ int main() { sf::Vector3i map_dim(100, 100, 100); Map* map = new Map(map_dim); + map->setVoxel(sf::Vector3i(77, 50, 85), 5); + cl_mem map_buff = clCreateBuffer( c.getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(char) * map_dim.x * map_dim.y * map_dim.z, map->list, NULL @@ -129,6 +131,9 @@ int main() { } } + int ind = 367; + printf("%i === %f, %f, %f\n", ind, view_matrix[ind * 4 + 0], view_matrix[ind * 4 + 1], view_matrix[ind * 4 + 2]); + cl_mem view_matrix_buff = clCreateBuffer( c.getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * 3 * view_res.x * view_res.y, view_matrix, NULL @@ -142,7 +147,7 @@ int main() { ); - sf::Vector3f cam_pos(50, 50, 50); + sf::Vector3f cam_pos(55, 50, 50); cl_mem cam_pos_buff = clCreateBuffer( c.getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * 4, &cam_pos, NULL @@ -194,15 +199,18 @@ int main() { c.set_kernel_arg("min_kern", 5, "cam_pos_buffer"); c.set_kernel_arg("min_kern", 6, "image_buffer"); + const int size = 100 * 100; + c.run_kernel("min_kern", size); - c.run_kernel("min_kern"); + + clFinish(c.getCommandQueue()); error = clEnqueueReleaseGLObjects(c.getCommandQueue(), 1, &image_buff, 0, NULL, NULL); if (c.assert(error, "clEnqueueReleaseGLObjects")) return -1; s.setTexture(t); - + // The step size in milliseconds between calls to Update() // Lets set it to 16.6 milliseonds (60FPS) float step_size = 0.0166f; @@ -224,9 +232,6 @@ int main() { // State values - - - sf::Vector3f cam_vec(0, 0, 0); RayCaster ray_caster(map, map_dim, view_res);