From 594c4e9fa72bcc7af062bf4e5b74406054321bc0 Mon Sep 17 00:00:00 2001
From: felix <felix.li@amd.com>
Date: Tue, 28 Oct 2025 14:37:24 +0000
Subject: [PATCH] set nt load

---
 .../ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp | 4 ++--
 include/ck_tile/ops/rmsnorm2d/kernel/rmsnorm2d_fwd_kernel.hpp | 4 ++--
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp b/include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp
index 6998b358d8..ee53b177d9 100644
--- a/include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp
+++ b/include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp
@@ -194,7 +194,7 @@ struct Layernorm2dFwd
         const auto iM = get_block_id() * Block_M;
 
         const auto x_window = [&]() {
-            const auto tmp_ = make_naive_tensor_view<address_space_enum::global>(
+            const auto tmp_ = make_naive_tensor_view<address_space_enum::global, memory_operation_enum::set, amd_buffer_coherence_enum::slc>(
                 static_cast<const XDataType*>(kargs.p_x),
                 make_tuple(kargs.m, kargs.n),
                 make_tuple(kargs.x_stride, 1),
@@ -283,7 +283,7 @@ struct Layernorm2dFwd
         }();
 
         auto y_window = [&]() {
-            auto tmp_ = make_naive_tensor_view<address_space_enum::global>(
+            auto tmp_ = make_naive_tensor_view<address_space_enum::global,  memory_operation_enum::set, amd_buffer_coherence_enum::slc>(
                 static_cast<YDataType*>(kargs.p_y),
                 make_tuple(kargs.m, kargs.n),
                 make_tuple(kargs.y_stride, 1),
diff --git a/include/ck_tile/ops/rmsnorm2d/kernel/rmsnorm2d_fwd_kernel.hpp b/include/ck_tile/ops/rmsnorm2d/kernel/rmsnorm2d_fwd_kernel.hpp
index e7f4ce0ba8..fd6e72e2a5 100644
--- a/include/ck_tile/ops/rmsnorm2d/kernel/rmsnorm2d_fwd_kernel.hpp
+++ b/include/ck_tile/ops/rmsnorm2d/kernel/rmsnorm2d_fwd_kernel.hpp
@@ -202,7 +202,7 @@ struct Rmsnorm2dFwd
             if constexpr(kFusedAdd == Rmsnorm2dFusedAddEnum::PRE_ADD ||
                          kFusedAdd == Rmsnorm2dFusedAddEnum::PRE_ADD_STORE)
             {
-                const auto tmp_ = make_naive_tensor_view<address_space_enum::global>(
+                const auto tmp_ = make_naive_tensor_view<address_space_enum::global, memory_operation_enum::set, amd_buffer_coherence_enum::slc>(
                     static_cast<const XResidualDataType*>(kargs.p_x_residual),
                     make_tuple(kargs.m, kargs.n),
                     make_tuple(kargs.xr_stride, 1),
@@ -236,7 +236,7 @@ struct Rmsnorm2dFwd
         }();
 
         auto y_window = [&]() {
-            auto tmp_ = make_naive_tensor_view<address_space_enum::global>(
+            auto tmp_ = make_naive_tensor_view<address_space_enum::global, memory_operation_enum::set, amd_buffer_coherence_enum::slc>(
                 static_cast<YDataType*>(kargs.p_y),
                 make_tuple(kargs.m, kargs.n),
                 make_tuple(kargs.y_stride, 1),