pub const CULL: &str = "// GPU-driven draw cull. One thread per draw record: tests the record\'s AABB\n// against the frustum, its cull distance and the previous frame\'s Hi-Z pyramid,\n// then writes either a real indirect draw or an instance_count-0 no-op into the\n// command buffer the indirect draw reads.\n//\n// Three variants out of one entry point, selected by define exactly as the\n// hand-written sources were:\n//\n// (none) phase 1: frustum + distance + Hi-Z, and records the outcome in\n// `cull_status` so phase 2 knows what to reconsider\n// CULL_PHASE2 re-tests only phase 1\'s Hi-Z candidates against the rebuilt\n// pyramid, for two-pass occlusion\n// SHADOW_CULL light-frustum only, one dispatch per cascade\n//\n// Single source for all three hosts. On Vulkan and DirectX the kernel writes\n// the indirect command record itself; on Metal it writes only `cull_status`,\n// and the hand-written `metal/shaders/cull_encode.metal` turns that status into\n// the indirect command buffer. Metal has no multi-draw-indirect and its ICB\n// encoding (`render_command`, `array<command_buffer, N>`) is a declaration\n// Slang cannot express, but nothing in the decision needs it, so the decision\n// is what lives here and the encoding is what stays per backend.\n//\n// DXIL_ABI and METAL_BINDINGS select a *host\'s* constant shape, a host\n// difference rather than a target one. Vulkan hands the kernel a push constant for the\n// frustum and bucket routing plus a set-1 uniform buffer for the Hi-Z\n// reprojection, because the Hi-Z half changes on a different schedule and is\n// owned by `vulkan/hiz.rs`. DirectX fuses both into one b0 root-constant block\n// (directx/cull.rs), so the same fields arrive as one struct there. The\n// accessor defines below are what let the body read them without caring.\n//\n// The command record itself genuinely differs and cannot be unified: a Vulkan\n// indirect draw is a bare `VkDrawIndexedIndirectCommand` whose `first_instance`\n// carries the object id, while a D3D12 ExecuteIndirect command signature\n// prepends the b0 root constant the shaders read the object id from, so its\n// record is a word longer and starts with that id.\n\n{OBJECT_COMMON}\n\n// `cull_status` values. STATUS_HIZ_CANDIDATE is the only outcome phase 2\n// re-tests against the rebuilt pyramid; the others are settled by phase 1.\nstatic const uint STATUS_DRAWN = 0u; // visible in phase 1 -> never re-tested\nstatic const uint STATUS_HIZ_CANDIDATE = 1u; // Hi-Z-occluded in phase 1 -> candidate\nstatic const uint STATUS_CULLED = 2u; // frustum/distance/disabled -> settled\nstatic const uint STATUS_REDRAW = 3u; // candidate found visible in phase 2\nstatic const uint STATUS_HIZ_CULLED = 4u; // candidate still occluded in phase 2\n\n// ---- Per-host constants ----\n\n#if defined(METAL_BINDINGS)\n\n// Layout matches `metal::uniforms::CullUniforms` (208 B), the inline constant\n// block at buffer(2). `cam_pos` is a whole float4 lane (w unused): MSL sizes a\n// constant-buffer float3 at 16 bytes, so a scalar after it would not pack.\n// `skinned_base` and `cascade_base` are read by the encode kernel\'s host, not\n// here; they ride the one block both dispatches share.\nstruct CullParams\n{\n float4 planes[6];\n float4 cam_pos;\n float4x4 prev_view_proj;\n float2 hiz_size;\n uint hiz_mip_count;\n uint hiz_enabled;\n uint object_count;\n uint skinned_base;\n uint cascade_base;\n uint bucket_count;\n};\n\n[[vk::push_constant]] ConstantBuffer<CullParams> cull : register(b2);\n\n#define CULL_PREV_VP cull.prev_view_proj\n#define CULL_HIZ_SIZE cull.hiz_size\n#define CULL_HIZ_MIPS cull.hiz_mip_count\n#define CULL_HIZ_ENABLED cull.hiz_enabled\n\n#elif defined(DXIL_ABI)\n\n// Layout matches `directx::uniforms::CullParams` (208 B), the b0 root\n// constants. `cam_pos` shares its 16-byte row with `object_count`.\nstruct CullParams\n{\n float4 planes[6];\n float3 cam_pos;\n uint object_count;\n float4x4 prev_view_proj;\n float2 hiz_size;\n uint hiz_mip_count;\n uint hiz_enabled;\n uint bucket_count;\n uint bucket_stride;\n uint _pad0;\n uint _pad1;\n};\n\nConstantBuffer<CullParams> cull : register(b0);\n\n#define CULL_PREV_VP cull.prev_view_proj\n#define CULL_HIZ_SIZE cull.hiz_size\n#define CULL_HIZ_MIPS cull.hiz_mip_count\n#define CULL_HIZ_ENABLED cull.hiz_enabled\n\n#else\n\n// Layout matches `vulkan::uniforms::CullParams` (120 B), the push constant.\nstruct CullParams\n{\n float4 planes[6];\n float3 cam_pos;\n uint object_count;\n uint bucket_count;\n uint bucket_stride;\n};\n\n// Layout matches `vulkan::uniforms::CullHizParams` (80 B), set 1 binding 1.\nstruct CullHizParams\n{\n float4x4 prev_view_proj;\n float2 hiz_size;\n uint hiz_mip_count;\n uint hiz_enabled;\n};\n\n[[vk::push_constant]] ConstantBuffer<CullParams> cull;\n\n#ifndef SHADOW_CULL\n[[vk::binding(1, 1)]] ConstantBuffer<CullHizParams> hiz;\n#endif\n\n#define CULL_PREV_VP hiz.prev_view_proj\n#define CULL_HIZ_SIZE hiz.hiz_size\n#define CULL_HIZ_MIPS hiz.hiz_mip_count\n#define CULL_HIZ_ENABLED hiz.hiz_enabled\n\n#endif\n\n// ---- Buffers ----\n//\n// Every register is pinned rather than left to declaration order, because the\n// order here is not the root signature\'s: `cull_status` has to be declared\n// before `DrawCommand` (the shadow variant declares no status, and the struct\n// differs per host), which would otherwise hand it u0 and give `commands` u1 --\n// the reverse of what directx/cull.rs binds. The `[[vk::binding]]`s spell\n// Vulkan\'s own set/binding numbers on the same declarations, which is also what\n// keeps slangc from warning about a bare `register()`.\n\n[[vk::binding(0, 0)]] StructuredBuffer<GpuObjectData> objects : register(t0);\n[[vk::binding(1, 0)]] StructuredBuffer<GpuDrawArgs> draw_args : register(t1);\n\n#ifdef METAL_BINDINGS\n\n// Metal\'s one buffer index space: the numbers are the slots `metal/cull.rs`\n// binds, shared with the encode kernel that reads the status back. The status\n// is the kernel\'s only output here, so every variant declares it, the shadow\n// one at the cascade\'s region of the shadow status buffer. The Hi-Z pyramid\n// takes its texture index from declaration order (the only texture), which\n// `assert_slang_metal_abi` locks to texture(0).\n[[vk::binding(3, 0)]] RWStructuredBuffer<uint> cull_status : register(u5);\n#ifndef SHADOW_CULL\nTexture2D<float> hiz_tex;\n#endif\n\n#else\n\n// The GPU-driven shadow cull is frustum-only: sun cascades have no light-space\n// depth pyramid, so there is no Hi-Z test and no two-pass occlusion. It omits\n// the status buffer and the whole set-1 Hi-Z block, which is what leaves its\n// Vulkan pipeline layout the lean three-binding set its host creates.\n#ifndef SHADOW_CULL\n// Per-object outcome of phase-1 cull. Always bound and written so phase 2 sees\n// valid data; under single-pass occlusion the values are simply ignored.\n[[vk::binding(3, 0)]] RWStructuredBuffer<uint> cull_status : register(u1);\n#endif\n\n#ifdef DXIL_ABI\n\n// One ExecuteIndirect command: the b0 object-id root constant followed by\n// D3D12_DRAW_INDEXED_ARGUMENTS. 24 bytes; matches the command signature in\n// directx/init/pipelines.rs.\nstruct DrawCommand\n{\n uint object_id;\n uint index_count;\n uint instance_count;\n uint start_index;\n int base_vertex;\n uint start_instance;\n};\n\nDrawCommand make_command(uint i, GpuDrawArgs a)\n{\n DrawCommand c;\n c.object_id = i;\n c.index_count = a.index_count;\n c.instance_count = 1u;\n c.start_index = a.index_offset;\n c.base_vertex = int(a.base_vertex);\n c.start_instance = 0u;\n return c;\n}\n\n#else\n\n// Matches VkDrawIndexedIndirectCommand (20 bytes). `first_instance` delivers\n// the record id to the depth-only and bindless vertex stages, which read it\n// back through `object_instance_index`.\nstruct DrawCommand\n{\n uint index_count;\n uint instance_count;\n uint first_index;\n int vertex_offset;\n uint first_instance;\n};\n\nDrawCommand make_command(uint i, GpuDrawArgs a)\n{\n DrawCommand c;\n c.index_count = a.index_count;\n c.instance_count = 1u;\n c.first_index = a.index_offset;\n c.vertex_offset = int(a.base_vertex);\n c.first_instance = i;\n return c;\n}\n\n#endif\n\n[[vk::binding(2, 0)]] RWStructuredBuffer<DrawCommand> commands : register(u0);\n\n#ifndef SHADOW_CULL\n// Read by texel coordinate only, never sampled. Vulkan\'s cull-read set binds a\n// combined image sampler, so this is a `Sampler2D` there; DirectX binds a plain\n// SRV through the root signature\'s descriptor table.\n#ifdef DXIL_ABI\nTexture2D<float> hiz_tex : register(t2);\n#else\n[[vk::binding(0, 1)]] Sampler2D<float> hiz_tex;\n#endif\n#endif\n\n#endif // !METAL_BINDINGS\n\n// AABB entirely behind any plane -> outside the frustum. Negation of\n// gfx::frustum::Frustum::intersects_aabb (the p-vertex test).\nbool frustum_culled(float3 bb_min, float3 bb_max)\n{\n [unroll] for (uint i = 0u; i < 6u; ++i)\n {\n float3 n = cull.planes[i].xyz;\n float3 farthest = float3(\n n.x >= 0.0 ? bb_max.x : bb_min.x,\n n.y >= 0.0 ? bb_max.y : bb_min.y,\n n.z >= 0.0 ? bb_max.z : bb_min.z);\n if (dot(n, farthest) + cull.planes[i].w < 0.0)\n {\n return true;\n }\n }\n return false;\n}\n\n// Squared distance from the camera to the closest point on the AABB; 0 when the\n// camera is inside. Mirrors gfx::frustum::aabb_distance_sq.\nfloat aabb_distance_sq(float3 bb_min, float3 bb_max)\n{\n float3 d = max(max(bb_min - cull.cam_pos.xyz, cull.cam_pos.xyz - bb_max), float3(0.0, 0.0, 0.0));\n return dot(d, d);\n}\n\n#ifndef SHADOW_CULL\n\n// Project the eight AABB corners through the reprojection matrix and reduce to\n// a screen-space rect (NDC.xy in [-1, 1]) plus the AABB\'s closest NDC depth.\n// Returns false if any corner ended up behind the camera (w <= 0), in which\n// case the caller conservatively keeps the object.\nbool project_aabb(\n float3 bb_min,\n float3 bb_max,\n out float2 ndc_min,\n out float2 ndc_max,\n out float min_depth)\n{\n ndc_min = float2( 1.0, 1.0);\n ndc_max = float2(-1.0, -1.0);\n min_depth = 1.0;\n [unroll] for (uint i = 0u; i < 8u; ++i)\n {\n float3 corner = float3(\n (i & 1u) != 0u ? bb_max.x : bb_min.x,\n (i & 2u) != 0u ? bb_max.y : bb_min.y,\n (i & 4u) != 0u ? bb_max.z : bb_min.z);\n float4 clip = mul(CULL_PREV_VP, float4(corner, 1.0));\n if (clip.w <= 0.0)\n {\n return false;\n }\n float3 ndc = clip.xyz / clip.w;\n ndc_min = min(ndc_min, ndc.xy);\n ndc_max = max(ndc_max, ndc.xy);\n min_depth = min(min_depth, ndc.z);\n }\n return true;\n}\n\n// True when the AABB is fully occluded by the Hi-Z pyramid. Conservative: any\n// uncertain case returns false (keep the object). The NDC y-flip matches the\n// main pass\'s viewport, and depth is the [0, 1] range both APIs use, so\n// MAX-reduced texels store the farthest occluder.\nbool hiz_occluded(float3 bb_min, float3 bb_max)\n{\n float2 ndc_min, ndc_max;\n float aabb_min_depth;\n if (!project_aabb(bb_min, bb_max, ndc_min, ndc_max, aabb_min_depth))\n {\n return false;\n }\n // Clip to NDC bounds so the UV math stays sane (an AABB straddling the\n // viewport on both sides of an axis is already kept by the frustum test).\n ndc_min = max(ndc_min, float2(-1.0, -1.0));\n ndc_max = min(ndc_max, float2( 1.0, 1.0));\n if (any(ndc_min > ndc_max))\n {\n return false;\n }\n // Standard depth: nearest point of the AABB at NDC.z near 0. Behind-near or\n // behind-far means we conservatively keep the AABB.\n if (aabb_min_depth < 0.0 || aabb_min_depth > 1.0)\n {\n return false;\n }\n // Map NDC -> UV (y flips because NDC y is up, UV v is down).\n float2 uv_min = float2(ndc_min.x * 0.5 + 0.5, 0.5 - ndc_max.y * 0.5);\n float2 uv_max = float2(ndc_max.x * 0.5 + 0.5, 0.5 - ndc_min.y * 0.5);\n // Size of the rect at mip 0, in texels.\n float2 size_tex = (uv_max - uv_min) * CULL_HIZ_SIZE;\n float max_dim = max(size_tex.x, size_tex.y);\n // Pick the mip whose texels are roughly the rect size so a 2x2 footprint\n // covers the rect (the standard Hi-Z 4-tap pattern).\n int mip = int(ceil(log2(max(max_dim, 1.0))));\n mip = clamp(mip, 0, int(CULL_HIZ_MIPS) - 1);\n float2 mip_dim = max(CULL_HIZ_SIZE / float(1u << uint(mip)), float2(1.0, 1.0));\n int2 lo = int2(floor(uv_min * mip_dim));\n int2 hi = int2(floor(uv_max * mip_dim));\n int2 max_xy = int2(mip_dim) - int2(1, 1);\n lo = clamp(lo, int2(0, 0), max_xy);\n hi = clamp(hi, int2(0, 0), max_xy);\n float d0 = hiz_tex.Load(int3(lo.x, lo.y, mip));\n float d1 = hiz_tex.Load(int3(hi.x, lo.y, mip));\n float d2 = hiz_tex.Load(int3(lo.x, hi.y, mip));\n float d3 = hiz_tex.Load(int3(hi.x, hi.y, mip));\n float occluder_depth = max(max(d0, d1), max(d2, d3));\n // If the AABB\'s closest projected depth is strictly behind the farthest\n // previously-rasterised surface in this region, the whole AABB is hidden.\n return aabb_min_depth > occluder_depth;\n}\n\n#ifndef METAL_BINDINGS\n// Write one record\'s command into its own shader bucket\'s region and a no-op\n// into every other region. `c` already carries instance_count 0 when the record\n// was culled, so a culled record resets every region. The reset sweep matters\n// because a freed draw slot can be reused by a record of a DIFFERENT bucket,\n// which would otherwise leave the old bucket\'s command stale and still\n// executing.\nvoid write_bucket_commands(uint i, DrawCommand c, uint flags)\n{\n uint bucket = min((flags >> DRAW_BUCKET_SHIFT) & DRAW_BUCKET_MASK,\n cull.bucket_count - 1u);\n DrawCommand noop = c;\n noop.instance_count = 0u;\n for (uint b = 0u; b < cull.bucket_count; ++b)\n {\n if (b == bucket)\n {\n commands[b * cull.bucket_stride + i] = c;\n }\n else\n {\n commands[b * cull.bucket_stride + i] = noop;\n }\n }\n}\n#endif // !METAL_BINDINGS\n\n#endif // !SHADOW_CULL\n\n[shader(\"compute\")]\n[numthreads(64, 1, 1)]\nvoid cull_kernel(uint3 tid : SV_DispatchThreadID)\n{\n uint i = tid.x;\n if (i >= cull.object_count)\n {\n return;\n }\n\n#ifdef SHADOW_CULL\n GpuDrawArgs a = draw_args[i];\n // Light-frustum only, against this cascade\'s `planes`. Deliberately NO\n // per-object distance cull: the cascade light frustum already bounds the\n // shadow draw distance through its extents, and the per-object view\n // `cull_distance` is a view-LOD-fade concept that must not silence shadows\n // (the legacy CPU shadow pass drew every caster; an off-screen caster\n // beyond cull_distance can still cast into the visible scene).\n bool draw = (a.flags & DRAW_ENABLED) != 0u;\n if (draw && (a.flags & DRAW_CULLABLE) != 0u)\n {\n GpuObjectData obj = objects[i];\n draw = !frustum_culled(obj.bb_min_cull_distance.xyz, obj.bb_max_alpha_cutoff.xyz);\n }\n#ifdef METAL_BINDINGS\n cull_status[cull.cascade_base + i] = draw ? STATUS_DRAWN : STATUS_CULLED;\n#else\n DrawCommand c = make_command(i, a);\n if (!draw)\n {\n c.instance_count = 0u;\n }\n commands[i] = c;\n#endif\n#elif defined(CULL_PHASE2)\n // Runs after the pyramid was rebuilt from this frame\'s phase-1 depth.\n // Re-tests only the objects phase 1 marked STATUS_HIZ_CANDIDATE against the\n // fresh pyramid (projected through this frame\'s VP, carried in\n // `prev_view_proj` exactly as phase 1 used the previous frame\'s). A\n // candidate still occluded by this frame\'s actual depth stays culled; one\n // now visible is redrawn. Everything else phase 1 settled and is left as it\n // was.\n bool candidate = cull_status[i] == STATUS_HIZ_CANDIDATE;\n bool redraw = candidate;\n if (candidate)\n {\n GpuObjectData obj = objects[i];\n if (CULL_HIZ_ENABLED != 0u\n && hiz_occluded(obj.bb_min_cull_distance.xyz, obj.bb_max_alpha_cutoff.xyz))\n {\n redraw = false;\n }\n cull_status[i] = redraw ? STATUS_REDRAW : STATUS_HIZ_CULLED;\n }\n#ifndef METAL_BINDINGS\n GpuDrawArgs a = draw_args[i];\n DrawCommand c = make_command(i, a);\n if (!redraw)\n {\n c.instance_count = 0u;\n }\n write_bucket_commands(i, c, a.flags);\n#endif\n#else\n // Phase 1. Records the outcome in `cull_status` for two-pass occlusion: a\n // Hi-Z cull is the only outcome phase 2 reconsiders against the rebuilt\n // pyramid; everything else is settled here.\n GpuDrawArgs a = draw_args[i];\n uint status = STATUS_DRAWN;\n if ((a.flags & DRAW_ENABLED) == 0u)\n {\n status = STATUS_CULLED;\n }\n else if ((a.flags & DRAW_CULLABLE) != 0u)\n {\n GpuObjectData obj = objects[i];\n float3 bb_min = obj.bb_min_cull_distance.xyz;\n float3 bb_max = obj.bb_max_alpha_cutoff.xyz;\n float cull_distance = obj.bb_min_cull_distance.w;\n if (frustum_culled(bb_min, bb_max))\n {\n status = STATUS_CULLED;\n }\n else if (cull_distance > 0.0\n && aabb_distance_sq(bb_min, bb_max) > cull_distance * cull_distance)\n {\n status = STATUS_CULLED;\n }\n else if (CULL_HIZ_ENABLED != 0u && hiz_occluded(bb_min, bb_max))\n {\n // Cull when the AABB is fully behind the previous frame\'s depth\n // pyramid. Skipped on the first frame and after a resize\n // (`hiz_enabled == 0`), where no valid pyramid exists yet.\n status = STATUS_HIZ_CANDIDATE;\n }\n }\n cull_status[i] = status;\n#ifndef METAL_BINDINGS\n DrawCommand c = make_command(i, a);\n if (status != STATUS_DRAWN)\n {\n c.instance_count = 0u;\n }\n write_bucket_commands(i, c, a.flags);\n#endif\n#endif\n}\n";Expand description
cull.slang.