From File to Pixel · Part 04: The Shader

pavelzosim:~/atlas_SYS.ONLINE / UTC+3

07 From File to Pixel — Compilation Pipeline

A shader file is text. A pixel on screen is the result of running compiled machine code on a GPU. Between those two points lies a compilation pipeline with 6–8 stages, each with its own failure modes, performance implications, and platform differences. Understanding this pipeline is the prerequisite for diagnosing shader stutter, variant bloat, and cross-platform discrepancies.

Compilation Stages: Source → ISA

1 Source (HLSL/GLSL/MSL) Text file on disk. Preprocessor expands #include, #define, #if. Variant keywords selected. Output: expanded source.
2 Frontend Compiler DXC (HLSL→DXIL/SPIR-V), glslang (GLSL→SPIR-V), Metal compiler (MSL→AIR). Parses, type-checks, produces IR. Happens at build time or on first use.
3 IR / Bytecode DXIL (DX12), SPIR-V (Vulkan/OpenGL), AIR (Metal). Platform-independent intermediate representation. Stored in shader bundle / cache. Inspectable with spirv-dis, dxc /Fc.
4 Backend / Driver JIT Driver compiles IR → GPU ISA (machine code). NVIDIA: PTX → SASS. AMD: SPIR-V → GCN/RDNA ISA. Apple: AIR → GPU binary. This is the stutter source on first draw — always warm your PSO cache.
5 PSO Creation Pipeline State Object bundles compiled shaders + render state (blend, depth, raster). DX12: ID3D12PipelineState. Vulkan: VkPipeline. Metal: MTLRenderPipelineState. Creation is expensive — do it at load time, never mid-frame.
6 PSO Cache Serialized compiled GPU binaries. DX12: ID3D12PipelineLibrary. Vulkan: VkPipelineCache. Metal: MTLBinaryArchive. Load on startup → zero JIT stutter on first draw. Ship in build or generate on first run.
7 Execution Compiled ISA runs on SM/CU. Registers allocated. Warps scheduled. Fragment output written through ROPs to framebuffer. This is the pixel.

Shader stutter = PSO created mid-frame. Always create PSOs during loading screens or scene transitions, never in response to a runtime condition. Use async PSO compilation (DX12 CreatePipelineState with D3D12_PIPELINE_STATE_FLAG_NONE, Vulkan VK_PIPELINE_CREATE_FAIL_ON_PIPELINE_COMPILE_REQUIRED_BIT) to fall back to a simpler variant while the real PSO compiles in a background thread.

Shader Variants — Combinatorial Explosion

Every #pragma multi_compile A B C multiplies your total shader count. Three such pragmas with 3 keywords each produce 27 variants. Five pragmas produce 243. Unity compiles all variants unless you strip them explicitly — this is why shader compilation takes minutes on large projects.

