perf: reuse persistent atomic buffer in S4 DPAS dispatch
Replace per-call sycl::malloc_device/free of the 1-int32 work-group
counter in the prefill and decode S4 DPAS grouped-GEMM dispatchers with
a persistent per-queue buffer. Each malloc/free forces a queue sync,
which is pure overhead on the decode hot path. The kernel self-zeroes
the counter at launch and every launch is synchronous, so a single
reused buffer per queue is safe.
Co-authored-by: a32543254 <53296245+a32543254@users.noreply.github.com>