From d7579321c5ce49986662ef67157611750f6d009e Mon Sep 17 00:00:00 2001
From: Muhammed Emin Ozturk <meozturk@t008-002.hpcfund>
Date: Wed, 26 Feb 2025 19:31:38 -0600
Subject: [PATCH] bf16 compilation error is fixed

---
 ...hreadwise_tensor_slice_transfer_v6r1r2.hpp |  3 ++
 include/ck/utility/dynamic_buffer.hpp         | 53 +++++++++++++++----
 2 files changed, 46 insertions(+), 10 deletions(-)

diff --git a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp
index bba2eb8dde..3a096df4c8 100755
--- a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp
+++ b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1r2.hpp
@@ -189,6 +189,9 @@ struct ThreadwiseTensorSliceTransfer_v6r1r2
             }
 #endif
 
+
+
+
             // copy data from dst_vector into dst_buf
             dst_buf.template Update<DstInMemOp, dst_vector_t>(
                 dst_coord_.GetOffset(),
diff --git a/include/ck/utility/dynamic_buffer.hpp b/include/ck/utility/dynamic_buffer.hpp
index 6de17a6152..5b035337d6 100755
--- a/include/ck/utility/dynamic_buffer.hpp
+++ b/include/ck/utility/dynamic_buffer.hpp
@@ -134,7 +134,8 @@ struct DynamicBuffer
     template <InMemoryDataOperationEnum Op,
               typename X,
               typename enable_if<is_same<typename scalar_type<remove_cvref_t<X>>::type,
-                                         typename scalar_type<remove_cvref_t<T>>::type>::value,
+                                         typename scalar_type<remove_cvref_t<T>>::type>::value ||
+                                     !is_native_type<X>(),
                                  bool>::type = false>
     __host__ __device__ void Update(index_t i, bool is_valid_element, const X& x)
     {
@@ -154,14 +155,43 @@ struct DynamicBuffer
         {
             auto tmp       = this->template Get<X>(i, is_valid_element);
             using scalar_t = typename scalar_type<remove_cvref_t<T>>::type;
-            // handle bfloat addition
-            if constexpr(is_same_v<scalar_t, bhalf_t>)
+
+            // // handle bfloat addition
+            // if constexpr(is_same_v<scalar_t, bhalf_t>)
+            // {
+            //     if constexpr(is_scalar_type<X>::value)
+            //     {
+            //         // Scalar type
+            //         auto result =
+            //             type_convert<X>(type_convert<float>(x) + type_convert<float>(tmp));
+            //         this->template Set<X>(i, is_valid_element, result);
+            //     }
+            //     else
+            //     {
+            //         // Vector type
+            //         constexpr auto vector_size = scalar_type<remove_cvref_t<X>>::vector_size;
+            //         const vector_type<scalar_t, vector_size> a_vector{tmp};
+            //         const vector_type<scalar_t, vector_size> b_vector{x};
+            //         static_for<0, vector_size, 1>{}([&](auto idx) {
+            //             auto result = type_convert<scalar_t>(
+            //                 type_convert<float>(a_vector.template AsType<scalar_t>()[idx]) +
+            //                 type_convert<float>(b_vector.template AsType<scalar_t>()[idx]));
+            //             this->template Set<scalar_t>(i + idx, is_valid_element, result);
+            //         });
+            //     }
+            // }
+            // else
+            // {
+            //     this->template Set<X>(i, is_valid_element, x + tmp);
+            // }
+
+            // Properly handle addition for all low-precision types
+            if constexpr(is_same_v<scalar_t, bhalf_t> || is_same_v<scalar_t, half_t>)
             {
                 if constexpr(is_scalar_type<X>::value)
                 {
-                    // Scalar type
-                    auto result =
-                        type_convert<X>(type_convert<float>(x) + type_convert<float>(tmp));
+                    // Scalar type: Convert to float, add, convert back
+                    auto result = type_convert<X>(type_convert<float>(x) + type_convert<float>(tmp));
                     this->template Set<X>(i, is_valid_element, result);
                 }
                 else
@@ -170,6 +200,8 @@ struct DynamicBuffer
                     constexpr auto vector_size = scalar_type<remove_cvref_t<X>>::vector_size;
                     const vector_type<scalar_t, vector_size> a_vector{tmp};
                     const vector_type<scalar_t, vector_size> b_vector{x};
+                    
+                    // Process each element of the vector in higher precision
                     static_for<0, vector_size, 1>{}([&](auto idx) {
                         auto result = type_convert<scalar_t>(
                             type_convert<float>(a_vector.template AsType<scalar_t>()[idx]) +
@@ -178,10 +210,11 @@ struct DynamicBuffer
                     });
                 }
             }
-            else
-            {
-                this->template Set<X>(i, is_valid_element, x + tmp);
-            }
+            // else
+            // {
+            //     // // For other types, direct addition is fine
+            //     // this->template Set<X>(i, is_valid_element, x + tmp);
+            // }
         }
     }