Hi! I'm porting Cascaded Shadow Maps(CSM) from HLSL Direct3D12/Vulkan I'm using GPU Culling, indirect Call + Vertex Amplification for shadow Casters for creating Shadow map texture. Shadow map texture - Texture2DArray
Shadow castes are rendering using Instancing and IndirectDraw, VS shader are checking [amplification_id]] and view mask
CS shader are finding shadow caster using InstanceId and view mask for each shadow cascade:
// dataIndex - Instance index for featching Instance Matrix
// viewMask - for VS [[amplification_id]]
uint PackInstanceData(uint viewMask, uint dataIndex) {
return (viewMask << 28u) | dataIndex;
}
kernel void CullObjects(uint2 DTid [[thread_position_in_grid]], constant ArgBuf& argBuf [[buffer(0)]], uint groupIndex [[thread_index_in_threadgroup]]) {
threadgroup atomic_uint sharedNumObjects;
threadgroup uint sharedInstanceId[32];
threadgroup uint globalBufferIndex;
constant auto& cbData = *argBuf.cbData;
uint index = DTid.x;
if (!groupIndex) {
atomic_store_explicit(&sharedNumObjects, 0u, memory_order_relaxed);
globalBufferIndex = 0u;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
bool visible = false;
uint viewMask = 0;
if (index < cbData.numInstanceObjects) {
const device CullingObjInfo& cullingObjectInfo = argBuf.cullingObjInfo[index];
// find shadow casters for each shadow cascade
for (uint i = 0u; i < NUM_FRUSTUMS; ++i) {
constant Frustum& frustum = cbData.frustums[i];
visible = SphereInFrustum(frustum, cullingObjectInfo.boundingSphere);
if (visible) {
visible = BoundingBoxInFrustum(frustum, cullingObjectInfo.boundingBox.minPoint, cullingObjectInfo.boundingBox.maxPoint);
}
if (visible) {
viewMask |= (1u << i);
}
}
visible = bool(viewMask);
}
}
if (visible) {
uint bufferIndex = atomic_fetch_add_explicit(&sharedNumObjects, 1u, memory_order_relaxed);
sharedInstanceId[bufferIndex] = PackInstanceData(viewMask, index);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
uint localNumObjects = atomic_load_explicit(&sharedNumObjects, memory_order_relaxed);
if (!groupIndex && localNumObjects) {
globalBufferIndex = atomic_fetch_add_explicit(&argBuf.numObjects->instanceCount, localNumObjects, memory_order_relaxed);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
if (groupIndex < localNumObjects) {
argBuf.instanceId[globalBufferIndex + groupIndex] = sharedInstanceId[groupIndex];
}
}
VS shader for depth shadow map rendering:
//
struct PSOutput {
float4 pos [[position]];
float2 Depth [[user(locn0)]];
float CullDist [[clip_distance]];
uint layer [[render_target_array_index]];
};
struct VS {
float4 pos [[attribute(0)]];
uint InstanceId [[attribute(9)]];
};
struct b {
float4x4 lightViewProj[NUM_SM_CASCADES];
};
struct Arg {
const device float3x4* InstanceData [[id(0)]];
constant b* cb [[id(1)]];
};
uint UnpackViewInstancingMask(uint packedData) {
return packedData >> 28u;
}
uint UnpackInstanceIndex(uint packedData) {
return packedData & 0x0FFFFFFFu;
}
vertex PSOutput vsMain(VS In [[stage_in]], constant Arg& arg [[buffer(VSArgumentBufferIndex)]], ushort ampID [[amplification_id]]) {
constant auto& b0 = *arg.cb;
PSOutput output;
float4 WorldPos = float4(In.pos.xyz, 1.0f);
uint viewMask = UnpackViewInstancingMask(In.InstanceId);
output.layer = ampID;
// check shadow caster rendering for each cascade using ampID and viewMask
if (viewMask & (1u << ampID)) {
float3x4 instanceMatrix = arg.InstanceData[UnpackInstanceIndex(In.InstanceId)];
WorldPos = float4((WorldPos * instanceMatrix).xyz, 1.0f);
output.CullDist = 1.0f;
}
else {
output.pos = float4(0.0f, 0.0f, 2.0f, 1.0f);
// skeep unvisible object for this cascade
output.CullDist = -1.0f;
return output;
}
output.pos = b0.lightViewProj[ampID] * WorldPos;
output.Depth.xy = output.pos.zz;
return output;
}
I have result left part: CPU Version - separate rendering for each layer for texture2Darray shadow map right part: GPU culling + Indirect + Vertex Amplification
Got to fix error ? Why 1 and 3 cascade are empty? 0 cascade has duplicate data from other cascades.
Thanks for the side-by-side image. It made the pattern easy to match. On the right, layer 2 shows what the CPU version draws for cascade 1, and layer 0 shows cascades 0, 2, and 3 together. In my testing, setting the layer in only one place, either your shader or the view mapping offsets, puts each cascade in its own layer.
That pattern appears when both your shader and the view mappings supply the layer. The renderTargetArrayIndexOffset documentation says: "If you are using vertex amplification, Metal calculates the index for each amplified vertex by adding the index offset for the vertex to the value returned from your shader." Your vertex shader writes output.layer = ampID. If your setVertexAmplificationCount(_:viewMappings:) call also passes view mappings with offsets 0, 1, 2, and 3, each cascade lands at twice its index. Cascades 0 through 3 go to layers 0, 2, 4, and 6.
I reproduced this with a small test on a Mac with M5 Max running macOS 27.0.1. It drew with an amplification count of 4 into a four-layer array texture, with offsets 0 through 3 and the shader writing ampID. Each layer received these amplification IDs:
- Layer 0: 0, 2, and 3
- Layer 1: none
- Layer 2: 1
- Layer 3: none
The out-of-range layers 4 and 6 ended up in layer 0, which matches the extra content in your cascade 0. The Metal Shading Language Specification describes this: if [[render_target_array_index]] "is out-of-bounds (that is, greater than or equal to renderTargetArrayLength), the hardware interprets this value as 0."
Either of these changes gave each layer exactly one amplification ID in the same test:
- Keep the view mapping offsets, and write
output.layer = 0in the vertex shader. - Keep
output.layer = ampID, and pass view mappings with zero offsets, or passnilfor the view mappings.
If your host code doesn't pass view mapping offsets, please post your setVertexAmplificationCount call and render pass setup, and I'll take another look.
