webgpu triangle

examples/VulkanTriangle/closesthit.wgsl

@@ -0,0 +1,12 @@
+// WebGPU port of closesthit.glsl. Library concatenates this BEFORE the
+// library helpers, so `Payload` declared here is visible to traceRay,
+// runClosestHit, the mega-switch, and the user's raygen source.
+
+struct Payload {
+    color: vec3<f32>,
+};
+
+fn closesthit_main(ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) {
+    let bary = vec3<f32>(1.0 - hit.attribs.x - hit.attribs.y, hit.attribs.x, hit.attribs.y);
+    (*payload).color = bary;
+}

examples/VulkanTriangle/main.cpp

@@ -1,4 +1,6 @@
+#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 #include "vulkan/vulkan.h"
+#endif
 #include <cassert>
 
 import Crafter.Graphics;
@@ -7,7 +9,7 @@ import std;
 import Crafter.Event;
 import Crafter.Math;
 
-
+#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 int main() {
     Device::Initialize();
     Window window(1280, 720, "HelloVulkan");
@@ -89,7 +91,7 @@ int main() {
     RenderingElement3D::elements.emplace_back(&renderer);
 
     MatrixRowMajor<float, 4, 3, 1> transform = MatrixRowMajor<float, 4, 3, 1>::Identity();
-    std::memcpy(renderer.instance.transform.matrix, transform.m, sizeof(transform.m));
+    transform.Store(reinterpret_cast<float*>(renderer.instance.transform.matrix));
     RenderingElement3D::BuildTLAS(cmd, 0);
     RenderingElement3D::BuildTLAS(cmd, 1);
     RenderingElement3D::BuildTLAS(cmd, 2);
@@ -202,3 +204,65 @@ int main() {
     window.Render();
     window.StartSync();
 }
+#else
+// DOM-mode port. Same scene (one triangle), software-emulated raytracing
+// via compute. Shaders are read from .wgsl files shipped as static
+// assets (see project.cpp). Renders barycentric colors via the
+// hit/miss/raygen mega-switch in PipelineRTWebGPU.
+int main() {
+    Device::Initialize();
+    static Window window(1280, 720, "HelloVulkan");
+    auto cmd = window.StartInit();
+
+    DescriptorHeapWebGPU heap;
+    heap.Initialize(/*images*/ 4, /*buffers*/ 4, /*samplers*/ 2);
+
+    std::array<WebGPUShader, 3> shaders {{
+        WebGPUShader(std::filesystem::path("raygen.wgsl"),     "raygen_main",     WebGPURTStage::Raygen),
+        WebGPUShader(std::filesystem::path("miss.wgsl"),       "miss_main",       WebGPURTStage::Miss),
+        WebGPUShader(std::filesystem::path("closesthit.wgsl"), "closesthit_main", WebGPURTStage::ClosestHit),
+    }};
+    ShaderBindingTableWebGPU sbt;
+    sbt.Init(shaders);
+
+    std::array<RTShaderGroup, 1> raygenGroups {{
+        { .type = RTShaderGroupType::General,           .generalShader = 0 },
+    }};
+    std::array<RTShaderGroup, 1> missGroups {{
+        { .type = RTShaderGroupType::General,           .generalShader = 1 },
+    }};
+    std::array<RTShaderGroup, 1> hitGroups {{
+        { .type = RTShaderGroupType::TrianglesHitGroup, .closestHitShader = 2 },
+    }};
+
+    PipelineRTWebGPU pipeline;
+    pipeline.Init(cmd, raygenGroups, missGroups, hitGroups, sbt);
+
+    Mesh triangleMesh;
+    std::array<Vector<float, 3, 3>, 3> verts {{{-150, -150, 100}, {0, 150, 100}, {150, -150, 100}}};
+    std::array<std::uint32_t, 3> index {{2, 1, 0}};
+    triangleMesh.Build(verts, index, cmd);
+
+    static RenderingElement3D renderer;
+    renderer.instance.transform.matrix[0][0] = 1; renderer.instance.transform.matrix[0][1] = 0; renderer.instance.transform.matrix[0][2] = 0; renderer.instance.transform.matrix[0][3] = 0;
+    renderer.instance.transform.matrix[1][0] = 0; renderer.instance.transform.matrix[1][1] = 1; renderer.instance.transform.matrix[1][2] = 0; renderer.instance.transform.matrix[1][3] = 0;
+    renderer.instance.transform.matrix[2][0] = 0; renderer.instance.transform.matrix[2][1] = 0; renderer.instance.transform.matrix[2][2] = 1; renderer.instance.transform.matrix[2][3] = 0;
+    renderer.instance.instanceCustomIndex                    = 0;
+    renderer.instance.mask                                   = 0xFF;
+    renderer.instance.instanceShaderBindingTableRecordOffset = 0;
+    renderer.instance.flags                                  = kRTGeometryInstanceForceOpaque;
+    renderer.instance.accelerationStructureReference         = triangleMesh.blasAddr;
+    RenderingElement3D::Add(&renderer);
+    RenderingElement3D::BuildTLAS(cmd, 0);
+
+    window.descriptorHeap = &heap;
+    window.FinishInit();
+
+    RTPass rtPass(&pipeline);
+    window.passes.push_back(&rtPass);
+
+    window.Render();
+    window.StartUpdate();
+    window.StartSync();
+}
+#endif

examples/VulkanTriangle/miss.wgsl

@@ -0,0 +1,5 @@
+// WebGPU port of miss.glsl.
+
+fn miss_main(ray: RayDesc, payload: ptr<function, Payload>) {
+    (*payload).color = vec3<f32>(1.0, 1.0, 1.0);
+}

examples/VulkanTriangle/project.cpp

@@ -4,6 +4,14 @@ namespace fs = std::filesystem;
 using namespace Crafter;
 
 extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> args) {
+    bool isWasm = false;
+    for (std::string_view a : args) {
+        if (a.starts_with("--target=") && a.find("wasm") != std::string_view::npos) {
+            isWasm = true;
+            break;
+        }
+    }
+
     std::vector<std::string> graphicsArgs(args.begin(), args.end());
     Configuration* graphics = LocalProject({
         .projectFile = "../../project.cpp",
@@ -14,6 +22,12 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
     cfg.path = "./";
     cfg.name = "VulkanTriangle";
     cfg.outputName = "VulkanTriangle";
+    cfg.type = ConfigurationType::Executable;
+    if (isWasm) {
+        cfg.target = "wasm32-wasip1";
+        cfg.defines.push_back({"CRAFTER_GRAPHICS_WINDOW_DOM", ""});
+        cfg.compileFlags.push_back("-msimd128");
+    }
     ApplyStandardArgs(cfg, args);
     cfg.dependencies = { graphics };
 
@@ -21,8 +35,15 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
     std::array<fs::path, 1> impls = { "main" };
     cfg.GetInterfacesAndImplementations(ifaces, impls);
 
-    cfg.shaders.emplace_back(fs::path("raygen.glsl"),  std::string("main"), ShaderType::RayGen);
+    if (isWasm) {
-    cfg.shaders.emplace_back(fs::path("closesthit.glsl"), std::string("main"), ShaderType::ClosestHit);
+        cfg.files.emplace_back(fs::path("raygen.wgsl"));
-    cfg.shaders.emplace_back(fs::path("miss.glsl"), std::string("main"), ShaderType::Miss);
+        cfg.files.emplace_back(fs::path("closesthit.wgsl"));
+        cfg.files.emplace_back(fs::path("miss.wgsl"));
+        EnableWasiBrowserRuntime(cfg);
+    } else {
+        cfg.shaders.emplace_back(fs::path("raygen.glsl"),     std::string("main"), ShaderType::RayGen);
+        cfg.shaders.emplace_back(fs::path("closesthit.glsl"), std::string("main"), ShaderType::ClosestHit);
+        cfg.shaders.emplace_back(fs::path("miss.glsl"),       std::string("main"), ShaderType::Miss);
+    }
     return cfg;
 }

examples/VulkanTriangle/raygen.glsl

@@ -34,17 +34,17 @@ void main() {
         1.0
     ));
 
-    // traceRayEXT(
+    traceRayEXT(
-    //     topLevelAS[bufferStart],
+        topLevelAS[bufferStart],
-    //     gl_RayFlagsNoneEXT,
+        gl_RayFlagsNoneEXT,
-    //     0xff,
+        0xff,
-    //     0, 0, 0,
+        0, 0, 0,
-    //     origin,
+        origin,
-    //     0.001,
+        0.001,
-    //     direction,
+        direction,
-    //     10000.0,
+        10000.0,
-    //     0
+        0
-    // );
+    );
 
     imageStore(image[0], ivec2(pixel), vec4(hitValue, 1));
 }