[ VARIANT_CONTROL // shader_feature vs multi_compile ]
PRAGMACOMPILED WHENSTRIPPED WHENUSE CASE
multi_compile A B Always — all variants Never automatically Global keywords that change at runtime (fog, shadows, instancing). Must always exist.
shader_feature A B Only if material uses keyword If no material in build uses it Per-material toggles (normal map, emission). Safe to use liberally.
multi_compile_local A B Per-shader scope only Never automatically Reduces keyword namespace pollution. Use when keyword is shader-specific.
shader_feature_local A B Per-material, local scope If no material uses it Best choice for optional per-material features in URP/HDRP.
// Unity — IPreprocessShaders: strip variants at build time
// Runs in Editor before build — removes unused variants from build output
class MyShaderStripper : IPreprocessShaders
{
    public int callbackOrder => 0;

    public void OnProcessShader(
        Shader shader,
        ShaderSnippetData snippet,
        IList<ShaderCompilerData> data)
    {
        // Strip all variants that have _EXPENSIVE_FEATURE_ON
        // unless the current build target needs them
        for (int i = data.Count - 1; i >= 0; i--)
        {
            if (data[i].shaderKeywordSet.IsEnabled(
                new ShaderKeyword("_EXPENSIVE_FEATURE_ON"))
                && !NeedsExpensiveFeature(snippet))
            {
                data.RemoveAt(i); // stripped — not compiled into build
            }
        }
    }
}

// ── DXC: compile HLSL → SPIR-V from command line ─────────
// dxc -spirv -T vs_6_5 -E vert shader.hlsl -Fo shader.vert.spv
// dxc -spirv -T ps_6_5 -E frag shader.hlsl -Fo shader.frag.spv
// Add -Od for debug info, -O3 for release optimization

08 GPU Instancing Patterns

Instancing is the primary mechanism for rendering thousands of identical (or near-identical) meshes with minimal CPU overhead. The right pattern depends on whether data is static or dynamic, how frequently transforms update, and whether you're in a legacy or modern pipeline.

[ INSTANCING_PATTERNS // UNITY ]
PATTERNAPICPU COSTDYNAMIC DATABEST FOR
SRP BatcherAutomatic (CBUFFER)Very lowPer-material onlyStatic/semi-static scene objects with varying materials
GPU InstancingUNITY_INSTANCING_* macrosLowPer-instance via instanced propertiesMedium counts (<1K), per-instance color/scale
DrawMeshInstancedC# API, managed arrayMedium (GC risk)Yes — upload array per frameDynamic objects, runtime spawned
DrawMeshInstancedIndirectComputeBuffer argsNear-zeroGPU-driven (compute writes args)GPU culling pipelines, >10K instances
DOTS / Entities GraphicsECS + BRGMinimal (Burst)Yes — NativeArray/chunk dataMassive counts, physics-driven, procedural

Pattern Details

SRP Batcher — automatic, zero shader changes

The SRP Batcher is not instancing — it batches SetPass calls, not draw calls. Each object still gets its own draw call, but the GPU state (shader + material) is cached. It activates automatically when your shader wraps per-material properties in CBUFFER_START(UnityPerMaterial).

// SRP Batcher compatible shader — ONE requirement
CBUFFER_START(UnityPerMaterial)
    float4 _BaseColor;
    float  _Roughness;
    float  _Metallic;
    // ALL per-material properties must be inside this CBUFFER
    // SRP Batcher caches this block between draws — no re-upload
CBUFFER_END

// Per-object built-in data is automatically in UnityPerDraw CBUFFER
// (unity_ObjectToWorld, unity_WorldToObject, etc.)
// Both must be CBUFFERed — then Batcher activates

GPU Instancing — per-instance varying data

// Shader — declare instanced property
UNITY_INSTANCING_BUFFER_START(InstanceProps)
    UNITY_DEFINE_INSTANCED_PROP(float4, _Color)
    UNITY_DEFINE_INSTANCED_PROP(float,  _Scale)
UNITY_INSTANCING_BUFFER_END(InstanceProps)

Varyings vert(Attributes input, UNITY_VERTEX_INPUT_INSTANCE_ID)
{
    UNITY_SETUP_INSTANCE_ID(input);
    float4 color = UNITY_ACCESS_INSTANCED_PROP(InstanceProps, _Color);
    // use color in vertex or pass to fragment via Varyings
    ...
}

// C# — set per-instance data via MaterialPropertyBlock
var block = new MaterialPropertyBlock();
for (int i = 0; i < count; i++) {
    block.SetColor("_Color", colors[i]);
    renderer.SetPropertyBlock(block, i); // per-instance block
}

DrawMeshInstancedIndirect — GPU-driven, zero CPU per instance

// C# — setup indirect args buffer
// Args: [indexCount, instanceCount, startIndex, baseVertex, startInstance]
uint[] args = new uint[] { mesh.GetIndexCount(0), 0, 0, 0, 0 };
argsBuffer = new ComputeBuffer(1, args.Length * sizeof(uint),
    ComputeBufferType.IndirectArguments);
argsBuffer.SetData(args);

// Compute shader writes instance count after GPU culling:
// RWBuffer<uint> _ArgsBuffer; → _ArgsBuffer[1] = visibleCount;

// C# — draw call (zero per-instance CPU work)
Graphics.DrawMeshInstancedIndirect(
    mesh, 0, material, bounds, argsBuffer);

// HLSL shader — read instance data from StructuredBuffer
StructuredBuffer<float4x4> _InstanceMatrices;

Varyings vert(Attributes input, uint instanceID : SV_InstanceID)
{
    float4x4 mat = _InstanceMatrices[instanceID];
    output.posCS = mul(mat, float4(input.posOS, 1.0));
    ...
}

When to use which: Static scene objects with material variety → SRP Batcher. Hundreds of objects with per-instance color/scale → GPU Instancing. Thousands of runtime objects → DrawMeshInstancedIndirect with compute culling. Tens of thousands with physics → DOTS Entities Graphics.

09 HLSL / GLSL / MSL — Cross-Language Reference

All three shader languages express the same GPU concepts — vertex transformation, texture sampling, compute dispatch — with different syntax and slightly different semantics. If you know one, the others are an afternoon of pattern matching. The table below is that pattern match.

Syntax Equivalence

[ SYNTAX_EQUIVALENCE // HLSL · GLSL · MSL ]
CONCEPTHLSL (DX12/Unity)GLSL (Vulkan/GL)MSL (Metal/Apple)
Float 2/3/4 vectorfloat2/3/4vec2/3/4float2/3/4
Int vectorint2/3/4ivec2/3/4int2/3/4
Matrix 4×4float4x4mat4float4x4
Matrix multiplymul(mat, vec)mat * vecmat * vec
Texture 2D typeTexture2D<T>sampler2D / texture2Dtexture2d<T>
Sample texturetex.Sample(samp, uv)texture(sampler2D, uv)tex.sample(samp, uv)
Sample explicit LODtex.SampleLevel(s,uv,lod)textureLod(s,uv,lod)tex.sample(s,uv,level(lod))
Vertex position outSV_POSITION semanticgl_Position built-in[[position]] attribute
Fragment color out: SV_Targetlayout(location=0) out[[color(0)]] return
Instance IDSV_InstanceIDgl_InstanceIndex[[instance_id]]
Vertex IDSV_VertexIDgl_VertexIndex[[vertex_id]]
Clip positionSV_POSITIONgl_Position[[position]]
Saturate (clamp 0–1)saturate(x)clamp(x,0.0,1.0)saturate(x)
Fused multiply-addmad(a,b,c)fma(a,b,c)fma(a,b,c)
Partial derivative ∂/∂xddx(x)dFdx(x)dfdx(x)
Partial derivative ∂/∂yddy(x)dFdy(x)dfdy(x)
Atomics (shared mem)InterlockedAdd(dst, val)atomicAdd(dst, val)atomic_fetch_add_explicit()
Group memory barrierGroupMemoryBarrierWithGroupSync()barrier() + memoryBarrierShared()threadgroup_barrier(mem_flags::mem_threadgroup)
Discard fragmentdiscard / clip(val)discarddiscard_fragment()
Constant buffercbuffer / CBUFFER_STARTlayout(std140) uniform Block {}constant T& [[buffer(N)]]
Compute thread IDSV_DispatchThreadIDgl_GlobalInvocationID[[thread_position_in_grid]]
Group thread IDSV_GroupThreadIDgl_LocalInvocationID[[thread_position_in_threadgroup]]
Group shared memorygroupshared T var[]shared T var[]threadgroup T var[]

Precision Qualifiers Across Languages

PRECISIONHLSLGLSL / VulkanMSLMOBILE IMPACT
32-bit floatfloathighp floatfloatFull cost — use for positions, world-space
16-bit floathalf / min16floatmediump floathalf2× faster ALU on Adreno/Mali — use for colors, normals, UVs
10-bit floatmin10float (DX only)lowp floatRare; useful for 0–1 values on very low-end
16-bit intmin16intmediump intshortFast on mobile; use for indices, flags

Precision on desktop vs mobile: On PC and console, half and float compile to the same 32-bit operations — the hint is ignored. On Mali/Adreno/Apple GPU, half genuinely executes at 16-bit precision with 2× throughput. This makes mobile-targeted shaders meaningfully faster when written with half for non-position data — but you must verify that 16-bit precision is sufficient for the calculation.

10 Compute Shaders

Compute shaders run outside the graphics pipeline — no vertex stage, no rasterization, no fixed render targets. They dispatch arbitrary parallel workloads on the GPU: particle simulation, physics, GPU culling, procedural generation, post-processing. Every modern GPU supports them. They are the primary tool for GPU-driven rendering.

Execution Model · Thread Groups · Shared Memory

A compute dispatch is a 3D grid of thread groups. Each thread group is a 3D block of threads. The total thread count is groups.x × groups.y × groups.z × numthreads.x × numthreads.y × numthreads.z.

CONCEPTHLSLGLSLMSLHARDWARE
Thread group size[numthreads(X,Y,Z)]layout(local_size_x=X,...)kernel void f([[ ... ]])Must be ≤ 1024 total threads
Global thread IDSV_DispatchThreadIDgl_GlobalInvocationID[[thread_position_in_grid]]gid = groupID * groupSize + localID
Local thread IDSV_GroupThreadIDgl_LocalInvocationID[[thread_position_in_threadgroup]]0..numthreads-1 per group
Group indexSV_GroupIndexgl_LocalInvocationIndex[[thread_index_in_threadgroup]]Flattened: z*XY + y*X + x
Shared memorygroupshared T[]shared T[]threadgroup T[]On-chip SRAM — 32–100 KB
Sync barrierGroupMemoryBarrierWithGroupSync()barrier()threadgroup_barrier()All threads in group reach barrier

Thread Group Size Selection

WORKLOADRECOMMENDED SIZEREASON
1D data (particles, bones)[256, 1, 1]Maps to NVIDIA warp×8 or AMD wave×4. Good occupancy.
2D image / texture[8, 8, 1]64 threads — one NVIDIA warp. Matches 2D cache locality.
2D with shared memory[16, 16, 1]256 threads. Good for tiled algorithms (blur, lighting).
Wave-level reduction[32, 1, 1] or [64, 1, 1]Match wave size exactly — one wave, no cross-wave sync.

Common Patterns

Pattern 1: Parallel Reduction (sum / max / min)

// HLSL — parallel reduction using wave intrinsics (DXC SM6.0+)
// Step 1: each thread reduces its element
// Step 2: wave intrinsic reduces across the wave in one instruction
// Step 3: one thread per wave writes to groupshared
// Step 4: first wave reduces groupshared values

RWBuffer<float> _Input;
RWBuffer<float> _Output;
groupshared float gs_waveResults[32]; // max 32 waves per group

[numthreads(256, 1, 1)]
void CSReduce(uint tid : SV_DispatchThreadID,
              uint lid : SV_GroupIndex,
              uint3 gid : SV_GroupID)
{
    float val = (tid < _Count) ? _Input[tid] : 0.0f;

    // Wave-level reduction — one hardware instruction on NVIDIA/AMD
    val = WaveActiveSum(val); // sum across all active lanes in wave

    // One thread per wave writes to groupshared
    if (WaveIsFirstLane())
        gs_waveResults[lid / WaveGetLaneCount()] = val;

    GroupMemoryBarrierWithGroupSync(); // wait for all waves

    // First wave reduces the per-wave results
    if (lid < WaveGetLaneCount()) {
        float waveVal = (lid < ceil(256.0/WaveGetLaneCount()))
                        ? gs_waveResults[lid] : 0.0f;
        waveVal = WaveActiveSum(waveVal);
        if (lid == 0) _Output[gid.x] = waveVal;
    }
}

Pattern 2: Tiled Algorithm (screen-space lighting)

// HLSL — Tiled Forward lighting: cull lights per tile in compute
// Classic approach: 16×16 tile = 256 threads = one dispatch per frame

#define TILE_SIZE 16
#define MAX_LIGHTS_PER_TILE 256

groupshared uint  gs_minDepth;
groupshared uint  gs_maxDepth;
groupshared uint  gs_lightCount;
groupshared uint  gs_lightIndices[MAX_LIGHTS_PER_TILE];

[numthreads(TILE_SIZE, TILE_SIZE, 1)]
void CSLightCull(
    uint2 gid  : SV_GroupID,
    uint2 tid  : SV_DispatchThreadID,
    uint  lid  : SV_GroupIndex)
{
    // Step 1: init tile depth range in groupshared
    if (lid == 0) {
        gs_minDepth  = 0x7F7FFFFF; // max float as uint
        gs_maxDepth  = 0;
        gs_lightCount = 0;
    }
    GroupMemoryBarrierWithGroupSync();

    // Step 2: each thread contributes its pixel's depth
    float depth = _DepthBuffer[tid].r;
    uint  uDepth = asuint(depth);
    InterlockedMin(gs_minDepth, uDepth);
    InterlockedMax(gs_maxDepth, uDepth);
    GroupMemoryBarrierWithGroupSync();

    // Step 3: distribute light culling across threads
    // Each thread tests one light against tile frustum
    for (uint i = lid; i < _LightCount; i += TILE_SIZE * TILE_SIZE)
    {
        if (LightIntersectsTile(_Lights[i], gid, gs_minDepth, gs_maxDepth))
        {
            uint idx;
            InterlockedAdd(gs_lightCount, 1, idx);
            if (idx < MAX_LIGHTS_PER_TILE)
                gs_lightIndices[idx] = i;
        }
    }
    GroupMemoryBarrierWithGroupSync();

    // Step 4: write tile light list to global buffer
    uint tileIdx = gid.y * _TilesX + gid.x;
    _TileLightCounts[tileIdx] = gs_lightCount;
    for (uint j = lid; j < gs_lightCount; j += TILE_SIZE * TILE_SIZE)
        _TileLightIndices[tileIdx * MAX_LIGHTS_PER_TILE + j] = gs_lightIndices[j];
}

Key compute rules: always synchronise shared memory writes with GroupMemoryBarrierWithGroupSync() before reading in another thread. Use wave intrinsics (WaveActiveSum, WaveActiveBallot) instead of groupshared when operating within a single wave — they're hardware instructions with no sync cost. On mobile, prefer smaller workgroup sizes (64–128) to stay within tile SRAM budget.

// FROM_FILE_TO_PIXEL // PART 04 OF 05 // CHAPTERS 07–10 // EOF