如何代替stmatrix指令完成存回比较好？

加速计算 CUDA
1

在使用ptx mma指令mma.sync.aligned.m16n8k64，在一个warp的寄存器中布局，应当满足以下布局

image
image1178×669 77.8 KB

显然 在同一个寄存器中的值，对应的逻辑位置的内存地址是不连续的，所以当我尝试写回共享内存时，我的写法如下：

      for(int k = 0;k <2;++k){
        half *ptr = &smem[i * smem_ldm * M + j * N];
        // half2 * ptr_half2 = &smem[E2S(i * smem_ldm * M + j * N)];
        // half2 value = __floats2half2_rn(c_frag[i * WARP_ROW_TILES * 4 + j * 4 + 0 + 2*k],c_frag[i * WARP_ROW_TILES * 4 + j * 4 + 1 + 2*k]);
        // ptr_half2[(tj+8*k) * E2S(smem_ldm) + ti ] = value;
        ptr[(tj+8*k) * smem_ldm + ti * 2 + 0] = __float2half(c_frag[i * WARP_ROW_TILES * 4 + j * 4 + 0 + 2*k]);
        ptr[(tj+8*k) * smem_ldm + ti * 2 + 1] = __float2half(c_frag[i * WARP_ROW_TILES * 4 + j * 4 + 1 + 2*k]);
        __syncthreads();
      }

WARP_ROW_TILES代表是多个输入矩阵的索引，在循环时， 我按照8x的间隔去写回，但是在ncu的报告中一直提示非合并的内存访问，虽然 ptx中存在stmatrix指令去实现以上过程，但是目前我无法使用sm90以上的设备，那么该如何在sm90以下设备去写回共享内存比较好呢？