From c8a3df5820f7f0f6a2ef7215368294c766f7ca1b Mon Sep 17 00:00:00 2001 From: Giulio Eulisse <10544+ktf@users.noreply.github.com> Date: Sun, 15 Nov 2026 09:37:49 +0100 Subject: [PATCH] GPU Metal: fence only threadgroup memory in the work-group collectives The Metal work-group scan, reduction and broadcast exchange data through threadgroup memory only, so their barriers need not order device memory. GPUbarrier() fences both, which on Apple GPUs makes each barrier wait for the outstanding device-memory accesses. With threadgroup-only fences the stream-compaction scanStart takes 12% less GPU time (30 kHz Pb-Pb, 1/4 TF, no --debug, 3 interleaved rounds); results unchanged. --- GPU/Common/GPUCommonAlgorithm.h | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/GPU/Common/GPUCommonAlgorithm.h b/GPU/Common/GPUCommonAlgorithm.h index fddfdd2f889da..477bc3c34591b 100644 --- a/GPU/Common/GPUCommonAlgorithm.h +++ b/GPU/Common/GPUCommonAlgorithm.h @@ -508,7 +508,9 @@ GPUdi() T warp_reduce_max(T v) // consecutive threads, of GPUCA_WARP_SIZE each. The SIMD scans take 32-bit // lanes; the narrower types are counts and are widened. Sums are returned // promoted, as v + 0 would be, since the callers pass !!pred: a bool would -// truncate them to 0 / 1. +// truncate them to 0 / 1. They only exchange data through smem, so their +// barriers fence threadgroup memory only: a device-memory fence, as +// GPUbarrier() has, costs the stream compaction ~10%. template GPUdi() auto warp_scan_inclusive_add(T v) -> decltype(+v) @@ -567,18 +569,18 @@ GPUdi() auto work_group_scan_inclusive_add_FUNC(T v, S& smem, int32_t iThread, i if (iThread % GPUCA_WARP_SIZE == GPUCA_WARP_SIZE - 1 || iThread == nThreads - 1) { smem.warpPredicateSum[iWarp] = x; } - GPUbarrier(); + threadgroup_barrier(mem_flags::mem_threadgroup); if (iWarp == 0) { const int32_t s = simd_prefix_inclusive_sum(iThread < nWarps ? smem.warpPredicateSum[iThread] : 0); if (iThread < nWarps) { smem.warpPredicateSum[iThread] = s; } } - GPUbarrier(); + threadgroup_barrier(mem_flags::mem_threadgroup); if (iWarp > 0) { x += smem.warpPredicateSum[iWarp - 1]; } - GPUbarrier(); + threadgroup_barrier(mem_flags::mem_threadgroup); return x; } @@ -589,9 +591,9 @@ GPUdi() T work_group_broadcast_FUNC(T v, int32_t i, S& smem, int32_t iThread) if (iThread == i) { smem.tmpBroadcast = v; } - GPUbarrier(); + threadgroup_barrier(mem_flags::mem_threadgroup); T retVal = smem.tmpBroadcast; - GPUbarrier(); + threadgroup_barrier(mem_flags::mem_threadgroup); return retVal; }