Commit ea724623 authored by Po-Yen, Chen's avatar Po-Yen, Chen
Browse files

Default enable intrinsic function calls

parent 1e1e8fd9
...@@ -100,6 +100,7 @@ struct GridwiseGemmPipeline_v2 ...@@ -100,6 +100,7 @@ struct GridwiseGemmPipeline_v2
// global read i + 2 // global read i + 2
b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf); b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf);
#if 1
__builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
...@@ -126,6 +127,7 @@ struct GridwiseGemmPipeline_v2 ...@@ -126,6 +127,7 @@ struct GridwiseGemmPipeline_v2
__builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read __builtin_amdgcn_sched_group_barrier(0x020, 1, 0); // VMEM read
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
__builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS write __builtin_amdgcn_sched_group_barrier(0x200, 1, 0); // DS write
#endif
++i; ++i;
} while(i < (num_loop - 2)); } while(i < (num_loop - 2));
......
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