4e42d663a6
Replace the megakernel @compute entry with five wavefront kernels sharing one module, connected by GPU ray/hit/payload buffers and a GPU-driven indirect bounce loop: GENERATE -> (PREP -> TRACE -> SHADE) x maxDepth -> RESOLVE - TRACE contains zero user code (pure _rtwTraverseTlas/Blas, opaque-only). - PREP publishes dispatchWorkgroupsIndirect args from the live ray count; the indirect-args buffer lives in its own bind group so it is never bound read-write in the same dispatch that consumes it as INDIRECT. - New emit/accumulate API: rtEmitPrimaryRay / rtEmitRay / rtAccumulate, plus an optional user Resolve stage (tonemap hook; identity by default). - Per-pass WfParams via a dynamic-offset uniform ring (curIsA/bounce vary between passes within one submit). - Payload-typed wfPayload binding emitted in the codegen region after the user's struct Payload; payload travels with each ray (2*W*H slots). - Request maxBufferSize / maxStorageBufferBindingSize / maxComputeWorkgroups PerDimension so the W*H-sized work buffers fit past the 128MB baseline. VulkanTriangle ported to the new API and renders bit-identical to the megakernel baseline at maxDepth=1. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
101 lines
3.8 KiB
C++
101 lines
3.8 KiB
C++
/*
|
||
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;
|
||
|
||
This library is distributed in the hope that it will be useful,
|
||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
|
||
Lesser General Public License for more details.
|
||
|
||
You should have received a copy of the GNU Lesser General Public
|
||
License along with this library; if not, write to the Free Software
|
||
Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
|
||
*/
|
||
module;
|
||
#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
|
||
#include "vulkan/vulkan.h"
|
||
#endif // !CRAFTER_GRAPHICS_WINDOW_DOM
|
||
export module Crafter.Graphics:RTPass;
|
||
#ifndef CRAFTER_GRAPHICS_WINDOW_DOM
|
||
import std;
|
||
import :RenderPass;
|
||
import :Window;
|
||
import :Device;
|
||
import :PipelineRTVulkan;
|
||
|
||
export namespace Crafter {
|
||
struct RTPass : RenderPass {
|
||
PipelineRTVulkan* pipeline;
|
||
|
||
RTPass(PipelineRTVulkan* p) : pipeline(p) {}
|
||
|
||
void Record(VkCommandBuffer cmd, std::uint32_t /*frameIdx*/, Window& window) override {
|
||
vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_RAY_TRACING_KHR, pipeline->pipeline);
|
||
Device::vkCmdTraceRaysKHR(cmd,
|
||
&pipeline->raygenRegion,
|
||
&pipeline->missRegion,
|
||
&pipeline->hitRegion,
|
||
&pipeline->callableRegion,
|
||
window.width, window.height, 1);
|
||
}
|
||
};
|
||
}
|
||
#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;
|
||
// Resolved WebGPU resource handles for each user binding the
|
||
// pipeline was loaded with, in declaration order. The example
|
||
// owns the storage (typically a small std::array of u32). Null /
|
||
// 0 means "no user bindings".
|
||
const void* handlesPtr = nullptr;
|
||
std::uint32_t handlesCount = 0;
|
||
// Wavefront bounce budget: number of (TRACE; SHADE) iterations.
|
||
// 1 = primary rays only; 2 = primary + one continuation/shadow
|
||
// bounce; etc. The library unrolls GENERATE; (PREP; TRACE; SHADE)
|
||
// ×maxDepth; RESOLVE.
|
||
std::uint32_t maxDepth = 1;
|
||
|
||
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),
|
||
handlesPtr,
|
||
static_cast<std::int32_t>(handlesCount),
|
||
static_cast<std::int32_t>(maxDepth));
|
||
}
|
||
};
|
||
}
|
||
#endif // CRAFTER_GRAPHICS_WINDOW_DOM
|