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.
Topic:
Graphics & Games
SubTopic:
Metal
2
0
740