From 1620f40d021d61ed861369e06ec6a56c228edb2d Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Fri, 23 Sep 2016 14:08:01 -0700 Subject: [PATCH] Fixed the camera movement so it now take into account sub-voxel movement Also fixed the distortion around the XY axes --- include/Map.h | 5 +++- ...minimal_kernel.cl => ray_caster_kernel.cl} | 24 ++++++++++++++----- src/CL_Wrapper.cpp | 7 ++---- src/Camera.cpp | 2 +- src/main.cpp | 13 ++++------ 5 files changed, 29 insertions(+), 22 deletions(-) rename kernels/{minimal_kernel.cl => ray_caster_kernel.cl} (85%) diff --git a/include/Map.h b/include/Map.h index 0efcef2..7f65f64 100644 --- a/include/Map.h +++ b/include/Map.h @@ -198,7 +198,10 @@ public: if (height_map[x + y * dim.x] > 0) { int z = height_map[x + y * dim.x]; - list[x + dim.x * (y + dim.z * z)] = 5; + while (z > 0){ + list[x + dim.x * (y + dim.z * z)] = 5; + z--; + } } } diff --git a/kernels/minimal_kernel.cl b/kernels/ray_caster_kernel.cl similarity index 85% rename from kernels/minimal_kernel.cl rename to kernels/ray_caster_kernel.cl index 7ed7ec3..f48c323 100644 --- a/kernels/minimal_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -45,10 +45,6 @@ __kernel void min_kern( int3 voxel_step = {1, 1, 1}; 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 int3 voxel = convert_int3(*cam_pos); @@ -56,9 +52,25 @@ __kernel void min_kern( // traverse an integer split float3 delta_t = fabs(1.0f / ray_dir); - // Intersection T is the collection of the next intersection points + // offset is how far we are into a voxel, enables sub voxel movement + float3 offset = ((*cam_pos) - floor(*cam_pos)) * convert_float3(voxel_step); + + //offset.x += delta_t.x * convert_float((voxel_step.x < 0)); + //offset -= delta_t * floor(offset / delta_t); + + // Intersection T is the collection of the next intersection points // for all 3 axis XYZ. - float3 intersection_t = delta_t; + float3 intersection_t = delta_t * offset; + + 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; + } int2 randoms = { 3, 14 }; uint seed = randoms.x + id; diff --git a/src/CL_Wrapper.cpp b/src/CL_Wrapper.cpp index 7fdbe47..b2ba706 100644 --- a/src/CL_Wrapper.cpp +++ b/src/CL_Wrapper.cpp @@ -291,18 +291,16 @@ int CL_Wrapper::run_kernel(std::string kernel_name, const int work_size){ if (assert(error, "clEnqueueNDRangeKernel")) return -1; + clFinish(getCommandQueue()); + // What if errors out and gl objects are never released? error = clEnqueueReleaseGLObjects(getCommandQueue(), 1, &buffer_map.at("image_buffer"), 0, NULL, NULL); if (assert(error, "clEnqueueReleaseGLObjects")) return -1; - - return 1; } - - cl_device_id CL_Wrapper::getDeviceID(){ return device_id; }; cl_platform_id CL_Wrapper::getPlatformID(){ return platform_id; }; cl_context CL_Wrapper::getContext(){ return context; }; @@ -316,7 +314,6 @@ bool CL_Wrapper::assert(int error_code, std::string function_name){ switch (error_code) { - case CL_SUCCESS: return false; diff --git a/src/Camera.cpp b/src/Camera.cpp index 06419a0..1ace366 100644 --- a/src/Camera.cpp +++ b/src/Camera.cpp @@ -68,7 +68,7 @@ int Camera::update(double delta_time) { // so vector multiplication broke? // have to do it component wise - double multiplier = 50; + double multiplier = 10; position.x += movement.x * delta_time * multiplier; position.y += movement.y * delta_time * multiplier; diff --git a/src/main.cpp b/src/main.cpp index 10ab7b2..ba76946 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -68,12 +68,8 @@ int main() { //Map m(sf::Vector3i (50, 50, 50)); //return 1; - - - sf::RenderWindow window(sf::VideoMode(WINDOW_X, WINDOW_Y), "SFML"); - // Setup CL, instantiate and pass in values to the kernel CL_Wrapper c; query_platform_devices(); @@ -81,9 +77,9 @@ int main() { c.create_shared_context(); c.create_command_queue(); - //c.compile_kernel("../kernels/kernel.cl", true, "hello"); - if (c.compile_kernel("../kernels/minimal_kernel.cl", true, "min_kern") < 0) { + if (c.compile_kernel("../kernels/ray_caster_kernel.cl", true, "min_kern") < 0) { std::cin.get(); + return -1; } std::cout << "map..."; @@ -277,7 +273,7 @@ int main() { camera.add_relative_impulse(Camera::DIRECTION::RIGHT); } if (sf::Keyboard::isKeyPressed(sf::Keyboard::T)) { - camera.set_position(sf::Vector3f(20, 20, 20)); + camera.set_position(sf::Vector3f(50, 50, 50)); } camera.add_static_impulse(cam_vec); @@ -313,8 +309,7 @@ int main() { // Run the raycast c.run_kernel("min_kern", WORK_SIZE); - clFinish(c.getCommandQueue()); - + // ==== RENDER ==== window.clear(sf::Color::Black);