在使用ds_read加mma指令进行运算时我发现这种情况下`__sycnthreads()`并不能同步线程 开始时代码如下 ```c++ #pragma unroll for (int r = 0; r < 2; r++) { int input_mma_load_offset = ((tid % 64) / 8 + r * 8) * 64 + (tid % 8) * 4; int input_lds_addr = input_mma_load_inner_offset + input_mma_load_offset; input_lds_addr = input_lds_addr * sizeof(float); PREFETCH_LDS(_b, input_lds_addr); int weight_mma_load_offset = ((tid % 64) / 8 + r * 8) * 64 + (tid % 8) * 4; int weight_lds_addr = weight_mma_load_inner_offset + weight_mma_load_offset + weight_dsread_offset; weight_lds_addr = weight_lds_addr * sizeof(float); PREFETCH_LDS(_a, weight_lds_addr); __syncthreads(); _d[0] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[0], _b.f2[0], _d[0]); _d[1] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[0], _b.f2[1], _d[1]); _d[2] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[1], _b.f2[0], _d[2]); _d[3] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[1], _b.f2[1], _d[3]); } ``` 在上面的代码里，我只用了一个`__sycnthreads()`来确保mma在lds之后进行，但是代码编译运行后出错，表现为：在不同的位置处，计算出错误的值，而这个错误的值每次运行都会变。很明显这是因为同步竞争。 查阅资料后我发现了`s_waitcnt lgkmcnt(0)`和`s_waitcnt vmcnt(0)`两条同步指令，在将我的同步方法换成用这两条指令后，同步竞争导致的错误消失了。 ```c++ #pragma unroll for (int r = 0; r < 2; r++) { int input_mma_load_offset = ((tid % 64) / 8 + r * 8) * 64 + (tid % 8) * 4; int input_lds_addr = input_mma_load_inner_offset + input_mma_load_offset; // DEBUG_PRINT("input_lds_addr: %d\n", input_lds_addr); input_lds_addr = input_lds_addr * sizeof(float); PREFETCH_LDS(_b, input_lds_addr); int weight_mma_load_offset = ((tid % 64) / 8 + r * 8) * 64 + (tid % 8) * 4; int weight_lds_addr = weight_mma_load_inner_offset + weight_mma_load_offset + weight_dsread_offset; weight_lds_addr = weight_lds_addr * sizeof(float); PREFETCH_LDS(_a, weight_lds_addr); asm volatile("s_waitcnt lgkmcnt(0)\n\t"); _d[0] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[0], _b.f2[0], _d[0]); _d[1] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[0], _b.f2[1], _d[1]); _d[2] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[1], _b.f2[0], _d[2]); _d[3] = __builtin_amdgcn_mmac_f32_16x16x8f32(_a.f2[1], _b.f2[1], _d[3]); asm volatile("s_waitcnt vmcnt(0)\n\t"); __syncthreads(); } ``` 我的问题是：在使用`ds_read`和`mma`的时候，是否确如我所说需要使用`s_waitcnt lgkmcnt(0)`和`s_waitcnt vmcnt(0)`来进行同步而不能只使用`__sycnthreads()`？又是什么原因导致的？