[clr] Out-of-bounds kernarg write in single-grid cooperative launch

[neoblizz]

Problem Description

The single-grid cooperative-groups launch path in CLR under-allocates the kernarg
buffer that backs a Device::MGSyncInfo, then writes the full struct into it. The
write overruns the reservation by ~40 bytes into the shared kernarg pool's free
region.

In projects/clr/rocclr/device/device.hpp:

static constexpr size_t kMGSyncDataSize = sizeof(MGSyncData);   // 8
static constexpr size_t kSGInfoSize     = kMGSyncDataSize;      // 8  <-- wrong

MGSyncInfo is 48 bytes (pointer 8 + grid_id 4 + num_grids 4 + prev_sum 8 +
all_sum 8 + sgs 8 + num_wg 4 + 4 tail padding), not 8. Note the multi-grid
constant right above it gets this right:
kMGInfoSizePerDevice = kMGSyncDataSize + sizeof(MGSyncInfo).

In projects/clr/rocclr/device/rocm/rocvirtual.cpp (HiddenMultiGridSync,
~line 4655) the singleGridSync branch allocates with the bad constant and then
populates every field:

syncInfo = reinterpret_cast<Device::MGSyncInfo*>(allocKernArg(Device::kSGInfoSize, 64));
syncInfo->mgs = nullptr;
...
syncInfo->sgs = {0};
syncInfo->grid_id   = vcmd->gridId();
syncInfo->num_grids = vcmd->numGrids();
syncInfo->prev_sum  = vcmd->prevGridSum();
syncInfo->all_sum   = vcmd->allGridSum();
syncInfo->num_wg    = vcmd->numWorkgroups();

allocKernArg -> ManagedBuffer::Acquire is a bump allocator over a single
pre-allocated pool and advances its cursor by exactly the requested size, so the trailing bytes land in memory the allocator still considers free.

Impact

Latent memory corruption, currently masked by allocator slack:

• The next kernarg allocation from the same pool can be handed memory that the
  sync-info write already clobbered, or can itself overwrite the barrier fields
  before the cooperative kernel reads them.
• Near a chunk boundary the overrun can straddle the flush/wrap point, silently
  dropping prev_sum / all_sum / num_wg.
• Whether it manifests depends on allocation order, alignment, and pool size —
  so it presents as rare, non-deterministic hangs or wrong results in
  this_grid().sync() on a single device, not a consistent failure.

Expected Behavior

kSGInfoSize reserves at least sizeof(MGSyncInfo) bytes, so the write stays
inside its own allocation.
so it presents as rare, non-deterministic hangs or wrong results in
this_grid().sync() on a single device, not a consistent failure.

Fixed by #8748.

Operating System

Ubuntu 24.04

CPU

Any

GPU

MI300X

ROCm Version

ROCm 7

ROCm Component

clr

Steps to Reproduce

1. Launch a single-device cooperative kernel via hipLaunchCooperativeKernel
   that calls cooperative_groups::this_grid().sync().
2. Precede/follow it with other kernel launches so the kernarg bump allocator's
   cursor lands unfavorably. The corruption is allocation-order dependent and
   will not reproduce on every run.
3. Under ASAN/valgrind on the host-visible kernarg pool, or by instrumenting
   ManagedBuffer::Acquire to poison the free region, the 8-byte reservation
   followed by a 48-byte write is directly observable.

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

rocminfo --support output

Paste output here

Additional Information

No response