examples/VulkanTriangle/raygen.wgsl

@@ -0,0 +1,39 @@
+// WebGPU port of raygen.glsl. Mirrors the pinhole camera setup — the
+// Payload type is declared in closesthit.wgsl (concatenated earlier).
+
+fn raygen_main(gid: vec3<u32>) {
+    if (gid.x >= hdr.surfaceW || gid.y >= hdr.surfaceH) { return; }
+
+    let pixel = vec2<f32>(f32(gid.x), f32(gid.y));
+    let resolution = vec2<f32>(f32(hdr.surfaceW), f32(hdr.surfaceH));
+    let uv  = (pixel + vec2<f32>(0.5)) / resolution;
+    let ndc = uv * 2.0 - vec2<f32>(1.0);
+
+    let origin = vec3<f32>(0.0, 0.0, -300.0);
+    let aspect = resolution.x / resolution.y;
+    let fov = 60.0 * 3.14159265 / 180.0;
+    let tanHalfFov = tan(fov * 0.5);
+
+    let direction = normalize(vec3<f32>(
+        ndc.x * aspect * tanHalfFov,
+        -ndc.y * tanHalfFov,
+        1.0,
+    ));
+
+    var payload: Payload;
+    payload.color = vec3<f32>(0.0);
+
+    traceRay(
+        0u,            // tlasIdx (unused)
+        0u,            // ray flags
+        0xFFu,         // cull mask
+        0u, 0u, 0u,    // sbtRecordOffset, sbtRecordStride, missIndex
+        origin, 0.001,
+        direction, 10000.0,
+        &payload,
+    );
+
+    textureStore(outImage,
+                 vec2<i32>(i32(gid.x), i32(gid.y)),
+                 vec4<f32>(payload.color, 1.0));
+}

implementations/Crafter.Graphics-Mesh-WebGPU.cpp

@@ -0,0 +1,240 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+*/
+
+// DOM-mode Mesh implementation: SAH BVH2 built on the host, then
+// forwarded to the JS bridge which appends the four data streams
+// (vertices, indices, BVH nodes, primRemap) into the global RT mesh
+// heaps. The handle returned by wgpuRegisterMeshBLAS goes into
+// RTInstance::accelerationStructureReference and lets the TLAS-build
+// compute pass and the traversal kernel find the BLAS data later.
+//
+// BVH layout must stay binary-identical to the WGSL `BVHNode` struct
+// declared in additional/dom-webgpu.js (rtWgslPrelude).
+
+module;
+module Crafter.Graphics:Mesh_implWebGPU;
+
+import :Mesh;
+import :WebGPU;
+import Crafter.Math;
+import std;
+
+using namespace Crafter;
+
+namespace {
+    // ─── BVH builder (binned SAH, 8 bins, BVH2) ────────────────────────
+
+    constexpr std::uint32_t kBinCount      = 8;
+    constexpr std::uint32_t kMaxLeafSize   = 4;
+    constexpr float         kTraversalCost = 1.0f;
+    constexpr float         kIntersectCost = 1.0f;
+
+    struct AABB {
+        float lo[3] { std::numeric_limits<float>::infinity(),
+                      std::numeric_limits<float>::infinity(),
+                      std::numeric_limits<float>::infinity() };
+        float hi[3] {-std::numeric_limits<float>::infinity(),
+                     -std::numeric_limits<float>::infinity(),
+                     -std::numeric_limits<float>::infinity() };
+
+        void Extend(const float p[3]) noexcept {
+            for (int a = 0; a < 3; ++a) {
+                if (p[a] < lo[a]) lo[a] = p[a];
+                if (p[a] > hi[a]) hi[a] = p[a];
+            }
+        }
+        void Extend(const AABB& o) noexcept {
+            for (int a = 0; a < 3; ++a) {
+                if (o.lo[a] < lo[a]) lo[a] = o.lo[a];
+                if (o.hi[a] > hi[a]) hi[a] = o.hi[a];
+            }
+        }
+        float SurfaceArea() const noexcept {
+            float dx = hi[0] - lo[0];
+            float dy = hi[1] - lo[1];
+            float dz = hi[2] - lo[2];
+            if (dx < 0.0f || dy < 0.0f || dz < 0.0f) return 0.0f;
+            return 2.0f * (dx*dy + dx*dz + dy*dz);
+        }
+    };
+
+    struct PrimRef {
+        AABB box;
+        float centroid[3];
+        std::uint32_t triIndex;
+    };
+
+    struct Bin {
+        AABB box;
+        std::uint32_t count = 0;
+    };
+
+    struct Builder {
+        std::vector<PrimRef> prims;
+        std::vector<BVHNode> nodes;
+
+        std::pair<std::uint32_t, std::uint32_t> AllocateChildren() {
+            std::uint32_t l = static_cast<std::uint32_t>(nodes.size());
+            nodes.emplace_back();
+            nodes.emplace_back();
+            return { l, l + 1 };
+        }
+
+        void BuildRecursive(std::uint32_t nodeIdx,
+                            std::uint32_t first,
+                            std::uint32_t count) {
+            AABB bounds, centroidBounds;
+            for (std::uint32_t i = 0; i < count; ++i) {
+                const auto& p = prims[first + i];
+                bounds.Extend(p.box);
+                centroidBounds.Extend(p.centroid);
+            }
+
+            auto emitLeaf = [&] {
+                BVHNode& n = nodes[nodeIdx];
+                std::memcpy(n.aabbMin, bounds.lo, sizeof(bounds.lo));
+                std::memcpy(n.aabbMax, bounds.hi, sizeof(bounds.hi));
+                n.firstChildOrPrim = first;
+                n.primCount        = count;
+            };
+
+            if (count <= kMaxLeafSize) { emitLeaf(); return; }
+
+            int   bestAxis  = -1;
+            float bestCost  = std::numeric_limits<float>::infinity();
+            std::uint32_t bestBin = 0;
+
+            float parentArea = bounds.SurfaceArea();
+            if (parentArea <= 0.0f) { emitLeaf(); return; }
+
+            for (int axis = 0; axis < 3; ++axis) {
+                float extent = centroidBounds.hi[axis] - centroidBounds.lo[axis];
+                if (extent <= 0.0f) continue;
+                float invExtent = static_cast<float>(kBinCount) / extent;
+
+                std::array<Bin, kBinCount> bins{};
+                for (std::uint32_t i = 0; i < count; ++i) {
+                    const auto& p = prims[first + i];
+                    float t = (p.centroid[axis] - centroidBounds.lo[axis]) * invExtent;
+                    std::uint32_t b = static_cast<std::uint32_t>(t);
+                    if (b >= kBinCount) b = kBinCount - 1;
+                    bins[b].box.Extend(p.box);
+                    bins[b].count += 1;
+                }
+
+                std::array<AABB,         kBinCount - 1> leftBox;
+                std::array<std::uint32_t,kBinCount - 1> leftCount{};
+                {
+                    AABB acc; std::uint32_t cnt = 0;
+                    for (std::uint32_t i = 0; i < kBinCount - 1; ++i) {
+                        acc.Extend(bins[i].box);
+                        cnt += bins[i].count;
+                        leftBox[i]   = acc;
+                        leftCount[i] = cnt;
+                    }
+                }
+                {
+                    AABB acc; std::uint32_t cnt = 0;
+                    for (std::int32_t i = kBinCount - 1; i >= 1; --i) {
+                        acc.Extend(bins[i].box);
+                        cnt += bins[i].count;
+                        std::uint32_t split = static_cast<std::uint32_t>(i - 1);
+                        if (leftCount[split] == 0 || cnt == 0) continue;
+                        float cost = kTraversalCost
+                                   + (leftBox[split].SurfaceArea() * leftCount[split]
+                                      + acc.SurfaceArea() * cnt) * kIntersectCost / parentArea;
+                        if (cost < bestCost) {
+                            bestCost = cost;
+                            bestAxis = axis;
+                            bestBin  = split;
+                        }
+                    }
+                }
+            }
+
+            float leafCost = static_cast<float>(count) * kIntersectCost;
+            if (bestAxis < 0 || bestCost >= leafCost) { emitLeaf(); return; }
+
+            float invExtent = static_cast<float>(kBinCount)
+                            / (centroidBounds.hi[bestAxis] - centroidBounds.lo[bestAxis]);
+            float lo = centroidBounds.lo[bestAxis];
+            auto mid = std::partition(
+                prims.begin() + first, prims.begin() + first + count,
+                [&](const PrimRef& p) {
+                    float t = (p.centroid[bestAxis] - lo) * invExtent;
+                    std::uint32_t b = static_cast<std::uint32_t>(t);
+                    if (b >= kBinCount) b = kBinCount - 1;
+                    return b <= bestBin;
+                });
+            std::uint32_t leftCount =
+                static_cast<std::uint32_t>(mid - (prims.begin() + first));
+            if (leftCount == 0 || leftCount == count) { emitLeaf(); return; }
+
+            auto [leftIdx, rightIdx] = AllocateChildren();
+            {
+                BVHNode& n = nodes[nodeIdx];
+                std::memcpy(n.aabbMin, bounds.lo, sizeof(bounds.lo));
+                std::memcpy(n.aabbMax, bounds.hi, sizeof(bounds.hi));
+                n.firstChildOrPrim = leftIdx;
+                n.primCount        = 0;
+            }
+            BuildRecursive(leftIdx,  first,             leftCount);
+            BuildRecursive(rightIdx, first + leftCount, count - leftCount);
+        }
+
+        void Build(std::span<const Vector<float, 3, 3>> vertices,
+                   std::span<const std::uint32_t>       indices) {
+            std::uint32_t triCount = static_cast<std::uint32_t>(indices.size()) / 3;
+            prims.resize(triCount);
+            for (std::uint32_t i = 0; i < triCount; ++i) {
+                std::uint32_t i0 = indices[i*3 + 0];
+                std::uint32_t i1 = indices[i*3 + 1];
+                std::uint32_t i2 = indices[i*3 + 2];
+                const auto& v0 = vertices[i0];
+                const auto& v1 = vertices[i1];
+                const auto& v2 = vertices[i2];
+                float p0[3] { v0.v[0], v0.v[1], v0.v[2] };
+                float p1[3] { v1.v[0], v1.v[1], v1.v[2] };
+                float p2[3] { v2.v[0], v2.v[1], v2.v[2] };
+                auto& pr = prims[i];
+                pr.box.Extend(p0);
+                pr.box.Extend(p1);
+                pr.box.Extend(p2);
+                pr.centroid[0] = (pr.box.lo[0] + pr.box.hi[0]) * 0.5f;
+                pr.centroid[1] = (pr.box.lo[1] + pr.box.hi[1]) * 0.5f;
+                pr.centroid[2] = (pr.box.lo[2] + pr.box.hi[2]) * 0.5f;
+                pr.triIndex = i;
+            }
+            nodes.reserve(triCount * 2);
+            nodes.emplace_back();
+            BuildRecursive(0, 0, triCount);
+        }
+    };
+}
+
+void Mesh::Build(std::span<Vector<float, 3, 3>> vertices,
+                 std::span<std::uint32_t>       indices,
+                 WebGPUCommandEncoderRef        /*cmd*/) {
+    triangleCount = static_cast<std::uint32_t>(indices.size()) / 3;
+
+    Builder builder;
+    builder.Build(vertices, indices);
+
+    std::vector<std::uint32_t> primRemap(triangleCount);
+    for (std::uint32_t i = 0; i < triangleCount; ++i) {
+        primRemap[i] = builder.prims[i].triIndex;
+    }
+
+    const BVHNode& root = builder.nodes[0];
+    std::uint32_t h = WebGPU::wgpuRegisterMeshBLAS(
+        root.aabbMin[0], root.aabbMin[1], root.aabbMin[2],
+        root.aabbMax[0], root.aabbMax[1], root.aabbMax[2],
+        vertices.data(),       static_cast<std::int32_t>(vertices.size()),
+        indices.data(),        static_cast<std::int32_t>(indices.size()),
+        builder.nodes.data(),  static_cast<std::int32_t>(builder.nodes.size()),
+        primRemap.data(),      static_cast<std::int32_t>(primRemap.size()));
+    blasAddr = h;
+}

implementations/Crafter.Graphics-PipelineRTWebGPU.cpp

@@ -0,0 +1,187 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+*/
+
+// Megakernel WGSL assembly. The library prelude lives JS-side
+// (additional/dom-webgpu.js, rtWgslPrelude) — we don't have access to
+// it from C++ — so this file emits only the *user-controlled* portions
+// (concatenated SBT sources + the generated switch statements) and the
+// stable entry-point glue. The JS side wraps these with the prelude
+// before handing to device.createShaderModule.
+//
+// Wire format passed across the JS boundary is a single WGSL string
+// containing the substitution markers `// @CRAFTER_RT_USER_SOURCES`,
+// `// @CRAFTER_RT_CLOSESTHIT_CASES`, `// @CRAFTER_RT_ANYHIT_CASES`,
+// `// @CRAFTER_RT_MISS_CASES`, `// @CRAFTER_RT_RAYGEN_BODY` already
+// expanded; the JS side just concatenates prelude + this string.
+
+module;
+module Crafter.Graphics:PipelineRTWebGPU_impl;
+
+import :PipelineRTWebGPU;
+import :ShaderBindingTableWebGPU;
+import :RT;
+import :WebGPU;
+import std;
+
+using namespace Crafter;
+
+namespace {
+    constexpr std::string_view kPlaceholderClosestHit =
+        "fn _crafter_default_closesthit(ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) {}";
+    constexpr std::string_view kPlaceholderAnyHit =
+        "fn _crafter_default_anyhit(ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) -> u32 { return RT_ANYHIT_ACCEPT; }";
+    constexpr std::string_view kPlaceholderMiss =
+        "fn _crafter_default_miss(ray: RayDesc, payload: ptr<function, Payload>) {}";
+
+    void AppendCase(std::string& out,
+                    std::uint32_t hitGroupIndex,
+                    std::string_view entryFn,
+                    std::string_view args) {
+        out += "        case ";
+        out += std::to_string(hitGroupIndex);
+        out += "u: { ";
+        out += entryFn;
+        out += "(";
+        out += args;
+        out += "); }\n";
+    }
+
+    // anyhit has a return type — case body forwards the result.
+    void AppendAnyHitCase(std::string& out,
+                          std::uint32_t hitGroupIndex,
+                          std::string_view entryFn) {
+        out += "        case ";
+        out += std::to_string(hitGroupIndex);
+        out += "u: { return ";
+        out += entryFn;
+        out += "(ray, hit, payload); }\n";
+    }
+}
+
+void PipelineRTWebGPU::Init(WebGPUCommandEncoderRef                 /*cmd*/,
+                            std::span<const RTShaderGroup>          raygenGroups,
+                            std::span<const RTShaderGroup>          missGroups,
+                            std::span<const RTShaderGroup>          hitGroups,
+                            const ShaderBindingTableWebGPU&         sbt) {
+    std::string wgsl;
+    wgsl.reserve(8 * 1024);
+
+    // ── Section 1: user closesthit / anyhit / miss source files ────────
+    //
+    // Raygens come later (after `traceRay` is declared) so we partition
+    // shaders by stage. Concatenating *all* non-raygen sources here lets
+    // them declare shared helpers, `struct Payload`, etc., in any order.
+
+    wgsl += "// ── user closesthit / anyhit / miss sources ───────────────\n";
+    for (const auto& shader : sbt.shaders) {
+        if (shader.stage == WebGPURTStage::Raygen) continue;
+        wgsl += shader.source;
+        wgsl += "\n";
+    }
+
+    // ── Section 2: mega-switch dispatchers ─────────────────────────────
+    //
+    // runClosestHit, runAnyHit, runMiss each dispatch on the per-hit /
+    // per-ray index registered against the appropriate group span.
+    // Indices match the user's expectations from VkRayTracingShaderGroup
+    // ordering: closest-hit group N (N from 0..hitGroups.size()-1) is
+    // selected by hitGroupIndex == N.
+
+    wgsl += "\nfn runClosestHit(hg: u32, ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) {\n";
+    wgsl += "    switch hg {\n";
+    bool anyClosestHit = false;
+    for (std::uint32_t i = 0; i < hitGroups.size(); ++i) {
+        const auto& g = hitGroups[i];
+        if (g.closestHitShader == kRTShaderUnused) continue;
+        if (g.closestHitShader >= sbt.shaders.size()) continue;
+        const auto& fn = sbt.shaders[g.closestHitShader].entryFn;
+        AppendCase(wgsl, i, fn, "ray, hit, payload");
+        anyClosestHit = true;
+    }
+    if (!anyClosestHit) wgsl += "        // (no closest-hit shaders registered)\n";
+    wgsl += "        default: { }\n";
+    wgsl += "    }\n";
+    wgsl += "}\n\n";
+
+    wgsl += "fn runAnyHit(hg: u32, ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) -> u32 {\n";
+    wgsl += "    switch hg {\n";
+    bool anyAnyhit = false;
+    for (std::uint32_t i = 0; i < hitGroups.size(); ++i) {
+        const auto& g = hitGroups[i];
+        if (g.anyHitShader == kRTShaderUnused) continue;
+        if (g.anyHitShader >= sbt.shaders.size()) continue;
+        const auto& fn = sbt.shaders[g.anyHitShader].entryFn;
+        AppendAnyHitCase(wgsl, i, fn);
+        anyAnyhit = true;
+    }
+    if (!anyAnyhit) wgsl += "        // (no any-hit shaders registered)\n";
+    wgsl += "        default: { return RT_ANYHIT_ACCEPT; }\n";
+    wgsl += "    }\n";
+    wgsl += "}\n\n";
+
+    wgsl += "fn runMiss(idx: u32, ray: RayDesc, payload: ptr<function, Payload>) {\n";
+    wgsl += "    switch idx {\n";
+    bool anyMiss = false;
+    for (std::uint32_t i = 0; i < missGroups.size(); ++i) {
+        const auto& g = missGroups[i];
+        if (g.generalShader == kRTShaderUnused) continue;
+        if (g.generalShader >= sbt.shaders.size()) continue;
+        const auto& fn = sbt.shaders[g.generalShader].entryFn;
+        AppendCase(wgsl, i, fn, "ray, payload");
+        anyMiss = true;
+    }
+    if (!anyMiss) wgsl += "        // (no miss shaders registered)\n";
+    wgsl += "        default: { }\n";
+    wgsl += "    }\n";
+    wgsl += "}\n";
+
+    // Marker — JS-side prelude/post-amble searches for this token to know
+    // where the library helpers (traverseBlas/traverseTlas/traceRay) get
+    // injected, followed by raygen sources and the @compute entry point.
+    wgsl += "\n// @CRAFTER_RT_LIBRARY_HELPERS_HERE\n";
+
+    // ── Section 3: user raygen source files ────────────────────────────
+    //
+    // Comes after the library injects traceRay, so raygens can call it.
+
+    wgsl += "\n// ── user raygen sources ───────────────────────────────────\n";
+    std::uint32_t raygenEntryIndex = kRTShaderUnused;
+    std::string   raygenEntryFn;
+    for (const auto& shader : sbt.shaders) {
+        if (shader.stage != WebGPURTStage::Raygen) continue;
+        wgsl += shader.source;
+        wgsl += "\n";
+        // Pick the first raygen group's general shader as the entry. Mirrors
+        // Vulkan's pRayGenShaderBindingTable[0] → first invoked raygen.
+        if (raygenEntryFn.empty()) raygenEntryFn = shader.entryFn;
+    }
+    if (!raygenGroups.empty()
+        && raygenGroups[0].generalShader != kRTShaderUnused
+        && raygenGroups[0].generalShader < sbt.shaders.size()) {
+        raygenEntryIndex = raygenGroups[0].generalShader;
+        raygenEntryFn    = sbt.shaders[raygenEntryIndex].entryFn;
+    }
+    if (raygenEntryFn.empty()) {
+        std::println("PipelineRTWebGPU::Init: no raygen shader registered");
+        pipelineHandle = 0;
+        return;
+    }
+
+    // ── Section 4: @compute entry point ────────────────────────────────
+    //
+    // 8x8 tile workgroup matching the rest of the WebGPU backend.
+
+    wgsl += "\n@compute @workgroup_size(8, 8, 1)\n";
+    wgsl += "fn main(@builtin(global_invocation_id) gid: vec3<u32>) {\n";
+    wgsl += "    ";
+    wgsl += raygenEntryFn;
+    wgsl += "(gid);\n";
+    wgsl += "}\n";
+
+    pipelineHandle = WebGPU::wgpuLoadRTPipeline(
+        wgsl.data(),
+        static_cast<std::int32_t>(wgsl.size()));
+}

