From 1f342dc8846adbf50bbb7a029bd1066477eb6c47 Mon Sep 17 00:00:00 2001 From: James Date: Mon, 29 Jan 2018 13:45:04 +0000 Subject: [PATCH] fun with raytracing --- fluid.cl | 54 ++++++++++++++++++++++++++++++++++++++++++++++++++++ lighting.hpp | 48 ++++++++++++++++++++++++++++++++++++++++++++++ main.cpp | 16 ++++++++++++---- 3 files changed, 114 insertions(+), 4 deletions(-) create mode 100644 lighting.hpp diff --git a/fluid.cl b/fluid.cl index 6461478..6a25b5a 100644 --- a/fluid.cl +++ b/fluid.cl @@ -449,3 +449,57 @@ void wavelet_upscale(__read_only image2d_t w_of_in, __read_only image2d_t veloci ///so we wanna advect in this kernel. Its very expensive to store intermediate data ///So: Go back in time, and generate the velocity there as well so we can store that? } + +__kernel +void lighting_raytrace_point(float2 point, float radius, int num_tracers, __read_only image2d_t fluid, __write_only image2d_t screen) +{ + sampler_t sam = CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR; + + int gid = get_global_id(0); + + if(gid >= num_tracers) + return; + + int sw = get_image_width(screen); + int sh = get_image_height(screen); + + float angle = 2 * M_PI * (float)gid / num_tracers; + + float2 dir = (float2){cos(angle), sin(angle)}; + + float2 finish = point + dir * radius; + + //float max_dir = max(fabs(dir.x), fabs(dir.y)); + + //dir = dir / max_dir; + + //int num = max(fabs(dir.x * radius), fabs(dir.y * radius)); + + int num = radius; + + float2 cur = point; + + float brightness = 1.f; + + ///after a distance of radius, we should be fully absorbed + ///in... a uniform cloud..? + for(int i=0; i < num; i++, cur += dir) + { + float density = read_imagef(fluid, sam, cur + 0.5f).x; + + float amount_reflected = brightness * density * 0.01f; + + brightness -= amount_reflected; + + float distance_curve = 1.f - (float)i / num; + + float extra_bright = 200; + + if(cur.x >= 0 && cur.y >= 0 && cur.x < sw && cur.y < sh) + { + write_imagef(screen, convert_int2(cur), (float4)(amount_reflected*distance_curve*extra_bright, 0, 0, 1)); + } + } +} diff --git a/lighting.hpp b/lighting.hpp new file mode 100644 index 0000000..148c02e --- /dev/null +++ b/lighting.hpp @@ -0,0 +1,48 @@ +#ifndef LIGHTING_HPP_INCLUDED +#define LIGHTING_HPP_INCLUDED + +struct lighting_manager +{ + cl::buffer* colour; + + vec2i saved_dim; + + void init(cl::context& ctx, cl::buffer_manager& buffers, cl::program& program, cl::command_queue& cqueue, vec2i screen_dim) + { + colour = buffers.fetch(ctx, nullptr); + + std::vector colour_buf; + + for(int y=0; y < screen_dim.y(); y++) + { + for(int x=0; x < screen_dim.x(); x++) + { + colour_buf.push_back({0,0,0,1}); + } + } + + colour->alloc_img(cqueue, colour_buf, screen_dim, CL_RGBA, CL_FLOAT); + + saved_dim = screen_dim; + } + + void tick(cl::cl_gl_interop_texture* interop, cl::buffer_manager& buffers, cl::program& program, cl::command_queue& cqueue, vec2f mpos, cl::buffer* fluid) + { + mpos.y() = saved_dim.y() - mpos.y(); + + float rad = 800.f; + + int num_tracers = ceil(2 * M_PI * rad)*2.5; + + cl::args raytrace_args; + raytrace_args.push_back(mpos); + raytrace_args.push_back(rad); + raytrace_args.push_back(num_tracers); + raytrace_args.push_back(fluid); + raytrace_args.push_back(interop); + + cqueue.exec(program, "lighting_raytrace_point", raytrace_args, {num_tracers}, {128}); + } +}; + +#endif // LIGHTING_HPP_INCLUDED diff --git a/main.cpp b/main.cpp index b5727a2..7c7b609 100644 --- a/main.cpp +++ b/main.cpp @@ -4,6 +4,7 @@ #include #include #include "fluid.hpp" +#include "lighting.hpp" int main() { @@ -28,7 +29,7 @@ int main() cl::buffer_manager buffer_manage; - cl::buffer* buf = buffer_manage.fetch(ctx, nullptr); + /*cl::buffer* buf = buffer_manage.fetch(ctx, nullptr); std::vector data; @@ -48,25 +49,30 @@ int main() cl::buffer* image = buffer_manage.fetch(ctx, nullptr); - image->alloc_img(cqueue, idata, window_size); + image->alloc_img(cqueue, idata, window_size);*/ cl::cl_gl_interop_texture* interop = buffer_manage.fetch(ctx, nullptr, win.getSize().x, win.getSize().y); interop->acquire(cqueue); - cl::args none; + /*cl::args none; //none.push_back(buf); none.push_back(interop); none.push_back(image); cqueue.exec(program, "fluid_test", none, {128}, {16}); - cqueue.block(); + cqueue.block();*/ vec2i screen_dim = {win.getSize().x, win.getSize().y}; fluid_manager fluid_manage; fluid_manage.init(ctx, buffer_manage, program, cqueue, screen_dim, screen_dim, screen_dim*2); + + lighting_manager lighting_manage; + lighting_manage.init(ctx, buffer_manage, program, cqueue, screen_dim); + + sf::Clock clk; sf::Keyboard key; @@ -107,6 +113,8 @@ int main() fluid_manage.tick(interop, buffer_manage, program, cqueue); + lighting_manage.tick(interop, buffer_manage, program, cqueue, cur_mouse, fluid_manage.dye[fluid_manage.which_dye]); + interop->gl_blit_me(0, cqueue);