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<cl::buffer>(ctx, nullptr);
+
+        std::vector<vec4f> 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 <ocl/logging.hpp>
 #include <SFML/Graphics.hpp>
 #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<cl::buffer>(ctx, nullptr);
+    /*cl::buffer* buf = buffer_manage.fetch<cl::buffer>(ctx, nullptr);
 
     std::vector<int> data;
 
@@ -48,25 +49,30 @@ int main()
 
     cl::buffer* image = buffer_manage.fetch<cl::buffer>(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<cl::cl_gl_interop_texture>(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);