From 237363809d5fc97f717eff3df8fc8cb33b5d17f0 Mon Sep 17 00:00:00 2001
From: zanzhang <zanzhang@amd.com>
Date: Wed, 29 Oct 2025 20:35:25 +0800
Subject: [PATCH] update coherence

---
 .../ck_tile/core/arch/amd_buffer_addressing.hpp  | 13 +++++++++++++
 .../core/arch/amd_buffer_addressing_builtins.hpp | 11 +++++++++++
 .../ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp  | 16 ++++++++++++----
 3 files changed, 36 insertions(+), 4 deletions(-)

diff --git a/include/ck_tile/core/arch/amd_buffer_addressing.hpp b/include/ck_tile/core/arch/amd_buffer_addressing.hpp
index 6e0fb4110d..f8fc1f9a16 100644
--- a/include/ck_tile/core/arch/amd_buffer_addressing.hpp
+++ b/include/ck_tile/core/arch/amd_buffer_addressing.hpp
@@ -1326,6 +1326,19 @@ enum struct amd_buffer_coherence_enum
     glc               = 1,
     slc               = 2,
     glc_slc           = 3,
+
+    // s[1:0] System Cache Level: 0=warp, 1=group, 2=device, 3=system
+    // bit0 = sc0, bit1 = nt, bit2 = swz?, bit4 = sc1
+    //
+
+    WAVE_NT0 = 0,
+    WAVE_NT1 = 2,
+    GROUP_NT0 = 1,
+    GROUP_NT1 = 3,
+    DEVICE_NT0 = 8,
+    DEVICE_NT1 = 10,
+    SYSTEM_NT0 = 9,
+    SYSTEM_NT1 = 11,
 };
 
 template <index_t N,
diff --git a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp
index bf95111d78..a1eaa14774 100644
--- a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp
+++ b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp
@@ -1186,6 +1186,17 @@ enum struct amd_buffer_coherence_enum
     glc               = 1,
     slc               = 2,
     glc_slc           = 3,
+    // s[1:0] System Cache Level: 0=warp, 1=group, 2=device, 3=system
+    // bit0 = sc0, bit1 = nt, bit2 = swz?, bit4 = sc1
+    //
+    WAVE_NT0 = 0,
+    WAVE_NT1 = 2,
+    GROUP_NT0 = 1,
+    GROUP_NT1 = 3,
+    DEVICE_NT0 = 8,
+    DEVICE_NT1 = 10,
+    SYSTEM_NT0 = 9,
+    SYSTEM_NT1 = 11,
 };
 
 template <index_t N,
diff --git a/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp b/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp
index ede07e221c..98517e66b1 100755
--- a/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp
+++ b/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp
@@ -569,7 +569,9 @@ struct FlatmmKernel
         const auto& a_tensor_view = [&]() {
             if constexpr(std::is_same_v<ALayout, tensor_layout::gemm::RowMajor>)
             {
-                return make_naive_tensor_view<address_space_enum::global>(
+                return make_naive_tensor_view<address_space_enum::global,
+                                              memory_operation_enum::set,
+                                              amd_buffer_coherence_enum::SYSTEM_NT1>(
                     a_ptr,
                     make_tuple(kargs.M, splitk_batch_offset.splitted_k),
                     make_tuple(kargs.stride_A, 1),
@@ -578,7 +580,9 @@ struct FlatmmKernel
             }
             else
             {
-                return make_naive_tensor_view<address_space_enum::global>(
+                return make_naive_tensor_view<address_space_enum::global,
+                                              memory_operation_enum::set,
+                                              amd_buffer_coherence_enum::SYSTEM_NT1>(
                     a_ptr,
                     make_tuple(splitk_batch_offset.splitted_k, kargs.M),
                     make_tuple(kargs.stride_A, 1),
@@ -628,7 +632,9 @@ struct FlatmmKernel
         const auto& e_tensor_view = [&]() {
             if constexpr(std::is_same_v<ELayout, tensor_layout::gemm::RowMajor>)
             {
-                return make_naive_tensor_view<address_space_enum::global, DstInMemOp>(
+                return make_naive_tensor_view<address_space_enum::global,
+                                              DstInMemOp,
+                                              amd_buffer_coherence_enum::SYSTEM_NT1>(
                     e_ptr,
                     make_tuple(kargs.M, kargs.N),
                     make_tuple(kargs.stride_E, 1),
@@ -637,7 +643,9 @@ struct FlatmmKernel
             }
             else
             {
-                return make_naive_tensor_view<address_space_enum::global, DstInMemOp>(
+                return make_naive_tensor_view<address_space_enum::global,
+                                              DstInMemOp,
+                                              amd_buffer_coherence_enum::SYSTEM_NT1>(
                     e_ptr,
                     make_tuple(kargs.N, kargs.M),
                     make_tuple(kargs.stride_E, 1),