Skip to content

Commit b2e2125

Browse files
committed
save code
1 parent e717474 commit b2e2125

1 file changed

Lines changed: 5 additions & 3 deletions

File tree

csrc/xpu_cutlass_fusion.cpp

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -525,6 +525,7 @@ printf("src_compress_size = %d, dst_compress_size = %d, src_vec_size = %d, dst_v
525525
for (int v = 0; v < src_vec_size; v++) {
526526
src_compress_type src_value = src[v];
527527
int dst_base_idx = l * src_vec_size * src_compress_size + v * src_compress_size;
528+
//int dst_base_idx = 0;
528529
#pragma unroll
529530
for (int c = 0; c < src_compress_size; c++) {
530531
uint8_t bit_value = (src_value >> (4 * (((c + 1) & 1) + (c >> 1) * 2))) & 0xF;
@@ -537,10 +538,11 @@ printf("src_compress_size = %d, dst_compress_size = %d, src_vec_size = %d, dst_v
537538
}
538539

539540
// #pragma unroll
540-
// for (int l = 0; l < dst_loop_num; l++) {
541+
// for (int l = 0; l < dst_loop_num / 4; l++) {
541542
// //reinterpret_cast<sycl::vec<dst_compress_type, dst_vec_size>*>(cute::raw_pointer_cast(mma_B.data()))[n * dst_loop_num + l] = reinterpret_cast<sycl::vec<dst_compress_type, dst_vec_size>*>(dst)[l];
542-
// reinterpret_cast<dst_compress_type*>(cute::raw_pointer_cast(mma_B.data()))[n * dst_loop_num + l] = reinterpret_cast<dst_compress_type*>(dst)[l];
543-
543+
// //reinterpret_cast<dst_compress_type*>(cute::raw_pointer_cast(mma_B.data()))[n * dst_loop_num + l] = reinterpret_cast<dst_compress_type*>(dst)[l];
544+
// reinterpret_cast<sycl::vec<dst_compress_type, 4>*>(cute::raw_pointer_cast(mma_B.data()))[n*dst_loop_num + l] = reinterpret_cast<sycl::vec<dst_compress_type, 4>*>(dst)[l];
545+
//
544546
// }
545547
}
546548
};

0 commit comments

Comments
 (0)