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);