[https://nvbugs/5970614][fix] Sync CTA before PDL trigger in quantize_with_block_size

Signed-off-by: Tianyu Xiong <117647511+tianyuxbear@users.noreply.github.com>

Author: tianyuxbear

cpp/tensorrt_llm/kernels/quantization.cuh

@@ -897,6 +897,15 @@ quantize_with_block_size(
             }
         }
     }
+    // Fix for nvbugs/5970614 (https://nvbugspro.nvidia.com/bug/5970614).
+    // PDL completion is reported when every CTA has either exited or called
+    // this function at least once (per CUDA Programming Guide). Without a
+    // CTA-wide barrier, an early-finishing warp can trigger completion while
+    // other warps in the same CTA are still writing sf_out / out, allowing the
+    // downstream NVF4 GEMM consumer to read partial data once
+    // wait_on_dependent_grids returns. Drain the CTA's stores before trigger.
+    __syncthreads();
+    __threadfence();
     cudaTriggerProgrammaticLaunchCompletion();
 #endif
 }