Fix correctness of streamk by adding a required barrier (#745)

Author: AlexAUT

python/perf-kernels/streamk/streamk_kernel.py

@@ -162,6 +162,7 @@ def streamk_gemm(
             rn1 = tl.max_contiguous(tl.multiple_of(rn1, BLOCK_SIZE_N), BLOCK_SIZE_N)
             P_ = P + pid * BLOCK_SIZE_M * BLOCK_SIZE_N + rm1[:, None] * BLOCK_SIZE_N + rn1[None, :]
             tl.store(P_, acc, cache_modifier=".wt")
+            tl.debug_barrier()
             tl.store(locks + pid, 1, cache_modifier=".wt")
         #   tl.store(P_, acc)
         #   tl.debug_barrier()