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.
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.
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.
| PRAGMA | COMPILED WHEN | STRIPPED WHEN | USE 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
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.
| PATTERN | API | CPU COST | DYNAMIC DATA | BEST FOR |
|---|---|---|---|---|
| SRP Batcher | Automatic (CBUFFER) | Very low | Per-material only | Static/semi-static scene objects with varying materials |
| GPU Instancing | UNITY_INSTANCING_* macros | Low | Per-instance via instanced properties | Medium counts (<1K), per-instance color/scale |
| DrawMeshInstanced | C# API, managed array | Medium (GC risk) | Yes — upload array per frame | Dynamic objects, runtime spawned |
| DrawMeshInstancedIndirect | ComputeBuffer args | Near-zero | GPU-driven (compute writes args) | GPU culling pipelines, >10K instances |
| DOTS / Entities Graphics | ECS + BRG | Minimal (Burst) | Yes — NativeArray/chunk data | Massive counts, physics-driven, procedural |
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
// 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
}
// 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.
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.
| CONCEPT | HLSL (DX12/Unity) | GLSL (Vulkan/GL) | MSL (Metal/Apple) |
|---|---|---|---|
| Float 2/3/4 vector | float2/3/4 | vec2/3/4 | float2/3/4 |
| Int vector | int2/3/4 | ivec2/3/4 | int2/3/4 |
| Matrix 4×4 | float4x4 | mat4 | float4x4 |
| Matrix multiply | mul(mat, vec) | mat * vec | mat * vec |
| Texture 2D type | Texture2D<T> | sampler2D / texture2D | texture2d<T> |
| Sample texture | tex.Sample(samp, uv) | texture(sampler2D, uv) | tex.sample(samp, uv) |
| Sample explicit LOD | tex.SampleLevel(s,uv,lod) | textureLod(s,uv,lod) | tex.sample(s,uv,level(lod)) |
| Vertex position out | SV_POSITION semantic | gl_Position built-in | [[position]] attribute |
| Fragment color out | : SV_Target | layout(location=0) out | [[color(0)]] return |
| Instance ID | SV_InstanceID | gl_InstanceIndex | [[instance_id]] |
| Vertex ID | SV_VertexID | gl_VertexIndex | [[vertex_id]] |
| Clip position | SV_POSITION | gl_Position | [[position]] |
| Saturate (clamp 0–1) | saturate(x) | clamp(x,0.0,1.0) | saturate(x) |
| Fused multiply-add | mad(a,b,c) | fma(a,b,c) | fma(a,b,c) |
| Partial derivative ∂/∂x | ddx(x) | dFdx(x) | dfdx(x) |
| Partial derivative ∂/∂y | ddy(x) | dFdy(x) | dfdy(x) |
| Atomics (shared mem) | InterlockedAdd(dst, val) | atomicAdd(dst, val) | atomic_fetch_add_explicit() |
| Group memory barrier | GroupMemoryBarrierWithGroupSync() | barrier() + memoryBarrierShared() | threadgroup_barrier(mem_flags::mem_threadgroup) |
| Discard fragment | discard / clip(val) | discard | discard_fragment() |
| Constant buffer | cbuffer / CBUFFER_START | layout(std140) uniform Block {} | constant T& [[buffer(N)]] |
| Compute thread ID | SV_DispatchThreadID | gl_GlobalInvocationID | [[thread_position_in_grid]] |
| Group thread ID | SV_GroupThreadID | gl_LocalInvocationID | [[thread_position_in_threadgroup]] |
| Group shared memory | groupshared T var[] | shared T var[] | threadgroup T var[] |
| PRECISION | HLSL | GLSL / Vulkan | MSL | MOBILE IMPACT |
|---|---|---|---|---|
| 32-bit float | float | highp float | float | Full cost — use for positions, world-space |
| 16-bit float | half / min16float | mediump float | half | 2× faster ALU on Adreno/Mali — use for colors, normals, UVs |
| 10-bit float | min10float (DX only) | lowp float | — | Rare; useful for 0–1 values on very low-end |
| 16-bit int | min16int | mediump int | short | Fast 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.
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.
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.
| CONCEPT | HLSL | GLSL | MSL | HARDWARE |
|---|---|---|---|---|
| Thread group size | [numthreads(X,Y,Z)] | layout(local_size_x=X,...) | kernel void f([[ ... ]]) | Must be ≤ 1024 total threads |
| Global thread ID | SV_DispatchThreadID | gl_GlobalInvocationID | [[thread_position_in_grid]] | gid = groupID * groupSize + localID |
| Local thread ID | SV_GroupThreadID | gl_LocalInvocationID | [[thread_position_in_threadgroup]] | 0..numthreads-1 per group |
| Group index | SV_GroupIndex | gl_LocalInvocationIndex | [[thread_index_in_threadgroup]] | Flattened: z*XY + y*X + x |
| Shared memory | groupshared T[] | shared T[] | threadgroup T[] | On-chip SRAM — 32–100 KB |
| Sync barrier | GroupMemoryBarrierWithGroupSync() | barrier() | threadgroup_barrier() | All threads in group reach barrier |
| WORKLOAD | RECOMMENDED SIZE | REASON |
|---|---|---|
| 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. |
// 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;
}
}
// 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.