Commit 17c97f58 authored by danyao12's avatar danyao12
Browse files

solve lds read&write conflicts

parent 0178da6f
...@@ -258,6 +258,7 @@ struct BlockFmhaBwdDQDKDVPipelineKRKTRVR ...@@ -258,6 +258,7 @@ struct BlockFmhaBwdDQDKDVPipelineKRKTRVR
block_sync_lds(); block_sync_lds();
v_reg_tensor = load_tile(v_lds_read_window); v_reg_tensor = load_tile(v_lds_read_window);
block_sync_lds();
//---------------------------- Loop Load in ----------------------------// //---------------------------- Loop Load in ----------------------------//
// Q: HBM ->Reg ->LDS // Q: HBM ->Reg ->LDS
auto q_dram_window = auto q_dram_window =
......
...@@ -505,6 +505,7 @@ struct BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP ...@@ -505,6 +505,7 @@ struct BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP
/* /*
* Store prefetched data into LDS * Store prefetched data into LDS
*/ */
block_sync_lds();
store_tile(q_lds_window, q_block_tile); store_tile(q_lds_window, q_block_tile);
shuffle_tile(qt_block_tile, q_block_tile); shuffle_tile(qt_block_tile, q_block_tile);
store_tile(qt_lds_write_window, qt_block_tile); store_tile(qt_lds_write_window, qt_block_tile);
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment