summaryrefslogtreecommitdiff
path: root/kernel.cl
diff options
context:
space:
mode:
Diffstat (limited to 'kernel.cl')
-rw-r--r--kernel.cl886
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
generated by cgit v1.2.1 (git 2.18.0) at 2024-12-23 08:45:10 +0000