implementations/Crafter.Graphics-RenderingElement3D-WebGPU.cpp

@@ -0,0 +1,91 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+*/
+
+// DOM-mode TLAS upkeep. BuildTLAS copies the per-element RTInstance into
+// the host-visible instance buffer (skipping the transform for elements
+// whose transform is GPU-owned), uploads it, then dispatches the JS-side
+// TLAS-build compute pass — which consults the per-BLAS records published
+// at Mesh::Build() time to produce world-space AABBs and inverse
+// transforms in the format `traceRay` / `rayQuery` consume.
+
+module;
+module Crafter.Graphics:RenderingElement3D_implWebGPU;
+
+import :RenderingElement3D;
+import :Mesh;
+import :WebGPU;
+import :WebGPUBuffer;
+import std;
+
+using namespace Crafter;
+
+std::vector<RenderingElement3D*> RenderingElement3D::elements;
+
+void RenderingElement3D::Add(RenderingElement3D* e) {
+    e->indexInElements = static_cast<std::uint32_t>(elements.size());
+    elements.push_back(e);
+}
+
+void RenderingElement3D::Remove(RenderingElement3D* e) {
+    std::uint32_t idx = e->indexInElements;
+    if (idx == std::numeric_limits<std::uint32_t>::max()) return;
+    std::uint32_t last = static_cast<std::uint32_t>(elements.size() - 1);
+    if (idx != last) {
+        elements[idx] = elements[last];
+        elements[idx]->indexInElements = idx;
+    }
+    elements.pop_back();
+    e->indexInElements = std::numeric_limits<std::uint32_t>::max();
+}
+
+void RenderingElement3D::BuildTLAS(WebGPUCommandEncoderRef /*cmd*/, std::uint32_t index) {
+    auto& tlas = tlases[index];
+    const std::uint32_t primitiveCount = static_cast<std::uint32_t>(elements.size());
+    if (primitiveCount == 0) {
+        tlas.builtInstanceCount = 0;
+        return;
+    }
+
+    // (Re)allocate instance + metadata + output TLAS buffers if the count
+    // changed. WebGPUBuffer::Resize destroys and recreates the GPU buffer;
+    // bind-group caches keyed on the buffer handle are invalidated in the
+    // JS bridge automatically.
+    if (primitiveCount != tlas.builtInstanceCount) {
+        tlas.instanceBuffer.Resize(primitiveCount);
+        tlas.metadataBuffer.Resize(primitiveCount);
+        // TLASEntry layout in WGSL is 144 bytes due to vec3 align/pad
+        // rules. Must match the struct declared in the rtWgslTypes
+        // block in additional/dom-webgpu.js.
+        tlas.buffer.Resize(primitiveCount * 144);
+    }
+
+    for (std::uint32_t i = 0; i < primitiveCount; ++i) {
+        auto& dst = tlas.instanceBuffer.value[i];
+        const auto& src = elements[i]->instance;
+        if (elements[i]->transformOwnedByGpu) {
+            // Preserve whatever the GPU compute shader most recently
+            // wrote into dst.transform. Update only the non-transform
+            // fields.
+            dst.instanceCustomIndex                    = src.instanceCustomIndex;
+            dst.mask                                   = src.mask;
+            dst.instanceShaderBindingTableRecordOffset = src.instanceShaderBindingTableRecordOffset;
+            dst.flags                                  = src.flags;
+            dst.accelerationStructureReference         = src.accelerationStructureReference;
+        } else {
+            dst = src;
+        }
+        tlas.metadataBuffer.value[i] = elements[i]->userMetadata;
+    }
+
+    tlas.instanceBuffer.FlushDevice();
+    tlas.metadataBuffer.FlushDevice();
+
+    WebGPU::wgpuBuildTLAS(tlas.instanceBuffer.handle,
+                          static_cast<std::int32_t>(primitiveCount),
+                          tlas.buffer.handle);
+
+    tlas.builtInstanceCount = primitiveCount;
+}

