Initial commitHEADmaster

1 files changed, 886 insertions, 0 deletions

diff --git a/kernel.cl b/kernel.cl
new file mode 100644
index 0000000..fb24c29
--- /dev/null
+++ b/kernel.cl
@@ -0,0 +1,886 @@
+#ifndef EPSILON
+#define EPSILON 1E-6f
+#endif
+
+#ifndef SPLT_EPS
+#define SPLT_EPS 4 * EPSILON
+#endif
+
+#ifndef LGT_EPS
+#define LGT_EPS 5E-5f
+#endif
+
+#ifndef REFR_EPS
+#define REFR_EPS 1E-4f
+#endif
+
+#ifndef NORM_EPS
+#define NORM_EPS 1E-12f
+#endif
+
+constant uchar PRIMITIVE_TYPE_BOX = 0;
+constant uchar PRIMITIVE_TYPE_INFINITEPLANE = 1;
+constant uchar PRIMITIVE_TYPE_SPHERE = 2;
+constant uchar PRIMITIVE_TYPE_TRIANGLE = 3;
+
+constant uchar SHADER_TYPE_BRDF = 0;
+constant uchar SHADER_TYPE_COOKTORRANCE = 1;
+constant uchar SHADER_TYPE_FLAT = 2;
+constant uchar SHADER_TYPE_LAMBERT = 3;
+constant uchar SHADER_TYPE_MATERIAL = 4;
+constant uchar SHADER_TYPE_MIRROR = 5;
+constant uchar SHADER_TYPE_NORMAL = 6;
+constant uchar SHADER_TYPE_PHONG = 7;
+constant uchar SHADER_TYPE_REFRACTION = 8;
+constant uchar SHADER_TYPE_SIMPLESHADOW = 9;
+
+constant uchar LIGHT_TYPE_AMBIENT = 0;
+constant uchar LIGHT_TYPE_POINT = 1;
+constant uchar LIGHT_TYPE_SPOT = 2;
+
+typedef float3 color_t;
+
+typedef struct __attribute__ ((packed)) _shader {
+    const uchar shader_type;
+} shader_t;
+
+typedef struct __attribute__ ((packed)) _shader_cooktorrance {
+    shader_t base;
+    color_t diffuse_color;
+    color_t ct_color;
+    float F0;
+    float m;
+} shader_cooktorrance_t;
+
+typedef struct __attribute__ ((packed)) _shader_flat {
+    shader_t base;
+    color_t object_color;
+} shader_flat_t;
+
+typedef struct __attribute__ ((packed)) _shader_lambert {
+    shader_t base;
+    color_t diffuse_color;
+} shader_lambert_t;
+
+typedef struct __attribute__ ((packed)) _shader_mirror {
+    shader_t base;
+} shader_mirror_t;
+
+typedef struct __attribute__ ((packed)) _shader_normal {
+    shader_t base;
+} shader_normal_t;
+
+typedef struct __attribute__ ((packed)) _shader_phong {
+    shader_t base;
+    color_t diffuse_color;
+    float diffuse_coefficient;
+    color_t specular_color;
+    float specular_coefficient;
+    float shininess_exponent;
+} shader_phong_t;
+
+typedef struct __attribute__ ((packed)) _shader_refraction {
+    shader_t base;
+    float index_inside;
+    float index_outside;
+} shader_refraction_t;
+
+typedef struct __attribute__ ((packed)) _shader_simpleshadow {
+    shader_t base;
+    color_t object_color;
+} shader_simpleshadow_t;
+
+typedef struct __attribute__ ((packed)) _primitive {
+    const uchar primitive_type;
+} primitive_t;
+
+typedef struct __attribute__ ((packed)) _primitive_box {
+    primitive_t base;
+    float3 center;
+    float3 size;
+} primitive_box_t;
+
+typedef struct __attribute__ ((packed)) _primitive_infiniteplane {
+    primitive_t base;
+    float3 origin;
+    float3 normal;
+} primitive_infiniteplane_t;
+
+typedef struct __attribute__ ((packed)) _primitive_sphere {
+    primitive_t base;
+    float3 center;
+    float radius;
+} primitive_sphere_t;
+
+typedef struct __attribute__ ((packed)) _primitive_triangle {
+    primitive_t base;
+    float3 vertices[3];
+    float3 normals[3];
+    float3 tangents[3];
+    float3 bitangents[3];
+    float2 surfaces[3];
+} primitive_triangle_t;
+
+typedef struct __attribute__ ((packed)) _ray {
+    float3 origin;
+    float3 direction;
+    float length;
+    float3 normal;
+    float2 surface;
+    float3 tangent;
+    float3 bitangent;
+    uint primitive_index;
+    uint bounces;
+} ray_t;
+
+typedef struct __attribute__ ((packed)) _light {
+    uchar light_type;
+    color_t color;
+    float intensity;
+} light_t;
+
+typedef struct __attribute__ ((packed)) _light_ambient {
+    light_t base;
+} light_ambient_t;
+
+typedef struct __attribute__ ((packed)) _light_point {
+    light_t base;
+    float3 position;
+} light_point_t;
+
+typedef struct __attribute__ ((packed)) _light_spot {
+    light_t base;
+    float3 position;
+    float3 direction;
+    float alpha_min;
+    float alpha_max;
+} light_spot_t;
+
+typedef struct __attribute__ ((packed)) _illumination {
+    color_t color;
+    float3 direction;
+} illumination_t;
+
+typedef struct _scene {
+    const color_t background_color;
+    __global const primitive_t *primitives;
+    const uint n_primitives;
+    __global const light_t *lights;
+    const uint n_lights;
+    __global const shader_t *shaders;
+} scene_t;
+
+bool intersect_box(__global const primitive_box_t *box, ray_t *r);
+bool intersect_infiniteplane(__global const primitive_infiniteplane_t *plane, ray_t *r);
+bool intersect_sphere(__global const primitive_sphere_t *sphere, ray_t *r);
+bool intersect_triangle(__global const primitive_triangle_t *triangle, ray_t *r);
+bool is_transparent(__global const shader_t *shader);
+bool find_occlusion(const scene_t *scene, ray_t *ray);
+void illuminate_ambient(illumination_t *illum, __global const light_ambient_t *light, const scene_t *scene, ray_t *ray);
+void illuminate_point(illumination_t *illum, __global const light_point_t *light, const scene_t *scene, ray_t *ray);
+void illuminate_spot(illumination_t *illum, __global const light_spot_t *light, const scene_t *scene, ray_t *ray);
+color_t shade_simpleshadow(__global const shader_simpleshadow_t *shader, const scene_t *scene, ray_t *ray);
+float ct_d(__global const shader_cooktorrance_t *shader, float nxh);
+float ct_f(__global const shader_cooktorrance_t *shader, float vxh);
+float ct_g(__global const shader_cooktorrance_t *shader, float nxh, float nxv, float vxh, float nxl);
+color_t shade_cooktorrance(__global const shader_cooktorrance_t *shader, const scene_t *scene, ray_t *ray);
+color_t shade_flat(__global const shader_flat_t *shader, const scene_t *scene, ray_t *ray);
+color_t shade_lambert(__global const shader_lambert_t *shader, const scene_t *scene, ray_t *ray);
+color_t shade_mirror(__global const shader_mirror_t *shader, const scene_t *scene, ray_t *ray);
+color_t shade_normal(__global const shader_normal_t *shader, const scene_t *scene, ray_t *ray);
+color_t shade_phong(__global const shader_phong_t *shader, const scene_t *scene, ray_t *ray);
+color_t shade_refraction(__global const shader_refraction_t *shader, const scene_t *scene, ray_t *ray);
+void trace_ray(const scene_t *scene, color_t *target_color, ray_t *target_ray);
+__global primitive_t *get_primitive(const scene_t *scene, uint n);
+__global shader_t *get_shader(const scene_t *scene, uint n);
+
+__global primitive_t *get_primitive(const scene_t *scene, uint n)
+{
+    __global const primitive_t *base = scene->primitives;
+
+    for(uint i = 0; i < n; i++) {
+        switch(base->primitive_type) {
+        case PRIMITIVE_TYPE_BOX:
+            base = (__global const primitive_t *) (((__global const primitive_box_t *) base) + 1);
+            break;
+        case PRIMITIVE_TYPE_INFINITEPLANE:
+            base = (__global const primitive_t *) (((__global const primitive_infiniteplane_t *) base) + 1);
+            break;
+        case PRIMITIVE_TYPE_SPHERE:
+            base = (__global const primitive_t *) (((__global const primitive_sphere_t *) base) + 1);
+            break;
+        case PRIMITIVE_TYPE_TRIANGLE:
+            base = (__global const primitive_t *) (((__global const primitive_triangle_t *) base) + 1);
+            break;
+        }
+    }
+
+    return base;
+}
+
+__global shader_t *get_shader(const scene_t *scene, uint n)
+{
+    __global const shader_t *base = scene->shaders;
+
+    for(uint i = 0; i < n; i++) {
+        switch(base->shader_type) {
+        case SHADER_TYPE_COOKTORRANCE:
+            base = (__global const primitive_t *) (((__global const shader_cooktorrance_t *) base) + 1);
+            break;
+        case SHADER_TYPE_FLAT:
+            base = (__global const primitive_t *) (((__global const shader_flat_t *) base) + 1);
+            break;
+        case SHADER_TYPE_LAMBERT:
+            base = (__global const primitive_t *) (((__global const shader_lambert_t *) base) + 1);
+            break;
+        case SHADER_TYPE_MIRROR:
+            base = (__global const primitive_t *) (((__global const shader_mirror_t *) base) + 1);
+            break;
+        case SHADER_TYPE_NORMAL:
+            base = (__global const primitive_t *) (((__global const shader_normal_t *) base) + 1);
+            break;
+        case SHADER_TYPE_PHONG:
+            base = (__global const primitive_t *) (((__global const shader_phong_t *) base) + 1);
+            break;
+        case SHADER_TYPE_REFRACTION:
+            base = (__global const primitive_t *) (((__global const shader_refraction_t *) base) + 1);
+            break;
+        default:
+        case SHADER_TYPE_SIMPLESHADOW:
+            base = (__global const primitive_t *) (((__global const shader_simpleshadow_t *) base) + 1);
+            break;
+        }
+    }
+
+    return base;
+}
+
+bool intersect_box(__global const primitive_box_t *box, ray_t *r)
+{
+    float3 const minBounds = box->center - box->size / 2;
+    float3 const maxBounds = box->center + box->size / 2;
+
+    float3 t1 = (minBounds - r->origin) / r->direction;
+    float3 t2 = (maxBounds - r->origin) / r->direction;
+
+    float tNear = -INFINITY;
+    float tFar = +INFINITY;
+    int tNearIndex = 0;
+    int tFarIndex = 0;
+
+    for (int d = 0; d < 3; ++d) {
+        if (r->direction[d] == 0 && (r->origin[d] < minBounds[d] || r->origin[d] > maxBounds[d]))
+            return false;
+
+        if (t1[d] > t2[d]) {
+            float3 temp = t1;
+            t1 = t2;
+            t2 = temp;
+        }
+
+        if (t1[d] > tNear) {
+            tNear = t1[d];
+            tNearIndex = d;
+        }
+
+        if (t2[d] < tFar) {
+            tFar = t2[d];
+            tFarIndex = d;
+        }
+
+        if (tFar < 0 || tNear > tFar)
+            return false;
+    }
+
+    float const t = (tNear >= 0 ? tNear : tFar);
+    int const tIndex = tNear >= 0 ? tNearIndex : tFarIndex;
+
+    if (r->length < t)
+        return false;
+
+    r->normal = (float3) (0.0f, 0.0f, 0.0f);
+    r->normal[tIndex] = 1.0f * ((r->direction[tIndex] < 0.0f) ? -1.0f : +1.0f);
+    r->normal[tIndex] *= ((tNear < 0.0f) ? +1.0f : -1.0f);
+
+    float3 const target = r->origin + t * r->direction;
+    float3 const surface = (target - minBounds) / (maxBounds - minBounds);
+
+    if (tIndex == 0) {
+        r->surface = (float2)(surface[2], surface[1]);
+        r->tangent = (float3)(0, 0, 1);
+    } else if (tIndex == 1) {
+        r->surface = (float2)(surface[0], surface[2]);
+        r->tangent = (float3)(1, 0, 0);
+    } else {
+        r->surface = (float2)(surface[0], surface[1]);
+        r->tangent = (float3)(1, 0, 0);
+    }
+
+    r->length = t;
+
+    return true;
+}
+
+bool intersect_infiniteplane(__global const primitive_infiniteplane_t *plane, ray_t *r)
+{
+    float const cosine = dot(r->direction, plane->normal);
+
+    if(cosine > 0)
+        return false;
+
+    float const t = dot(plane->origin - r->origin, plane->normal) / cosine;
+
+    if (t < EPSILON || r->length < t)
+        return false;
+
+    r->normal = plane->normal;
+    r->length = t;
+
+    return true;
+}
+
+bool intersect_sphere(__global const primitive_sphere_t *sphere, ray_t *r)
+{
+    float3 const difference = r->origin - sphere->center;
+
+    float const a = 1.0f;
+    float const b = 2.0f * dot(r->direction, difference);
+    float const c = dot(difference, difference) - sphere->radius * sphere->radius;
+    float const discriminant = b * b - 4 * a * c;
+
+    if (discriminant < 0)
+        return false;
+
+    float const root = sqrt(discriminant);
+
+    float const q = -0.5f * (b < 0 ? (b - root) : (b + root));
+    float const t0 = q / a;
+    float const t1 = c / q;
+    float t = min(t0, t1);
+    if (t < EPSILON)
+        t = max(t0, t1);
+
+    if (t < EPSILON || r->length < t)
+        return false;
+
+    float3 const hitPoint = r->origin + t * r->direction;
+    r->normal = normalize(hitPoint - sphere->center);
+
+    float const phi = acos(r->normal.y);
+    float const rho = atan2(r->normal.z, r->normal.x) + M_PI_F;
+    r->surface = (float2)(rho / (2 * M_PI_F), phi / M_PI_F);
+    r->tangent = (float3)(sin(rho), 0, cos(rho));
+    r->bitangent = normalize(cross(r->normal, r->tangent));
+
+    r->length = t;
+
+    return true;
+}
+
+bool intersect_triangle(__global const primitive_triangle_t *triangle, ray_t *r)
+{
+    float3 const edge1 = triangle->vertices[1] - triangle->vertices[0];
+    float3 const edge2 = triangle->vertices[2] - triangle->vertices[0];
+
+    float3 const pVec = cross(r->direction, edge2);
+
+    float const det = dot(edge1, pVec);
+    if (fabs(det) < EPSILON)
+        return false;
+    float const inv_det = 1.0f / det;
+
+    float3 const tVec = r->origin - triangle->vertices[0];
+    float const u = dot(tVec, pVec) * inv_det;
+    if (0.0f > u || u > 1.0f)
+        return false;
+
+    float3 const qVec = cross(tVec, edge1);
+    float const v = dot(r->direction, qVec) * inv_det;
+    if (0.0f > v || u + v > 1.0f)
+        return false;
+
+    float const t = dot(edge2, qVec) * inv_det;
+    if (t < EPSILON || r->length < t)
+        return false;
+
+    if (length(triangle->normals[0]) * length(triangle->normals[1]) * length(triangle->normals[2]) > EPSILON)
+        r->normal = normalize(u * triangle->normals[1] + v * triangle->normals[2] + (1 - u - v) * triangle->normals[0]);
+    else
+        r->normal = normalize(cross(edge1, edge2));
+    r->tangent = normalize(u * triangle->tangents[1] + v * triangle->tangents[2] + (1 - u - v) * triangle->tangents[0]);
+    r->bitangent = normalize(u * triangle->bitangents[1] + v * triangle->bitangents[2] + (1 - u - v) * triangle->bitangents[0]);
+
+    r->surface = u * triangle->surfaces[1] + v * triangle->surfaces[2] + (1 - u - v) * triangle->surfaces[0];
+
+    r->length = t;
+
+    return true;
+}
+
+bool is_transparent(__global const shader_t *shader)
+{
+    switch(shader->shader_type) {
+    case SHADER_TYPE_REFRACTION:
+        return true;
+    default:
+        return false;
+    }
+}
+
+bool find_occlusion(const scene_t *scene, ray_t *ray)
+{
+    __global const primitive_t *base = scene->primitives;
+    __global const primitive_box_t *box = NULL;
+    __global const primitive_infiniteplane_t *plane = NULL;
+    __global const primitive_sphere_t *sphere = NULL;
+    __global const primitive_triangle_t *triangle = NULL;
+
+    for(uint i = 0; i < scene->n_primitives; i++) {
+        switch(base->primitive_type) {
+        case PRIMITIVE_TYPE_BOX:
+            box = (__global const primitive_box_t *) base;
+            if(intersect_box((__global const primitive_box_t *) base, ray)) {
+                ray->primitive_index = i;
+                if(is_transparent(get_shader(base->shader_index)))
+                    return true;
+            }
+            base = (__global const primitive_t *) (box + 1);
+            break;
+        case PRIMITIVE_TYPE_INFINITEPLANE:
+            plane = (__global const primitive_infiniteplane_t *) base;
+            if(intersect_infiniteplane((__global const primitive_infiniteplane_t *) base, ray)) {
+                ray->primitive_index = i;
+                if(is_transparent(get_shader(base->shader_index)))
+                    return true;
+            }
+            base = (__global const primitive_t *) (plane + 1);
+            break;
+        case PRIMITIVE_TYPE_SPHERE:
+            sphere = (__global const primitive_sphere_t *) base;
+            if(intersect_sphere((__global const primitive_sphere_t *) base, ray)) {
+                ray->primitive_index = i;
+                if(is_transparent(get_shader(base->shader_index)))
+                    return true;
+            }
+            base = (__global const primitive_t *) (sphere + 1);
+            break;
+        case PRIMITIVE_TYPE_TRIANGLE:
+            triangle = (__global const primitive_triangle_t *) base;
+            if(intersect_triangle((__global const primitive_triangle_t *) base, ray)) {
+                ray->primitive_index = i;
+                if(is_transparent(get_shader(base->shader_index)))
+                    return true;
+            }
+            base = (__global const primitive_t *) (triangle + 1);
+            break;
+        }
+    }
+
+    return false;
+}
+
+void illuminate_ambient(illumination_t *illum, __global const light_ambient_t *light, const scene_t *scene, ray_t *ray)
+{
+    illum->color = light->base.color * light->base.intensity;
+    illum->direction = ray->normal * -1.0f;
+}
+
+void illuminate_point(illumination_t *illum, __global const light_point_t *light, const scene_t *scene, ray_t *ray)
+{
+    float3 const target = ray->origin + (ray->length - LGT_EPS) * ray->direction;
+
+    illum->direction = normalize(target - light->position);
+
+    float const distance = length(target - light->position);
+
+    ray_t light_ray;
+    light_ray.origin = target;
+    light_ray.direction = illum->direction * -1.0f;
+    light_ray.length = distance - LGT_EPS;
+
+    if(!find_occlusion(scene, &light_ray))
+        illum->color = 1.0f / (distance * distance) * light->base.color * light->base.intensity;
+}
+
+void illuminate_spot(illumination_t *illum, __global const light_spot_t *light, const scene_t *scene, ray_t *ray)
+{
+    float3 const target = ray->origin + (ray->length - LGT_EPS) * ray->direction;
+
+    illum->direction = normalize(target - light->position);
+
+    float const distance = length(target - light->position);
+
+    ray_t light_ray;
+    light_ray.origin = target;
+    light_ray.direction = illum->direction * -1.0f;
+    light_ray.length = distance - LGT_EPS;
+
+    float const alpha = fabs(acos(dot(illum->direction, light->direction)) * 180.0f / M_PI_F);
+
+    if (light->alpha_max > alpha) {
+        if (!find_occlusion(scene, &light_ray)) {
+            illum->color = 1.0f / (distance * distance) * light->base.color * light->base.intensity;
+            if (light->alpha_min < alpha)
+                illum->color *= 1.0f - (alpha - light->alpha_min) / (light->alpha_max - light->alpha_min);
+        }
+    }
+}
+
+color_t shade_simpleshadow(__global const shader_simpleshadow_t *shader, const scene_t *scene, ray_t *ray)
+{
+    color_t fragment_color;
+
+    __global const light_t *base = scene->lights;
+    __global const light_ambient_t *ambient = NULL;
+    __global const light_point_t *point = NULL;
+    __global const light_spot_t *spot = NULL;
+
+    for(uint i = 0; i < scene->n_lights; i++) {
+        illumination_t illum;
+        switch(base->light_type) {
+        case LIGHT_TYPE_AMBIENT:
+            ambient = (__global const light_ambient_t *) base;
+            illuminate_ambient(&illum, ambient, scene, ray);
+            base = (__global const light_t *) (ambient + 1);
+            break;
+        case LIGHT_TYPE_POINT:
+            point = (__global const light_point_t *) base;
+            illuminate_point(&illum, point, scene, ray);
+            base = (__global const light_t *) (point + 1);
+            break;
+        case LIGHT_TYPE_SPOT:
+            spot = (__global const light_spot_t *) base;
+            illuminate_spot(&illum, spot, scene, ray);
+            base = (__global const light_t *) (spot + 1);
+            break;
+        }
+        fragment_color += illum.color;
+    }
+
+    return fragment_color * shader->object_color;
+}
+
+float ct_d(__global const shader_cooktorrance_t *shader, float nxh)
+{
+    float const r2 = shader->m * shader->m;
+    float const nxh2 = nxh * nxh;
+    return exp((nxh2 - 1.0f) / (r2 * nxh2)) / (4.0f * r2 * pow(nxh, 4.0f));
+}
+
+float ct_f(__global const shader_cooktorrance_t *shader, float vxh)
+{
+    return shader->F0 + (1.0f - shader->F0) * pow(1.0f - vxh, 5);
+}
+
+float ct_g(__global const shader_cooktorrance_t *shader, float nxh, float nxv, float vxh, float nxl)
+{
+    return min(1.0f, min(2.0f * nxh * nxv / vxh, 2.0f * nxh * nxl / vxh));
+}
+
+color_t shade_cooktorrance(__global const shader_cooktorrance_t *shader, const scene_t *scene, ray_t *ray)
+{
+    color_t fragment_color;
+
+    if (shader->m >= 0.0f) {
+        __global const light_t *base = scene->lights;
+        __global const light_ambient_t *ambient = NULL;
+        __global const light_point_t *point = NULL;
+        __global const light_spot_t *spot = NULL;
+
+        for(uint i = 0; i < scene->n_lights; i++) {
+            illumination_t illum;
+            switch(base->light_type) {
+            case LIGHT_TYPE_AMBIENT:
+                ambient = (__global const light_ambient_t *) base;
+                illuminate_ambient(&illum, ambient, scene, ray);
+                base = (__global const light_t *) (ambient + 1);
+                break;
+            case LIGHT_TYPE_POINT:
+                point = (__global const light_point_t *) base;
+                illuminate_point(&illum, point, scene, ray);
+                base = (__global const light_t *) (point + 1);
+                break;
+            case LIGHT_TYPE_SPOT:
+                spot = (__global const light_spot_t *) base;
+                illuminate_spot(&illum, spot, scene, ray);
+                base = (__global const light_t *) (spot + 1);
+                break;
+            }
+
+            float const nxl = max(0.0f, dot(illum.direction * -1.0f, ray->normal));
+            if (nxl <= 0.0f)
+                continue;
+
+            color_t const diffuse = shader->diffuse_color / M_PI_F;
+            fragment_color += diffuse * nxl * illum.color;
+
+            float3 const h = normalize((illum.direction * -1.0f) - ray->direction);
+            float const nxh = max(0.0f, dot(ray->normal, h));
+            float const nxv = max(0.0f, dot(ray->normal, ray->direction * -1.0f));
+            float const vxh = max(0.0f, dot(ray->direction * -1.0f, h));
+            
+            if (nxv * nxl > EPSILON) {
+                color_t const specular = shader->ct_color * (ct_f(shader, vxh) * ct_d(shader, nxh) * ct_g(shader, nxh, nxv, vxh, nxl)) / (M_PI_F * nxv * nxl);
+
+                fragment_color += specular * nxl * illum.color;
+            }
+        }
+    }
+
+    return fragment_color;
+}
+
+color_t shade_flat(__global const shader_flat_t *shader, const scene_t *scene, ray_t *ray)
+{
+    return shader->object_color;
+}
+
+color_t shade_lambert(__global const shader_lambert_t *shader, const scene_t *scene, ray_t *ray)
+{
+    color_t fragment_color;
+
+    __global const light_t *base = scene->lights;
+    __global const light_ambient_t *ambient = NULL;
+    __global const light_point_t *point = NULL;
+    __global const light_spot_t *spot = NULL;
+
+    for(uint i = 0; i < scene->n_lights; i++) {
+        illumination_t illum;
+        switch(base->light_type) {
+        case LIGHT_TYPE_AMBIENT:
+            ambient = (__global const light_ambient_t *) base;
+            illuminate_ambient(&illum, ambient, scene, ray);
+            base = (__global const light_t *) (ambient + 1);
+            break;
+        case LIGHT_TYPE_POINT:
+            point = (__global const light_point_t *) base;
+            illuminate_point(&illum, point, scene, ray);
+            base = (__global const light_t *) (point + 1);
+            break;
+        case LIGHT_TYPE_SPOT:
+            spot = (__global const light_spot_t *) base;
+            illuminate_spot(&illum, spot, scene, ray);
+            base = (__global const light_t *) (spot + 1);
+            break;
+        }
+        color_t const diffuse = shader->diffuse_color * max(dot(illum.direction * -1.0f, ray->normal), 0.0f);
+        fragment_color += diffuse * illum.color;
+    }
+
+    return fragment_color;
+}
+
+color_t shade_mirror(__global const shader_mirror_t *shader, const scene_t *scene, ray_t *ray)
+{
+    float3 const reflection = ray->direction - 2 * dot(ray->normal, ray->direction) * ray->normal;
+
+    ray_t reflection_ray = *ray;
+    reflection_ray.origin = ray->origin + (ray->length - REFR_EPS) * ray->direction;
+    reflection_ray.direction = normalize(reflection);
+    reflection_ray.length = INFINITY;
+    reflection_ray.primitive_index = 0;
+
+    color_t target_color;
+    trace_ray(scene, &target_color, &reflection_ray);
+    return target_color;
+}
+
+color_t shade_normal(__global const shader_normal_t *shader, const scene_t *scene, ray_t *ray)
+{
+    return (color_t)((ray->normal.x + 1.0f) / 2.0f, (ray->normal.y + 1.0f) / 2.0f, (ray->normal.z + 1.0f) / 2.0f);
+}
+
+color_t shade_phong(__global const shader_phong_t *shader, const scene_t *scene, ray_t *ray)
+{
+    color_t fragment_color;
+
+    float3 const reflection = ray->direction - 2 * dot(ray->normal, ray->direction) * ray->normal;
+
+    __global const light_t *base = scene->lights;
+    __global const light_ambient_t *ambient = NULL;
+    __global const light_point_t *point = NULL;
+    __global const light_spot_t *spot = NULL;
+
+    for(uint i = 0; i < scene->n_lights; i++) {
+        illumination_t illum;
+        switch(base->light_type) {
+        case LIGHT_TYPE_AMBIENT:
+            ambient = (__global const light_ambient_t *) base;
+            illuminate_ambient(&illum, ambient, scene, ray);
+            base = (__global const light_t *) (ambient + 1);
+            break;
+        case LIGHT_TYPE_POINT:
+            point = (__global const light_point_t *) base;
+            illuminate_point(&illum, point, scene, ray);
+            base = (__global const light_t *) (point + 1);
+            break;
+        case LIGHT_TYPE_SPOT:
+            spot = (__global const light_spot_t *) base;
+            illuminate_spot(&illum, spot, scene, ray);
+            base = (__global const light_t *) (spot + 1);
+            break;
+        }
+
+        color_t const diffuse = shader->diffuse_coefficient * shader->diffuse_color * max(dot(illum.direction * -1.0f, ray->normal), 0.0f);
+        fragment_color += diffuse * illum.color;
+
+        float const cosine = dot(illum.direction * -1.0f, reflection);
+        if(cosine > 0.0f) {
+            color_t const specular = shader->specular_coefficient * shader->specular_color * pow(cosine, shader->shininess_exponent);
+            fragment_color += specular * illum.color;
+        }
+    }
+
+    return fragment_color;
+}
+
+color_t shade_refraction(__global const shader_refraction_t *shader, const scene_t *scene, ray_t *ray)
+{
+    if (ray->bounces > 0) {
+        float3 normal_vector = ray->normal;
+
+        float refractiveIndex = shader->index_outside / shader->index_inside;
+        
+        if (dot(normal_vector, ray->direction) > 0) {
+            normal_vector = -normal_vector;
+            refractiveIndex = shader->index_inside / shader->index_outside;
+        }
+
+        float cosineTheta = dot(normal_vector, -ray->direction);
+        float cosinePhi = sqrt(1 + refractiveIndex * refractiveIndex * (cosineTheta * cosineTheta - 1));
+        
+        float3 t = refractiveIndex * ray->direction + (refractiveIndex * cosineTheta - cosinePhi) * normal_vector;
+
+        ray_t refractionRay = *ray;
+        
+        refractionRay.length = INFINITY;
+        refractionRay.primitive_index = 0;
+
+        if (dot(t, normal_vector) <= 0.0f) {
+            refractionRay.origin = ray->origin + (ray->length + REFR_EPS) * ray->direction;
+            refractionRay.direction = normalize(t);
+        } else {
+            refractionRay.origin = ray->origin + (ray->length - REFR_EPS) * ray->direction;
+            
+            float3 const reflectionVector = ray->direction - 2.0f * dot(normal_vector, ray->direction) * normal_vector;
+
+            refractionRay.direction = normalize(reflectionVector);
+        }
+
+        color_t target_color;
+        trace_ray(scene, &target_color, &refractionRay);
+        return target_color;
+    }
+
+    return (color_t)(0.0f, 0.0f, 0.0f);
+}
+
+void trace_ray(const scene_t *scene, float3 *target_color, ray_t *target_ray)
+{
+    bool found = false;
+
+    __global const primitive_t *base = scene->primitives;
+    __global const primitive_box_t *box = NULL;
+    __global const primitive_infiniteplane_t *plane = NULL;
+    __global const primitive_sphere_t *sphere = NULL;
+    __global const primitive_triangle_t *triangle = NULL;
+
+    for(uint i = 0; i < scene->n_primitives; i++) {
+        uint pp = 0;
+        switch(base->primitive_type) {
+        case PRIMITIVE_TYPE_BOX:
+            box = (__global const primitive_box_t *) base;
+            if(intersect_box(box, target_ray)) {
+                found = true;
+                target_ray->primitive_index = i;
+            }
+            base = (__global const primitive_t *) (box + 1);
+            break;
+        case PRIMITIVE_TYPE_INFINITEPLANE:
+            plane = (__global const primitive_infiniteplane_t *) base;
+            if(intersect_infiniteplane((__global const primitive_infiniteplane_t *) base, target_ray)) {
+                found = true;
+                target_ray->primitive_index = i;
+            }
+            base = (__global const primitive_t *) (plane + 1);
+            break;
+        case PRIMITIVE_TYPE_SPHERE:
+            sphere = (__global const primitive_sphere_t *) base;
+            if(intersect_sphere((__global const primitive_sphere_t *) base, target_ray)) {
+                found = true;
+                target_ray->primitive_index = i;
+            }
+            base = (__global const primitive_t *) (sphere + 1);
+            break;
+        case PRIMITIVE_TYPE_TRIANGLE:
+            triangle = (__global const primitive_triangle_t *) base;
+            if(intersect_triangle((__global const primitive_triangle_t *) base, target_ray)) {
+                found = true;
+                target_ray->primitive_index = i;
+            }
+            base = (__global const primitive_t *) (triangle + 1);
+            break;
+        }
+    }
+
+    if(found && ((target_ray->bounces--) > 0)) {
+        __global const shader_t *shader = get_shader(get_primitive(target_ray->primitive_index)->shader_index);
+        switch(shader->shader_type) {
+        case SHADER_TYPE_COOKTORRANCE:
+            *target_color = shade_cooktorrance((__global const shader_cooktorrance_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_FLAT:
+            *target_color = shade_flat((__global const shader_flat_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_LAMBERT:
+            *target_color = shade_lambert((__global const shader_lambert_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_MIRROR:
+            *target_color = shade_mirror((__global const shader_mirror_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_NORMAL:
+            *target_color = shade_normal((__global const shader_normal_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_PHONG:
+            *target_color = shade_phong((__global const shader_phong_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_REFRACTION:
+            *target_color = shade_refraction((__global const shader_refraction_t *) shader, scene, target_ray);
+            break;
+        case SHADER_TYPE_SIMPLESHADOW:
+        default:
+            *target_color = shade_simpleshadow((__global const shader_simpleshadow_t *) shader, scene, target_ray);
+            break;
+        }
+    } else {
+        *target_color = scene->background_color;
+    }
+}
+
+void __kernel raytrace(__global float3 *buf_color, __global ray_t *buf_ray, const uint n_rays,
+    const uint offset_x, const uint offset_y, const uint width, const uint height,
+    __global const primitive_t *primitives, const uint n_primitives,
+    __global const light_t *lights, const uint n_lights,
+    __global const shader_t *shaders, const color_t background_color)
+{
+    uint target = ((offset_x + get_global_id(0)) * height) + offset_y + get_global_id(1);
+
+    if(target < n_rays) {
+        float3 target_color = buf_color[target];
+        ray_t target_ray = buf_ray[target];
+
+        const scene_t scene = {
+            .background_color = background_color,
+            .primitives = primitives,
+            .n_primitives = n_primitives,
+            .lights = lights,
+            .n_lights = n_lights,
+            .shaders = shaders
+        };
+
+        trace_ray(&scene, &target_color, &target_ray);
+
+        buf_color[target] = target_color;
+        buf_ray[target] = target_ray;
+    }
+}
\ No newline at end of file