From 3c96f553e06bf18526bb3b219eb0f0c71e1e302a Mon Sep 17 00:00:00 2001 From: Leonard Kugis Date: Thu, 2 Feb 2023 04:48:33 +0100 Subject: Initial commit --- .gitignore | 131 +++++++++ kernel.cl | 886 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 1017 insertions(+) create mode 100644 .gitignore create mode 100644 kernel.cl diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..0b4ca40 --- /dev/null +++ b/.gitignore @@ -0,0 +1,131 @@ +# Created by https://www.toptal.com/developers/gitignore/api/visualstudiocode,c++,windows,linux,macos +# Edit at https://www.toptal.com/developers/gitignore?templates=visualstudiocode,c++,windows,linux,macos + +### C++ ### +# Prerequisites +*.d + +# Compiled Object files +*.slo +*.lo +*.o +*.obj + +# Precompiled Headers +*.gch +*.pch + +# Compiled Dynamic libraries +*.so +*.dylib +*.dll + +# Fortran module files +*.mod +*.smod + +# Compiled Static libraries +*.lai +*.la +*.a +*.lib + +# Executables +*.exe +*.out +*.app + +### Linux ### +*~ + +# temporary files which can be created if a process still has a handle open of a deleted file +.fuse_hidden* + +# KDE directory preferences +.directory + +# Linux trash folder which might appear on any partition or disk +.Trash-* + +# .nfs files are created when an open file is removed but is still being accessed +.nfs* + +### macOS ### +# General +.DS_Store +.AppleDouble +.LSOverride + +# Icon must end with two \r +Icon + + +# Thumbnails +._* + +# Files that might appear in the root of a volume +.DocumentRevisions-V100 +.fseventsd +.Spotlight-V100 +.TemporaryItems +.Trashes +.VolumeIcon.icns +.com.apple.timemachine.donotpresent + +# Directories potentially created on remote AFP share +.AppleDB +.AppleDesktop +Network Trash Folder +Temporary Items +.apdisk + +### macOS Patch ### +# iCloud generated files +*.icloud + +### VisualStudioCode ### +.vscode/* +!.vscode/settings.json +!.vscode/tasks.json +!.vscode/launch.json +!.vscode/extensions.json +!.vscode/*.code-snippets + +# Local History for Visual Studio Code +.history/ + +# Built Visual Studio Code Extensions +*.vsix + +### VisualStudioCode Patch ### +# Ignore all local history of files +.history +.ionide + +### Windows ### +# Windows thumbnail cache files +Thumbs.db +Thumbs.db:encryptable +ehthumbs.db +ehthumbs_vista.db + +# Dump file +*.stackdump + +# Folder config file +[Dd]esktop.ini + +# Recycle Bin used on file shares +$RECYCLE.BIN/ + +# Windows Installer files +*.cab +*.msi +*.msix +*.msm +*.msp + +# Windows shortcuts +*.lnk + +# End of https://www.toptal.com/developers/gitignore/api/visualstudiocode,c++,windows,linux,macos 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 -- cgit v1.2.1