moved headers from headers folder to new kernels folder

src/kernels/master_kernel.h

@@ -0,0 +1,272 @@
+const char *kernel_source_master = " \
+int is_intersecting(float a, float b, float c, float d, float p, float q, float r, float s) { \
+    float det, gamma, lambda; \
+    \
+    det = (c - a) * (s - q) - (r - p) * (d - b); \
+    \
+    if (det == 0) { \
+        return 1; \
+    }  \
+    else { \
+        lambda = ((s - q) * (r - a) + (p - r) * (s - b)) / det; \
+        gamma = ((b - d) * (r - a) + (c - a) * (s - b)) / det; \
+        return (0 < lambda && lambda < 1) && (0 < gamma && gamma < 1); \
+    } \
+} \
+\
+void norm(float dest[2], float a[2]){ \
+    float magnitude = sqrt(a[0] * a[0] + a[1] * a[1]); \
+    \
+    dest[0] = a[0] / magnitude; \
+    dest[1] = a[1] / magnitude; \
+    } \
+    \
+    void sub(float dest[2], float a[2], float b[2]){ \
+    dest[0] = a[0] - b[0]; \
+    dest[1] = a[1] - b[1]; \
+    } \
+    \
+    void add(float dest[2], float a[2], float b[2]){ \
+    dest[0] = a[0] + b[0]; \
+    dest[1] = a[1] + b[1]; \
+} \
+\
+__kernel void raster_kernel(__global int* objects, __global float* verticies, __global float* normals, __global float* uvs, __global int* triangles, __global uint* frame_buffer, __global uchar* textures, __global int* texture_info, int object_count, int width, int height, int show_buffer){ \
+    int id_x = get_global_id(0) - width / 2; \
+    int id_y = get_global_id(1) - height / 2; \
+    \
+    float z_pixel = 0; \
+    uint frame_pixel = 0x22222222; \
+    \
+    float highest_z = 800;\
+    float lowest_z = 0;\
+    \
+    int has_normals = 1;\
+    int has_uvs = 1;\
+    \
+    float w0;\
+    float w1;\
+    float w2;\
+    \
+\
+    for (int object = 0; object < object_count; object++){ \
+        int base = object * 15;\
+        \
+        int object_x =   objects[base + 0]; \
+        int object_y =   objects[base + 1]; \
+        int object_z =   objects[base + 2]; \
+        int object_r_x = objects[base + 3]; \
+        int object_r_y = objects[base + 4]; \
+        int object_r_z = objects[base + 5]; \
+        int object_s_x = objects[base + 6]; \
+        int object_s_y = objects[base + 7]; \
+        int object_s_z = objects[base + 8]; \
+        \
+        int triangle_count = objects[base + 9];\
+        int triangle_index = objects[base + 10];\
+        int vertex_index = objects[base + 11];\
+        int normal_index = objects[base + 12];\
+        int uv_index = objects[base + 13];\
+        int texture_index = objects[base + 14];\
+        \
+        for (int triangle = 0; triangle < triangle_count; triangle++){\
+            int triangle_base = (triangle + triangle_index) * 9; \
+\
+            int i0 = (vertex_index + triangles[triangle_base + 0]) * 3;\
+            int i1 = (vertex_index + triangles[triangle_base + 1]) * 3;\
+            int i2 = (vertex_index + triangles[triangle_base + 2]) * 3;\
+\
+            int i3 = (normal_index + triangles[triangle_base + 3]) * 3;\
+            int i4 = (normal_index + triangles[triangle_base + 4]) * 3;\
+            int i5 = (normal_index + triangles[triangle_base + 5]) * 3;\
+\
+            int i6 = (uv_index + triangles[triangle_base + 6]) * 3;\
+            int i7 = (uv_index + triangles[triangle_base + 7]) * 3;\
+            int i8 = (uv_index + triangles[triangle_base + 8]) * 3;\
+            \
+            float z0 = verticies[i0 + 2] * object_s_z + object_z;\
+            float x0 = (verticies[i0 + 0] * object_s_x + object_x);\
+            float y0 = (verticies[i0 + 1] * object_s_y + object_y);\
+            \
+            float z1 = verticies[i1 + 2] * object_s_z + object_z;\
+            float x1 = (verticies[i1 + 0] * object_s_x + object_x);\
+            float y1 = (verticies[i1 + 1] * object_s_y + object_y);\
+            \
+            float z2 = verticies[i2 + 2] * object_s_z + object_z;\
+            float x2 = (verticies[i2 + 0] * object_s_x + object_x);\
+            float y2 = (verticies[i2 + 1] * object_s_y + object_y);\
+            \
+            \
+            if (i3 < 0 || i4 < 0 || i5 < 0){\
+                has_normals = 0;\
+            }\
+            if (i6 < 0 || i7 < 0 || i8 < 0){\
+                has_uvs = 0;\
+            }\
+            \
+            if (isinf(x0) || isinf(y0) || isinf(z0) || isinf(x1) || isinf(y1) || isinf(z1) || isinf(x2) || isinf(y2) || isinf(z2)){\
+                continue;\
+            }\
+            \
+            float smallest_x = x0; \
+            float largest_x = x0; \
+            float smallest_y = y0; \
+            float largest_y = y0; \
+            \
+            if (x0 > largest_x) largest_x = x0;\
+            if (x1 > largest_x) largest_x = x1;\
+            if (x2 > largest_x) largest_x = x2;\
+            \
+            if (x0 < smallest_x) smallest_x = x0;\
+            if (x1 < smallest_x) smallest_x = x1;\
+            if (x2 < smallest_x) smallest_x = x2;\
+            \
+            if (y0 > largest_y) largest_y = y0;\
+            if (y1 > largest_y) largest_y = y1;\
+            if (y2 > largest_y) largest_y = y2;\
+            \
+            if (y0 < smallest_y) smallest_y = y0;\
+            if (y1 < smallest_y) smallest_y = y1;\
+            if (y2 < smallest_y) smallest_y = y2;\
+            \
+            smallest_x = fmin(smallest_x, 0); \
+            largest_x = fmax(largest_x, width); \
+            smallest_y = fmin(smallest_y, 0); \
+            largest_y = fmax(largest_y, height); \
+            \
+            if (id_x >= smallest_x && id_x <= largest_x && id_y >= smallest_y && id_y <= largest_y){ \
+                int intersections = 0; \
+                \
+                intersections += is_intersecting(id_x, id_y, 10000, 100000, x0, y0, x1, y1); \
+                intersections += is_intersecting(id_x, id_y, 10000, 100000, x1, y1, x2, y2); \
+                intersections += is_intersecting(id_x, id_y, 10000, 100000, x2, y2, x0, y0); \
+                \
+                if (intersections % 2 == 0){ \
+                    continue; \
+                } \
+                \
+                float denominator = (y1 - y2) * (x0 - x2) + (x2 - x1) * (y0 - y2); \
+                \
+                if (denominator <= 0) { \
+                    continue; \
+                } \
+                w0 = ((y1 - y2) * (id_x - x2) + (x2 - x1) * (id_y - y2)) / denominator; \
+                w1 = ((y2 - y0) * (id_x - x0) + (x0 - x2) * (id_y - y0)) / denominator; \
+                w2 = 1.0 - w0 - w1; \
+                \
+                float z = w0 * z0 + w1 * z1 + w2 * z2; \
+                \
+                if (z > z_pixel){ \
+                    z_pixel = z; \
+                    \
+                    float n_x0 = normals[i3 + 0];\
+                    float n_y0 = normals[i3 + 1];\
+                    float n_z0 = normals[i3 + 2];\
+                    \
+                    float n_x1 = normals[i4 + 0];\
+                    float n_y1 = normals[i4 + 1];\
+                    float n_z1 = normals[i4 + 2];\
+                    \
+                    float n_x2 = normals[i5 + 0];\
+                    float n_y2 = normals[i5 + 1];\
+                    float n_z2 = normals[i5 + 2];\
+                    \
+                    float u_x0 = uvs[i6 + 0];\
+                    float u_y0 = uvs[i6 + 1];\
+                    float u_z0 = uvs[i6 + 2];\
+                    \
+                    float u_x1 = uvs[i7 + 0];\
+                    float u_y1 = uvs[i7 + 1];\
+                    float u_z1 = uvs[i7 + 2];\
+                    \
+                    float u_x2 = uvs[i8 + 0];\
+                    float u_y2 = uvs[i8 + 1];\
+                    float u_z2 = uvs[i8 + 2];\
+                    \
+                    switch (show_buffer){\
+                        case 0:{\
+                            float ux = w0 * u_x0 + w1 * u_x1 + w2 * u_x2;\
+                            float uy = w0 * u_y0 + w1 * u_y1 + w2 * u_y2;\
+                            \
+                            int texture_width = texture_info[texture_index * 2];\
+                            int texture_height = texture_info[texture_index * 2 + 1];\
+                            \
+                            int ix = clamp((int)ux * texture_width, 0, texture_width - 1);\
+                            int iy = clamp((int)uy * texture_height, 0, texture_height - 1);\
+                            \
+                            int pixel_offset = (iy * texture_width + ix) * 4;\
+                            \
+                            uchar r = textures[0 + 0];\
+                            uchar g = textures[0 + 1];\
+                            uchar b = textures[0 + 2];\
+                            uchar a = textures[0 + 3];\
+                            \
+                            frame_pixel = (a << 24) | (r << 16) | (g << 8) | b;\
+                            \
+                            break;}\
+                        case 1:{\
+                            float z = clamp(z_pixel, 0.0f, highest_z);\
+                            \
+                            float norm_z = z / highest_z;\
+                            \
+                            uchar intensity = (uchar)(norm_z * 255.0f);\
+                            \
+                            frame_pixel = 0xFF000000 | (intensity << 16) | (intensity << 8) | intensity;\
+                            \
+                            break;}\
+                        case 2:{\
+                            float nx = w0 * n_x0 + w1 * n_x1 + w2 * n_x2;\
+                            float ny = w0 * n_y0 + w1 * n_y1 + w2 * n_y2;\
+                            float nz = w0 * n_z0 + w1 * n_z1 + w2 * n_z2;\
+                            nx = clamp((nx * 0.5f + 0.5f) * 255.0f, 0.0f, 255.0f);\
+                            ny = clamp((ny * 0.5f + 0.5f) * 255.0f, 0.0f, 255.0f);\
+                            nz = clamp((nz * 0.5f + 0.5f) * 255.0f, 0.0f, 255.0f);\
+                            \
+                            uchar r = (uchar)nx;\
+                            uchar g = (uchar)ny;\
+                            uchar b = (uchar)nz;\
+                            \
+                            if (!has_normals){\
+                                r = 20;\
+                                g = 20;\
+                                b = 20;\
+                            }\
+                            \
+                            frame_pixel = 0xFF000000 | (r << 16) | (g << 8) | b;\
+                            \
+                            break;}\
+                        case 3:{\
+                            float ux = w0 * u_x0 + w1 * u_x1 + w2 * u_x2;\
+                            float uy = w0 * u_y0 + w1 * u_y1 + w2 * u_y2;\
+                            \
+                            uchar r = (uchar)clamp(ux * 255.0f, 0.0f, 255.0f);\
+                            uchar g = (uchar)clamp(uy * 255.0f, 0.0f, 255.0f);\
+                            uchar b = 0;\
+                            \
+                            if (!has_uvs){\
+                                r = 20;\
+                                g = 20;\
+                                b = 20;\
+                            }\
+                            \
+                            frame_pixel = 0xFF000000 | (r << 16) | (g << 8) | b;\
+                            \
+                            break;}\
+                        default:{\
+                            frame_pixel = 0xFF00FFFF;\
+                            \
+                            break;}\
+                    }\
+                } \
+            }\
+        }\
+    }\
+    \
+    int pixel_coord = (height * 0.5 - id_y) * width + id_x + width * 0.5;\
+    \
+    if (pixel_coord >= width * height || pixel_coord < 0){\
+        return;\
+    }\
+    frame_buffer[pixel_coord] = frame_pixel; \
+    \
+}\n";

