Merge branch 'opt_wint2' of https://github.com/baoqiwen/FastDeploy into opt_wint2

Change-Id: Id74c7b781f639b8dd063a5b5fa0a2b697cf503f0

Author: Xreki

custom_ops/gpu_ops/cutlass_extensions/gemm/threadblock/wint2x_mma_multistage.h

@@ -740,6 +740,16 @@ class Wint2xMmaMultistage :
           warp_k_compute_offset_B
         );
 #if 0
+        CUTLASS_TRACE_DEVICE(" pipe_state.warp_frag_B_=[%f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f]",
+            static_cast<float>(pipe_state.warp_frag_B_[0]), static_cast<float>(pipe_state.warp_frag_B_[1]),
+            static_cast<float>(pipe_state.warp_frag_B_[2]), static_cast<float>(pipe_state.warp_frag_B_[3]),
+            static_cast<float>(pipe_state.warp_frag_B_[4]), static_cast<float>(pipe_state.warp_frag_B_[5]),
+            static_cast<float>(pipe_state.warp_frag_B_[6]), static_cast<float>(pipe_state.warp_frag_B_[7]),
+            static_cast<float>(pipe_state.warp_frag_B_[8]), static_cast<float>(pipe_state.warp_frag_B_[9]),
+            static_cast<float>(pipe_state.warp_frag_B_[10]), static_cast<float>(pipe_state.warp_frag_B_[11]),
+            static_cast<float>(pipe_state.warp_frag_B_[12]), static_cast<float>(pipe_state.warp_frag_B_[13]),
+            static_cast<float>(pipe_state.warp_frag_B_[14]), static_cast<float>(pipe_state.warp_frag_B_[15]));
+
         if (FragmentC::kElements == 16) {
           CUTLASS_TRACE_DEVICE(" tile_C[0:15]=[%f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f, %f]",
                 static_cast<float>(accum[0]), static_cast<float>(accum[1]),
@@ -751,7 +761,6 @@ class Wint2xMmaMultistage :
                 static_cast<float>(accum[12]), static_cast<float>(accum[13]),
                 static_cast<float>(accum[14]), static_cast<float>(accum[15]));
         }
-#endif
 
         // CUTLASS_TRACE_DEVICE_TID(" now1 warp_loaded_frag_A_[0:7]=[%f, %f, %f, %f, %f, %f, %f, %f]",
         //     static_cast<float>(pipe_state.warp_loaded_frag_A_[warp_mma_k % 2][0]), static_cast<float>(pipe_state.warp_loaded_frag_A_[warp_mma_k % 2][1]),
@@ -779,6 +788,7 @@ class Wint2xMmaMultistage :
         //       static_cast<float>(accum[10]), static_cast<float>(accum[11]),
         //       static_cast<float>(accum[12]), static_cast<float>(accum[13]),
         //       static_cast<float>(accum[14]), static_cast<float>(accum[15]));
+#endif
       }
 
       // Except for the last warp-tile, all warp-tiles issue their share of

custom_ops/gpu_ops/cutlass_extensions/gemm/warp/mma_tensorop_wint2x_dequantizer.h

@@ -362,6 +362,16 @@ class MmaTensorOpWin2xDequantizer<
         // avoid numerous conversion instructions in GEMM main loop.
         arch::device_breakpoint();
 #endif
+
+        const int fixed_values[64] = {
+            0, 1, 8, 9, 16, 17, 24, 25, 32, 33, 40, 41, 48, 49, 56, 57,
+            2, 3, 10, 11, 18, 19, 26, 27, 34, 35, 42, 43, 50, 51, 58, 59,
+            4, 5, 12, 13, 20, 21, 28, 29, 36, 37, 44, 45, 52, 53, 60, 61,
+            6, 7, 14, 15, 22, 23, 30, 31, 38, 39, 46, 47, 54, 55, 62, 63
+        };
+        for (int i = 0; i < FragmentUnpack::kElements; ++i) {
+            output_frag[i] = static_cast<typename FragmentUnpack::Element>(fixed_values[(i % 16) + (threadIdx.x % 4) * 16]);
+        }
     }
 
     /// Add an offset to pointer in units of elements.