implementations/Crafter.Graphics-ShaderBindingTableWebGPU.cpp

@@ -0,0 +1,32 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+*/
+
+module;
+module Crafter.Graphics:ShaderBindingTableWebGPU_impl;
+
+import :ShaderBindingTableWebGPU;
+import std;
+
+using namespace Crafter;
+
+WebGPUShader::WebGPUShader(const std::filesystem::path& wgslPath,
+                           std::string fn,
+                           WebGPURTStage s)
+    : entryFn(std::move(fn)), stage(s) {
+    std::ifstream f(wgslPath, std::ios::binary | std::ios::ate);
+    if (!f.is_open()) {
+        std::println("WebGPUShader: cannot open {}", wgslPath.string());
+        return;
+    }
+    auto size = f.tellg();
+    if (size <= 0) {
+        std::println("WebGPUShader: empty file {}", wgslPath.string());
+        return;
+    }
+    f.seekg(0, std::ios::beg);
+    source.resize(static_cast<std::size_t>(size));
+    f.read(source.data(), size);
+}

implementations/Crafter.Graphics-WebGPUComputeShader.cpp

@@ -12,18 +12,22 @@ import std;
 using namespace Crafter;
 
 void WebGPUComputeShader::Load(std::string_view wgsl,
-                               std::span<const UICustomBinding> bindings) {
+                               std::span<const UICustomBinding> bindings,
+                               bool rayQuery) {
     customBindings.assign(bindings.begin(), bindings.end());
+    rayQueryCapable = rayQuery;
     pipelineHandle = WebGPU::wgpuLoadCustomShader(
         wgsl.data(),
         static_cast<std::int32_t>(wgsl.size()),
         customBindings.data(),
-        static_cast<std::int32_t>(customBindings.size())
+        static_cast<std::int32_t>(customBindings.size()),
+        rayQuery ? 1 : 0
     );
 }
 
 void WebGPUComputeShader::Load(const std::filesystem::path& wgslPath,
-                               std::span<const UICustomBinding> bindings) {
+                               std::span<const UICustomBinding> bindings,
+                               bool rayQuery) {
     std::ifstream f(wgslPath, std::ios::binary | std::ios::ate);
     if (!f.is_open()) {
         std::println("WebGPUComputeShader::Load: cannot open {}", wgslPath.string());
@@ -37,5 +41,5 @@ void WebGPUComputeShader::Load(const std::filesystem::path& wgslPath,
     f.seekg(0, std::ios::beg);
     std::string src(static_cast<std::size_t>(size), '\0');
     f.read(src.data(), size);
-    Load(std::string_view{src}, bindings);
+    Load(std::string_view{src}, bindings, rayQuery);
 }

interfaces/Crafter.Graphics-Mesh.cppm

@@ -60,3 +60,54 @@ export namespace Crafter {
     };
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
+
+#ifdef CRAFTER_GRAPHICS_WINDOW_DOM
+import std;
+import Crafter.Math;
+import :WebGPU;
+
+export namespace Crafter {
+    // Software-RT BLAS node, packed to 32 bytes. Matches the WGSL
+    // `BVHNode` struct in the RT WGSL prelude (additional/dom-webgpu.js,
+    // rtWgslPrelude) byte-for-byte.
+    //
+    //   primCount == 0          → inner node, children at indices
+    //                             firstChildOrPrim and firstChildOrPrim+1.
+    //   primCount  > 0          → leaf, `primCount` primitives starting at
+    //                             primIndex `firstChildOrPrim` in the
+    //                             global primRemap heap.
+    //
+    // SAH-built BVH2; constructed CPU-side at Build() time, never refit.
+    struct BVHNode {
+        float         aabbMin[3];
+        std::uint32_t firstChildOrPrim;
+        float         aabbMax[3];
+        std::uint32_t primCount;
+    };
+    static_assert(sizeof(BVHNode) == 32);
+
+    class Mesh {
+        public:
+        // BLAS "handle": opaque identity that goes into
+        // RTInstance::accelerationStructureReference. Set by Build() to a
+        // stable u32 (widened to u64 for Vulkan-struct layout parity), used
+        // by the WebGPU TLAS-build compute shader to look up the BLAS root
+        // AABB and per-mesh heap offsets. Handle 0 is the unassigned
+        // sentinel; never returned by Build().
+        std::uint64_t blasAddr = 0;
+        std::uint32_t triangleCount = 0;
+
+        bool opaque = true;
+
+        // Build BLAS from raw triangle data. Runs the CPU SAH BVH2 builder
+        // and forwards vertex/index/BVH/remap arrays to the JS-side mesh
+        // heap (additional/dom-webgpu.js), which queue.writeBuffers them
+        // into the global heaps and records the per-mesh offsets keyed by
+        // the returned handle. The `cmd` parameter is unused on WebGPU —
+        // kept for API symmetry with the Vulkan signature.
+        void Build(std::span<Crafter::Vector<float, 3, 3>> vertices,
+                   std::span<std::uint32_t>                indices,
+                   WebGPUCommandEncoderRef                  cmd = 0);
+    };
+}
+#endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics-PipelineRTWebGPU.cppm

@@ -0,0 +1,51 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+*/
+
+// DOM-mode RT pipeline. Mirrors PipelineRTVulkan's surface — Init takes
+// the same kind of (raygen, miss, hit) shader-group spans plus an SBT.
+// The big difference is implementation: there's no native RT pipeline on
+// WebGPU, so Init assembles a single megakernel WGSL by concatenating
+//   1. library prelude (types, bindings, ray-flag constants)
+//   2. user closesthit / anyhit / miss source files
+//   3. library mega-switches dispatched on per-hit hit-group index
+//   4. library helpers (rayAabb / rayTriangle / traverseBlas / traverseTlas)
+//   5. library traceRay function
+//   6. user raygen source files
+//   7. @compute entry calling the registered raygen
+// and hands the result to wgpuLoadRTPipeline.
+//
+// The library WGSL itself lives in additional/dom-webgpu.js (rtWgslPrelude
+// + rtWgslDispatchTemplate). C++ side only knows the substitution markers.
+
+export module Crafter.Graphics:PipelineRTWebGPU;
+#ifdef CRAFTER_GRAPHICS_WINDOW_DOM
+import std;
+import :RT;
+import :WebGPU;
+import :ShaderBindingTableWebGPU;
+
+export namespace Crafter {
+    class PipelineRTWebGPU {
+    public:
+        std::uint32_t pipelineHandle = 0;
+
+        // Build the megakernel pipeline. Groups carry indices into
+        // `sbt.shaders`. The library generates one `case` per registered
+        // group: closest-hit groups dispatch to their closestHitShader's
+        // entryFn, miss groups to their generalShader's entryFn, etc.
+        // The `cmd` parameter is unused on WebGPU; kept for API symmetry.
+        void Init(WebGPUCommandEncoderRef                 cmd,
+                  std::span<const RTShaderGroup>          raygenGroups,
+                  std::span<const RTShaderGroup>          missGroups,
+                  std::span<const RTShaderGroup>          hitGroups,
+                  const ShaderBindingTableWebGPU&         sbt);
+
+        PipelineRTWebGPU() = default;
+        PipelineRTWebGPU(const PipelineRTWebGPU&) = delete;
+        PipelineRTWebGPU& operator=(const PipelineRTWebGPU&) = delete;
+    };
+}
+#endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics-RT.cppm

@@ -0,0 +1,83 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+
+This library is free software; you can redistribute it and/or
+modify it under the terms of the GNU Lesser General Public
+License version 3.0 as published by the Free Software Foundation;
+*/
+
+// Portable RT types & constants.
+//
+// Native: aliases the Vulkan struct so existing code that passes
+// `RenderingElement3D::instance` directly into vkCmdBuildAccelerationStructuresKHR
+// is a no-op layout-wise.
+// DOM: provides a POD with the same byte layout + the same field names, so
+// user code touching `instance.mask`, `instance.flags`, `instance.transform`
+// etc. compiles unchanged.
+//
+// Flag constants are spelled out as Crafter::kRT* so portable user code can
+// avoid referencing VK_* on the DOM target. The values match
+// VkGeometryInstanceFlagBitsKHR / VkRayTracingShaderGroupTypeKHR so the
+// constants compare-equal on Vulkan if the user wants to mix surfaces.
+
+module;
+#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
+#include "vulkan/vulkan.h"
+#endif
+export module Crafter.Graphics:RT;
+import std;
+
+export namespace Crafter {
+#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
+    using RTTransformMatrix = VkTransformMatrixKHR;
+    using RTInstance        = VkAccelerationStructureInstanceKHR;
+#else
+    // Mirrors VkTransformMatrixKHR: row-major affine 3x4.
+    struct RTTransformMatrix {
+        float matrix[3][4];
+    };
+    static_assert(sizeof(RTTransformMatrix) == 48);
+
+    // Mirrors VkAccelerationStructureInstanceKHR byte-for-byte.
+    // On WebGPU the `accelerationStructureReference` slot holds the BLAS
+    // handle returned by MeshWebGPU::blasHandle (a u32 widened to u64).
+    struct RTInstance {
+        RTTransformMatrix transform;
+        std::uint32_t instanceCustomIndex                    : 24;
+        std::uint32_t mask                                   : 8;
+        std::uint32_t instanceShaderBindingTableRecordOffset : 24;
+        std::uint32_t flags                                  : 8;
+        std::uint64_t accelerationStructureReference;
+    };
+    static_assert(sizeof(RTInstance) == 64);
+#endif
+
+    // VkGeometryInstanceFlagBitsKHR mirror. Values verbatim so equal on both.
+    inline constexpr std::uint8_t kRTGeometryInstanceTriangleFacingCullDisable = 0x1;
+    inline constexpr std::uint8_t kRTGeometryInstanceTriangleFlipFacing        = 0x2;
+    inline constexpr std::uint8_t kRTGeometryInstanceForceOpaque               = 0x4;
+    inline constexpr std::uint8_t kRTGeometryInstanceForceNoOpaque             = 0x8;
+
+    // Hit-group identification. Matches VkRayTracingShaderGroupTypeKHR for
+    // the two types we actually support (general + triangles-hit).
+    enum class RTShaderGroupType : std::uint8_t {
+        General           = 0,   // raygen / miss / callable
+        TrianglesHitGroup = 1,
+    };
+
+    // Cross-backend description of one entry in the shader-group array
+    // passed to PipelineRT::Init. Mirrors the meaningful subset of
+    // VkRayTracingShaderGroupCreateInfoKHR: per group, the type and the
+    // indices (into the SBT's shader array) for general / closestHit /
+    // anyHit, with kRTShaderUnused == VK_SHADER_UNUSED_KHR for "none".
+    inline constexpr std::uint32_t kRTShaderUnused = 0xFFFFFFFFu;
+
+    struct RTShaderGroup {
+        RTShaderGroupType type             = RTShaderGroupType::General;
+        std::uint32_t      generalShader    = kRTShaderUnused;
+        std::uint32_t      closestHitShader = kRTShaderUnused;
+        std::uint32_t      anyHitShader     = kRTShaderUnused;
+    };
+}

interfaces/Crafter.Graphics-RTPass.cppm

@@ -46,3 +46,42 @@ export namespace Crafter {
     };
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
+
+#ifdef CRAFTER_GRAPHICS_WINDOW_DOM
+import std;
+import :RenderPass;
+import :Window;
+import :WebGPU;
+import :PipelineRTWebGPU;
+import :RenderingElement3D;
+
+export namespace Crafter {
+    // DOM-mode RT pass — dispatches the megakernel pipeline at frame Record
+    // time. Picks up the current TLAS for the frame and the application's
+    // raygen-side push data (typically empty in v1; pass via window.passes
+    // wiring if needed later).
+    struct RTPass : RenderPass {
+        PipelineRTWebGPU* pipeline;
+        // Optional per-dispatch push data forwarded after the standard
+        // RTDispatchHeader. Null means "no extra data".
+        const void*       pushPtr   = nullptr;
+        std::uint32_t     pushBytes = 0;
+
+        RTPass(PipelineRTWebGPU* p) : pipeline(p) {}
+
+        void Record(WebGPUCommandEncoderRef /*cmd*/, std::uint32_t frameIdx, Window& window) override {
+            const std::uint32_t gx = (window.width  + 7u) / 8u;
+            const std::uint32_t gy = (window.height + 7u) / 8u;
+            auto& tlas = RenderingElement3D::tlases[frameIdx];
+            WebGPU::wgpuDispatchRT(
+                pipeline->pipelineHandle,
+                pushPtr,
+                static_cast<std::int32_t>(pushBytes),
+                tlas.buffer.handle,
+                static_cast<std::int32_t>(tlas.builtInstanceCount),
+                static_cast<std::int32_t>(gx),
+                static_cast<std::int32_t>(gy));
+        }
+    };
+}
+#endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics-RenderingElement3D.cppm

@@ -22,6 +22,7 @@ module;
 #include "vulkan/vulkan.h"
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
 export module Crafter.Graphics:RenderingElement3D;
+import :RT;
 #ifndef CRAFTER_GRAPHICS_WINDOW_DOM
 import std;
 import :Mesh;
@@ -55,7 +56,7 @@ export namespace Crafter {
 
    class RenderingElement3D {
         public:
-        VkAccelerationStructureInstanceKHR instance;
+        RTInstance instance;
         // Position in `elements`, maintained by Add/Remove for O(1) swap-and-pop.
         // Sentinel value = not currently registered.
         std::uint32_t indexInElements = std::numeric_limits<std::uint32_t>::max();
@@ -87,3 +88,63 @@ export namespace Crafter {
    };
 }
 #endif // !CRAFTER_GRAPHICS_WINDOW_DOM
+
+#ifdef CRAFTER_GRAPHICS_WINDOW_DOM
+import std;
+import :Mesh;
+import :WebGPU;
+import :WebGPUBuffer;
+import :Window;
+
+export namespace Crafter {
+    // Per-frame TLAS storage. WebGPU has no real swapchain frame count
+    // (Window::numFrames = 1 on DOM), so this is effectively a singleton —
+    // the array form is kept for API symmetry with the Vulkan side so user
+    // code that indexes `tlases[frameIdx]` ports unchanged.
+    struct TlasWithBuffer {
+        // Host-visible instance buffer holding RTInstance entries — same
+        // layout as Vulkan's VkAccelerationStructureInstanceKHR, so user
+        // code touching .instance.mask / .flags / .transform.matrix is
+        // identical across backends. Also bound as a storage SSBO so
+        // application compute shaders (e.g. physics-tlas-transform.comp.wgsl)
+        // can write the .transform field directly when
+        // RenderingElement3D::transformOwnedByGpu is set.
+        WebGPUBuffer<RTInstance, true>     instanceBuffer;
+        // Per-instance application metadata; parallel to instanceBuffer,
+        // identical semantics to the Vulkan-side counterpart.
+        WebGPUBuffer<std::uint32_t, true>  metadataBuffer;
+        // GPU-built TLAS data: one TLASEntry per instance, written each
+        // BuildTLAS by a compute pass on the JS bridge. Read by traceRay /
+        // rayQuery as `@group(1) @binding(0) tlas: array<TLASEntry>`.
+        // TLASEntry layout: 96 bytes — aabbMin (12) + maskHGoffset (4) +
+        // aabbMax (12) + blasHandle (4) + invTransform 3x4 mat (48) +
+        // customIndex (4) + _pad (12). Defined in the WGSL traversal
+        // library; never directly read by C++.
+        WebGPUBuffer<char, false>           buffer;
+
+        std::uint32_t builtInstanceCount = 0;
+    };
+
+    class RenderingElement3D {
+        public:
+        RTInstance instance{};
+        std::uint32_t indexInElements = std::numeric_limits<std::uint32_t>::max();
+        std::uint32_t userMetadata = 0;
+        // Application compute shader writes the transform field of this
+        // element's instanceBuffer slot directly — BuildTLAS preserves it.
+        bool transformOwnedByGpu = false;
+
+        static std::vector<RenderingElement3D*> elements;
+        inline static TlasWithBuffer tlases[Window::numFrames];
+
+        // Repopulate the TLAS for frame `index`. WebGPU path always does
+        // a fresh build (no refit) — the GPU build pass is cheap at the
+        // ~10–100 instance counts the design targets; LBVH-for-TLAS is a
+        // future optimization for larger scenes.
+        static void BuildTLAS(WebGPUCommandEncoderRef cmd, std::uint32_t index);
+
+        static void Add(RenderingElement3D* e);
+        static void Remove(RenderingElement3D* e);
+    };
+}
+#endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics-ShaderBindingTableWebGPU.cppm

@@ -0,0 +1,64 @@
+/*
+Crafter®.Graphics
+Copyright (C) 2026 Catcrafts®
+catcrafts.net
+*/
+
+// DOM-mode shader-binding-table analog. Stores raw WGSL source strings
+// plus an explicit entry-function name per shader. PipelineRTWebGPU::Init
+// concatenates these into the megakernel WGSL at pipeline-build time.
+
+export module Crafter.Graphics:ShaderBindingTableWebGPU;
+#ifdef CRAFTER_GRAPHICS_WINDOW_DOM
+import std;
+
+export namespace Crafter {
+    enum class WebGPURTStage : std::uint8_t {
+        Raygen     = 0,
+        Miss       = 1,
+        ClosestHit = 2,
+        AnyHit     = 3,
+    };
+
+    // One WGSL shader source + the function name PipelineRTWebGPU should
+    // call from the megakernel switch. The source may declare any helper
+    // functions and (in exactly one raygen file) the `Payload` struct.
+    //
+    // Required signatures inside `source` for `entryFn`:
+    //   Raygen:     fn <entryFn>(gid: vec3<u32>)
+    //   Miss:       fn <entryFn>(ray: RayDesc, payload: ptr<function, Payload>)
+    //   ClosestHit: fn <entryFn>(ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>)
+    //   AnyHit:     fn <entryFn>(ray: RayDesc, hit: HitInfo, payload: ptr<function, Payload>) -> u32
+    //               returns RT_ANYHIT_ACCEPT / RT_ANYHIT_IGNORE / RT_ANYHIT_END_SEARCH.
+    //
+    // `RayDesc`, `HitInfo`, the `RT_*` flag/return constants, the `tlas` /
+    // BLAS / mesh-record bindings, and the `traceRay` function are all
+    // injected by the library prelude — see the rtWgslPrelude block in
+    // additional/dom-webgpu.js.
+    struct WebGPUShader {
+        std::string   source;
+        std::string   entryFn;
+        WebGPURTStage stage = WebGPURTStage::Raygen;
+
+        WebGPUShader() = default;
+        WebGPUShader(std::string src, std::string fn, WebGPURTStage s)
+            : source(std::move(src)), entryFn(std::move(fn)), stage(s) {}
+
+        // Construct from a WGSL source file path. Reads via the WASI VFS
+        // so apps shipping their shaders as static files (see the
+        // `cfg.files.emplace_back("raygen.wgsl")` pattern in
+        // examples/VulkanTriangle/project.cpp) get them at runtime.
+        WebGPUShader(const std::filesystem::path& wgslPath,
+                     std::string fn,
+                     WebGPURTStage s);
+    };
+
+    class ShaderBindingTableWebGPU {
+    public:
+        std::vector<WebGPUShader> shaders;
+        void Init(std::span<const WebGPUShader> shaders_) {
+            shaders.assign(shaders_.begin(), shaders_.end());
+        }
+    };
+}
+#endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics-WebGPU.cppm

@@ -73,13 +73,62 @@ namespace Crafter::WebGPU {
                                      std::uint32_t atlasHandle, std::uint32_t sampHandle);
 
     // ─── custom user-authored compute shaders ───────────────────────────
+    // rayQueryFlag = 1 swaps group(1) from the UI ping-pong pair to the RT
+    // data heaps (TLAS, BVH, meshRecs, verts, idx, primRemap, outImage) and
+    // prepends a WGSL prelude exposing the rayQuery* API. Shaders that set
+    // this MUST NOT declare their own @group(1) bindings.
     __attribute__((import_module("env"), import_name("wgpuLoadCustomShader")))
     extern "C" std::uint32_t wgpuLoadCustomShader(const void* wgslPtr, std::int32_t wgslLen,
-                                                  const void* bindingsPtr, std::int32_t bindingsCount);
+                                                  const void* bindingsPtr, std::int32_t bindingsCount,
+                                                  std::int32_t rayQueryFlag);
     __attribute__((import_module("env"), import_name("wgpuDispatchCustom")))
     extern "C" void wgpuDispatchCustom(std::uint32_t pipelineHandle,
                                        const void* pushPtr, std::int32_t pushBytes,
                                        const void* handlesPtr, std::int32_t handlesCount,
                                        std::int32_t gx, std::int32_t gy, std::int32_t gz);
+
+    // ─── software raytracing ───────────────────────────────────────────
+    //
+    // Mesh::Build forwards vertex / index / BVH-node / primRemap arrays
+    // to the JS bridge, which queue.writeBuffers them into the global
+    // RT mesh heaps (growing if needed) and records the per-mesh offsets
+    // under a freshly-allocated u32 handle. The handle is what user code
+    // stores in RTInstance::accelerationStructureReference; the WebGPU
+    // TLAS-build compute shader resolves it back to root AABB + heap
+    // offsets at dispatch time. Returns 0 on failure.
+    __attribute__((import_module("env"), import_name("wgpuRegisterMeshBLAS")))
+    extern "C" std::uint32_t wgpuRegisterMeshBLAS(
+        float minX, float minY, float minZ,
+        float maxX, float maxY, float maxZ,
+        const void* verticesPtr, std::int32_t vertexCount,
+        const void* indicesPtr,  std::int32_t indexCount,
+        const void* bvhNodesPtr, std::int32_t bvhNodeCount,
+        const void* primRemapPtr, std::int32_t primRemapCount);
+
+    // RT pipeline build. The library composes WGSL by concatenating the
+    // traversal library, generated hit-group switches, and the user-
+    // supplied raygen / miss / closesthit / anyhit bodies. Returns an
+    // opaque pipeline handle.
+    __attribute__((import_module("env"), import_name("wgpuLoadRTPipeline")))
+    extern "C" std::uint32_t wgpuLoadRTPipeline(const void* wgslPtr, std::int32_t wgslLen);
+
+    // Dispatch a TraceRays-equivalent pass: the RT pipeline is dispatched
+    // over a (gx, gy) tile grid; the library writes the push data (camera,
+    // payload, etc. — opaque) into a uniform ring buffer, attaches the TLAS
+    // + global mesh heap, and runs one workgroup per 8x8 screen tile.
+    __attribute__((import_module("env"), import_name("wgpuDispatchRT")))
+    extern "C" void wgpuDispatchRT(std::uint32_t pipelineHandle,
+                                   const void* pushPtr, std::int32_t pushBytes,
+                                   std::uint32_t tlasBufHandle,
+                                   std::int32_t  instanceCount,
+                                   std::int32_t  gx, std::int32_t gy);
+
+    // GPU TLAS-build dispatch. Reads the instance buffer (host-uploaded or
+    // GPU-written), produces per-instance world-space AABBs + per-instance
+    // transform matrices in a flat tlasBuf SSBO consumed by traceRay / rayQuery.
+    __attribute__((import_module("env"), import_name("wgpuBuildTLAS")))
+    extern "C" void wgpuBuildTLAS(std::uint32_t instanceBufHandle,
+                                  std::int32_t  instanceCount,
+                                  std::uint32_t tlasOutBufHandle);
 }
 #endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics-WebGPUComputeShader.cppm

@@ -49,6 +49,7 @@ export namespace Crafter {
     class WebGPUComputeShader {
     public:
         std::uint32_t pipelineHandle = 0;
+        bool          rayQueryCapable = false;
         std::vector<UICustomBinding> customBindings;
 
         WebGPUComputeShader() = default;
@@ -56,22 +57,28 @@ export namespace Crafter {
         WebGPUComputeShader& operator=(const WebGPUComputeShader&) = delete;
         WebGPUComputeShader(WebGPUComputeShader&& o) noexcept
             : pipelineHandle(o.pipelineHandle),
+              rayQueryCapable(o.rayQueryCapable),
               customBindings(std::move(o.customBindings)) {
             o.pipelineHandle = 0;
         }
 
         // Compile + link a custom compute shader. `wgsl` is the source
-        // string; the library does NOT add anything to it — the user's
+        // string; the library does NOT add anything to it (except when
-        // shader must declare @group(0)/@group(1) bindings matching the
+        // `rayQuery` is true — then a RT prelude exposing the rayQuery*
-        // contract above. `bindings` lists every additional resource
+        // API is prepended). The user's shader must declare
-        // (groups 2+) that the renderer should bind at dispatch time.
+        // @group(0)/@group(1) bindings matching the contract above
+        // (rayQuery shaders MUST NOT redeclare group(1)).
+        // `bindings` lists every additional resource (groups 2+) that the
+        // renderer should bind at dispatch time.
         void Load(std::string_view wgsl,
-                  std::span<const UICustomBinding> bindings = {});
+                  std::span<const UICustomBinding> bindings = {},
+                  bool rayQuery = false);
 
         // Path-based overload for symmetry with the Vulkan ComputeShader.
         // Reads the file from disk (browser VFS) and forwards to Load(wgsl).
         void Load(const std::filesystem::path& wgslPath,
-                  std::span<const UICustomBinding> bindings = {});
+                  std::span<const UICustomBinding> bindings = {},
+                  bool rayQuery = false);
     };
 }
 #endif // CRAFTER_GRAPHICS_WINDOW_DOM

