From 73026bc65d84fe4ff6fc9cd21f0c0268095de6f2 Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Wed, 18 Jan 2017 22:06:33 -0800 Subject: [PATCH] Very very close to pixel perfect shadows. Having some of that axis switch malarchy slowing things down --- include/Hardware_Caster.h | 1 + include/util.hpp | 4 ++ kernels/ray_caster_kernel.cl | 123 ++++++++++++++--------------------- src/Camera.cpp | 2 +- src/Hardware_Caster.cpp | 16 ++++- src/Old_Map.cpp | 12 ++++ src/main.cpp | 16 +++-- 7 files changed, 93 insertions(+), 81 deletions(-) diff --git a/include/Hardware_Caster.h b/include/Hardware_Caster.h index 5a5c863..97b24cb 100644 --- a/include/Hardware_Caster.h +++ b/include/Hardware_Caster.h @@ -61,6 +61,7 @@ public: void draw(sf::RenderWindow* window) override; + int debug_quick_recompile(); void test_edit_viewport(int width, int height, float v_fov, float h_fov); private: diff --git a/include/util.hpp b/include/util.hpp index b8d1bbc..8c6bf9b 100644 --- a/include/util.hpp +++ b/include/util.hpp @@ -158,6 +158,10 @@ inline float AngleBetweenVectors(sf::Vector3f a, sf::Vector3f b){ return acos(DotProduct(a, b) / (Magnitude(a) * Magnitude(b))); } +inline float DistanceBetweenPoints(sf::Vector3f a, sf::Vector3f b) { + return sqrt(DotProduct(a, b)); +} + inline float DegreesToRadians(float in) { return static_cast(in * PI / 180.0f); } diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index 13cb6f4..235e255 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -47,7 +47,9 @@ int rand(int* seed) // 1 <= *seed < m return(*seed); } - +float DistanceBetweenPoints(float3 a, float3 b) { + return sqrt(pow(a.x - b.x, 2) + pow(a.y - b.y, 2) + pow(a.z - b.z, 2)); +} // =================================== Boolean ray intersection ============================ // ========================================================================================= @@ -62,6 +64,8 @@ bool cast_light_intersection_ray( ){ + float distance_to_light = DistanceBetweenPoints(ray_pos, (float3)(lights[4], lights[5], lights[6])); + // Setup the voxel step based on what direction the ray is pointing int3 voxel_step = { 1, 1, 1 }; voxel_step *= (ray_dir > 0) - (ray_dir < 0); @@ -80,29 +84,14 @@ bool cast_light_intersection_ray( // for all 3 axis XYZ. float3 intersection_t = delta_t *offset; - // for negative values, wrap around the delta_t, rather not do this - // component wise, but it doesn't appear to want to work - if (intersection_t.x < 0) { - intersection_t.x += delta_t.x; - } - if (intersection_t.y < 0) { - intersection_t.y += delta_t.y; - } - if (intersection_t.z < 0) { - intersection_t.z += delta_t.z; - } - - // Hard cut-off for how far the ray can travel - int max_dist = 800; - int dist = 0; + // for negative values, wrap around the delta_t + intersection_t += delta_t * -convert_float3(isless(intersection_t, 0)); int3 face_mask = { 0, 0, 0 }; // Andrew Woo's raycasting algo do { - - // Fancy no branch version of the logic step face_mask = intersection_t.xyz <= min(intersection_t.yzx, intersection_t.zxy); intersection_t += delta_t * fabs(convert_float3(face_mask.xyz)); @@ -112,10 +101,8 @@ bool cast_light_intersection_ray( int3 overshoot = voxel < *map_dim; int3 undershoot = voxel >= 0; - if (overshoot.x == 0 || overshoot.y == 0 || overshoot.z == 0 || undershoot.x == 0 || undershoot.y == 0) { - return false; - } - if (undershoot.z == 0) { + if (any(overshoot == (int3)(0, 0, 0)) || + any(undershoot == (int3)(0, 0, 0))) { return false; } @@ -123,14 +110,14 @@ bool cast_light_intersection_ray( int index = voxel.x + (*map_dim).x * (voxel.y + (*map_dim).z * (voxel.z)); int voxel_data = map[index]; - if (voxel_data != 0) { + if (voxel_data != 0) return true; - } - - dist++; - - } while (dist < 700); + //} while (any(isless(intersection_t, (float3)(distance_to_light - 1)))); + } while (intersection_t.x < distance_to_light - 1 || + intersection_t.y < distance_to_light - 1 || + intersection_t.z < distance_to_light - 1 ); + return false; } @@ -284,19 +271,11 @@ __kernel void raycaster( // Intersection T is the collection of the next intersection points // for all 3 axis XYZ. - float3 intersection_t = delta_t * (offset); + float3 intersection_t = delta_t * offset; - // for negative values, wrap around the delta_t, rather not do this - // component wise, but it doesn't appear to want to work - if (intersection_t.x < 0) { - intersection_t.x += delta_t.x; - } - if (intersection_t.y < 0) { - intersection_t.y += delta_t.y; - } - if (intersection_t.z < 0) { - intersection_t.z += delta_t.z; - } + // for negative values, wrap around the delta_t + intersection_t += delta_t * -convert_float3(isless(intersection_t, 0)); + // Hard cut-off for how far the ray can travel int max_dist = 800; @@ -350,64 +329,62 @@ __kernel void raycaster( } // set to which face - float3 face_position = convert_float3(face_mask * voxel_step); + float3 face_position = (float)(0); //convert_float3(-face_mask * voxel_step);// convert_float3(face_mask * voxel_step) * -1; - - - if (face_mask.x * voxel_step.x == -1) { - float z_percent = (intersection_t.x - (intersection_t.z - delta_t.z)) / delta_t.z; - float y_percent = (intersection_t.x - (intersection_t.y - delta_t.y)) / delta_t.y; + if (face_mask.x == -1) { - face_position = (float3)(1.0f, y_percent, z_percent); + //float z_percent = ((intersection_t.z - delta_t.z) - intersection_t.x) / delta_t.z; + //float y_percent = ((intersection_t.y - delta_t.y) - intersection_t.x) / delta_t.y; + + float z_percent = (intersection_t.z - (intersection_t.x - delta_t.x)) / delta_t.z; + float y_percent = (intersection_t.y - (intersection_t.x - delta_t.x)) / delta_t.y; - if (face_mask.x == 1) - face_position *= -1; + //if (z_percent > 0 && z_percent > 1) + // face_position = (float3)(-1.0f, 1-y_percent, 1-z_percent); + face_position = (float3)(1.0f, y_percent, z_percent); } + else if (face_mask.y == -1) { - if (face_mask.y * voxel_step.y == -1) { + //float x_percent = ((intersection_t.x - delta_t.x) - intersection_t.y) / delta_t.x; + //float z_percent = ((intersection_t.z - delta_t.z) - intersection_t.y) / delta_t.z; - float x_percent = (intersection_t.y - (intersection_t.x - delta_t.x)) / delta_t.x; - float z_percent = (intersection_t.y - (intersection_t.z - delta_t.z)) / delta_t.z; + float x_percent = (intersection_t.x - (intersection_t.y - delta_t.y)) / delta_t.x; + float z_percent = (intersection_t.z - (intersection_t.y - delta_t.y)) / delta_t.z; face_position = (float3)(x_percent, 1.0f, z_percent); - if (face_mask.y == 1) - face_position *= -1; } - - if (face_mask.z * voxel_step.z == -1) { - - float vx = intersection_t.x - delta_t.x; - float vy = intersection_t.y - delta_t.y; - float vz = intersection_t.z - delta_t.z; - float x_percent = (intersection_t.z - (intersection_t.x - delta_t.x)) / delta_t.x; - float y_percent = (intersection_t.z - (intersection_t.y - delta_t.y)) / delta_t.y; + else if (face_mask.z == -1) { - face_position = (float3)(x_percent, y_percent, 1.0f); + //float x_percent = ((intersection_t.x - delta_t.x) - intersection_t.z) / delta_t.x; + //float y_percent = ((intersection_t.y - delta_t.y) - intersection_t.z) / delta_t.y; - if (face_mask.z == 1) - face_position *= -1; + float x_percent = (intersection_t.x - (intersection_t.z - delta_t.z)) / delta_t.x; + float y_percent = (intersection_t.y - (intersection_t.z - delta_t.z)) / delta_t.y; + + face_position = (float3)(x_percent, y_percent, 1.0f); } - - // set the xy for that face - //face_position += - // convert_float3(face_mask == (int3)(1,1,1)) - // convert_float3(face_mask == (int3)(0,0,0)) * (intersection_t - delta_t); + if (ray_dir.x > 0) + face_position.x = - face_position.x - 1; + + if (ray_dir.y > 0) + face_position.y = - face_position.y - 1; + if (ray_dir.z > 0) + face_position.z = - face_position.z - 1; - //face_position += convert_float3(face_mask == (int3)(0,0,0)) * (rand(&seed) % 10) / 50.0; if (cast_light_intersection_ray( map, map_dim, - normalize((float3)(lights[4], lights[5], lights[6]) - (convert_float3(voxel))), - (convert_float3(voxel) + face_position),//convert_float3(face_mask * voxel_step)),//face_position),// + normalize((float3)(lights[4], lights[5], lights[6]) - (convert_float3(voxel) + face_position)), + (convert_float3(voxel) + face_position), lights, light_count )) { diff --git a/src/Camera.cpp b/src/Camera.cpp index 760d7ef..2c37f6c 100644 --- a/src/Camera.cpp +++ b/src/Camera.cpp @@ -94,7 +94,7 @@ void Camera::recieve_event(VrEventPublisher* publisher, std::unique_ptr(event.get()); if (held_event->code == sf::Keyboard::LShift) { - default_impulse = 0.2f; + default_impulse = 0.1f; } if (held_event->code == sf::Keyboard::RShift) { default_impulse = 1.0f; diff --git a/src/Hardware_Caster.cpp b/src/Hardware_Caster.cpp index 3aa7c7f..02e72d8 100644 --- a/src/Hardware_Caster.cpp +++ b/src/Hardware_Caster.cpp @@ -197,6 +197,17 @@ void Hardware_Caster::draw(sf::RenderWindow* window) { window->draw(viewport_sprite); } +int Hardware_Caster::debug_quick_recompile() +{ + int error = compile_kernel("../kernels/ray_caster_kernel.cl", true, "raycaster"); + if (assert(error, "compile_kernel")) { + std::cin.get(); // hang the output window so we can read the error + return error; + } + validate(); + +} + void Hardware_Caster::test_edit_viewport(int width, int height, float v_fov, float h_fov) { sf::Vector2i view_res(width, height); @@ -439,6 +450,7 @@ int Hardware_Caster::compile_kernel(std::string kernel_source, bool is_path, std size_t kernel_source_size = strlen(source); // Load the source into CL's data structure + cl_program program = clCreateProgramWithSource( context, 1, &source, @@ -474,7 +486,9 @@ int Hardware_Caster::compile_kernel(std::string kernel_source, bool is_path, std if (assert(error, "clCreateKernel")) return OPENCL_ERROR; - kernel_map.emplace(std::make_pair(kernel_name, kernel)); + // Do I want these to overlap when repeated?? + kernel_map[kernel_name] = kernel; + //kernel_map.emplace(std::make_pair(kernel_name, kernel)); return 1; } diff --git a/src/Old_Map.cpp b/src/Old_Map.cpp index 3b7f19a..50a0744 100644 --- a/src/Old_Map.cpp +++ b/src/Old_Map.cpp @@ -158,6 +158,18 @@ void Old_Map::generate_terrain() { } } + for (int x = 30; x < 60; x++) { + //for (int y = 0; y < dimensions.y; y++) { + for (int z = 0; z < 25; z++) { + voxel_data[x + dimensions.x * (50 + dimensions.z * z)] = 6; + } + //} + } + + set_voxel(sf::Vector3i(45, 70, 5), 1); + set_voxel(sf::Vector3i(47, 70, 5), 1); + set_voxel(sf::Vector3i(100, 100, 50), 1); + } diff --git a/src/main.cpp b/src/main.cpp index c552d38..d1ecb24 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -176,18 +176,22 @@ int main() { input_handler.handle_held_keys(); input_handler.dispatch_events(); - if (sf::Keyboard::isKeyPressed(sf::Keyboard::Equal)) { - raycaster->test_edit_viewport(WINDOW_X, WINDOW_Y, w += 5, h += 5); - } - if (sf::Keyboard::isKeyPressed(sf::Keyboard::Dash)) { - raycaster->test_edit_viewport(WINDOW_X, WINDOW_Y, w -= 5, h -= 5); + + if (sf::Keyboard::isKeyPressed(sf::Keyboard::F11)) { + while (raycaster->debug_quick_recompile() != 0); } + //if (sf::Keyboard::isKeyPressed(sf::Keyboard::Equal)) { + // raycaster->test_edit_viewport(WINDOW_X, WINDOW_Y, w += 5, h += 5); + //} + //if (sf::Keyboard::isKeyPressed(sf::Keyboard::Dash)) { + // raycaster->test_edit_viewport(WINDOW_X, WINDOW_Y, w -= 5, h -= 5); + //} if (sf::Keyboard::isKeyPressed(sf::Keyboard::LAlt)) { light_vec.at(0).position = camera->get_position(); light_vec.at(0).direction_cartesian = SphereToCart(camera->get_direction()); } if (sf::Keyboard::isKeyPressed(sf::Keyboard::O)) { - light_vec.at(0).orbit_around_center(timer_accumulator += delta_time); + // light_vec.at(0).orbit_around_center(timer_accumulator += delta_time); } if (sf::Keyboard::isKeyPressed(sf::Keyboard::M)) {