mitchellhansen/voxel-raycaster
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

10 FPS average increase from changing global work size to 2D

Browse Source
This commit is contained in:
MitchellHansen
2017-03-19 00:42:54 -07:00
parent ec65ef7741
commit 7e5d4ef947
4 changed files with 17 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/raycaster/Hardware_Caster.h
View File

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

7
kernels/ray_caster_kernel.cl
View File

@@ -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) {

3
src/main.cpp
View File

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

13
src/raycaster/Hardware_Caster.cpp
View File

@@ -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<size_t>(work_size) };
//size_t global_work_size[2] = { static_cast<size_t>(work_dim_x), static_cast<size_t>(work_dim_y)};
size_t global_work_size[2] = { static_cast<size_t>(1440), static_cast<size_t>(900)};
//size_t global_work_size[1] = { static_cast<size_t>(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; };