[HGEMM] fix naive hgemm f16x8 LDST error (#157)

wjj19950828 · aiyi.wjj · web-flow · commit a6831457f03f · 2024-12-10T19:51:16.000+08:00
Co-authored-by: aiyi.wjj &lt;aiyi.wjj@alibaba-inc.com&gt;
diff --git a/kernels/hgemm/naive/hgemm.cu b/kernels/hgemm/naive/hgemm.cu
@@ -735,7 +735,7 @@ __global__ void hgemm_t_8x8_sliced_k_f16x8_pack_bcf_dbuf_kernel(
     s_a[smem_sel_next][load_a_smem_k + 1][load_a_smem_m] = r_load_a[1];
     s_a[smem_sel_next][load_a_smem_k + 2][load_a_smem_m] = r_load_a[2];
     s_a[smem_sel_next][load_a_smem_k + 3][load_a_smem_m] = r_load_a[3];
-    LDST128BITS(s_b[smem_sel_next][load_b_smem_k][load_b_smem_n]) = LDST128BITS(r_load_b[0]);
+    LDST64BITS(s_b[smem_sel_next][load_b_smem_k][load_b_smem_n]) = LDST64BITS(r_load_b[0]);
 
     __syncthreads();
   }

Original file line number	Diff line number	Diff line change
`@@ -735,7 +735,7 @@ __global__ void hgemm_t_8x8_sliced_k_f16x8_pack_bcf_dbuf_kernel(`
`735`	`735`	`s_a[smem_sel_next][load_a_smem_k + 1][load_a_smem_m] = r_load_a[1];`
`736`	`736`	`s_a[smem_sel_next][load_a_smem_k + 2][load_a_smem_m] = r_load_a[2];`
`737`	`737`	`s_a[smem_sel_next][load_a_smem_k + 3][load_a_smem_m] = r_load_a[3];`
`738`		`- LDST128BITS(s_b[smem_sel_next][load_b_smem_k][load_b_smem_n]) = LDST128BITS(r_load_b[0]);`
	`738`	`+ LDST64BITS(s_b[smem_sel_next][load_b_smem_k][load_b_smem_n]) = LDST64BITS(r_load_b[0]);`
`739`	`739`
`740`	`740`	`__syncthreads();`
`741`	`741`	`}`