diff --git a/assets/screenshot.PNG b/assets/screenshot.PNG index 0e1ca7b..c688dcb 100644 Binary files a/assets/screenshot.PNG and b/assets/screenshot.PNG differ diff --git a/assets/screenshot1.png b/assets/screenshot1.png index c688dcb..cd8c29d 100644 Binary files a/assets/screenshot1.png and b/assets/screenshot1.png differ diff --git a/assets/textures/minecraft_tiles.png b/assets/textures/minecraft_tiles.png new file mode 100644 index 0000000..d7045db Binary files /dev/null and b/assets/textures/minecraft_tiles.png differ diff --git a/assets/textures/spritesheet_tiles.png b/assets/textures/spritesheet_tiles.png new file mode 100644 index 0000000..b3ac99b Binary files /dev/null and b/assets/textures/spritesheet_tiles.png differ diff --git a/include/Hardware_Caster.h b/include/Hardware_Caster.h index 97b24cb..40997b9 100644 --- a/include/Hardware_Caster.h +++ b/include/Hardware_Caster.h @@ -49,12 +49,16 @@ public: // 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 *lights) 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; @@ -74,7 +78,7 @@ private: int check_cl_khr_gl_sharing(); - int create_image_buffer(std::string buffer_name, cl_uint size, void* data); + 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); diff --git a/kernels/ray_caster_kernel.cl b/kernels/ray_caster_kernel.cl index fdf4839..e496356 100644 --- a/kernels/ray_caster_kernel.cl +++ b/kernels/ray_caster_kernel.cl @@ -23,13 +23,13 @@ float4 white_light(float4 input, float3 light, int3 mask) { float4 view_light(float4 in_color, float3 light, float3 view, int3 mask) { float diffuse = max(dot(normalize(convert_float3(mask)), normalize(light)), 0.0f); - in_color += diffuse * 0.5; + in_color += diffuse * 0.2; if (dot(light, normalize(convert_float3(mask))) > 0.0) { float3 halfwayVector = normalize(normalize(light) + normalize(view)); float specTmp = max(dot(normalize(convert_float3(mask)), halfwayVector), 0.0f); - in_color += pow(specTmp, 1.0f) * 0.1; + in_color += pow(specTmp, 1.0f) * 0.5; } //in_color += 0.02; @@ -134,7 +134,10 @@ __kernel void raycaster( global float* lights, global int* light_count, __write_only image2d_t image, - global int* seed_memory + global int* seed_memory, + __read_only image2d_t texture_atlas, + global int2 *atlas_dim, + global int2 *tile_dim ){ int global_id = get_global_id(0); @@ -230,27 +233,23 @@ __kernel void raycaster( if (voxel_data != 0) { + // Determine where on the 2d plane the ray intersected - if (voxel_data == 6) { - voxel_color = (float4)(0.0, 0.239, 0.419, 0.3); - } - else if (voxel_data == 5) { - voxel_color = (float4)(0.25, 0.52, 0.30, 0.1); - } - else if (voxel_data == 1) { - voxel_color = (float4)(0.929, 0.957, 0.027, 0.7); - } - - // set to which face float3 face_position = (float)(0); + float2 texture_position = (float)(0); + // First determine the percent of the way the ray is towards the next intersection_t + // in relation to the xyz position on the plane if (face_mask.x == -1) { float z_percent = (intersection_t.z - (intersection_t.x - delta_t.x)) / delta_t.z; float y_percent = (intersection_t.y - (intersection_t.x - delta_t.x)) / delta_t.y; + // Since we intersected face x, we know that we are at the face (1.0) + // Not entirely sure what is causing the 1.0 vs 1.001 rendering bug face_position = (float3)(1.001f, y_percent, z_percent); + texture_position = (float2)(y_percent, z_percent); } else if (face_mask.y == -1) { @@ -258,7 +257,7 @@ __kernel void raycaster( float z_percent = (intersection_t.z - (intersection_t.y - delta_t.y)) / delta_t.z; face_position = (float3)(x_percent, 1.001f, z_percent); - + texture_position = (float2)(x_percent, z_percent); } else if (face_mask.z == -1) { @@ -267,28 +266,55 @@ __kernel void raycaster( float y_percent = (intersection_t.y - (intersection_t.z - delta_t.z)) / delta_t.y; face_position = (float3)(x_percent, y_percent, 1.001f); + texture_position = (float2)(x_percent, y_percent); } + // We now need to account for the ray wanting to skip the axis in which + // it flips its sign - if (ray_dir.x > 0) - face_position.x = - face_position.x + 1; + // TODO: improve this - if (ray_dir.x < 0) - face_position.x = face_position.x + 0; + if (ray_dir.x > 0) { + face_position.x = -face_position.x + 1; + texture_position.x = -texture_position.x + 1.0; + } + //if (ray_dir.x < 0) + // face_position.x = face_position.x + 0; - if (ray_dir.y > 0) + if (ray_dir.y > 0){ face_position.y = - face_position.y + 1; + texture_position.y = -texture_position.y + 1.0; + } - if (ray_dir.y < 0) - face_position.y = face_position.y + 0; + //if (ray_dir.y < 0) + // face_position.y = face_position.y + 0; - if (ray_dir.z > 0) + if (ray_dir.z > 0) { face_position.z = - face_position.z + 1; + texture_position.y = -texture_position.y + 1.0; + } - if (ray_dir.z < 0) - face_position.z = face_position.z + 0; + //if (ray_dir.z < 0) + // face_position.z = face_position.z + 0; + + + // Now either use the face position to retrieve a texture sample, or + // just a plain color for the voxel color + + if (voxel_data == 6) { + voxel_color = (float4)(0.0, 0.239, 0.419, 0.3); + } + else if (voxel_data == 5) { + float2 tile_size = convert_float2(*atlas_dim / *tile_dim); + voxel_color = read_imagef(texture_atlas, convert_int2(texture_position * tile_size) + convert_int2((float2)(3, 0) * tile_size)); + //voxel_color = (float4)(0.25, 0.52, 0.30, 0.1); + } + else if (voxel_data == 1) { + voxel_color = (float4)(0.929, 0.957, 0.027, 0.7); + } + // if (cast_light_intersection_ray( map, diff --git a/src/Hardware_Caster.cpp b/src/Hardware_Caster.cpp index 02e72d8..ffbe5dd 100644 --- a/src/Hardware_Caster.cpp +++ b/src/Hardware_Caster.cpp @@ -87,6 +87,9 @@ void Hardware_Caster::validate() 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(); } @@ -94,6 +97,18 @@ void Hardware_Caster::validate() } +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? @@ -177,7 +192,7 @@ void Hardware_Caster::create_viewport(int width, int height, float v_fov, float viewport_sprite.setTexture(viewport_texture); // Pass the buffer to opencl - create_image_buffer("image", sizeof(sf::Uint8) * width * height * 4, viewport_image); + create_image_buffer("image", sizeof(sf::Uint8) * width * height * 4, &viewport_texture); } @@ -511,7 +526,7 @@ int Hardware_Caster::set_kernel_arg( } -int Hardware_Caster::create_image_buffer(std::string buffer_name, cl_uint size, void* data) { +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 @@ -522,7 +537,7 @@ int Hardware_Caster::create_image_buffer(std::string buffer_name, cl_uint size, int error; cl_mem buff = clCreateFromGLTexture( getContext(), CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, - 0, viewport_texture.getNativeHandle(), &error); + 0, texture->getNativeHandle(), &error); if (assert(error, "clCreateFromGLTexture")) return OPENCL_ERROR; diff --git a/src/main.cpp b/src/main.cpp index 9df2995..36091fb 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -124,10 +124,17 @@ int main() { // *links* the lights to the GPU raycaster->assign_lights(&light_vec); - // Checks to see if proper data was uploaded, then sets the kernel args - raycaster->validate(); + + // Load in the spritesheet texture + sf::Texture spritesheet; + spritesheet.loadFromFile("../assets/textures/minecraft_tiles.png"); + spritesheet.getNativeHandle(); + raycaster->create_texture_atlas(&spritesheet, sf::Vector2i(16, 16)); + // Checks to see if proper data was uploaded, then sets the kernel args + // ALL DATA LOADING MUST BE FINISHED + raycaster->validate(); // ========== DEBUG ========== fps_counter fps;