CFT builds add about 2.6 µs to every thread-scope nvshmem_quiet() under TMA, even with logical endpoints off #138

With NVSHMEM_CFT_HANDLES_SUPPORT=ON, shmem_put_tma_bw is noticeably slower under NVSHMEM_TMA_POLICY=ENABLE than the same perftest built without CFT. The difference is a constant cost per iteration, about 2.6 µs, not a cost per put. Logical endpoints are not enabled in these runs (NVSHMEM_ENABLE_LOGICAL_ENDPOINT unset, POSIX FD heap handles), so no handle operation can be pending.

Measurements

B300, 2 PEs on one node over NVSwitch, devel 2dca770, library and perftests built for sm_103 with CUDA 13.3. The two builds differ only in -DNVSHMEM_CFT_HANDLES_SUPPORT=ON. Command: shmem_put_tma_bw -n 200 -w 20 -c 1 -t 128 -b 4096 -e 8388608, median of 5 runs, NVSHMEM_PERFTEST_VERIFY=1 (all 30 runs passed). GB/s:

FORCE matches ENABLE in both builds, and DISABLE is unaffected by CFT. Converted to time per iteration, CFT adds 2.58 µs at 4 KiB and 3.16 µs at 8 MiB, i.e. roughly the same amount whether an iteration issues 1 or 128 puts. On two RTX PRO 6000 (sm_120, PCIe) the same comparison showed CFT ENABLE about 2.4× slower at 4–32 KiB (single run, no verification).

Nsight Compute, one launch of bw_smem_tma at 4 KiB (1 CTA × 128 threads, 200 iterations), rank 0, one metric per run:

Where it comes from

The perftest calls nvshmem_quiet() from all 128 threads once per iteration. In a CFT build, the thread-scope quiet first runs the handle quiet (nvshmemi_common_device.cuh:1191-1194), which for NVSHMEMI_THREADGROUP_THREAD is nvshmemi_handle_quiet_all() (:752-756 → :693-706). Every calling thread looks up the TMA shared-memory registration (dependent volatile global loads) and then reads pending_handle_active of all 32 handle-barrier slots through a generic pointer. That matches the ~32 extra generic loads and ~3 extra global loads per warp per iteration above. The only gate is a valid TMA registration, which is why DISABLE is unaffected. Handle operations themselves are only issued when nvshmemi_is_le_implemented() holds, which requires a valid logical-endpoint id (logical_endpoint_device.cuh:193-206; callers at nvshmemi_common_device.cuh:1369, :1410, :1430). With logical endpoints off, no slot can have a pending handle, but the scan still runs.

Questions

• Could the scan be skipped when no handle operation can be pending, for example by checking the logical-endpoint state in the device state, or a per-registration flag that the handle put sets?
• Since thread-scope quiet must cover every warp, would scanning once per warp (one lane, then a shuffle) be acceptable instead of once per thread?

Separately, the perftest fills shared memory from offset 0 after nvshmemx_give_smem(), so it overwrites the region that nvshmemx_give_smem() reserves for mbarriers and handle-barrier slots (NVSHMEMI_SMEM_DATA_REGION_OFFSET). It happens to leave pending_handle_active at 0 because the stored int values are thread ids below 65536. Should the perftest use the data region only?