From 7e5d4ef947f7c62f4a637f226cb152e61714cf4d Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Sun, 19 Mar 2017 00:42:54 -0700 Subject: [PATCH] 10 FPS average increase from changing global work size to 2D --- include/raycaster/Hardware_Caster.h | 2 +- kernels/ray_caster_kernel.cl | 7 ++++++- src/main.cpp | 3 ++- src/raycaster/Hardware_Caster.cpp | 13 ++++++++----- 4 files changed, 17 insertions(+), 8 deletions(-) diff --git a/include/raycaster/Hardware_Caster.h b/include/raycaster/Hardware_Caster.h index d960c6b..e3073fe 100644 --- a/include/raycaster/Hardware_Caster.h +++ b/include/raycaster/Hardware_Caster.h @@ -136,7 +136,7 @@ private: // Run the kernel using a 1d work size // TODO: Test 2d worksize - int run_kernel(std::string kernel_name, const int work_size); + int run_kernel(std::string kernel_name, const int work_dim_x, const int work_dim_y); // Run a test kernel that prints out the kernel args void print_kernel_arguments(); diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index ef73dbb..e363bb8 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -149,6 +149,9 @@ __kernel void raycaster( ){ int global_id = get_global_id(0); + + int x = get_global_id(0); + int y = get_global_id(1); // Get and set the random seed from seed memory int seed = seed_memory[global_id]; @@ -156,7 +159,9 @@ __kernel void raycaster( seed_memory[global_id] = seed; // Get the pixel on the viewport, and find the view matrix ray that matches it - int2 pixel = { global_id % (*resolution).x, global_id / (*resolution).x}; + //int2 pixel = { global_id % (*resolution).x, global_id / (*resolution).x }; + int2 pixel = (int2)(x, y); + float3 ray_dir = projection_matrix[pixel.x + (*resolution).x * pixel.y]; //if (pixel.x == 960 && pixel.y == 540) { diff --git a/src/main.cpp b/src/main.cpp index 07b6635..8dab75f 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -104,6 +104,7 @@ int main() { window.setMouseCursorVisible(false); window.setKeyRepeatEnabled(false); window.setFramerateLimit(60); + window.setVerticalSyncEnabled(false); ImGui::SFML::Init(window); window.resetGLStates(); @@ -264,7 +265,7 @@ int main() { handle->set_rgbi(light); } - if (ImGui::SliderFloat4("Position", light_pos, 0, MAP_X)) { + if (ImGui::SliderFloat3("Position", light_pos, 0, MAP_X)) { sf::Vector3f light(light_pos[0], light_pos[1], light_pos[2]); handle->set_position(light); } diff --git a/src/raycaster/Hardware_Caster.cpp b/src/raycaster/Hardware_Caster.cpp index d473f06..c7fe33c 100644 --- a/src/raycaster/Hardware_Caster.cpp +++ b/src/raycaster/Hardware_Caster.cpp @@ -104,7 +104,7 @@ void Hardware_Caster::create_texture_atlas(sf::Texture *t, sf::Vector2i tile_dim void Hardware_Caster::compute() { // correlating work size with texture size? good, bad? - run_kernel("raycaster", viewport_texture.getSize().x * viewport_texture.getSize().y); + run_kernel("raycaster", viewport_texture.getSize().x, viewport_texture.getSize().y); } // There is a possibility that I would want to move this over to be all inside it's own @@ -663,9 +663,12 @@ int Hardware_Caster::store_buffer(cl_mem buffer, std::string buffer_name) { return 1; } -int Hardware_Caster::run_kernel(std::string kernel_name, const int work_size) { +int Hardware_Caster::run_kernel(std::string kernel_name, const int work_dim_x, const int work_dim_y) { - size_t global_work_size[1] = { static_cast(work_size) }; + //size_t global_work_size[2] = { static_cast(work_dim_x), static_cast(work_dim_y)}; + + size_t global_work_size[2] = { static_cast(1440), static_cast(900)}; + //size_t global_work_size[1] = { static_cast(1440*900) }; cl_kernel kernel = kernel_map.at(kernel_name); @@ -676,7 +679,7 @@ int Hardware_Caster::run_kernel(std::string kernel_name, const int work_size) { //error = clEnqueueTask(command_queue, kernel, 0, NULL, NULL); error = clEnqueueNDRangeKernel( command_queue, kernel, - 1, NULL, global_work_size, + 2, NULL, global_work_size, NULL, 0, NULL, NULL); if (vr_assert(error, "clEnqueueNDRangeKernel")) @@ -705,7 +708,7 @@ void Hardware_Caster::print_kernel_arguments() set_kernel_arg("printer", 7, "light_count"); set_kernel_arg("printer", 8, "image"); - run_kernel("printer", 1); + run_kernel("printer", 1, 1); } cl_device_id Hardware_Caster::getDeviceID() { return device_id; };