WORKING! Awesome! It now casts fully inside the gpu,

context is then switched to gl and then rendered via sfml.
It has no loop, no controls, and the aspect ratio is off,
but holy hell it works!

include/CL_Wrapper.h

@@ -39,7 +39,7 @@ public:
     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 store_buffer(cl_mem, std::string buffer_name);
-	int run_kernel(std::string kernel_name);
+	int run_kernel(std::string kernel_name, const int work_size);
 
 
 	bool assert(int error_code, std::string function_name);

include/Map.h

@@ -34,6 +34,12 @@ public:
 	char *list;
 	sf::Vector3i dimensions;
 
+	void setVoxel(sf::Vector3i position, int val){
+
+		list[position.x + dimensions.x * (position.y + dimensions.z * position.z)] = val;
+
+	};
+
 	void moveLight(sf::Vector2f in);
 	sf::Vector3f global_light;
 

kernels/minimal_kernel.c

@@ -10,12 +10,159 @@ __kernel void min_kern(
 
     size_t id = get_global_id(0);
 
-    float4 black = (float4)(.49, .68, .81, 1);
+    int2 pixel = {id % resolution->x, id / resolution->x};
 
-    int2 pixelcoord = (int2) (id, id);
 
-    write_imagef(image, pixelcoord, black);
+    float3 ray_dir = projection_matrix[pixel.x + resolution->x * pixel.y];
+    //printf("%i === %f, %f, %f\n", id, ray_dir.x, ray_dir.y, ray_dir.z);
 
+    // Y axis, pitch
+    //ray_dir.x = ray_dir.z * sin(cam_dir->y) + ray_dir.x * cos(cam_dir->y);
+    //ray_dir.y = ray_dir.y;
+    //ray_dir.z = ray_dir.z * cos(cam_dir->y) - ray_dir.x * sin(cam_dir->y);
+
+
+    ray_dir = (float3)(
+            ray_dir.z * sin(cam_dir->y) + ray_dir.x * cos(cam_dir->y),
+            ray_dir.y,
+            ray_dir.z * cos(cam_dir->y) - ray_dir.x * sin(cam_dir->y)
+    );
+
+    // Z axis, yaw
+    //ray_dir.x = ray_dir.x * cos(cam_dir->z) - ray_dir.y * sin(cam_dir->z);
+    //ray_dir.y = ray_dir.x * sin(cam_dir->z) + ray_dir.y * cos(cam_dir->z);
+    //ray_dir.z = ray_dir.z;
+
+    ray_dir = (float3)(
+            ray_dir.x * cos(cam_dir->z) - ray_dir.y * sin(cam_dir->z),
+            ray_dir.x * sin(cam_dir->z) + ray_dir.y * cos(cam_dir->z),
+            ray_dir.z
+    );
+
+
+    // Setup the voxel step based on what direction the ray is pointing
+    int3 voxel_step = {1, 1, 1};
+    voxel_step.x *= (ray_dir.x > 0) - (ray_dir.x < 0);
+    voxel_step.y *= (ray_dir.y > 0) - (ray_dir.y < 0);
+    voxel_step.z *= (ray_dir.z > 0) - (ray_dir.z < 0);
+
+    // Setup the voxel coords from the camera origin
+    int3 voxel = {
+            floorf(cam_pos->x),
+            floorf(cam_pos->y),
+            floorf(cam_pos->z)
+    };
+
+    // Delta T is the units a ray must travel along an axis in order to
+    // traverse an integer split
+    float3 delta_t = {
+            fabsf(1.0f / ray_dir.x),
+            fabsf(1.0f / ray_dir.y),
+            fabsf(1.0f / ray_dir.z)
+    };
+
+    // Intersection T is the collection of the next intersection points
+    // for all 3 axis XYZ.
+    float3 intersection_t = {
+            delta_t.x,
+            delta_t.y,
+            delta_t.z
+    };
+
+
+    int dist = 0;
+    int face = -1;
+    // X:0, Y:1, Z:2
+
+    // Andrew Woo's raycasting algo
+    do {
+        if ((intersection_t.x) < (intersection_t.y)) {
+            if ((intersection_t.x) < (intersection_t.z)) {
+
+                face = 0;
+                voxel.x += voxel_step.x;
+                intersection_t.x = intersection_t.x + delta_t.x;
+            } else {
+
+                face = 2;
+                voxel.z += voxel_step.z;
+                intersection_t.z = intersection_t.z + delta_t.z;
+            }
+        } else {
+            if ((intersection_t.y) < (intersection_t.z)) {
+
+                face = 1;
+                voxel.y += voxel_step.y;
+                intersection_t.y = intersection_t.y + delta_t.y;
+            } else {
+
+                face = 2;
+                voxel.z += voxel_step.z;
+                intersection_t.z = intersection_t.z + delta_t.z;
+            }
+        }
+
+        // If the ray went out of bounds
+        if (voxel.z >= map_dim->z) {
+            write_imagef(image, pixel, (float4)(.5, .50, .00, 1));
+            return;
+        }
+        if (voxel.x >= map_dim->x) {
+            write_imagef(image, pixel, (float4)(.00, .00, .99, 1));
+            return;
+        }
+        if (voxel.y >= map_dim->x) {
+            write_imagef(image, pixel, (float4)(.00, .99, .00, 1));
+            return;
+        }
+
+        if (voxel.x < 0) {
+            write_imagef(image, pixel, (float4)(.99, .00, .99, 1));
+            return;
+        }
+        if (voxel.y < 0) {
+            write_imagef(image, pixel, (float4)(.99, .99, .00, 1));
+            return;
+        }
+        if (voxel.z < 0) {
+            write_imagef(image, pixel, (float4)(.00, .99, .99, 1));
+            return;
+        }
+
+        // If we hit a voxel
+        int index = voxel.x + map_dim->x * (voxel.y + map_dim->z * voxel.z);
+        int voxel_data = map[index];
+
+        if (id == 100)
+            printf("%i, %i, %i\n", voxel.x, voxel.y, voxel.z);
+
+        switch (voxel_data) {
+            case 1:
+                write_imagef(image, pixel, (float4)(.50, .00, .00, 1));
+                return;
+            case 2:
+                write_imagef(image, pixel, (float4)(.00, .50, .00, 1.00));
+                return;
+            case 3:
+                write_imagef(image, pixel, (float4)(.00, .00, .50, 1.00));
+                return;
+            case 4:
+                write_imagef(image, pixel, (float4)(.25, .00, .25, 1.00));
+                return;
+            case 5:
+                write_imagef(image, pixel, (float4)(.10, .30, .80, 1.00));
+                return;
+            case 6:
+                write_imagef(image, pixel, (float4)(.30, .80, .10, 1.00));
+                return;
+        }
+
+        dist++;
+    } while (dist < 200);
+
+
+    write_imagef(image, pixel, (float4)(.00, .00, .00, .00));
+    return;
 
     //printf("%i %i -- ", id, map[id]);
     //printf("%i, %i, %i\n", map_dim->x, map_dim->y, map_dim->z);
@@ -23,7 +170,4 @@ __kernel void min_kern(
     //printf("%f, %f, %f\n", cam_dir->x, cam_dir->y, cam_dir->z);
     //printf("%f, %f, %f\n", cam_pos->x, cam_pos->y, cam_pos->z);
 
-
-
-
 }

src/CL_Wrapper.cpp

@@ -230,10 +230,10 @@ int CL_Wrapper::store_buffer(cl_mem buffer, std::string buffer_name){
     buffer_map.emplace(std::make_pair(buffer_name, buffer));
 }
 
-int CL_Wrapper::run_kernel(std::string kernel_name){
+int CL_Wrapper::run_kernel(std::string kernel_name, const int work_size){
 
     const int WORKER_SIZE = 10;
-    size_t global_work_size[1] = { WORKER_SIZE };
+    size_t global_work_size[1] = { static_cast<size_t>(work_size) };
 
     cl_kernel kernel = kernel_map.at(kernel_name);
 

src/main.cpp

@@ -29,8 +29,8 @@
 #include "RayCaster.h"
 #include "CL_Wrapper.h"
 
-const int WINDOW_X = 150;
+const int WINDOW_X = 100;
-const int WINDOW_Y = 150;
+const int WINDOW_Y = 100;
 
 
 
@@ -76,6 +76,8 @@ int main() {
     sf::Vector3i map_dim(100, 100, 100);
     Map* map = new Map(map_dim);
 
+    map->setVoxel(sf::Vector3i(77, 50, 85), 5);
+
     cl_mem map_buff = clCreateBuffer(
             c.getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
             sizeof(char) * map_dim.x * map_dim.y * map_dim.z, map->list, NULL
@@ -129,6 +131,9 @@ int main() {
         }
     }
 
+    int ind = 367;
+    printf("%i === %f, %f, %f\n", ind, view_matrix[ind * 4 + 0], view_matrix[ind * 4 + 1], view_matrix[ind * 4 + 2]);
+
     cl_mem view_matrix_buff = clCreateBuffer(
             c.getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
             sizeof(float) * 3 * view_res.x * view_res.y, view_matrix, NULL
@@ -142,7 +147,7 @@ int main() {
     );
 
 
-    sf::Vector3f cam_pos(50, 50, 50);
+    sf::Vector3f cam_pos(55, 50, 50);
     cl_mem cam_pos_buff = clCreateBuffer(
             c.getContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
             sizeof(float) * 4, &cam_pos, NULL
@@ -194,15 +199,18 @@ int main() {
     c.set_kernel_arg("min_kern", 5, "cam_pos_buffer");
     c.set_kernel_arg("min_kern", 6, "image_buffer");
 
+    const int size = 100 * 100;
+    c.run_kernel("min_kern", size);
 
-    c.run_kernel("min_kern");
+
+    clFinish(c.getCommandQueue());
 
     error = clEnqueueReleaseGLObjects(c.getCommandQueue(), 1, &image_buff, 0, NULL, NULL);
     if (c.assert(error, "clEnqueueReleaseGLObjects"))
         return -1;
 
     s.setTexture(t);
-    
+
     // The step size in milliseconds between calls to Update()
     // Lets set it to 16.6 milliseonds (60FPS)
     float step_size = 0.0166f;
@@ -224,9 +232,6 @@ int main() {
 
 	// State values
 
-
-
-
 	sf::Vector3f cam_vec(0, 0, 0);
 
 	RayCaster ray_caster(map, map_dim, view_res);