diff --git a/include/raycaster/Hardware_Caster.h b/include/raycaster/Hardware_Caster.h new file mode 100644 index 0000000..ee6bf4e --- /dev/null +++ b/include/raycaster/Hardware_Caster.h @@ -0,0 +1,112 @@ +#pragma once +#include +#include +#include +#include "util.hpp" +#include +#include + + +#ifdef linux +#include +#include +#include + +#elif defined _WIN32 +#include +#include +#include + +// Note: windows.h must be included before Gl/GL.h +#include +#include + +#elif defined TARGET_OS_MAC +# include +# include + +#endif + +struct device { + cl_device_id id; + cl_device_type type; + cl_uint clock_frequency; + char version[128]; + cl_platform_id platform; + cl_uint comp_units; +}; + +class Hardware_Caster : public RayCaster +{ +public: + Hardware_Caster(); + + virtual ~Hardware_Caster(); + + int init() override; + + // In interop mode, this will create a GL texture that we share + // Otherwise, it will create the pixel buffer and pass that in as an image, retrieving it each draw + // Both will create the view matrix, view res buffer + void create_viewport(int width, int height, float v_fov, float h_fov) override; + + void assign_lights(std::vector *data) override; + void assign_map(Old_Map *map) override; + void assign_camera(Camera *camera) override; + void validate() override; + + // TODO: Hoist this to the base class + void create_texture_atlas(sf::Texture *t, sf::Vector2i tile_dim); + + + // draw will abstract the gl sharing and software rendering + // methods of retrieving the screen buffer + void compute() override; + 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: + + + int acquire_platform_and_device(); + + int create_shared_context(); + + int create_command_queue(); + + int check_cl_khr_gl_sharing(); + + int create_image_buffer(std::string buffer_name, cl_uint size, sf::Texture* texture); + int create_buffer(std::string buffer_name, cl_uint size, void* data); + int create_buffer(std::string buffer_name, cl_uint size, void* data, cl_mem_flags flags); + int store_buffer(cl_mem, std::string buffer_name); + int release_buffer(std::string buffer_name); + + int compile_kernel(std::string kernel_source, bool is_path, std::string kernel_name); + + int set_kernel_arg(std::string kernel_name, int index, std::string buffer_name); + + int run_kernel(std::string kernel_name, const int work_size); + + void print_kernel_arguments(); + + bool assert(int error_code, std::string function_name); + + cl_device_id getDeviceID(); + cl_platform_id getPlatformID(); + cl_context getContext(); + cl_kernel getKernel(std::string kernel_name); + cl_command_queue getCommandQueue(); + + cl_platform_id platform_id; + cl_device_id device_id; + cl_context context; + cl_command_queue command_queue; + + std::map kernel_map; + std::map buffer_map; + +}; + diff --git a/include/raycaster/RayCaster.h b/include/raycaster/RayCaster.h new file mode 100644 index 0000000..a9c5d9e --- /dev/null +++ b/include/raycaster/RayCaster.h @@ -0,0 +1,52 @@ +#pragma once +#include +#include +#include +#include "Old_Map.h" +#include "Camera.h" +#include "LightController.h" + +class RayCaster { +public: + + enum ERROR_CODES { + SHARING_NOT_SUPPORTED = 800, + OPENCL_NOT_SUPPORTED = 801, + OPENCL_ERROR = 802, + ERR = 803 + }; + + RayCaster(); + virtual ~RayCaster(); + + virtual int init() = 0; + + virtual void assign_map(Old_Map *map) = 0; + virtual void assign_camera(Camera *camera) = 0; + virtual void create_viewport(int width, int height, float v_fov, float h_fov) = 0; + virtual void assign_lights(std::vector *data) = 0; + virtual void validate() = 0; + + // draw will abstract the gl sharing and software rendering + // methods of retrieving the screen buffer + virtual void compute() = 0; + virtual void draw(sf::RenderWindow* window) = 0; + +protected: + + sf::Sprite viewport_sprite; + sf::Texture viewport_texture; + + Old_Map * map = nullptr; + Camera *camera = nullptr; +// std::vector *lights; + std::vector *lights; + int light_count = 0; + sf::Uint8 *viewport_image = nullptr; + sf::Vector4f *viewport_matrix = nullptr; + sf::Vector2i viewport_resolution; + + int error = 0; + +}; + diff --git a/include/raycaster/Software_Caster.h b/include/raycaster/Software_Caster.h new file mode 100644 index 0000000..f2c28d4 --- /dev/null +++ b/include/raycaster/Software_Caster.h @@ -0,0 +1,35 @@ +#include "RayCaster.h" +#include + +class Software_Caster : public RayCaster +{ +public: + Software_Caster(); + + virtual ~Software_Caster(); + + int init() override; + + // In interop mode, this will create a GL texture that we share + // Otherwise, it will create the pixel buffer and pass that in as an image, retrieving it each draw + // Both will create the view matrix, view res buffer + void create_viewport(int width, int height, float v_fov, float h_fov) override; + + void assign_lights(std::vector *data) override; + void assign_map(Old_Map *map) override; + void assign_camera(Camera *camera) override; + void validate() override; + + // draw will abstract the gl sharing and software rendering + // methods of retrieving the screen buffer + void compute() override; + void draw(sf::RenderWindow* window) override; + +private: + + void cast_viewport(); + void cast_thread(int start_id, int end_id); + void cast_ray(int id); + void blit_pixel(sf::Color color, sf::Vector2i position, sf::Vector3i mask); + sf::Color global_light(sf::Color in, sf::Vector3i mask); +}; diff --git a/src/raycaster/Hardware_Caster.cpp b/src/raycaster/Hardware_Caster.cpp new file mode 100644 index 0000000..1ee91dd --- /dev/null +++ b/src/raycaster/Hardware_Caster.cpp @@ -0,0 +1,896 @@ +#include "Hardware_Caster.h" + + + +Hardware_Caster::Hardware_Caster() { + +} + + +Hardware_Caster::~Hardware_Caster() { +} + +int Hardware_Caster::init() { + + // Initialize opencl up to the point where we start assigning buffers + + error = acquire_platform_and_device(); + if(assert(error, "aquire_platform_and_device")) + return error; + + error = check_cl_khr_gl_sharing(); + if(assert(error, "check_cl_khr_gl_sharing")) + return error; + + error = create_shared_context(); + if (assert(error, "create_shared_context")) + return error; + + error = create_command_queue(); + if (assert(error, "create_command_queue")) + return error; + + 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; + } + + srand(time(NULL)); + + int *seed_memory = new int[1920*1080]; + + create_buffer("seed", sizeof(int) * 1920 * 1080, seed_memory); + + return 1; + +} + +void Hardware_Caster::assign_map(Old_Map *map) { + + this->map = map; + auto dimensions = map->getDimensions(); + + create_buffer("map", sizeof(char) * dimensions.x * dimensions.y * dimensions.z, map->get_voxel_data()); + create_buffer("map_dimensions", sizeof(int) * 3, &dimensions); + +} + +void Hardware_Caster::assign_camera(Camera *camera) { + + this->camera = camera; + + create_buffer("camera_direction", sizeof(float) * 4, (void*)camera->get_direction_pointer(), CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR); + create_buffer("camera_position", sizeof(float) * 4, (void*)camera->get_position_pointer(), CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR); +} + +void Hardware_Caster::validate() +{ + // Check to make sure everything has been entered; + if (camera == nullptr || + map == nullptr || + viewport_image == nullptr || + viewport_matrix == nullptr) { + + std::cout << "Raycaster.validate() failed, camera, map, or viewport not initialized"; + + } else { + + // Set all the kernel args + set_kernel_arg("raycaster", 0, "map"); + set_kernel_arg("raycaster", 1, "map_dimensions"); + set_kernel_arg("raycaster", 2, "viewport_resolution"); + set_kernel_arg("raycaster", 3, "viewport_matrix"); + set_kernel_arg("raycaster", 4, "camera_direction"); + set_kernel_arg("raycaster", 5, "camera_position"); + set_kernel_arg("raycaster", 6, "lights"); + set_kernel_arg("raycaster", 7, "light_count"); + set_kernel_arg("raycaster", 8, "image"); + set_kernel_arg("raycaster", 9, "seed"); + set_kernel_arg("raycaster", 10, "texture_atlas"); + set_kernel_arg("raycaster", 11, "atlas_dim"); + set_kernel_arg("raycaster", 12, "tile_dim"); + + //print_kernel_arguments(); + } + + +} + +void Hardware_Caster::create_texture_atlas(sf::Texture *t, sf::Vector2i tile_dim) { + + create_image_buffer("texture_atlas", t->getSize().x * t->getSize().x * 4 * sizeof(float), t); + + // create_buffer observes arg 3's + + sf::Vector2u v = t->getSize(); + create_buffer("atlas_dim", sizeof(sf::Vector2u) , &v); + + create_buffer("tile_dim", sizeof(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); +} + +// There is a possibility that I would want to move this over to be all inside it's own +// container to make it so it can be changed via CL_MEM_USE_HOST_PTR. But I doubt it +// would ever be called enough to warrent that +void Hardware_Caster::create_viewport(int width, int height, float v_fov, float h_fov) { + + // CL needs the screen resolution + sf::Vector2i view_res(width, height); + create_buffer("viewport_resolution", sizeof(int) * 2, &view_res); + + // And an array of vectors describing the way the "lens" of our + // camera works + + // This could be modified to make some odd looking camera lenses + + double y_increment_radians = DegreesToRadians(v_fov / view_res.y); + double x_increment_radians = DegreesToRadians(h_fov / view_res.x); + + viewport_matrix = new sf::Vector4f[width * height * 4]; + + for (int y = -view_res.y / 2; y < view_res.y / 2; y++) { + for (int x = -view_res.x / 2; x < view_res.x / 2; x++) { + + // The base ray direction to slew from + sf::Vector3f ray(1, 0, 0); + + // Y axis, pitch + ray = sf::Vector3f( + static_cast(ray.z * sin(y_increment_radians * y) + ray.x * cos(y_increment_radians * y)), + static_cast(ray.y), + static_cast(ray.z * cos(y_increment_radians * y) - ray.x * sin(y_increment_radians * y)) + ); + + // Z axis, yaw + ray = sf::Vector3f( + static_cast(ray.x * cos(x_increment_radians * x) - ray.y * sin(x_increment_radians * x)), + static_cast(ray.x * sin(x_increment_radians * x) + ray.y * cos(x_increment_radians * x)), + static_cast(ray.z) + ); + + // correct for the base ray pointing to (1, 0, 0) as (0, 0). Should equal (1.57, 0) + ray = sf::Vector3f( + static_cast(ray.z * sin(-1.57) + ray.x * cos(-1.57)), + static_cast(ray.y), + static_cast(ray.z * cos(-1.57) - ray.x * sin(-1.57)) + ); + + int index = (x + view_res.x / 2) + view_res.x * (y + view_res.y / 2); + ray = Normalize(ray); + + viewport_matrix[index] = sf::Vector4f( + ray.x, + ray.y, + ray.z, + 0 + ); + } + } + + create_buffer("viewport_matrix", sizeof(float) * 4 * view_res.x * view_res.y, viewport_matrix, CL_MEM_USE_HOST_PTR); + + // Create the image that opencl's rays write to + viewport_image = new sf::Uint8[width * height * 4]; + + for (int i = 0; i < width * height * 4; i += 4) { + + viewport_image[i] = 255; // R + viewport_image[i + 1] = 255; // G + viewport_image[i + 2] = 255; // B + viewport_image[i + 3] = 100; // A + } + + // Interop lets us keep a reference to it as a texture + viewport_texture.create(width, height); + viewport_texture.update(viewport_image); + viewport_sprite.setTexture(viewport_texture); + + // Pass the buffer to opencl + create_image_buffer("image", sizeof(sf::Uint8) * width * height * 4, &viewport_texture); + +} + +//void Hardware_Caster::assign_lights(std::vector *lights) { +// +// //this->lights = ; +// +// std::cout << sizeof(LightController); +// std::cout << sizeof(float); +// light_count = static_cast(lights->size()); +// +// //create_buffer("lights", sizeof(float) * 10 * light_count, this->lights->data(), CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR); +// +// create_buffer("light_count", sizeof(int), &light_count); +// +//} + +void Hardware_Caster::assign_lights(std::vector *data) { + + // Get a pointer to the packed light data +// this->lights = data; + + light_count = static_cast(lights->size()); + + size_t packed_size = sizeof(LightController::PackedData); + + create_buffer("lights", packed_size * light_count, lights->data(), CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR); + + create_buffer("light_count", sizeof(int), &light_count); + +} + +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(); + + return 1; +} + +void Hardware_Caster::test_edit_viewport(int width, int height, float v_fov, float h_fov) +{ + sf::Vector2i view_res(width, height); + + double y_increment_radians = DegreesToRadians(v_fov / view_res.y); + double x_increment_radians = DegreesToRadians(h_fov / view_res.x); + + for (int y = -view_res.y / 2; y < view_res.y / 2; y++) { + for (int x = -view_res.x / 2; x < view_res.x / 2; x++) { + + // The base ray direction to slew from + sf::Vector3f ray(1, 0, 0); + + // Y axis, pitch + ray = sf::Vector3f( + static_cast(ray.z * sin(y_increment_radians * y) + ray.x * cos(y_increment_radians * y)), + static_cast(ray.y), + static_cast(ray.z * cos(y_increment_radians * y) - ray.x * sin(y_increment_radians * y)) + ); + + // Z axis, yaw + ray = sf::Vector3f( + static_cast(ray.x * cos(x_increment_radians * x) - ray.y * sin(x_increment_radians * x)), + static_cast(ray.x * sin(x_increment_radians * x) + ray.y * cos(x_increment_radians * x)), + static_cast(ray.z) + ); + + // correct for the base ray pointing to (1, 0, 0) as (0, 0). Should equal (1.57, 0) + ray = sf::Vector3f( + static_cast(ray.z * sin(-1.57) + ray.x * cos(-1.57)), + static_cast(ray.y), + static_cast(ray.z * cos(-1.57) - ray.x * sin(-1.57)) + ); + + int index = (x + view_res.x / 2) + view_res.x * (y + view_res.y / 2); + ray = Normalize(ray); + + viewport_matrix[index] = sf::Vector4f( + ray.x, + ray.y, + ray.z, + 0 + ); + } + } +} + +int Hardware_Caster::acquire_platform_and_device() { + + // Get the number of platforms + cl_uint plt_cnt = 0; + clGetPlatformIDs(0, nullptr, &plt_cnt); + + // Fetch the platforms + std::map> plt_ids; + + // buffer before map init + std::vector plt_buf(plt_cnt); + clGetPlatformIDs(plt_cnt, plt_buf.data(), nullptr); + + // Map init + for (auto id : plt_buf) { + plt_ids.emplace(std::make_pair(id, std::vector())); + } + + // For each platform, populate its devices + for (unsigned int i = 0; i < plt_cnt; i++) { + + cl_uint deviceIdCount = 0; + error = clGetDeviceIDs(plt_buf[i], CL_DEVICE_TYPE_ALL, 0, nullptr, &deviceIdCount); + + // Check to see if we even have opencl on this machine + if (deviceIdCount == 0) { + std::cout << "There appears to be no platforms supporting opencl" << std::endl; + return OPENCL_NOT_SUPPORTED; + } + + // Get the device ids + std::vector deviceIds(deviceIdCount); + error = clGetDeviceIDs(plt_buf[i], CL_DEVICE_TYPE_ALL, deviceIdCount, deviceIds.data(), NULL); + + if (assert(error, "clGetDeviceIDs")) + return OPENCL_ERROR; + + for (unsigned int q = 0; q < deviceIdCount; q++) { + + device d; + + d.id = deviceIds[q]; + + clGetDeviceInfo(d.id, CL_DEVICE_PLATFORM, sizeof(cl_platform_id), &d.platform, NULL); + clGetDeviceInfo(d.id, CL_DEVICE_VERSION, sizeof(char) * 128, &d.version, NULL); + clGetDeviceInfo(d.id, CL_DEVICE_TYPE, sizeof(cl_device_type), &d.type, NULL); + clGetDeviceInfo(d.id, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(cl_uint), &d.clock_frequency, NULL); + clGetDeviceInfo(d.id, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &d.comp_units, NULL); + + plt_ids.at(d.platform).push_back(d); + } + } + + + // The devices how now been queried we want to shoot for a gpu with the fastest clock, + // falling back to the cpu with the fastest clock if we weren't able to find one + + device current_best_device; + current_best_device.type = 0; // Set this to 0 so the first run always selects a new device + current_best_device.clock_frequency = 0; + current_best_device.comp_units = 0; + + + for (auto kvp : plt_ids) { + + for (auto device : kvp.second) { + + // Gonna just split this up into cases. There are so many devices I cant test with + // that opencl supports. I'm not going to waste my time making a generic implimentation + + // Upon success of a condition, set the current best device values + + if (device.type == CL_DEVICE_TYPE_GPU && current_best_device.type != CL_DEVICE_TYPE_GPU) { + current_best_device = device; + } + else if (device.comp_units > current_best_device.comp_units) { + current_best_device = device; + } + else if (current_best_device.type != CL_DEVICE_TYPE_GPU && device.clock_frequency > current_best_device.clock_frequency) { + current_best_device = device; + } + } + } + + platform_id = current_best_device.platform; + device_id = current_best_device.id; + + return 1; +}; + +int Hardware_Caster::create_shared_context() { + + // Hurray for standards! + // Setup the context properties to grab the current GL context + +#ifdef linux + + cl_context_properties context_properties[] = { + CL_GL_CONTEXT_KHR, (cl_context_properties)glXGetCurrentContext(), + CL_GLX_DISPLAY_KHR, (cl_context_properties)glXGetCurrentDisplay(), + CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, + 0 + }; + +#elif defined _WIN32 + + HGLRC hGLRC = wglGetCurrentContext(); + HDC hDC = wglGetCurrentDC(); + cl_context_properties context_properties[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, + CL_GL_CONTEXT_KHR, (cl_context_properties)hGLRC, + CL_WGL_HDC_KHR, (cl_context_properties)hDC, + 0 + }; + + +#elif defined TARGET_OS_MAC + + CGLContextObj glContext = CGLGetCurrentContext(); + CGLShareGroupObj shareGroup = CGLGetShareGroup(glContext); + cl_context_properties context_properties[] = { + CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, + (cl_context_properties)shareGroup, + 0 + }; + +#endif + + // Create our shared context + context = clCreateContext( + context_properties, + 1, + &device_id, + nullptr, nullptr, + &error + ); + + if (assert(error, "clCreateContext")) + return OPENCL_ERROR; + + return 1; +} + +int Hardware_Caster::create_command_queue() { + + // If context and device_id have initialized + if (context && device_id) { + + command_queue = clCreateCommandQueue(context, device_id, 0, &error); + + if (assert(error, "clCreateCommandQueue")) + return OPENCL_ERROR; + + return 1; + } + else { + std::cout << "Failed creating the command queue. Context or device_id not initialized"; + return OPENCL_ERROR; + } +} + +int Hardware_Caster::check_cl_khr_gl_sharing() { + + // Test for sharing + size_t ext_str_size = 1024; + char *ext_str = new char[ext_str_size]; + clGetDeviceInfo(device_id, CL_DEVICE_EXTENSIONS, ext_str_size, ext_str, &ext_str_size); + + if (std::string(ext_str).find("cl_khr_gl_sharing") == std::string::npos) { + std::cout << "No support for the cl_khr_gl_sharing extension"; + delete ext_str; + return RayCaster::SHARING_NOT_SUPPORTED; + } + + delete ext_str; + return 1; +} + +int Hardware_Caster::compile_kernel(std::string kernel_source, bool is_path, std::string kernel_name) { + + const char* source; + std::string tmp; + + if (is_path) { + //Load in the kernel, and c stringify it + tmp = read_file(kernel_source); + source = tmp.c_str(); + } + else { + source = kernel_source.c_str(); + } + + size_t kernel_source_size = strlen(source); + + // Load the source into CL's data structure + + cl_program program = clCreateProgramWithSource( + context, 1, + &source, + &kernel_source_size, &error + ); + + // This is not for compilation, it only loads the source + if (assert(error, "clCreateProgramWithSource")) + return OPENCL_ERROR; + + + // Try and build the program + error = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); + + // Check to see if it errored out + if (assert(error, "clBuildProgram")) { + + // Get the size of the queued log + size_t log_size; + clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + char *log = new char[log_size]; + + // Grab the log + clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL); + + std::cout << log; + return OPENCL_ERROR; + } + + // Done initializing the kernel + cl_kernel kernel = clCreateKernel(program, kernel_name.c_str(), &error); + + if (assert(error, "clCreateKernel")) + return OPENCL_ERROR; + + // Do I want these to overlap when repeated?? + kernel_map[kernel_name] = kernel; + //kernel_map.emplace(std::make_pair(kernel_name, kernel)); + + return 1; +} + +int Hardware_Caster::set_kernel_arg( + std::string kernel_name, + int index, + std::string buffer_name) { + + error = clSetKernelArg( + kernel_map.at(kernel_name), + index, + sizeof(cl_mem), + (void *)&buffer_map.at(buffer_name)); + + if (assert(error, "clSetKernelArg")) + return OPENCL_ERROR; + + return 0; + +} + +int Hardware_Caster::create_image_buffer(std::string buffer_name, cl_uint size, sf::Texture* texture) { + + // I can imagine overwriting buffers will be common, so I think + // this is safe to overwrite / release old buffers quietly + if (buffer_map.count(buffer_name) > 0) { + release_buffer(buffer_name); + } + + int error; + cl_mem buff = clCreateFromGLTexture( + getContext(), CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, + 0, texture->getNativeHandle(), &error); + + if (assert(error, "clCreateFromGLTexture")) + return OPENCL_ERROR; + + store_buffer(buff, buffer_name); + + return 1; +} + +int Hardware_Caster::create_buffer(std::string buffer_name, cl_uint size, void* data, cl_mem_flags flags) { + + // I can imagine overwriting buffers will be common, so I think + // this is safe to overwrite / release old buffers quietly + if (buffer_map.count(buffer_name) > 0) { + release_buffer(buffer_name); + } + + cl_mem buff = clCreateBuffer( + getContext(), flags, + size, data, &error + ); + + if (assert(error, "clCreateBuffer")) + return OPENCL_ERROR; + + store_buffer(buff, buffer_name); + + return 1; + +} + +int Hardware_Caster::create_buffer(std::string buffer_name, cl_uint size, void* data) { + + // I can imagine overwriting buffers will be common, so I think + // this is safe to overwrite / release old buffers quietly + if (buffer_map.count(buffer_name) > 0) { + release_buffer(buffer_name); + } + + cl_mem buff = clCreateBuffer( + getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + size, data, &error + ); + + if (assert(error, "clCreateBuffer")) + return OPENCL_ERROR; + + store_buffer(buff, buffer_name); + + return 1; + +} + +int Hardware_Caster::release_buffer(std::string buffer_name) { + + if (buffer_map.count(buffer_name) > 0) { + + int error = clReleaseMemObject(buffer_map.at(buffer_name)); + + if (assert(error, "clReleaseMemObject")) { + std::cout << "Error releasing buffer : " << buffer_name; + std::cout << "Buffer not removed"; + return -1; + + } else { + buffer_map.erase(buffer_name); + } + + } else { + std::cout << "Error releasing buffer : " << buffer_name; + std::cout << "Buffer not found"; + return -1; + } + + return 1; + +} + +int Hardware_Caster::store_buffer(cl_mem buffer, std::string buffer_name) { + buffer_map.emplace(std::make_pair(buffer_name, buffer)); + return 1; +} + +int Hardware_Caster::run_kernel(std::string kernel_name, const int work_size) { + + size_t global_work_size[1] = { static_cast(work_size) }; + + cl_kernel kernel = kernel_map.at(kernel_name); + + error = clEnqueueAcquireGLObjects(getCommandQueue(), 1, &buffer_map.at("image"), 0, 0, 0); + if (assert(error, "clEnqueueAcquireGLObjects")) + return OPENCL_ERROR; + + //error = clEnqueueTask(command_queue, kernel, 0, NULL, NULL); + error = clEnqueueNDRangeKernel( + command_queue, kernel, + 1, NULL, global_work_size, + NULL, 0, NULL, NULL); + + if (assert(error, "clEnqueueNDRangeKernel")) + return OPENCL_ERROR; + + clFinish(getCommandQueue()); + + // What if errors out and gl objects are never released? + error = clEnqueueReleaseGLObjects(getCommandQueue(), 1, &buffer_map.at("image"), 0, NULL, NULL); + if (assert(error, "clEnqueueReleaseGLObjects")) + return OPENCL_ERROR; + + return 1; +} + +void Hardware_Caster::print_kernel_arguments() +{ + compile_kernel("../kernels/print_arguments.cl", true, "printer"); + set_kernel_arg("printer", 0, "map"); + set_kernel_arg("printer", 1, "map_dimensions"); + set_kernel_arg("printer", 2, "viewport_resolution"); + set_kernel_arg("printer", 3, "viewport_matrix"); + set_kernel_arg("printer", 4, "camera_direction"); + set_kernel_arg("printer", 5, "camera_position"); + set_kernel_arg("printer", 6, "lights"); + set_kernel_arg("printer", 7, "light_count"); + set_kernel_arg("printer", 8, "image"); + + run_kernel("printer", 1); +} + +cl_device_id Hardware_Caster::getDeviceID() { return device_id; }; +cl_platform_id Hardware_Caster::getPlatformID() { return platform_id; }; +cl_context Hardware_Caster::getContext() { return context; }; +cl_kernel Hardware_Caster::getKernel(std::string kernel_name) { return kernel_map.at(kernel_name); }; +cl_command_queue Hardware_Caster::getCommandQueue() { return command_queue; }; + +bool Hardware_Caster::assert(int error_code, std::string function_name) { + + // Just gonna do a little jump table here, just error codes so who cares + std::string err_msg = "Error : "; + + switch (error_code) { + + case CL_SUCCESS: + return false; + + case 1: + return false; + + case CL_DEVICE_NOT_FOUND: + err_msg += "CL_DEVICE_NOT_FOUND"; + break; + case CL_DEVICE_NOT_AVAILABLE: + err_msg = "CL_DEVICE_NOT_AVAILABLE"; + break; + case CL_COMPILER_NOT_AVAILABLE: + err_msg = "CL_COMPILER_NOT_AVAILABLE"; + break; + case CL_MEM_OBJECT_ALLOCATION_FAILURE: + err_msg = "CL_MEM_OBJECT_ALLOCATION_FAILURE"; + break; + case CL_OUT_OF_RESOURCES: + err_msg = "CL_OUT_OF_RESOURCES"; + break; + case CL_OUT_OF_HOST_MEMORY: + err_msg = "CL_OUT_OF_HOST_MEMORY"; + break; + case CL_PROFILING_INFO_NOT_AVAILABLE: + err_msg = "CL_PROFILING_INFO_NOT_AVAILABLE"; + break; + case CL_MEM_COPY_OVERLAP: + err_msg = "CL_MEM_COPY_OVERLAP"; + break; + case CL_IMAGE_FORMAT_MISMATCH: + err_msg = "CL_IMAGE_FORMAT_MISMATCH"; + break; + case CL_IMAGE_FORMAT_NOT_SUPPORTED: + err_msg = "CL_IMAGE_FORMAT_NOT_SUPPORTED"; + break; + case CL_BUILD_PROGRAM_FAILURE: + err_msg = "CL_BUILD_PROGRAM_FAILURE"; + break; + case CL_MAP_FAILURE: + err_msg = "CL_MAP_FAILURE"; + break; + case CL_MISALIGNED_SUB_BUFFER_OFFSET: + err_msg = "CL_MISALIGNED_SUB_BUFFER_OFFSET"; + break; + case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: + err_msg = "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; + break; + case CL_COMPILE_PROGRAM_FAILURE: + err_msg = "CL_COMPILE_PROGRAM_FAILURE"; + break; + case CL_LINKER_NOT_AVAILABLE: + err_msg = "CL_LINKER_NOT_AVAILABLE"; + break; + case CL_LINK_PROGRAM_FAILURE: + err_msg = "CL_LINK_PROGRAM_FAILURE"; + break; + case CL_DEVICE_PARTITION_FAILED: + err_msg = "CL_DEVICE_PARTITION_FAILED"; + break; + case CL_KERNEL_ARG_INFO_NOT_AVAILABLE: + err_msg = "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; + break; + case CL_INVALID_VALUE: + err_msg = "CL_INVALID_VALUE"; + break; + case CL_INVALID_DEVICE_TYPE: + err_msg = "CL_INVALID_DEVICE_TYPE"; + break; + case CL_INVALID_PLATFORM: + err_msg = "CL_INVALID_PLATFORM"; + break; + case CL_INVALID_DEVICE: + err_msg = "CL_INVALID_DEVICE"; + break; + case CL_INVALID_CONTEXT: + err_msg = "CL_INVALID_CONTEXT"; + break; + case CL_INVALID_QUEUE_PROPERTIES: + err_msg = "CL_INVALID_QUEUE_PROPERTIES"; + break; + case CL_INVALID_COMMAND_QUEUE: + err_msg = "CL_INVALID_COMMAND_QUEUE"; + break; + case CL_INVALID_HOST_PTR: + err_msg = "CL_INVALID_HOST_PTR"; + break; + case CL_INVALID_MEM_OBJECT: + err_msg = "CL_INVALID_MEM_OBJECT"; + break; + case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: + err_msg = "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; + break; + case CL_INVALID_IMAGE_SIZE: + err_msg = "CL_INVALID_IMAGE_SIZE"; + break; + case CL_INVALID_SAMPLER: + err_msg = "CL_INVALID_SAMPLER"; + break; + case CL_INVALID_BINARY: + err_msg = "CL_INVALID_BINARY"; + break; + case CL_INVALID_BUILD_OPTIONS: + err_msg = "CL_INVALID_BUILD_OPTIONS"; + break; + case CL_INVALID_PROGRAM: + err_msg = "CL_INVALID_PROGRAM"; + break; + case CL_INVALID_PROGRAM_EXECUTABLE: + err_msg = "CL_INVALID_PROGRAM_EXECUTABLE"; + break; + case CL_INVALID_KERNEL_NAME: + err_msg = "CL_INVALID_KERNEL_NAME"; + break; + case CL_INVALID_KERNEL_DEFINITION: + err_msg = "CL_INVALID_KERNEL_DEFINITION"; + break; + case CL_INVALID_KERNEL: + err_msg = "CL_INVALID_KERNEL"; + break; + case CL_INVALID_ARG_INDEX: + err_msg = "CL_INVALID_ARG_INDEX"; + break; + case CL_INVALID_ARG_VALUE: + err_msg = "CL_INVALID_ARG_VALUE"; + break; + case CL_INVALID_ARG_SIZE: + err_msg = "CL_INVALID_ARG_SIZE"; + break; + case CL_INVALID_KERNEL_ARGS: + err_msg = "CL_INVALID_KERNEL_ARGS"; + break; + case CL_INVALID_WORK_DIMENSION: + err_msg = "CL_INVALID_WORK_DIMENSION"; + break; + case CL_INVALID_WORK_GROUP_SIZE: + err_msg = "CL_INVALID_WORK_GROUP_SIZE"; + break; + case CL_INVALID_WORK_ITEM_SIZE: + err_msg = "CL_INVALID_WORK_ITEM_SIZE"; + break; + case CL_INVALID_GLOBAL_OFFSET: + err_msg = "CL_INVALID_GLOBAL_OFFSET"; + break; + case CL_INVALID_EVENT_WAIT_LIST: + err_msg = "CL_INVALID_EVENT_WAIT_LIST"; + break; + case CL_INVALID_EVENT: + err_msg = "CL_INVALID_EVENT"; + break; + case CL_INVALID_OPERATION: + err_msg = "CL_INVALID_OPERATION"; + break; + case CL_INVALID_GL_OBJECT: + err_msg = "CL_INVALID_GL_OBJECT"; + break; + case CL_INVALID_BUFFER_SIZE: + err_msg = "CL_INVALID_BUFFER_SIZE"; + break; + case CL_INVALID_MIP_LEVEL: + err_msg = "CL_INVALID_MIP_LEVEL"; + break; + case CL_INVALID_GLOBAL_WORK_SIZE: + err_msg = "CL_INVALID_GLOBAL_WORK_SIZE"; + break; + case CL_INVALID_PROPERTY: + err_msg = "CL_INVALID_PROPERTY"; + break; + case CL_INVALID_IMAGE_DESCRIPTOR: + err_msg = "CL_INVALID_IMAGE_DESCRIPTOR"; + break; + case CL_INVALID_COMPILER_OPTIONS: + err_msg = "CL_INVALID_COMPILER_OPTIONS"; + break; + case CL_INVALID_LINKER_OPTIONS: + err_msg = "CL_INVALID_LINKER_OPTIONS"; + break; + case CL_INVALID_DEVICE_PARTITION_COUNT: + err_msg = "CL_INVALID_DEVICE_PARTITION_COUNT"; + break; + case RayCaster::SHARING_NOT_SUPPORTED: + err_msg = "SHARING_NOT_SUPPORTED"; + break; + case RayCaster::OPENCL_NOT_SUPPORTED: + err_msg = "OPENCL_NOT_SUPPORTED"; + break; + case RayCaster::OPENCL_ERROR: + err_msg = "OPENCL_ERROR"; + break; + case RayCaster::ERR: + err_msg = "ERROR"; + break; + } + + std::cout << err_msg << " =at= " << function_name << std::endl; + return true; +} diff --git a/src/raycaster/RayCaster.cpp b/src/raycaster/RayCaster.cpp new file mode 100644 index 0000000..721b4cb --- /dev/null +++ b/src/raycaster/RayCaster.cpp @@ -0,0 +1,7 @@ +#include "RayCaster.h" + +RayCaster::RayCaster() { +} + +RayCaster::~RayCaster() { +} \ No newline at end of file diff --git a/src/raycaster/Software_Caster.cpp b/src/raycaster/Software_Caster.cpp new file mode 100644 index 0000000..5e664b4 --- /dev/null +++ b/src/raycaster/Software_Caster.cpp @@ -0,0 +1,343 @@ +#include "Software_Caster.h" + + + +Software_Caster::Software_Caster() +{ +} + + +Software_Caster::~Software_Caster() +{ +} + +int Software_Caster::init() +{ + return 1; +} + +void Software_Caster::create_viewport(int width, int height, float v_fov, float h_fov) +{ + // CL needs the screen resolution + viewport_resolution = sf::Vector2i(width, height); + + // And an array of vectors describing the way the "lens" of our + // camera works + // This could be modified to make some odd looking camera lenses + + double y_increment_radians = DegreesToRadians(v_fov / viewport_resolution.y); + double x_increment_radians = DegreesToRadians(h_fov / viewport_resolution.x); + + viewport_matrix = new sf::Vector4f[width * height * 4]; + + for (int y = -viewport_resolution.y / 2; y < viewport_resolution.y / 2; y++) { + for (int x = -viewport_resolution.x / 2; x < viewport_resolution.x / 2; x++) { + + // The base ray direction to slew from + sf::Vector3f ray(1, 0, 0); + + // Y axis, pitch + ray = sf::Vector3f( + static_cast(ray.z * sin(y_increment_radians * y) + ray.x * cos(y_increment_radians * y)), + static_cast(ray.y), + static_cast(ray.z * cos(y_increment_radians * y) - ray.x * sin(y_increment_radians * y)) + ); + + + // Z axis, yaw + ray = sf::Vector3f( + static_cast(ray.x * cos(x_increment_radians * x) - ray.y * sin(x_increment_radians * x)), + static_cast(ray.x * sin(x_increment_radians * x) + ray.y * cos(x_increment_radians * x)), + static_cast(ray.z) + ); + + int index = (x + viewport_resolution.x / 2) + viewport_resolution.x * (y + viewport_resolution.y / 2); + ray = Normalize(ray); + + viewport_matrix[index] = sf::Vector4f( + ray.x, + ray.y, + ray.z, + 0 + ); + } + } + + // Create the image that opencl's rays write to + viewport_image = new sf::Uint8[width * height * 4]; + + for (int i = 0; i < width * height * 4; i += 4) { + + viewport_image[i] = 255; // R + viewport_image[i + 1] = 255; // G + viewport_image[i + 2] = 255; // B + viewport_image[i + 3] = 255; // A + } + + // Interop lets us keep a reference to it as a texture + viewport_texture.create(width, height); + viewport_texture.update(viewport_image); + viewport_sprite.setTexture(viewport_texture); + + +} + +void Software_Caster::assign_lights(std::vector *data) { + +// this->lights = data; + + int light_count = static_cast(data->size()); +} + +void Software_Caster::assign_map(Old_Map * map) { + this->map = map; +} + +void Software_Caster::assign_camera(Camera * camera) { + this->camera = camera; +} + +void Software_Caster::validate() { + // Check to make sure everything has been entered; + if (camera == nullptr || + map == nullptr || + viewport_image == nullptr || + viewport_matrix == nullptr) { + + std::cout << "Raycaster.validate() failed, camera, map, or viewport not initialized"; + + } +} + +void Software_Caster::compute() { + cast_viewport(); +} + +void Software_Caster::draw(sf::RenderWindow * window) { + viewport_texture.update(viewport_image); + window->draw(viewport_sprite); +} + +void Software_Caster::cast_viewport() { + + std::vector threads; + for (int i = 0; i < 13; i++) { + int s = viewport_resolution.x * ((viewport_resolution.y / 13) * i); + int e = viewport_resolution.x * ((viewport_resolution.y / 13) * (i + 1)); + threads.push_back(new std::thread(&Software_Caster::cast_thread, this, s, e)); + } + + for (auto i : threads) { + i->join(); + delete i; + } +} + +void Software_Caster::cast_thread(int start_id, int end_id) { + + for (int i = start_id; i < end_id; i++) { + cast_ray(i); + } + +} + +void Software_Caster::cast_ray(int id) +{ + sf::Vector2i pixel = { id % viewport_resolution.x, id / viewport_resolution.x }; + + // 4f 3f ?? + sf::Vector4f ray_dir = viewport_matrix[pixel.x + viewport_resolution.x * pixel.y]; + + ray_dir = sf::Vector4f( + ray_dir.z * sin(camera->get_direction().x) + ray_dir.x * cos(camera->get_direction().x), + ray_dir.y, + ray_dir.z * cos(camera->get_direction().x) - ray_dir.x * sin(camera->get_direction().x), + 0 + ); + + ray_dir = sf::Vector4f( + ray_dir.x * cos(camera->get_direction().y) - ray_dir.y * sin(camera->get_direction().y), + ray_dir.x * sin(camera->get_direction().y) + ray_dir.y * cos(camera->get_direction().y), + ray_dir.z, + 0 + ); + + // Setup the voxel step based on what direction the ray is pointing + sf::Vector3i voxel_step = sf::Vector3i( + static_cast(1 * (abs(ray_dir.x) / ray_dir.x)), + static_cast(1 * (abs(ray_dir.y) / ray_dir.y)), + static_cast(1 * (abs(ray_dir.z) / ray_dir.z)) + ); + + // Setup the voxel coords from the camera origin + sf::Vector3i voxel = sf::Vector3i( + static_cast(camera->get_position().x), + static_cast(camera->get_position().y), + static_cast(camera->get_position().z) + ); + + // Delta T is the units a ray must travel along an axis in order to + // traverse an integer split + sf::Vector3f delta_t = sf::Vector3f( + fabs(1.0f / ray_dir.x), + fabs(1.0f / ray_dir.y), + fabs(1.0f / ray_dir.z) + ); + + // offset is how far we are into a voxel, enables sub voxel movement + sf::Vector3f offset = sf::Vector3f( + (camera->get_position().x - floor(camera->get_position().x)) * voxel_step.x, + (camera->get_position().y - floor(camera->get_position().y)) * voxel_step.y, + (camera->get_position().z - floor(camera->get_position().z)) * voxel_step.z + ); + + // Intersection T is the collection of the next intersection points + // for all 3 axis XYZ. + sf::Vector3f intersection_t = sf::Vector3f( + delta_t.x * offset.x, + delta_t.y * offset.y, + delta_t.z * offset.z + ); + + // 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; + } + + // use a ghetto ass rng to give rays a "fog" appearance + sf::Vector2i randoms = { 3, 14 }; + int seed = randoms.x + id; + int t = seed ^ (seed << 11); + int result = randoms.y ^ (randoms.y >> 19) ^ (t ^ (t >> 8)); + + int max_dist = 800 + result % 50; + int dist = 0; + + sf::Vector3i mask = { 0, 0, 0 }; + + // Andrew Woo's raycasting algo + do { + + if ((intersection_t.x) < (intersection_t.y)) { + if ((intersection_t.x) < (intersection_t.z)) { + + mask.x = 1; + voxel.x += voxel_step.x; + intersection_t.x = intersection_t.x + delta_t.x; + } + else { + + mask.z = 1; + voxel.z += voxel_step.z; + intersection_t.z = intersection_t.z + delta_t.z; + } + } + else { + if ((intersection_t.y) < (intersection_t.z)) { + + mask.y = 1; + voxel.y += voxel_step.y; + intersection_t.y = intersection_t.y + delta_t.y; + } + else { + + mask.z = 1; + voxel.z += voxel_step.z; + intersection_t.z = intersection_t.z + delta_t.z; + } + } + + + // If the ray went out of bounds + sf::Vector3i overshoot = sf::Vector3i( + voxel.x <= map->getDimensions().x, + voxel.y <= map->getDimensions().y, + voxel.z <= map->getDimensions().z + ); + + sf::Vector3i undershoot = sf::Vector3i( + voxel.x > 0, + voxel.y > 0, + voxel.z > 0 + ); + + if (overshoot.x == 0 || overshoot.y == 0 || overshoot.z == 0 || undershoot.x == 0 || undershoot.y == 0) { + blit_pixel(sf::Color::Yellow, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + } + if (undershoot.z == 0) { + blit_pixel(sf::Color::Yellow, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + } + + // If we hit a voxel + //int index = voxel.x * (*map_dim).y * (*map_dim).z + voxel.z * (*map_dim).z + voxel.y; + // Why the off by one on voxel.y? + int index = voxel.x + map->getDimensions().x * (voxel.y + map->getDimensions().z * (voxel.z - 1)); + int voxel_data = map->get_voxel_data()[index]; + + if (voxel_data != 0) { + switch (voxel_data) { + case 1: + blit_pixel(sf::Color::Green, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + case 2: + blit_pixel(sf::Color::Green, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + case 3: + blit_pixel(sf::Color::Green, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + case 4: + blit_pixel(sf::Color::Green, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + case 5: + blit_pixel(sf::Color(30, 10, 200, 100), sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + case 6: + blit_pixel(sf::Color::Green, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; + default: + //write_imagef(image, pixel, (float4)(.30, .2550, .2550, 255.00)); + return; + } + } + + dist++; + } while (dist < max_dist); + + blit_pixel(sf::Color::Red, sf::Vector2i{ pixel.x,pixel.y }, mask); + return; +} + +void Software_Caster::blit_pixel(sf::Color color, sf::Vector2i position, sf::Vector3i mask) { + + sf::Color t = global_light(color, mask); + viewport_image[(position.x + viewport_resolution.x * position.y) * 4 + 0] = t.r; + viewport_image[(position.x + viewport_resolution.x * position.y) * 4 + 1] = t.g; + viewport_image[(position.x + viewport_resolution.x * position.y) * 4 + 2] = t.b; + viewport_image[(position.x + viewport_resolution.x * position.y) * 4 + 3] = t.a; +} + +sf::Color Software_Caster::global_light(sf::Color in, sf::Vector3i mask) { + + // I think I may scrap this whole software fallback caster thing + + //sf::Vector3f mask_f(mask); + + //in.a = in.a + (int)acos( + // DotProduct( + // Normalize(lights->at(0).direction_cartesian), + // Normalize(mask_f) + // ) + // )/ 2; + + return in; + +} \ No newline at end of file