jayfoad wrote:
> Currently the default minimum waves/EU is function of the maximum flat work
> size but to me this is only a heuristic for what minimum occupancy we should
> strive to achieve (namely to fit a single workgroup of maximum size on a CU
> concurrently), not an ABI requirement.
Claude says this PR creates cases where a kernel can't be launched at all (i.e.
can't even achieve occupancy 1):
---
**Experiment.**
- I applied the PR's `AMDGPUSubtarget.{h,cpp}` changes onto `main`
(5758d8e3120a); they applied cleanly.
- The kernel: gfx942 (wave64, 4 SIMDs per CU),
`"amdgpu-flat-work-group-size"="1,1024"`,
`"amdgpu-waves-per-eu"="1"`, and about 300 live floats across a
barrier.
| | `main` | PR |
|---|---|---|
| `.max_flat_workgroup_size` | 1024 | 1024 |
| `.vgpr_count` | 128 | **301** |
| Occupancy | 4 | **1** |
| ScratchSize | 696 | 0 |
- A 1024-lane workgroup is 16 waves, and **all waves of a workgroup must
be resident on one CU at the same time** (barriers and LDS depend on
it). That needs 4 waves per SIMD, so at most 128 VGPRs on gfx942.
- With the PR, the kernel still advertises
`max_flat_workgroup_size: 1024` but uses 301 VGPRs. **Any launch with
a workgroup larger than 256 lanes cannot be scheduled**, and the
runtime rejects it or hangs.
- `main` correctly spills instead.
https://github.com/llvm/llvm-project/pull/138284
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
