faces now contain verts, norms, & uvs

src/headers/types.h

@@ -16,21 +16,35 @@ typedef struct { // A loaded texture file
 } RI_texture;
 
 typedef struct {
-    cl_uint position_0_index;
-    cl_uint position_1_index;
-    cl_uint position_2_index;
+    RI_vector_3 position_0;
+    RI_vector_3 position_1;
+    RI_vector_3 position_2;
 
-    cl_uint normal_0_index;
-    cl_uint normal_1_index;
-    cl_uint normal_2_index;
+    RI_vector_3 normal_0;
+    RI_vector_3 normal_1;
+    RI_vector_3 normal_2;
 
-    cl_uint uv_0_index;
-    cl_uint uv_1_index;
-    cl_uint uv_2_index;
+    RI_vector_2 uv_0;
+    RI_vector_2 uv_1;
+    RI_vector_2 uv_2;
 
     cl_uint should_render;
 } RI_face;
 
+typedef struct {
+    int position_0_index;
+    int position_1_index;
+    int position_2_index;
+
+    int normal_0_index;
+    int normal_1_index;
+    int normal_2_index;
+
+    int uv_0_index;
+    int uv_1_index;
+    int uv_2_index;
+} RI_temp_face;
+
 typedef struct {
     int face_count;
     int face_index;
@@ -116,9 +130,10 @@ typedef struct {
     cl_mem uvs_mem_buffer;
     RI_renderable_face *faces_to_render;
     RI_face *faces;
-    RI_vector_3 *vertecies;
-    RI_vector_3 *normals;
-    RI_vector_2 *uvs;
+    RI_temp_face *temp_faces;
+    RI_vector_3 *temp_vertecies;
+    RI_vector_3 *temp_normals;
+    RI_vector_2 *temp_uvs;
     int face_count;
     int vertex_count;
     int normal_count;

src/kernels/kernels.cl

@@ -29,17 +29,17 @@ typedef struct {
 } RI_renderable_face;
 
 typedef struct {
-    int position_0_index;
-    int position_1_index;
-    int position_2_index;
+    RI_vector_3 position_0;
+    RI_vector_3 position_1;
+    RI_vector_3 position_2;
 
-    int normal_0_index;
-    int normal_1_index;
-    int normal_2_index;
+    RI_vector_3 normal_0;
+    RI_vector_3 normal_1;
+    RI_vector_3 normal_2;
 
-    int uv_0_index;
-    int uv_1_index;
-    int uv_2_index;
+    RI_vector_2 uv_0;
+    RI_vector_2 uv_1;
+    RI_vector_2 uv_2;
 
     int should_render;
 } RI_face;
@@ -265,7 +265,7 @@ void global_quaternion_rotate(__global RI_vector_3 *position, RI_vector_4 rotati
     *position = (RI_vector_3){rotation.x, rotation.y, rotation.z};
 }
 
-__kernel void transformer(__global RI_face *faces, __global RI_vector_3 *vertecies, __global RI_vector_3 *normals, __global RI_vector_2 *uvs, __global RI_renderable_face *renderable_faces, double actor_x, double actor_y, double actor_z, double actor_r_w, double actor_r_x, double actor_r_y, double actor_r_z, double actor_s_x, double actor_s_y, double actor_s_z, int has_normals, int has_uvs, int face_array_offset_index, int face_count, int width, int height, double horizontal_fov_factor, double vertical_fov_factor, float min_clip, float max_clip, double camera_x, double camera_y, double camera_z, double camera_r_w, double camera_r_x, double camera_r_y, double camera_r_z, int renderable_face_offset, int face_sqrt){
+__kernel void transformer(__global RI_face *faces, __global RI_renderable_face *renderable_faces, double actor_x, double actor_y, double actor_z, double actor_r_w, double actor_r_x, double actor_r_y, double actor_r_z, double actor_s_x, double actor_s_y, double actor_s_z, int has_normals, int has_uvs, int face_array_offset_index, int face_count, int width, int height, double horizontal_fov_factor, double vertical_fov_factor, float min_clip, float max_clip, double camera_x, double camera_y, double camera_z, double camera_r_w, double camera_r_x, double camera_r_y, double camera_r_z, int renderable_face_offset, int face_sqrt){
     int face_index = get_global_id(1) * face_sqrt + get_global_id(0); if (face_index >= face_count) return;
 
     RI_vector_3 current_actor_position = (RI_vector_3){actor_x, actor_y, actor_z};
@@ -283,36 +283,24 @@ __kernel void transformer(__global RI_face *faces, __global RI_vector_3 *verteci
 
     cur_r_face->should_render = 1;
 
-    int vert_pos_0_index = cur_face->position_0_index;
-    int vert_pos_1_index = cur_face->position_1_index;
-    int vert_pos_2_index = cur_face->position_2_index;
-    
-    int normal_0_index = cur_face->normal_0_index;
-    int normal_1_index = cur_face->normal_1_index;
-    int normal_2_index = cur_face->normal_2_index;
-
-    int uv_0_index = cur_face->uv_0_index;
-    int uv_1_index = cur_face->uv_1_index;
-    int uv_2_index = cur_face->uv_2_index;
-
     // cur_r_face->parent_actor = current_actor;
     cur_r_face->shrunk = 0;
     cur_r_face->split = 0;
 
     // cur_r_face->material = current_actor.material;
 
-    cur_r_face->position_0 = vertecies[vert_pos_0_index];
-    cur_r_face->position_1 = vertecies[vert_pos_1_index];
-    cur_r_face->position_2 = vertecies[vert_pos_2_index];
+    cur_r_face->position_0 = cur_face->position_0;
+    cur_r_face->position_1 = cur_face->position_1;
+    cur_r_face->position_2 = cur_face->position_2;
 
-    cur_r_face->normal_0 = normals[normal_0_index];
-    cur_r_face->normal_1 = normals[normal_1_index];
-    cur_r_face->normal_2 = normals[normal_2_index];
+    cur_r_face->normal_0 = cur_face->normal_0;
+    cur_r_face->normal_1 = cur_face->normal_1;
+    cur_r_face->normal_2 = cur_face->normal_2;
 
     if (has_uvs){
-        cur_r_face->uv_0 = uvs[uv_0_index];
-        cur_r_face->uv_1 = uvs[uv_1_index];
-        cur_r_face->uv_2 = uvs[uv_2_index];
+        cur_r_face->uv_0 = cur_face->uv_0;
+        cur_r_face->uv_1 = cur_face->uv_1;
+        cur_r_face->uv_2 = cur_face->uv_2;
     }
 
     // scale

src/launch program/main.c

@@ -25,17 +25,17 @@ int main(){
     RI_actor *triangle = RI_new_actor();
 
     RI_mesh *homer_mesh = RI_load_mesh("objects/homer.obj");
-    RI_mesh *teapot_mesh = RI_load_mesh("objects/teapot.obj");
+    RI_mesh *teapot_mesh = RI_load_mesh("objects/stanford_bunny.obj");
 
     cube->mesh = homer_mesh;
-    triangle->mesh = homer_mesh;
+    triangle->mesh = teapot_mesh;
 
     cube->scale = (RI_vector_3){80, 80, 80};
     cube->position = (RI_vector_3){-20, -30, 100};
     cube->rotation = (RI_vector_4){1, 0, 0, 0};
 
-    triangle->scale = (RI_vector_3){10, 10, 10};
-    triangle->position = (RI_vector_3){20, 0, 100};
+    triangle->scale = (RI_vector_3){300, 300, 300};
+    triangle->position = (RI_vector_3){20, -20, 100};
     triangle->rotation = (RI_vector_4){1, 0, 0, 0};
     
     scene->actors = malloc(sizeof(RI_actor) * 10);
@@ -48,6 +48,8 @@ int main(){
     long int start, end;
     double fps = 0;
 
+    float total_fps = 0;
+
     double delta_time = 0;
     double delta_min = 0.00001;
     double delta_max = 1;
@@ -68,7 +70,9 @@ int main(){
         delta_time = fmin(fmax((double)(end - start) / (double)(CLOCKS_PER_SEC), delta_min), delta_max);
         fps = 1.0 / delta_time;
 
-        printf("%f\n", fps);
+        total_fps += fps;
+
+        printf("fps: %f average fps: %f\n", fps, total_fps / context->current_frame);
     }
 
     free(scene->actors);

src/main/main.c

@@ -103,35 +103,47 @@ RI_mesh *RI_load_mesh(char *filename){
     
     int face_count = 0;
 
+    int object_face_count = 0;
+    int object_vertecies_count = 0;
+    int object_normals_count = 0;
+    int object_uvs_count = 0;
+
     while (fgets(line, sizeof(line), file)) {
         if (line[0] == 'f' && line[1] == ' ') { // face
             ++face_count;
             ++context.opencl.face_count;
+            ++object_face_count;
         }
         else if (line[0] == 'v'){
             if (line[1] == ' ') { // vertex
                 ++context.opencl.vertex_count;
+                ++object_vertecies_count;
             }
             else if (line[1] == 'n') { // normal
                 ++context.opencl.normal_count;
+                ++object_normals_count;
             }
             else if (line[1] == 't') { // UV
                 ++context.opencl.uv_count;
+                ++object_uvs_count;
             }
         }
     }
 
     rewind(file);
 
+    debug("%d faces %d vertecies %d normals %d uvs", object_face_count, object_vertecies_count, object_normals_count, object_uvs_count);
+
     context.opencl.faces = RI_realloc(context.opencl.faces, sizeof(RI_face) * context.opencl.face_count);
-    context.opencl.vertecies = RI_realloc(context.opencl.vertecies, sizeof(RI_vector_3) * context.opencl.vertex_count);
-    context.opencl.normals = RI_realloc(context.opencl.normals, sizeof(RI_vector_3) * context.opencl.normal_count);
-    context.opencl.uvs = RI_realloc(context.opencl.uvs, sizeof(RI_vector_2) * context.opencl.uv_count);
+    context.opencl.temp_faces = RI_malloc(sizeof(RI_face) * object_face_count);
+    context.opencl.temp_vertecies = RI_malloc(sizeof(RI_vector_3) * object_vertecies_count);
+    if (object_normals_count > 0) context.opencl.temp_normals = RI_malloc(sizeof(RI_vector_3) * object_normals_count);
+    if (object_uvs_count > 0) context.opencl.temp_uvs = RI_malloc(sizeof(RI_vector_2) * object_uvs_count);
 
-    int current_face_index = previous_face_count;
-    int current_vertex_index = previous_vertecies_count;
-    int current_normal_index = previous_normals_count;
-    int current_uv_index = previous_uvs_count;
+    int current_face_index = 0;
+    int current_vertex_index = 0;
+    int current_normal_index = 0;
+    int current_uv_index = 0;
 
     int has_normals, has_uvs;
     has_normals = has_uvs = 0;
@@ -177,17 +189,17 @@ RI_mesh *RI_load_mesh(char *filename){
                 has_normals = has_uvs = 1;
             }
 
-            context.opencl.faces[current_face_index].position_0_index = vertex_0_index - 1 + previous_vertecies_count;
-            context.opencl.faces[current_face_index].position_1_index = vertex_1_index - 1 + previous_vertecies_count;
-            context.opencl.faces[current_face_index].position_2_index = vertex_2_index - 1 + previous_vertecies_count;
+            context.opencl.temp_faces[current_face_index].position_0_index = vertex_0_index - 1;
+            context.opencl.temp_faces[current_face_index].position_1_index = vertex_1_index - 1;
+            context.opencl.temp_faces[current_face_index].position_2_index = vertex_2_index - 1;
 
-            context.opencl.faces[current_face_index].normal_0_index = normal_0_index - 1 + previous_normals_count;
-            context.opencl.faces[current_face_index].normal_1_index = normal_1_index - 1 + previous_normals_count;
-            context.opencl.faces[current_face_index].normal_2_index = normal_2_index - 1 + previous_normals_count;
+            context.opencl.temp_faces[current_face_index].normal_0_index = normal_0_index - 1;
+            context.opencl.temp_faces[current_face_index].normal_1_index = normal_1_index - 1;
+            context.opencl.temp_faces[current_face_index].normal_2_index = normal_2_index - 1;
             
-            context.opencl.faces[current_face_index].uv_0_index = uv_0_index - 1 + previous_uvs_count;
-            context.opencl.faces[current_face_index].uv_1_index = uv_1_index - 1 + previous_uvs_count;
-            context.opencl.faces[current_face_index].uv_2_index = uv_2_index - 1 + previous_uvs_count;
+            context.opencl.temp_faces[current_face_index].uv_0_index = uv_0_index - 1;
+            context.opencl.temp_faces[current_face_index].uv_1_index = uv_1_index - 1;
+            context.opencl.temp_faces[current_face_index].uv_2_index = uv_2_index - 1;
 
             context.opencl.faces[current_face_index].should_render = 1;
 
@@ -198,9 +210,9 @@ RI_mesh *RI_load_mesh(char *filename){
             
             sscanf(line, "v %lf %lf %lf", &x, &y, &z);
 
-            context.opencl.vertecies[current_vertex_index].x = x;
-            context.opencl.vertecies[current_vertex_index].y = y;
-            context.opencl.vertecies[current_vertex_index].z = z;
+            context.opencl.temp_vertecies[current_vertex_index].x = x;
+            context.opencl.temp_vertecies[current_vertex_index].y = y;
+            context.opencl.temp_vertecies[current_vertex_index].z = z;
 
             ++current_vertex_index;
         } 
@@ -209,9 +221,9 @@ RI_mesh *RI_load_mesh(char *filename){
             
             sscanf(line, "vn %lf %lf %lf", &x, &y, &z);
 
-            context.opencl.normals[current_normal_index].x = x;
-            context.opencl.normals[current_normal_index].y = y;
-            context.opencl.normals[current_normal_index].z = z;
+            context.opencl.temp_normals[current_normal_index].x = x;
+            context.opencl.temp_normals[current_normal_index].y = y;
+            context.opencl.temp_normals[current_normal_index].z = z;
 
             ++current_normal_index;
         }
@@ -220,14 +232,33 @@ RI_mesh *RI_load_mesh(char *filename){
 
             sscanf(line, "vt %lf %lf %lf", &x, &y, &z);
 
-            context.opencl.uvs[current_uv_index].x = x;
-            context.opencl.uvs[current_uv_index].y = y;
+            context.opencl.temp_uvs[current_uv_index].x = x;
+            context.opencl.temp_uvs[current_uv_index].y = y;
             // UVS are almost always 2D so we don't need Z (the type itself is a vector 2f, not 3f) 
 
             ++current_uv_index;
         } 
     }
 
+    for (int i = 0; i < object_face_count; ++i){
+        context.opencl.faces[i + previous_face_count].position_0 = context.opencl.temp_vertecies[context.opencl.temp_faces[i].position_0_index];
+        context.opencl.faces[i + previous_face_count].position_1 = context.opencl.temp_vertecies[context.opencl.temp_faces[i].position_1_index];
+        context.opencl.faces[i + previous_face_count].position_2 = context.opencl.temp_vertecies[context.opencl.temp_faces[i].position_2_index];
+
+        context.opencl.faces[i + previous_face_count].normal_0 = context.opencl.temp_normals[context.opencl.temp_faces[i].normal_0_index];
+        context.opencl.faces[i + previous_face_count].normal_1 = context.opencl.temp_normals[context.opencl.temp_faces[i].normal_1_index];
+        context.opencl.faces[i + previous_face_count].normal_2 = context.opencl.temp_normals[context.opencl.temp_faces[i].normal_2_index];
+
+        context.opencl.faces[i + previous_face_count].uv_0 = context.opencl.temp_uvs[context.opencl.temp_faces[i].uv_0_index];
+        context.opencl.faces[i + previous_face_count].uv_1 = context.opencl.temp_uvs[context.opencl.temp_faces[i].uv_1_index];
+        context.opencl.faces[i + previous_face_count].uv_2 = context.opencl.temp_uvs[context.opencl.temp_faces[i].uv_2_index];
+    }
+
+    free(context.opencl.temp_faces);
+    free(context.opencl.temp_vertecies);
+    if (object_normals_count > 0) free(context.opencl.temp_normals);
+    if (object_uvs_count > 0) free(context.opencl.temp_uvs);
+
     char* loading_mesh_notice_string;
 
     if (has_normals && !has_uvs) loading_mesh_notice_string = "normals";
@@ -239,7 +270,7 @@ RI_mesh *RI_load_mesh(char *filename){
     new_mesh->has_normals = has_normals;
     new_mesh->has_uvs = has_uvs;
 
-    new_mesh->face_count = context.opencl.face_count - previous_face_count;
+    new_mesh->face_count = object_face_count;
     new_mesh->face_index = previous_face_count;
 
     debug("[Mesh Loader] Loaded mesh \"%s\"! %d faces, %d verticies, %d normals, %d uvs", filename, current_face_index, current_vertex_index, current_normal_index, current_uv_index); 
@@ -260,7 +291,7 @@ RI_mesh *RI_load_mesh(char *filename){
 
             context.opencl.length_of_renderable_faces_array = context.opencl.face_count * 2;
             
-            debug("reallocating %f mb", sizeof(RI_renderable_face) * context.opencl.length_of_renderable_faces_array / 1048576.0);
+            debug("reallocating %f mb (%d renderable faces)", sizeof(RI_renderable_face) * context.opencl.length_of_renderable_faces_array / 1048576.0, context.opencl.length_of_renderable_faces_array);
 
             context.opencl.faces_to_render = RI_malloc(sizeof(RI_renderable_face) * context.opencl.length_of_renderable_faces_array);
 
@@ -284,36 +315,6 @@ RI_mesh *RI_load_mesh(char *filename){
         }
     }
 
-    if (previous_vertecies_count != context.opencl.vertex_count) {
-        if (context.opencl.vertecies_mem_buffer) clReleaseMemObject(context.opencl.vertecies_mem_buffer);
-
-        context.opencl.vertecies_mem_buffer = clCreateBuffer(context.opencl.context, CL_MEM_READ_WRITE, sizeof(RI_vector_3) * context.opencl.vertex_count, NULL, NULL);
-        
-        clEnqueueWriteBuffer(context.opencl.queue, context.opencl.vertecies_mem_buffer, CL_TRUE, 0, sizeof(RI_vector_3) * context.opencl.vertex_count, context.opencl.vertecies, 0, NULL, NULL);
-
-        clSetKernelArg(context.opencl.transformation_kernel, 1, sizeof(cl_mem), &context.opencl.vertecies_mem_buffer);
-    }
-
-    if (previous_normals_count != context.opencl.normal_count) {
-        if (context.opencl.normals_mem_buffer) clReleaseMemObject(context.opencl.normals_mem_buffer);
-
-        context.opencl.normals_mem_buffer = clCreateBuffer(context.opencl.context, CL_MEM_READ_WRITE, sizeof(RI_vector_3) * context.opencl.normal_count, NULL, NULL);
-        
-        clEnqueueWriteBuffer(context.opencl.queue, context.opencl.normals_mem_buffer, CL_TRUE, 0, sizeof(RI_vector_3) * context.opencl.normal_count, context.opencl.normals, 0, NULL, NULL);
-
-        clSetKernelArg(context.opencl.transformation_kernel, 2, sizeof(cl_mem), &context.opencl.normals_mem_buffer);
-    }
-
-    if (previous_uvs_count != context.opencl.uv_count) {
-        if (context.opencl.uvs_mem_buffer) clReleaseMemObject(context.opencl.uvs_mem_buffer);
-
-        context.opencl.uvs_mem_buffer = clCreateBuffer(context.opencl.context, CL_MEM_READ_WRITE, sizeof(RI_vector_2) * context.opencl.uv_count, NULL, NULL);
-        
-        clEnqueueWriteBuffer(context.opencl.queue, context.opencl.uvs_mem_buffer, CL_TRUE, 0, sizeof(RI_vector_2) * context.opencl.uv_count, context.opencl.uvs, 0, NULL, NULL);
-
-        clSetKernelArg(context.opencl.transformation_kernel, 3, sizeof(cl_mem), &context.opencl.uvs_mem_buffer);
-    }
-
     fclose(file);
     
     end_time = clock();
@@ -341,30 +342,30 @@ void RI_render(RI_texture *target_texture, RI_scene *scene){
     // kernel args    
     
     // 21, double horizontal_fov_factor
-    clSetKernelArg(context.opencl.transformation_kernel, 21, sizeof(double), &horizontal_fov_factor);
+    clSetKernelArg(context.opencl.transformation_kernel, 18, sizeof(double), &horizontal_fov_factor);
     // 22, double vertical_fov_factor
-    clSetKernelArg(context.opencl.transformation_kernel, 22, sizeof(double), &vertical_fov_factor);
+    clSetKernelArg(context.opencl.transformation_kernel, 19, sizeof(double), &vertical_fov_factor);
 
     // 23, double min_clip
-    clSetKernelArg(context.opencl.transformation_kernel, 23, sizeof(float), &scene->camera.min_clip);
+    clSetKernelArg(context.opencl.transformation_kernel, 20, sizeof(float), &scene->camera.min_clip);
     // 24, double max_clip
-    clSetKernelArg(context.opencl.transformation_kernel, 24, sizeof(float), &scene->camera.max_clip);
+    clSetKernelArg(context.opencl.transformation_kernel, 21, sizeof(float), &scene->camera.max_clip);
 
     // 25, double camera_x
-    clSetKernelArg(context.opencl.transformation_kernel, 25, sizeof(double), &scene->camera.position.x);
+    clSetKernelArg(context.opencl.transformation_kernel, 22, sizeof(double), &scene->camera.position.x);
     // 26, double camera_y
-    clSetKernelArg(context.opencl.transformation_kernel, 26, sizeof(double), &scene->camera.position.y);
+    clSetKernelArg(context.opencl.transformation_kernel, 23, sizeof(double), &scene->camera.position.y);
     // 27, double camera_z
-    clSetKernelArg(context.opencl.transformation_kernel, 27, sizeof(double), &scene->camera.position.z);
+    clSetKernelArg(context.opencl.transformation_kernel, 24, sizeof(double), &scene->camera.position.z);
 
     // 28, double camera_r_w
-    clSetKernelArg(context.opencl.transformation_kernel, 28, sizeof(double), &scene->camera.rotation.w);
+    clSetKernelArg(context.opencl.transformation_kernel, 25, sizeof(double), &scene->camera.rotation.w);
     // 29, double camera_r_x
-    clSetKernelArg(context.opencl.transformation_kernel, 29, sizeof(double), &scene->camera.rotation.x);
+    clSetKernelArg(context.opencl.transformation_kernel, 26, sizeof(double), &scene->camera.rotation.x);
     // 30, double camera_r_y
-    clSetKernelArg(context.opencl.transformation_kernel, 30, sizeof(double), &scene->camera.rotation.y);
+    clSetKernelArg(context.opencl.transformation_kernel, 27, sizeof(double), &scene->camera.rotation.y);
     // 31, double camera_r_z
-    clSetKernelArg(context.opencl.transformation_kernel, 31, sizeof(double), &scene->camera.rotation.z);
+    clSetKernelArg(context.opencl.transformation_kernel, 28, sizeof(double), &scene->camera.rotation.z);
 
 
     int local_group_size_x = 16;
@@ -411,7 +412,7 @@ void RI_render(RI_texture *target_texture, RI_scene *scene){
     // set faces_to_render to zero
     memset(context.opencl.faces_to_render, 0, sizeof(RI_renderable_face) * scene->face_count * 2);
 
-    clSetKernelArg(context.opencl.transformation_kernel, 4, sizeof(cl_mem), &context.opencl.renderable_faces_mem_buffer);
+    clSetKernelArg(context.opencl.transformation_kernel, 1, sizeof(cl_mem), &context.opencl.renderable_faces_mem_buffer);
 
     context.current_renderable_face_index = 0;
     context.current_split_renderable_face_index = 0;
@@ -444,46 +445,48 @@ void RI_render(RI_texture *target_texture, RI_scene *scene){
         debug("(%d extra work items; %d items (%dx%d) - %d faces)", t_global_work_size[0] * t_global_work_size[1] - scene->actors[actor_index]->mesh->face_count, t_global_work_size[0] * t_global_work_size[1], t_global_work_size[0], t_global_work_size[1], scene->actors[actor_index]->mesh->face_count);
 
         // 5, double actor_x
-        clSetKernelArg(context.opencl.transformation_kernel, 5, sizeof(double), &actor->position.x);
+        clSetKernelArg(context.opencl.transformation_kernel, 2, sizeof(double), &actor->position.x);
         // 6, double actor_y
-        clSetKernelArg(context.opencl.transformation_kernel, 6, sizeof(double), &actor->position.y);
+        clSetKernelArg(context.opencl.transformation_kernel, 3, sizeof(double), &actor->position.y);
         // 7, double actor_z
-        clSetKernelArg(context.opencl.transformation_kernel, 7, sizeof(double), &actor->position.z);
+        clSetKernelArg(context.opencl.transformation_kernel, 4, sizeof(double), &actor->position.z);
 
         // 8, double actor_r_w
-        clSetKernelArg(context.opencl.transformation_kernel, 8, sizeof(double), &actor->rotation.w);
+        clSetKernelArg(context.opencl.transformation_kernel, 5, sizeof(double), &actor->rotation.w);
         // 9, double actor_r_x
-        clSetKernelArg(context.opencl.transformation_kernel, 9, sizeof(double), &actor->rotation.x);
+        clSetKernelArg(context.opencl.transformation_kernel, 6, sizeof(double), &actor->rotation.x);
         // 10, double actor_r_y
-        clSetKernelArg(context.opencl.transformation_kernel, 10, sizeof(double), &actor->rotation.y);
+        clSetKernelArg(context.opencl.transformation_kernel, 7, sizeof(double), &actor->rotation.y);
         // 11, double actor_r_z
-        clSetKernelArg(context.opencl.transformation_kernel, 11, sizeof(double), &actor->rotation.z);
+        clSetKernelArg(context.opencl.transformation_kernel, 8, sizeof(double), &actor->rotation.z);
 
         // 12, double actor_s_x
-        clSetKernelArg(context.opencl.transformation_kernel, 12, sizeof(double), &actor->scale.x);
+        clSetKernelArg(context.opencl.transformation_kernel, 9, sizeof(double), &actor->scale.x);
         // 13, double actor_s_y
-        clSetKernelArg(context.opencl.transformation_kernel, 13, sizeof(double), &actor->scale.y);
+        clSetKernelArg(context.opencl.transformation_kernel, 10, sizeof(double), &actor->scale.y);
         // 14, double actor_s_z
-        clSetKernelArg(context.opencl.transformation_kernel, 14, sizeof(double), &actor->scale.z);
+        clSetKernelArg(context.opencl.transformation_kernel, 11, sizeof(double), &actor->scale.z);
 
         // 15, int has_normals
-        clSetKernelArg(context.opencl.transformation_kernel, 15, sizeof(int), &actor->mesh->has_normals);
+        clSetKernelArg(context.opencl.transformation_kernel, 12, sizeof(int), &actor->mesh->has_normals);
         // 16, int has_uvs
-        clSetKernelArg(context.opencl.transformation_kernel, 16, sizeof(int), &actor->mesh->has_uvs);
+        clSetKernelArg(context.opencl.transformation_kernel, 13, sizeof(int), &actor->mesh->has_uvs);
         // 17, int face_array_offset_index
-        clSetKernelArg(context.opencl.transformation_kernel, 17, sizeof(int), &actor->mesh->face_index);
+        clSetKernelArg(context.opencl.transformation_kernel, 14, sizeof(int), &actor->mesh->face_index);
         // 18, int face_count
-        clSetKernelArg(context.opencl.transformation_kernel, 18, sizeof(int), &actor->mesh->face_count);
+        clSetKernelArg(context.opencl.transformation_kernel, 15, sizeof(int), &actor->mesh->face_count);
 
         // 32, int renderable_face_offset
-        clSetKernelArg(context.opencl.transformation_kernel, 32, sizeof(int), &renderable_face_index);
+        clSetKernelArg(context.opencl.transformation_kernel, 29, sizeof(int), &renderable_face_index);
         
         // 33, int face_sqrt
-        clSetKernelArg(context.opencl.transformation_kernel, 33, sizeof(int), &face_sqrt);
+        clSetKernelArg(context.opencl.transformation_kernel, 30, sizeof(int), &face_sqrt);
 
         debug("running kernel...");
 
-        clEnqueueNDRangeKernel(context.opencl.queue, context.opencl.transformation_kernel, 2, NULL, t_global_work_size, t_local_work_size, 0, NULL, &event);
+        cl_int error = clEnqueueNDRangeKernel(context.opencl.queue, context.opencl.transformation_kernel, 2, NULL, t_global_work_size, t_local_work_size, 0, NULL, &event);
+
+        if (error != CL_SUCCESS) debug("error enqueing kernel (%d)", error);
         clFinish(context.opencl.queue);
 
         clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(start), &start, NULL);
@@ -784,84 +787,80 @@ int RI_init(){
     // double camera_x, double camera_y, double camera_z, 
     // double camera_r_w, double camera_r_x, double camera_r_y, double camera_r_z
 
-    // // 0, __global RI_face *faces
+    // // 0: __global RI_face *faces
     // clSetKernelArg(context.opencl.transformation_kernel, 0, sizeof(cl_mem), &context.opencl.faces_mem_buffer);
-    // // 1, __global RI_vector_3 *vertecies
-    // clSetKernelArg(context.opencl.transformation_kernel, 1, sizeof(cl_mem), &context.opencl.vertecies_mem_buffer);
-    // // 2, __global RI_vector_3 *normals
-    // clSetKernelArg(context.opencl.transformation_kernel, 2, sizeof(cl_mem), &context.opencl.normals_mem_buffer);
-    // // 3, __global RI_vector_2 *uvs
-    // clSetKernelArg(context.opencl.transformation_kernel, 3, sizeof(cl_mem), &context.opencl.uvs_mem_buffer);
-    // // 4, __global RI_renderable_face *renderable_faces
-    // clSetKernelArg(context.opencl.transformation_kernel, 4, sizeof(cl_mem), &context.opencl.renderable_faces_mem_buffer);
-
-    // // 5, double actor_x
-    // clSetKernelArg(context.opencl.transformation_kernel, 5, sizeof(double), &actor_x);
-    // // 6, double actor_y
-    // clSetKernelArg(context.opencl.transformation_kernel, 6, sizeof(double), &actor_y);
-    // // 7, double actor_z
-    // clSetKernelArg(context.opencl.transformation_kernel, 7, sizeof(double), &actor_z);
-
-    // // 8, double actor_r_w
-    // clSetKernelArg(context.opencl.transformation_kernel, 8, sizeof(double), &actor_r_w);
-    // // 9, double actor_r_x
-    // clSetKernelArg(context.opencl.transformation_kernel, 9, sizeof(double), &actor_r_x);
-    // // 10, double actor_r_y
-    // clSetKernelArg(context.opencl.transformation_kernel, 10, sizeof(double), &actor_r_y);
-    // // 11, double actor_r_z
-    // clSetKernelArg(context.opencl.transformation_kernel, 11, sizeof(double), &actor_r_z);
-
-    // // 12, double actor_s_x
-    // clSetKernelArg(context.opencl.transformation_kernel, 12, sizeof(double), &actor_s_x);
-    // // 13, double actor_s_y
-    // clSetKernelArg(context.opencl.transformation_kernel, 13, sizeof(double), &actor_s_y);
-    // // 14, double actor_s_z
-    // clSetKernelArg(context.opencl.transformation_kernel, 14, sizeof(double), &actor_s_z);
-
-    // // 15, int has_normals
-    // clSetKernelArg(context.opencl.transformation_kernel, 15, sizeof(int), &has_normals);
-    // // 16, int has_uvs
-    // clSetKernelArg(context.opencl.transformation_kernel, 16, sizeof(int), &has_uvs);
-    // // 17, int face_array_offset_index
-    // clSetKernelArg(context.opencl.transformation_kernel, 17, sizeof(int), &face_array_offset_index);
-    // // 18, int face_count
-    // clSetKernelArg(context.opencl.transformation_kernel, 18, sizeof(int), &face_count);
-
-    // // 19, int width
-    clSetKernelArg(context.opencl.transformation_kernel, 19, sizeof(int), &context.window.width);
-    // // 20, int height
-    clSetKernelArg(context.opencl.transformation_kernel, 20, sizeof(int), &context.window.height);
-
-    // // 21, double horizontal_fov_factor
-    // clSetKernelArg(context.opencl.transformation_kernel, 21, sizeof(double), &horizontal_fov_factor);
-    // // 22, double vertical_fov_factor
-    // clSetKernelArg(context.opencl.transformation_kernel, 22, sizeof(double), &vertical_fov_factor);
-
-    // // 23, double min_clip
-    // clSetKernelArg(context.opencl.transformation_kernel, 23, sizeof(double), &min_clip);
-    // // 24, double max_clip
-    // clSetKernelArg(context.opencl.transformation_kernel, 24, sizeof(double), &max_clip);
-
-    // // 25, double camera_x
-    // clSetKernelArg(context.opencl.transformation_kernel, 25, sizeof(double), &camera_x);
-    // // 26, double camera_y
-    // clSetKernelArg(context.opencl.transformation_kernel, 26, sizeof(double), &camera_y);
-    // // 27, double camera_z
-    // clSetKernelArg(context.opencl.transformation_kernel, 27, sizeof(double), &camera_z);
-
-    // // 28, double camera_r_w
-    // clSetKernelArg(context.opencl.transformation_kernel, 28, sizeof(double), &camera_r_w);
-    // // 29, double camera_r_x
-    // clSetKernelArg(context.opencl.transformation_kernel, 29, sizeof(double), &camera_r_x);
-    // // 30, double camera_r_y
-    // clSetKernelArg(context.opencl.transformation_kernel, 30, sizeof(double), &camera_r_y);
-    // // 31, double camera_r_z
-    // clSetKernelArg(context.opencl.transformation_kernel, 31, sizeof(double), &camera_r_z);
-    
-    // // 32, int renderable_face_offset
-    // clSetKernelArg(context.opencl.transformation_kernel, 32, sizeof(int), &renderable_face_offset);
-    // // 33, int face_sqrt
-    // clSetKernelArg(context.opencl.transformation_kernel, 33, sizeof(int), &face_sqrt);
+
+    // // 1: __global RI_renderable_face *renderable_faces
+    // clSetKernelArg(context.opencl.transformation_kernel, 1, sizeof(cl_mem), &context.opencl.renderable_faces_mem_buffer);
+
+    // // 2: double actor_x
+    // clSetKernelArg(context.opencl.transformation_kernel, 2, sizeof(double), &actor_x);
+    // // 3: double actor_y
+    // clSetKernelArg(context.opencl.transformation_kernel, 3, sizeof(double), &actor_y);
+    // // 4: double actor_z
+    // clSetKernelArg(context.opencl.transformation_kernel, 4, sizeof(double), &actor_z);
+
+    // // 5: double actor_r_w
+    // clSetKernelArg(context.opencl.transformation_kernel, 5, sizeof(double), &actor_r_w);
+    // // 6: double actor_r_x
+    // clSetKernelArg(context.opencl.transformation_kernel, 6, sizeof(double), &actor_r_x);
+    // // 7: double actor_r_y
+    // clSetKernelArg(context.opencl.transformation_kernel, 7, sizeof(double), &actor_r_y);
+    // // 8: double actor_r_z
+    // clSetKernelArg(context.opencl.transformation_kernel, 8, sizeof(double), &actor_r_z);
+
+    // // 9: double actor_s_x
+    // clSetKernelArg(context.opencl.transformation_kernel, 9, sizeof(double), &actor_s_x);
+    // // 10: double actor_s_y
+    // clSetKernelArg(context.opencl.transformation_kernel, 10, sizeof(double), &actor_s_y);
+    // // 11: double actor_s_z
+    // clSetKernelArg(context.opencl.transformation_kernel, 11, sizeof(double), &actor_s_z);
+
+    // // 12: int has_normals
+    // clSetKernelArg(context.opencl.transformation_kernel, 12, sizeof(int), &has_normals);
+    // // 13: int has_uvs
+    // clSetKernelArg(context.opencl.transformation_kernel, 13, sizeof(int), &has_uvs);
+
+    // // 14: int face_array_offset_index
+    // clSetKernelArg(context.opencl.transformation_kernel, 14, sizeof(int), &face_array_offset_index);
+    // // 15: int face_count
+    // clSetKernelArg(context.opencl.transformation_kernel, 15, sizeof(int), &face_count);
+
+    // // 16: int width
+    clSetKernelArg(context.opencl.transformation_kernel, 16, sizeof(int), &context.window.width);
+    // // 17: int height
+    clSetKernelArg(context.opencl.transformation_kernel, 17, sizeof(int), &context.window.height);
+
+    // // 18: double horizontal_fov_factor
+    // clSetKernelArg(context.opencl.transformation_kernel, 18, sizeof(double), &horizontal_fov_factor);
+    // // 19: double vertical_fov_factor
+    // clSetKernelArg(context.opencl.transformation_kernel, 19, sizeof(double), &vertical_fov_factor);
+
+    // // 20: float min_clip
+    // clSetKernelArg(context.opencl.transformation_kernel, 20, sizeof(float), &min_clip_f);
+    // // 21: float max_clip
+    // clSetKernelArg(context.opencl.transformation_kernel, 21, sizeof(float), &max_clip_f);
+
+    // // 22: double camera_x
+    // clSetKernelArg(context.opencl.transformation_kernel, 22, sizeof(double), &camera_x);
+    // // 23: double camera_y
+    // clSetKernelArg(context.opencl.transformation_kernel, 23, sizeof(double), &camera_y);
+    // // 24: double camera_z
+    // clSetKernelArg(context.opencl.transformation_kernel, 24, sizeof(double), &camera_z);
+
+    // // 25: double camera_r_w
+    // clSetKernelArg(context.opencl.transformation_kernel, 25, sizeof(double), &camera_r_w);
+    // // 26: double camera_r_x
+    // clSetKernelArg(context.opencl.transformation_kernel, 26, sizeof(double), &camera_r_x);
+    // // 27: double camera_r_y
+    // clSetKernelArg(context.opencl.transformation_kernel, 27, sizeof(double), &camera_r_y);
+    // // 28: double camera_r_z
+    // clSetKernelArg(context.opencl.transformation_kernel, 28, sizeof(double), &camera_r_z);
+
+    // // 29: int renderable_face_offset
+    // clSetKernelArg(context.opencl.transformation_kernel, 29, sizeof(int), &renderable_face_offset);
+    // // 30: int face_sqrt
+    // clSetKernelArg(context.opencl.transformation_kernel, 30, sizeof(int), &face_sqrt);
 
     context.defaults.default_actor->mesh = RI_load_mesh("objects/cube.obj");