From ff866f6bb6fe5bfae27f4ad3a5121348af3ea0b9 Mon Sep 17 00:00:00 2001
From: "PoYen, Chen" <PoYen.Chen@amd.com>
Date: Wed, 12 Jun 2024 02:52:07 +0000
Subject: [PATCH] Support providing invalid element for tensor view

---
 include/ck_tile/core/tensor/tensor_view.hpp | 30 +++++++++++++++++++++
 1 file changed, 30 insertions(+)

diff --git a/include/ck_tile/core/tensor/tensor_view.hpp b/include/ck_tile/core/tensor/tensor_view.hpp
index 656309532e..137180d33a 100644
--- a/include/ck_tile/core/tensor/tensor_view.hpp
+++ b/include/ck_tile/core/tensor/tensor_view.hpp
@@ -222,6 +222,36 @@ make_naive_tensor_view(DataType* p,
     return tensor_view<decltype(buffer_view), decltype(desc), DstInMemOp>{buffer_view, desc};
 }
 
+template <address_space_enum BufferAddressSpace = address_space_enum::generic,
+          memory_operation_enum DstInMemOp      = memory_operation_enum::set,
+          typename DataType,
+          typename... Lengths,
+          typename... Strides,
+          typename X,
+          index_t GuaranteedLastDimensionVectorLength = -1,
+          index_t GuaranteedLastDimensionVectorStride = -1,
+          typename std::enable_if<sizeof...(Lengths) == sizeof...(Strides) &&
+                                      std::is_same_v<remove_cvref_t<DataType>, remove_cvref_t<X>>,
+                                  bool>::type         = false>
+CK_TILE_HOST_DEVICE constexpr auto
+make_naive_tensor_view(DataType* p,
+                       const tuple<Lengths...>& lengths,
+                       const tuple<Strides...>& strides,
+                       X invalid_element_value,
+                       number<GuaranteedLastDimensionVectorLength> = number<-1>{},
+                       number<GuaranteedLastDimensionVectorStride> = number<-1>{})
+{
+    auto desc = make_naive_tensor_descriptor(lengths,
+                                             strides,
+                                             number<GuaranteedLastDimensionVectorLength>{},
+                                             number<GuaranteedLastDimensionVectorStride>{});
+
+    auto buffer_view = make_buffer_view<BufferAddressSpace>(
+        p, desc.get_element_space_size(), invalid_element_value);
+
+    return tensor_view<decltype(buffer_view), decltype(desc), DstInMemOp>{buffer_view, desc};
+}
+
 template <address_space_enum BufferAddressSpace = address_space_enum::generic,
           typename DataType,
           typename... Lengths,