simply vector type At

composable_kernel/include/tensor_operation/threadwise_generic_tensor_slice_copy_v2.hpp

@@ -197,7 +197,7 @@ struct ThreadwiseGenericTensorSliceCopy_v5
                     ThreadBufferDesc::CalculateOffset(to_multi_index(long_vector_data_begin_id)) /
                     long_vector_size;
 
-                thread_buff.template At<SrcDataPerRead>()(Number<buff_off>{}) = src_buff;
+                thread_buff.At(Number<SrcDataPerRead>{})(Number<buff_off>{}) = src_buff;
             });
     }
 
@@ -225,7 +225,7 @@ struct ThreadwiseGenericTensorSliceCopy_v5
                     ThreadBufferDesc::CalculateOffset(to_multi_index(long_vector_data_begin_id)) /
                     long_vector_size;
 
-                auto src_buff = thread_buff.template At<DstDataPerWrite>()[Number<buff_off>{}];
+                auto src_buff = thread_buff.At(Number<DstDataPerWrite>{})[Number<buff_off>{}];
 
                 const auto dst_coord = mDstSliceOrigin + to_multi_index(long_vector_data_begin_id);
 

composable_kernel/include/utility/float_type.amd.hpp.in

@@ -37,22 +37,22 @@ union float_vec4_t
     __host__ __device__ constexpr float_vec4_t() {}
 
     template <index_t vs>
-    __host__ __device__ auto& At();
+    __host__ __device__ auto& At(Number<vs>);
 
     template <>
-    __host__ __device__ auto& At<1>()
+    __host__ __device__ auto& At(Number<1>)
     {
         return s1;
     }
 
     template <>
-    __host__ __device__ auto& At<2>()
+    __host__ __device__ auto& At(Number<2>)
     {
         return s2;
     }
 
     template <>
-    __host__ __device__ auto& At<4>()
+    __host__ __device__ auto& At(Number<4>)
     {
         return s4;
     }
@@ -67,28 +67,28 @@ union float_vec8_t
     __host__ __device__ constexpr float_vec8_t() {}
 
     template <index_t vs>
-    __host__ __device__ auto& At();
+    __host__ __device__ auto& At(Number<vs>);
 
     template <>
-    __host__ __device__ auto& At<1>()
+    __host__ __device__ auto& At(Number<1>)
     {
         return s1;
     }
 
     template <>
-    __host__ __device__ auto& At<2>()
+    __host__ __device__ auto& At(Number<2>)
     {
         return s2;
     }
 
     template <>
-    __host__ __device__ auto& At<4>()
+    __host__ __device__ auto& At(Number<4>)
     {
         return s4;
     }
 
     template <>
-    __host__ __device__ auto& At<8>()
+    __host__ __device__ auto& At(Number<8>)
     {
         return s8;
     }
@@ -104,34 +104,34 @@ union float_vec16_t
     __host__ __device__ constexpr float_vec16_t() {}
 
     template <index_t vs>
-    __host__ __device__ auto& At();
+    __host__ __device__ auto& At(Number<vs>);
 
     template <>
-    __host__ __device__ auto& At<1>()
+    __host__ __device__ auto& At(Number<1>)
     {
         return s1;
     }
 
     template <>
-    __host__ __device__ auto& At<2>()
+    __host__ __device__ auto& At(Number<2>)
     {
         return s2;
     }
 
     template <>
-    __host__ __device__ auto& At<4>()
+    __host__ __device__ auto& At(Number<4>)
     {
         return s4;
     }
 
     template <>
-    __host__ __device__ auto& At<8>()
+    __host__ __device__ auto& At(Number<8>)
     {
         return s8;
     }
 
     template <>
-    __host__ __device__ auto& At<16>()
+    __host__ __device__ auto& At(Number<16>)
     {
         return s16;
     }