[all-commits] [llvm/llvm-project] ea3f10: [OpenMP] Reserve the main thread's warp for generi...

[Larry Meadows via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: ea3f107a84eaa2459f6eb5d1c8dbeefea7576689
      https://github.com/llvm/llvm-project/commit/ea3f107a84eaa2459f6eb5d1c8dbeefea7576689
  Author: Larry Meadows <Lawrence.Meadows at amd.com>
  Date:   2026-09-09 (Wed, 09 Sep 2026)

  Changed paths:
    M clang/test/OpenMP/amdgcn-attributes.cpp
    M clang/test/OpenMP/amdgcn_target_device_vla.cpp
    M clang/test/OpenMP/nvptx_SPMD_codegen.cpp
    M clang/test/OpenMP/nvptx_target_teams_codegen.cpp
    M clang/test/OpenMP/ompx_attributes_codegen.cpp
    M clang/test/OpenMP/target_num_teams_num_threads_attributes.cpp
    M clang/test/OpenMP/target_teams_reduction_codegen.cpp
    M llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
    M mlir/test/Target/LLVMIR/omptarget-region-device-llvm.mlir
    M mlir/test/Target/LLVMIR/openmp-target-launch-device.mlir
    M offload/plugins-nextgen/common/src/RecordReplay.cpp
    A offload/test/offloading/thread_limit_below_wavefront.c

  Log Message:
  -----------
  [OpenMP] Reserve the main thread's warp for generic mode kernels (#218790)

In generic mode one warp of the block runs the main thread and the rest
run the
user's threads, so everything needing the team size computes
`BlockSize - WarpSize`: the device runtime's `getMaxTeamThreads`, and
the state
machine OpenMPOpt builds. The bound written into the kernel was the
program's
`thread_limit` exactly, so on gfx90a `thread_limit(4)` launched a block
of 4
threads and the team size came out as `4 - 64`. That subtraction is
unsigned, so
the runtime saw four billion workers rather than none: the parallel
region never
ran, and on current trunk the launch dies with a memory access fault.
Any
`thread_limit` below one warp was affected, so anything under 64 on a
wave64
target.

Add the warp where the bound is written, in `createTargetInit`, for
generic mode
only, so `thread_limit(4)` writes 68 and the team gets its 4, clamped to
the
target's maximum work group size because an out-of-range bound is
dropped rather
than clamped. A fixed 64 is reserved rather than the target's own
wavefront: the
bound is only a ceiling, and the runtime still adds the real warp size
and
clamps to it, so the block that launches is unchanged. Where the warp is
narrower the ceiling is 32 too high, which costs the backend some
register
budget and nothing else; NVPTX with `thread_limit(20)` now gets
`maxntid=84`
where it would have got 52, and still launches 52.

The record-replay recorder, which wrote `thread_limit` as the largest
launch a
replay would accept, now records the larger block that was actually
launched.
`thread_limit_below_wavefront.c` is the reproducer, and the remaining
files are
regenerated clang and MLIR expectations.

## Testing

Tested on `main` at 68622e02c9e5: `clang/test/OpenMP`,
`mlir/test/Target/LLVMIR`
and `llvm/test/Transforms/OpenMP` give 2114 passed, 14 unsupported, 4
failed.
The four are `mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared*`, which
fail on
`main` independently of this change: their RUN lines hand `llc` an
`-mxcoff-roptr` it only accepts on AIX.

`check-offload` on an MI210 has nothing failing: 1134 passed, 4
expectedly
failed, 363 unsupported, 12 skipped. That covers the new test,
`thread_limit_below_wavefront.c`, and the six `omp-kernel-replay` tests.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications