解读

国内Unity项目做GPU并行优化时，面试官常把“Warp占用率”与“LDS冲突”放在一起问，目的是快速判断候选人是否真正在移动/PC GPU上做过Compute Shader调优。Warp（NVIDIA术语，AMD称Wavefront）是GPU调度的最小线程束，Warp级占用率直接决定SM/CU里有没有足够的活跃Warp去掩盖访存延迟；而LDS（Local Data Share，即芯片上的共享内存）冲突一旦发生，会让整个Wavefront/Warp在LDS Bank仲裁上串行化，瞬间把并行度打回串行，两者叠加就是Compute性能“雪崩”的典型场景。答不出“Bank Conflict 等于把32线程变成1线程”这一层，基本会被判定为“只写过Demo，没上过真机”。

知识点

1. Warp/Wavefront：32（NVIDIA）或64（AMD）线程硬调度单位，一个Warp内执行完全SIMT，任意一条线程阻塞＝整个Warp阻塞。
2. Warp占用率：Active Warp数 / 理论最大Warp数，受寄存器配额、LDS大小、线程块尺寸三重限制；Unity Compute Shader里可通过[numthreads(x,y,z)]、shared memory size、pragma kernel_max_registers手动干预。
3. LDS Bank结构：移动端Mali/GPU通常16 Bank，桌面RDNA 32 Bank，相邻线程访问同一Bank不同地址即冲突，延迟=冲突线程数×Bank周转周期。
4. 冲突量化：理想无冲突时一个Wavefront 1 cycle返回数据；若32线程全部撞Bank，延迟放大32×，Wavefront吞吐量直接掉到1/32，占用率再高也白搭。
5. Unity实战：Compute Shader里声明groupshared float4 cache[256]；线程tid访问cache[tid]无事，若改为cache[tid*2]且线程步长=Bank数，则必冲突，真机Profiler会出现“LDS Stall 90%+”。

答案

Warp级占用率衡量的是GPU核心里同时活跃的Warp数量，越高越能把算术或访存延迟隐藏掉；但占用率只是“上限”，不代表一定能打满算力。当Compute Shader里用到groupshared（LDS）时，若线程索引对Bank取模后大量重复，就会触发LDS Bank Conflict，导致同一Warp内线程串行排队取数，一个时钟周期只能服务一条线程，相当于把并行度降到1/Bank冲突倍数。此时即使理论占用率100%，有效IPC也会被LDS Stall拖成个位数，整体Kernel时间呈数倍到数十倍放大。优化思路分三步：

1. 降低冲突：把groupshared数组按Bank宽度做pad，例如float4 cache[256+1]，让stride跨Bank；
2. 减LDS footprint：把每条线程的共享内存从float4降到float，腾出LDS容量，让调度器能同时多排几个Wavefront；
3. 重新分块：在Unity里把[numthreads(64,1,1)]改成[numthreads(8,8,1)]，寄存器压力下降，Warp占用率回升，再辅以Bank无冲突访问，即可在Adreno 730上把原本8 ms的CS Skinning压到1.2 ms，实测GPU利用率从23%提到71%。

拓展思考

国内项目越来越把Compute Shader当“性能救命稻草”，但真机瓶颈往往不在ALU而在Memory Hierarchy。下一步面试官通常会追问：

1. “如果把LDS数据改到Global UAV，还会快吗？”——需要答出“Global原子访问跨Workgroup竞争更大，且没有L1 Bank，延迟更高，除非做Wave-level tile-based同步否则得不偿失”；
2. “iOS A系列没有LDS，怎么办？”——要提到tile memory（苹果叫threadgroup memory）同样存在Bank冲突，只是Bank数=16，且与像素Tile共享带宽，需用simdgroup_barrier+quad_shuffle做跨线程数据交换；
3. “Unity SRP Batch/DR DrawProceduralIndirect会不会受Warp占用率影响？”——答出Indirect Draw的Compute Shader本身如果LDS冲突，会导致GPU端生成命令的延迟，进而把CPU等待GPU Frame的bubble放大，表现就是“SRP Batcher开了一堆，帧率还是掉”。把这三点连成“CPU-GPU端到端管线”视角，就能在资深/专家面里拿到高分。