From d2bc5e483aedd3a3668a951fcaaee74a86dd5412 Mon Sep 17 00:00:00 2001 From: mitchellhansen Date: Sat, 24 Feb 2018 21:04:29 -0800 Subject: [PATCH] The settings buffer is fully operational --- include/CLCaster.h | 21 ++++- kernels/ray_caster_kernel.cl | 155 +++-------------------------------- src/Application.cpp | 11 +-- src/CLCaster.cpp | 116 +++++++++++++++++--------- 4 files changed, 108 insertions(+), 195 deletions(-) diff --git a/include/CLCaster.h b/include/CLCaster.h index 86168ca..3469337 100644 --- a/include/CLCaster.h +++ b/include/CLCaster.h @@ -151,8 +151,13 @@ public: void save_config(); // Set a define - void setDefine(std::string name, std::string value); - void removeDefine(std::string name); + void set_define(std::string name, std::string value); + void remove_define(std::string name); + + bool create_settings_buffer(); + bool release_settings_buffer(); + bool add_to_settings_buffer(std::string setting_name, std::string define_accessor_name, int64_t *value); + bool remove_from_settings_buffer(std::string setting_name); // ================================== DEBUG ======================================= @@ -286,9 +291,18 @@ private: // Containers holding the kernels and buffers std::map kernel_map; std::map buffer_map; - std::map defines_map; std::unordered_map>> image_map; + const unsigned int SETTINGS_BUFFER_SIZE = 64; + unsigned int settings_buffer_position = 0; + int64_t* settings_buffer = nullptr; + + // name of setting, position in the settings buffer + std::map settings_buffer_indices; + + // name of define, value + std::map defines_map; + // Hardware caster holds and renders its own textures sf::Sprite viewport_sprite; sf::Texture viewport_texture; @@ -302,7 +316,6 @@ private: std::vector *lights; int light_count = 0; - int error = 0; diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index 06e91fd..a7a3a31 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -47,6 +47,11 @@ constant float4 overshoot_color_2 = { 0.00f, 0.00f, 0.00f, 0.00f }; // ========================================================================= // ========================================================================= + +#define setting(name) settings_buffer[name] + + + // ========================================================================= // ========================= HELPER FUNCTIONS ============================== @@ -120,7 +125,7 @@ bool get_oct_vox( struct TraversalState ts; // push the root node to the parent stack - ts.current_descriptor_index = *settings_buffer; + ts.current_descriptor_index = setting(OCTREE_ROOT_INDEX); ts.current_descriptor = octree_descriptor_buffer[ts.current_descriptor_index]; ts.scale = 0; ts.found = false; @@ -315,146 +320,6 @@ __kernel void raycaster( intersection_t += delta_t * fabs(convert_float3(face_mask.xyz)); voxel.xyz += voxel_step.xyz * face_mask.xyz; -// ======================================================================================================================================= -// ======================================================================================================================================= -// ======================================================================================================================================= - - // uchar prev_val = traversal_state.idx_stack[traversal_state.scale]; - // uint8_t this_face_mask = 0; - // - // // Check the voxel face that we traversed - // // and increment the idx in the idx stack - // if (face_mask.x) { - // this_face_mask = Octree::idx_set_x_mask; - // } - // else if (face_mask.y) { - // this_face_mask = Octree::idx_set_y_mask; - // } - // else if (face_mask.z) { - // this_face_mask = Octree::idx_set_z_mask; - // } - // - // traversal_state.idx_stack[traversal_state.scale] ^= this_face_mask; - // - // // Mask index is the 1D index'd value of the idx for interaction with the valid / leaf masks - // int mask_index = traversal_state.idx_stack[traversal_state.scale]; - // - // // Whether or not the next oct we want to enter in the current CD's valid mask is 1 or 0 - // bool is_valid = false; - // - // // TODO: Rework this logic so we don't have this bodgy if - // if (mask_index > prev_val) - // is_valid = (traversal_state.parent_stack[traversal_state.parent_stack_position] >> 16) & mask_8[mask_index]; - // - // // Check to see if the idx increased or decreased - // // If it decreased - // // Pop up the stack until the oct that the idx flip is valid and we landed on a valid oct - // while (mask_index < prev_val || !is_valid) { - // - // jump_power *= 2; - // - // // Keep track of the 0th edge of out current oct - // traversal_state.oct_pos.x = floor(voxel.x / 2) * jump_power; - // traversal_state.oct_pos.y = floor(voxel.y / 2) * jump_power; - // traversal_state.oct_pos.z = floor(voxel.z / 2) * jump_power; - // - // // Clear and pop the idx stack - // traversal_state.idx_stack[traversal_state.scale] = 0; - // - // // Scale is now set to the oct above. Be wary of this - // traversal_state.scale--; - // - // // Update the prev_val for our new idx - // prev_val = traversal_state.idx_stack[traversal_state.scale]; - // - // // Clear and pop the parent stack, maybe off by one error? - // traversal_state.parent_stack_index[traversal_state.parent_stack_position] = 0; - // traversal_state.parent_stack[traversal_state.parent_stack_position] = 0; - // traversal_state.parent_stack_position--; - // - // // Set the current CD to the one on top of the stack - // traversal_state.current_descriptor = - // traversal_state.parent_stack[traversal_state.parent_stack_position]; - // - // // Apply the face mask to the new idx for the while check - // traversal_state.idx_stack[traversal_state.scale] ^= this_face_mask; - // - // // Get the mask index of the new idx and check the valid status - // mask_index = traversal_state.idx_stack[traversal_state.scale]; - // is_valid = (traversal_state.parent_stack[traversal_state.parent_stack_position] >> 16) & mask_8[mask_index]; - // } - // - // // At this point parent_stack[position] is at the CD of an oct with a valid oct at the leaf indicated by the current - // // idx in the idx stack scale - // - // // While we haven't bottomed out and the oct we're looking at is valid - // while (jump_power > 1 && is_valid) { - // - // // If all went well and we found a valid non-leaf oct then we will traverse further down the hierarchy - // traversal_state.scale++; - // jump_power /= 2; - // - // // Count the number of valid octs that come before and add it to the index to get the position - // // Negate it by one as it counts itself - // int count = count_bits((uint8_t)(traversal_state.parent_stack[traversal_state.parent_stack_position] >> 16) & count_mask_8[mask_index]) - 1; - // - // // If this CD had the far bit set - // if (far_bit_mask & descriptor_buffer[traversal_state.parent_stack_index[traversal_state.parent_stack_position]]) { - // - // // access the far point at which the head points too. Determine it's value, and add - // // the count of the valid bits in the current CD to the index - // uint64_t far_pointer_index = - // traversal_state.parent_stack_index[traversal_state.parent_stack_position] + // current index + - // (traversal_state.parent_stack[traversal_state.parent_stack_position] & child_pointer_mask); // the relative prt to the far ptr - // - // // Get the absolute ptr from the far ptr and add the count to get the CD that we want - // traversal_state.parent_stack_index[traversal_state.parent_stack_position + 1] = descriptor_buffer[far_pointer_index] + count; - // } - // // If this CD doesn't have the far bit set, access the element at which head points to - // // and then add the specified number of indices to get to the correct child descriptor - // else { - // traversal_state.parent_stack_index[traversal_state.parent_stack_position + 1] = - // traversal_state.parent_stack_index[traversal_state.parent_stack_position] + // The current index to this CD - // (traversal_state.parent_stack[traversal_state.parent_stack_position] & child_pointer_mask) + count; // The relative dist + the number of bits that were valid - // } - // - // // Now that we have the index set we can increase our parent stack position to the next level and - // // retrieve the value of its CD - // traversal_state.parent_stack_position++; - // traversal_state.parent_stack[traversal_state.parent_stack_position] = descriptor_buffer[traversal_state.parent_stack_index[traversal_state.parent_stack_position]]; - // - // // Unlike the single shot DFS, it makes a bit more sense to have this at the tail of the while loop - // // Do the logic steps to find which sub oct we step down into - // if (voxel.x >= (jump_power / 2) + traversal_state.oct_pos.x) { - // - // // Set our voxel position to the (0,0) of the correct oct - // traversal_state.oct_pos.x += (jump_power / 2); - // - // // Set the idx to represent the move - // traversal_state.idx_stack[traversal_state.scale] |= idx_set_x_mask; - // - // } - // if (voxel.y >= (jump_power / 2) + traversal_state.oct_pos.y) { - // - // traversal_state.oct_pos.y += (jump_power / 2); - // traversal_state.idx_stack[traversal_state.scale] |= idx_set_y_mask; - // } - // if (voxel.z >= (jump_power / 2) + traversal_state.oct_pos.z) { - // - // traversal_state.oct_pos.z += (jump_power / 2); - // traversal_state.idx_stack[traversal_state.scale] |= idx_set_z_mask; - // } - // - // // Update the mask index with the new voxel we walked down to, and then check it's valid status - // mask_index = traversal_state.idx_stack[traversal_state.scale]; - // is_valid = (traversal_state.parent_stack[traversal_state.parent_stack_position] >> 16) & mask_8[mask_index]; - // - // } - -// ======================================================================================================================================= -// ======================================================================================================================================= -// ======================================================================================================================================= - // Test for out of bounds contions, add fog if (any(voxel >= *map_dim) || any(voxel < 0)){ voxel.xyz -= voxel_step.xyz * face_mask.xyz; @@ -462,11 +327,11 @@ __kernel void raycaster( color_accumulator.w *= 4; break; } + int vox_dim = setting(OCTDIM); - int vox_dim = OCTDIM; + // If we hit a voxel -// If we hit a voxel - // if (voxel.x < (*map_dim).x && voxel.y < (*map_dim).x && voxel.z < (*map_dim).x){ + if (voxel.x < (*map_dim).x && voxel.y < (*map_dim).x && voxel.z < (*map_dim).x){ // if (get_oct_vox( // voxel, // octree_descriptor_buffer, @@ -480,7 +345,7 @@ __kernel void raycaster( // } // } else { voxel_data = map[voxel.x + (*map_dim).x * (voxel.y + (*map_dim).z * (voxel.z))]; - //} + } diff --git a/src/Application.cpp b/src/Application.cpp index 59e8d30..758137a 100644 --- a/src/Application.cpp +++ b/src/Application.cpp @@ -23,18 +23,18 @@ Application::~Application() { else { Logger::log("Can't release window, shared_ptr count : " + window.use_count(), Logger::LogLevel::WARN); } - //light_handle->~LightHandle(); - //light_controller->~LightController(); } bool Application::init_clcaster() { // Start up the raycaster raycaster = std::make_shared(); - raycaster->setDefine("OCTDIM", std::to_string(MAP_X)); + if (!raycaster->init()) abort(); + raycaster->add_to_settings_buffer("octree_dimensions", "OCTDIM", (int64_t*)&MAP_X); + map = std::make_shared(MAP_X); // TODO: Implement this @@ -47,9 +47,6 @@ bool Application::init_clcaster() { raycaster->assign_octree(map); raycaster->assign_map(map); - - - camera = std::make_shared( sf::Vector3f(3.5f, 3.5f, 3.5f), // Starting position sf::Vector2f(1.57f, 0.0f), // Direction @@ -77,7 +74,7 @@ bool Application::init_clcaster() { Logger::log("Failed to load spritesheet from file", Logger::LogLevel::WARN); raycaster->create_texture_atlas(&spritesheet, sf::Vector2i(16, 16)); - // Checks to see if proper data was uploaded, then sets the kernel args + // Compiles the kernel, Checks to see if proper data was uploaded, then sets the kernel args // ALL DATA LOADING MUST BE FINISHED if (!raycaster->validate()) { abort(); diff --git a/src/CLCaster.cpp b/src/CLCaster.cpp index c4e7bda..39ef2e4 100644 --- a/src/CLCaster.cpp +++ b/src/CLCaster.cpp @@ -3,14 +3,6 @@ CLCaster::CLCaster() {} CLCaster::~CLCaster() { - // Causes sigabrt?? - //release_map(); - //release_camera(); - //release_octree(); - //clReleaseKernel(kernel_map.at("raycaster")); - //clReleaseProgram() - //release_viewport(); - delete[] viewport_matrix; delete[] viewport_image; @@ -64,12 +56,11 @@ bool CLCaster::init() { return false; } - if (!compile_kernel("../kernels/ray_caster_kernel.cl", true, "raycaster")) { - Logger::log("Failed to compile the kernel", Logger::LogLevel::ERROR, __LINE__, __FILE__); - std::cin.get(); // hang the output window so we can read the error - return false; - } - + if (!create_settings_buffer()) { + Logger::log("Failed to create settings buffer", Logger::LogLevel::ERROR, __LINE__, __FILE__); + return false; + } + srand(time(nullptr)); int *seed_memory = new int[1920*1080]; @@ -117,8 +108,8 @@ bool CLCaster::assign_octree(std::shared_ptr map) { return false; if (!create_buffer("octree_attachment_buffer", map->octree.buffer_size * sizeof(uint64_t), map->octree.attachment_buffer)) return false; - if (!create_buffer("settings_buffer", sizeof(uint64_t), &map->octree.root_index)) - return false; + + add_to_settings_buffer("octree_root_index", "OCTREE_ROOT_INDEX", (int64_t*)&map->octree.root_index); return true; } @@ -134,8 +125,6 @@ bool CLCaster::release_octree() return false; if (!release_buffer("octree_attachment_buffer")) return false; - if (!release_buffer("settings_buffer")) - return false; return true; } @@ -166,6 +155,12 @@ bool CLCaster::release_camera() { bool CLCaster::validate() { + if (!compile_kernel("../kernels/ray_caster_kernel.cl", true, "raycaster")) { + Logger::log("Failed to compile the kernel", Logger::LogLevel::ERROR, __LINE__, __FILE__); + std::cin.get(); // hang the output window so we can read the error + return false; + } + Logger::log("Validating OpenCL kernel args", Logger::LogLevel::INFO); // Check to make sure everything has been entered @@ -278,7 +273,7 @@ bool CLCaster::create_viewport(int width, int height, float v_fov, float h_fov) } } - if (!create_buffer("viewport_matrix", sizeof(float) * 4 * view_res.x * view_res.y, viewport_matrix, CL_MEM_USE_HOST_PTR)) + if (!create_buffer("viewport_matrix", sizeof(float) * 4 * view_res.x * view_res.y, viewport_matrix)) return false; // Create the image that opencl's rays write to @@ -635,7 +630,6 @@ bool CLCaster::create_shared_context() 0 }; - #elif defined TARGET_OS_MAC CGLContextObj glContext = CGLGetCurrentContext(); @@ -648,8 +642,6 @@ bool CLCaster::create_shared_context() #endif - - // Create our shared context context = clCreateContext( context_properties, @@ -718,7 +710,6 @@ bool CLCaster::compile_kernel(std::string kernel_source, bool is_path, std::stri Logger::log("Failed at clCreateProgramWithSource() :" + cl_err_lookup(error), Logger::LogLevel::ERROR, __LINE__, __FILE__); return false; } - std::stringstream build_string_stream; @@ -727,7 +718,6 @@ bool CLCaster::compile_kernel(std::string kernel_source, bool is_path, std::stri build_string_stream << " -D" << define.first << "=" << define.second; } - //build_string_stream << "-DOCTDIM=" << std::to_string(Application::MAP_X); build_string_stream << " -cl-finite-math-only -cl-fast-relaxed-math -cl-unsafe-math-optimizations"; std::string build_string = build_string_stream.str(); @@ -1177,14 +1167,63 @@ std::string CLCaster::cl_err_lookup(int error_code) { } -void CLCaster::setDefine(std::string name, std::string value) { +void CLCaster::set_define(std::string name, std::string value) { defines_map[name] = value; } -void CLCaster::removeDefine(std::string name) { +void CLCaster::remove_define(std::string name) { defines_map.erase(name); } +bool CLCaster::add_to_settings_buffer(std::string setting_name, std::string define_accessor_name, int64_t *value) { + + bool success = true; + + if (settings_buffer == nullptr){ + + Logger::log("Trying to push settings to an uninitialized settings buffer", Logger::LogLevel::ERROR, __LINE__, __FILE__); + success = false; + + } else if (defines_map.count(define_accessor_name)) { + + Logger::log("Define name already present in the defines map", Logger::LogLevel::ERROR, __LINE__, __FILE__); + success = false; + + } else { + + if (settings_buffer_position < SETTINGS_BUFFER_SIZE) { + defines_map[define_accessor_name] = std::to_string(settings_buffer_position); + settings_buffer[settings_buffer_position] = *value; + settings_buffer_position++; + } else { + Logger::log("Settings buffer has reached the maximum size of " + std::to_string(SETTINGS_BUFFER_SIZE) + " elements", Logger::LogLevel::ERROR, __LINE__, __FILE__); + success = false; + } + } + + return success; +} + +bool CLCaster::create_settings_buffer() { + + settings_buffer = new int64_t[SETTINGS_BUFFER_SIZE]; + if (!create_buffer("settings_buffer", sizeof(int64_t) * SETTINGS_BUFFER_SIZE, settings_buffer, CL_MEM_USE_HOST_PTR)) + return false; + return true; +} + +bool CLCaster::remove_from_settings_buffer(std::string setting_name) { + + Logger::log("remove_from_settings_buffer() not implimented", Logger::LogLevel::WARN, __LINE__, __FILE__); + return false; +} + +bool CLCaster::release_settings_buffer() { + if (!release_buffer("settings_buffer")) + return false; + return true; +} + CLCaster::device::device(cl_device_id device_id, cl_platform_id platform_id) { @@ -1237,26 +1276,25 @@ CLCaster::device::device(const device& d) { void CLCaster::device::print(std::ostream& stream) const { stream << "\n\tDevice ID : " << device_id << std::endl; - stream << "\tDevice Name : " << data.device_name << std::endl; + stream << "\tPlatform ID : " << platform_id << std::endl; - stream << "\tPlatform ID : " << platform_id << std::endl; - stream << "\tPlatform Name : " << data.platform_name << std::endl; + stream << "\tDevice Name : " << data.device_name << std::endl; + stream << "\tPlatform Name : " << data.platform_name << std::endl; + stream << "\tDevice Type : "; - stream << "\tOpenCL Version : " << data.opencl_version << std::endl; - stream << "\tSupports sharing : " << std::boolalpha << cl_gl_sharing << std::endl; - stream << "\tDevice Type : "; + if (data.device_type == CL_DEVICE_TYPE_CPU) + stream << "CPU" << std::endl; - if (data.device_type == CL_DEVICE_TYPE_CPU) - stream << "CPU" << std::endl; + else if (data.device_type == CL_DEVICE_TYPE_GPU) + stream << "GPU" << std::endl; - else if (data.device_type == CL_DEVICE_TYPE_GPU) - stream << "GPU" << std::endl; + else if (data.device_type == CL_DEVICE_TYPE_ACCELERATOR) + stream << "Accelerator" << std::endl; - else if (data.device_type == CL_DEVICE_TYPE_ACCELERATOR) - stream << "Accelerator" << std::endl; + stream << "\tOpenCL Version : " << data.opencl_version << std::endl; + stream << "\tSupports sharing : " << std::boolalpha << cl_gl_sharing << std::endl; stream << "\tIs Little Endian : " << std::boolalpha << is_little_endian << std::endl; - stream << "\tClock Frequency : " << data.clock_frequency << std::endl; stream << "\tCompute Units : " << data.compute_units << std::endl;