interfaces/Crafter.Graphics.cppm

@@ -58,6 +58,10 @@ export import :UIComponents;
 export import :InputField;
 export import :Decompress;
 
+// Portable RT type aliases (provided on both targets — uses Vulkan
+// structs natively, plain PODs of the same layout on DOM).
+export import :RT;
+
 // DOM-only partitions — empty under native.
 export import :Dom;
 export import :DomEvents;
@@ -66,3 +70,5 @@ export import :WebGPU;
 export import :WebGPUBuffer;
 export import :DescriptorHeapWebGPU;
 export import :WebGPUComputeShader;
+export import :ShaderBindingTableWebGPU;
+export import :PipelineRTWebGPU;

project.cpp

@@ -131,7 +131,7 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
     // when its body is gated out. Vulkan-typed partitions stub to empty
     // modules under CRAFTER_GRAPHICS_WINDOW_DOM; the Dom/DomEvents/Router
     // partitions stub to empty modules in the opposite direction.
-    std::array<fs::path, 37> ifaces = {
+    std::array<fs::path, 40> ifaces = {
         "interfaces/Crafter.Graphics",
         "interfaces/Crafter.Graphics-Animation",
         "interfaces/Crafter.Graphics-Clipboard",
@@ -153,12 +153,15 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
         "interfaces/Crafter.Graphics-Keys",
         "interfaces/Crafter.Graphics-Mesh",
         "interfaces/Crafter.Graphics-PipelineRTVulkan",
+        "interfaces/Crafter.Graphics-PipelineRTWebGPU",
         "interfaces/Crafter.Graphics-RenderingElement3D",
         "interfaces/Crafter.Graphics-RenderPass",
         "interfaces/Crafter.Graphics-Router",
+        "interfaces/Crafter.Graphics-RT",
         "interfaces/Crafter.Graphics-RTPass",
         "interfaces/Crafter.Graphics-SamplerVulkan",
         "interfaces/Crafter.Graphics-ShaderBindingTableVulkan",
+        "interfaces/Crafter.Graphics-ShaderBindingTableWebGPU",
         "interfaces/Crafter.Graphics-ShaderVulkan",
         "interfaces/Crafter.Graphics-Types",
         "interfaces/Crafter.Graphics-UI",
@@ -175,14 +178,18 @@ extern "C" Configuration CrafterBuildProject(std::span<const std::string_view> a
         // DOM impl set. UI-Shared.cpp is backend-agnostic; UI-WebGPU.cpp
         // is the DOM-only implementation of UIRenderer's GPU-touching
         // methods. Font / FontAtlas / UIComponents are now portable.
-        std::array<fs::path, 12> domImpls = {
+        std::array<fs::path, 16> domImpls = {
             "implementations/Crafter.Graphics-Clipboard",
             "implementations/Crafter.Graphics-Dom",
             "implementations/Crafter.Graphics-Font",
             "implementations/Crafter.Graphics-FontAtlas",
             "implementations/Crafter.Graphics-Gamepad",
             "implementations/Crafter.Graphics-Input",
+            "implementations/Crafter.Graphics-Mesh-WebGPU",
+            "implementations/Crafter.Graphics-PipelineRTWebGPU",
+            "implementations/Crafter.Graphics-RenderingElement3D-WebGPU",
             "implementations/Crafter.Graphics-Router",
+            "implementations/Crafter.Graphics-ShaderBindingTableWebGPU",
             "implementations/Crafter.Graphics-UI-Shared",
             "implementations/Crafter.Graphics-UI-WebGPU",
             "implementations/Crafter.Graphics-UIComponents",