src/kernels/non_master_kernel.h

@@ -0,0 +1,145 @@
+const char *kernel_source_non_master = " \
+int is_intersecting(float a, float b, float c, float d, float p, float q, float r, float s) { \
+    float det, gamma, lambda; \
+    \
+    det = (c - a) * (s - q) - (r - p) * (d - b); \
+    \
+    if (det == 0) { \
+        return 1; \
+    }  \
+    else { \
+        lambda = ((s - q) * (r - a) + (p - r) * (s - b)) / det; \
+        gamma = ((b - d) * (r - a) + (c - a) * (s - b)) / det; \
+        return (0 < lambda && lambda < 1) && (0 < gamma && gamma < 1); \
+    } \
+} \
+\
+void norm(float dest[2], float a[2]){ \
+    float magnitude = sqrt(a[0] * a[0] + a[1] * a[1]); \
+    \
+    dest[0] = a[0] / magnitude; \
+    dest[1] = a[1] / magnitude; \
+    } \
+    \
+    void sub(float dest[2], float a[2], float b[2]){ \
+    dest[0] = a[0] - b[0]; \
+    dest[1] = a[1] - b[1]; \
+    } \
+    \
+    void add(float dest[2], float a[2], float b[2]){ \
+    dest[0] = a[0] + b[0]; \
+    dest[1] = a[1] + b[1]; \
+} \
+\
+__kernel void raster_kernel(__global float* polygons, __global uint* frame_buffer, int polygon_count, int width, int height, int show_z_buffer, float highest_z){ \
+    int id_x = get_global_id(0); \
+    int id_y = get_global_id(1); \
+    \
+    float z_pixel = 0; \
+    uint frame_pixel = 0x22222222; \
+    \
+    float biggest_z = 0;\
+    \
+    for (int polygon = 0; polygon < polygon_count; polygon++){ \
+        int base = polygon * 9; \
+        float x0 = polygons[base]; \
+        float y0 = polygons[base + 1]; \
+        float z0 = polygons[base + 2]; \
+        float x1 = polygons[base + 3]; \
+        float y1 = polygons[base + 4]; \
+        float z1 = polygons[base + 5]; \
+        float x2 = polygons[base + 6]; \
+        float y2 = polygons[base + 7]; \
+        float z2 = polygons[base + 8]; \
+        \
+        if (isinf(x0) || isinf(y0) || isinf(z0) || isinf(x1) || isinf(y1) || isinf(z1) || isinf(x2) || isinf(y2) || isinf(z2)){\
+            return;\
+        }\
+        \
+        float smallest_x = x0; \
+        float largest_x = x0; \
+        float smallest_y = y0; \
+        float largest_y = y0; \
+        \
+        for (int point = 0; point < 3; point++){ \
+            float x = polygons[base + point * 3]; \
+            float y = polygons[base + point * 3 + 1]; \
+            \
+            if (x > largest_x){ \
+                largest_x = x; \
+            } \
+            \
+            if (x < smallest_x){ \
+                smallest_x = x; \
+            } \
+            \
+            if (y > largest_y){ \
+                largest_y = y; \
+            } \
+            \
+            if (y < smallest_y){\
+                smallest_y = y;\
+            } \
+        } \
+        \
+        smallest_x = fmin(smallest_x, 0); \
+        largest_x = fmax(largest_x, width); \
+        smallest_y = fmin(smallest_y, 0); \
+        largest_y = fmax(largest_y, height); \
+        \
+        if (id_x >= smallest_x && id_x <= largest_x && id_y >= smallest_y && id_y <= largest_y){ \
+            int intersections = 0; \
+            \
+            intersections += is_intersecting(id_x, id_y, 10000, 100000, x0, y0, x1, y1); \
+            intersections += is_intersecting(id_x, id_y, 10000, 100000, x1, y1, x2, y2); \
+            intersections += is_intersecting(id_x, id_y, 10000, 100000, x2, y2, x0, y0); \
+            \
+            if (intersections % 2 == 0){ \
+                continue; \
+            } \
+            \
+            float denominator = (y1 - y2) * (x0 - x2) + (x2 - x1) * (y0 - y2); \
+            float w0 = ((y1 - y2) * (id_x - x2) + (x2 - x1) * (id_y - y2)) / denominator; \
+            float w1 = ((y2 - y0) * (id_x - x0) + (x0 - x2) * (id_y - y2)) / denominator; \
+            float w2 = 1.0 - w0 - w1; \
+            \
+            if (denominator < 0) { \
+                w0 = -w0; \
+                w1 = -w1; \
+                w2 = -w2; \
+                denominator = -denominator; \
+            } \
+            \
+            float z = w0 * z0 + w1 * z1 + w2 * z2; \
+            \
+            if (z < 0){ \
+                z *= -1; \
+            } \
+            \
+            if (z > z_pixel){ \
+                z_pixel = z; \
+            } \
+             \
+            else { \
+                continue; \
+            } \
+            \
+            frame_pixel = 0xFFFFFFFF / polygon_count * (polygon + 1); \
+        } \
+    } \
+    \
+    if (id_y * width + id_x > width * height){\
+    return;\
+    }\
+    frame_buffer[id_y * width + id_x] = frame_pixel; \
+    \
+    if (!show_z_buffer){return;}\
+    \
+    float z = clamp(z_pixel, 0.0f, highest_z);\
+    \
+    float norm_z = z / highest_z;\
+    \
+    uchar intensity = (uchar)(norm_z * 255.0f);\
+    \
+    frame_buffer[id_y * width + id_x] = 0xFF000000 | (intensity << 16) | (intensity << 8) | intensity;\
+}\n";