path: root/kernel.cl

#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;
    }
}