add new method deqaunt

Author: xiaolil1

csrc/xpu_cutlass_fusion.cpp

@@ -283,6 +283,7 @@ class gemm_4bit_cutlass_kernel {
     }
 #endif
 #else
+#if 0    
     using format_type = uint32_t; //32
     static constexpr auto src_bits = sizeof_bits_v<SrcType>; //4
     static constexpr auto scalar = sizeof_bits_v<format_type> / src_bits; // 8
@@ -324,6 +325,27 @@ class gemm_4bit_cutlass_kernel {
         }
       }
     }
+#endif    
+    using compress_type = uint8_t;
+    static constexpr auto compress_ratio = sizeof_bits_v<compress_type> / sizeof_bits_v<SrcType>;
+    static constexpr auto K = decltype(size(out))::value / N;
+    auto s_tensor = make_tensor((compress_type*)(raw_pointer_cast(in.data())), Shape<Int<K/compress_ratio>, Int<N>>{});
+    auto d_tensor = make_tensor(out.data(), Shape<Int<K>, Int<N>>{});
+
+    #pragma unroll
+    for (int n = 0; n < N; n++) {
+      float ts = tCrS_input(n);
+      //auto& src = *(cute::array<SrcType, splits>*)(s_tensor(_, n).data());
+      //auto& dst = *(cute::array<DstType, splits>*)(d_tensor(_, n).data());
+      auto& src = s_tensor(_, n).data();
+      auto& dst = d_tensor(_, n).data();
+
+      #pragma unroll
+      for (int k = 0; k < K/compress_ratio/2; k++) {
+        dst[k * 2 + 1] = static_cast<DstType>(quant_map[src[k] >> 4] * ts);
+        dst[k * 2] = static_cast<DstType>(quant_map[src[k] & 0x0f] * ts);
+      }
+    }
 #endif     
   }