indicatrix 0.9.0

A physically-based spectral gemstone renderer that turns GemCAD-style cutting instructions into rendered output.
Documentation
// Phase 1 GPU struct-layout self-test kernels -- driven by
// `renderer::gpu::layout_check`'s `run_facet_plane`/`run_camera_params`/`run_ray`/
// `run_hit_record`, NOT physics kernels.
//
// Same purpose and mechanism as `layout_echo.wgsl` (see that file's own doc comment for
// the bug class this is built to catch): echo every named field of a new Phase-1
// storage/uniform struct straight through to an independent output buffer, so
// `layout_check` can compare the two buffers' raw bytes (padding included) and prove
// the hand-written `#[repr(C)]` offsets in `renderer::buffers` actually agree with what
// WGSL computes -- rather than trusting the offset comments there on their own.
//
// Four independent structs, four independent bind-group slots and entry points, all in
// one module (each entry point only touches its own pair of globals, so this needs no
// per-struct WGSL file).

struct FacetPlane {
    normal: vec3<f32>,
    d: f32,
}

struct GpuCameraParams {
    origin: vec3<f32>,
    fov_tan: f32,
    forward: vec3<f32>,
    width: f32,
    right: vec3<f32>,
    height: f32,
    up: vec3<f32>,
    num_samples: u32,
}

struct GpuRay {
    origin: vec3<f32>,
    dir: vec3<f32>,
}

struct GpuHitRecord {
    t: f32,
    facet_idx: i32,
    hit: u32,
    normal: vec3<f32>,
}

@group(0) @binding(0) var<storage, read> in_plane: FacetPlane;
@group(0) @binding(1) var<storage, read_write> out_plane: FacetPlane;

@group(0) @binding(2) var<storage, read> in_camera: GpuCameraParams;
@group(0) @binding(3) var<storage, read_write> out_camera: GpuCameraParams;

@group(0) @binding(4) var<storage, read> in_ray: GpuRay;
@group(0) @binding(5) var<storage, read_write> out_ray: GpuRay;

@group(0) @binding(6) var<storage, read> in_hit: GpuHitRecord;
@group(0) @binding(7) var<storage, read_write> out_hit: GpuHitRecord;

@compute @workgroup_size(1)
fn echo_plane() {
    out_plane.normal = in_plane.normal;
    out_plane.d = in_plane.d;
}

@compute @workgroup_size(1)
fn echo_camera() {
    out_camera.origin = in_camera.origin;
    out_camera.fov_tan = in_camera.fov_tan;
    out_camera.forward = in_camera.forward;
    out_camera.width = in_camera.width;
    out_camera.right = in_camera.right;
    out_camera.height = in_camera.height;
    out_camera.up = in_camera.up;
    out_camera.num_samples = in_camera.num_samples;
}

@compute @workgroup_size(1)
fn echo_ray() {
    out_ray.origin = in_ray.origin;
    out_ray.dir = in_ray.dir;
}

@compute @workgroup_size(1)
fn echo_hit() {
    out_hit.t = in_hit.t;
    out_hit.facet_idx = in_hit.facet_idx;
    out_hit.hit = in_hit.hit;
    out_hit.normal = in_hit.normal;
}

// Frosted girdle finish GPU port: renderer::buffers::facet_finish's
// `array<u32>` upload -- a runtime-sized array, unlike the four single-instance structs
// above, so it gets its own workgroup-parallel entry point (one thread per element)
// rather than reusing the `@workgroup_size(1)` single-instance echo pattern.
@group(0) @binding(8) var<storage, read> in_facet_finishes: array<u32>;
@group(0) @binding(9) var<storage, read_write> out_facet_finishes: array<u32>;

@compute @workgroup_size(64)
fn echo_facet_finishes(@builtin(global_invocation_id) gid: vec3<u32>) {
    let idx = gid.x;
    if (idx >= arrayLength(&in_facet_finishes)) {
        return;
    }
    out_facet_finishes[idx] = in_facet_finishes[idx];
}