From 2fdcb55cf12468738893ef5c30c5cc0a51c2f22c Mon Sep 17 00:00:00 2001
From: joye <joye@amd.com>
Date: Thu, 15 May 2025 14:36:02 +0800
Subject: [PATCH] fix some compiling errors

---
 include/ck_tile/core/arch/amd_buffer_addressing.hpp |  3 +++
 include/ck_tile/core/tensor/load_tile.hpp           | 12 ++++--------
 include/ck_tile/core/tensor/tile_window_linear.hpp  |  2 +-
 3 files changed, 8 insertions(+), 9 deletions(-)

diff --git a/include/ck_tile/core/arch/amd_buffer_addressing.hpp b/include/ck_tile/core/arch/amd_buffer_addressing.hpp
index 50f81ecf84..69443abd1c 100644
--- a/include/ck_tile/core/arch/amd_buffer_addressing.hpp
+++ b/include/ck_tile/core/arch/amd_buffer_addressing.hpp
@@ -77,6 +77,9 @@ struct buffer_store;
 template <index_t bytes>
 struct buffer_store_if;
 
+template <index_t bytes, bool pre_nop = false>
+struct async_buffer_load;
+
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Wundefined-reinterpret-cast"
 // TODO: strict aliasing rule seems fail when reinterpret_cast between vector type
diff --git a/include/ck_tile/core/tensor/load_tile.hpp b/include/ck_tile/core/tensor/load_tile.hpp
index 7168551abf..8b7541bf23 100644
--- a/include/ck_tile/core/tensor/load_tile.hpp
+++ b/include/ck_tile/core/tensor/load_tile.hpp
@@ -92,18 +92,14 @@ CK_TILE_DEVICE auto load_tile_raw(T& tile,
 template <typename LdsTileWindow_,
           typename TileWindow_,
           index_t i_access           = -1,
-          bool oob_conditional_check = true,
-          bool pre_nop               = false>
+          bool oob_conditional_check = true>
 CK_TILE_DEVICE auto async_load_tile(LdsTileWindow_&& lds_tile,
                                     const TileWindow_& tile_window,
                                     number<i_access>                     = {},
-                                    bool_constant<oob_conditional_check> = {},
-                                    bool_constant<pre_nop>               = {})
+                                    bool_constant<oob_conditional_check> = {})
 {
-    return tile_window.async_load(lds_tile,
-                                  number<i_access>{},
-                                  bool_constant<oob_conditional_check>{},
-                                  bool_constant<pre_nop>{});
+    return tile_window.async_load(
+        lds_tile, number<i_access>{}, bool_constant<oob_conditional_check>{});
 }
 
 template <typename LdsTileWindow_,
diff --git a/include/ck_tile/core/tensor/tile_window_linear.hpp b/include/ck_tile/core/tensor/tile_window_linear.hpp
index 66f86c669d..a4669e2515 100644
--- a/include/ck_tile/core/tensor/tile_window_linear.hpp
+++ b/include/ck_tile/core/tensor/tile_window_linear.hpp
@@ -748,7 +748,7 @@ struct tile_window_linear
             auto bottom_tensor_flag         = cached_flags_[IAccess];
 
             auto lds_bottom_tensor_thread_idx =
-                lds_tile.get_window_origin() + window_adaptor_thread_coord.get_bottom_index();
+                lds_tile.get_window_origin() + window_adaptor_coord.get_bottom_index();
 
             const auto lds_coord =
                 make_tensor_coordinate(lds_tile.get_bottom_tensor_view().get_tensor_descriptor(),