CUDA编程实现矩阵转置时对线程索引的疑问

利用共享内存完成矩阵转置时，对线程索引很不理解，仔细研究之后才完全弄懂具体的原理。废话不多说，直接进入主题。普通的矩阵转置我就不详细说了，很好理解，这方面的资料也很多，直接从共享内存完成矩阵转置开始。
干货（**先给出代码，然后讲述具体原理**）：
`__global__ void matrix_transpo(const float*a, float*c, const int n, const int m)
{//**a是n行m列，c是m行n列**
    __shared__ float tem[BLOCK_DIM][BLOCK_DIM];
     int xindex = threadIdx.x + blockIdx.x*blockDim.x;
     int yindex = threadIdx.y + blockIdx.y*blockDim.y;
    if (xindex<m&&yindex<n)
    tem[threadIdx.y][threadIdx.x] = a[yindex*m + xindex];
    __syncthreads();
    xindex = blockIdx.y*BLOCK_DIM + threadIdx.x;//观察与上面的不同
    yindex = blockIdx.x*BLOCK_DIM + threadIdx.y;//这两行代码很多人都难以理解，后面将详细讲解
    if (xindex < n&&yindex < m)
        c[yindex*n + xindex] = tem[threadIdx.x][threadIdx.y];//观察数组tem的索引
}`
具体原理（**先给出模型图，然后详细解析**）：

假如将线程分成了2*2的线程块，每个块的线程为4*4，块内的共享内存的大小为线程块的大小（tem[4][4]),我们将数组a中的数据按行优先读取，然后存储在共享内存中(tem数组中), 再从共享内存中按列读取，将数据按行存储到数组c中，当进行矩阵转置时，首先是块进行转置（块2和块3变换位置），然后块内元素进行转置（可以画图试一试），这都不难理解，难理解的是线程索引的表示，如代码中的 ：

 xindex = blockIdx.y*BLOCK_DIM + threadIdx.x;//观察与上面的不同
 yindex = blockIdx.x*BLOCK_DIM + threadIdx.y;//

这两行代码是建立数组c的元素索引，很多人想不明白，为什么后面加的东西（红色字体)与自己想的不一样，甚至怀疑代码是不是有问题（代码没问题，验证过了），现在就来给大家说说原理，刚才已经说了，线程块是要进行转置的，所以纵向的线程块所以将是转置后数组的横向线程块索引，横向的线程块索引将是转置后的数组的纵向索引，然后线程块转置后内部的线程也要完成转置，同理，纵向的维度将是转置后的数组的横向维度，看着可能有点晕，看下图就知道了

 xindex = blockIdx.y*BLOCK_DIM + threadIdx.x;//观察与上面的不同
 yindex = blockIdx.x*BLOCK_DIM + threadIdx.y;

上图也就是对红色代码的解释， c[yindex*n + xindex] = tem[threadIdx.x][threadIdx.y];我们可以观察一下，这行代码中的tem数组的索引，tem[threadIdx.x][threadIdx.y];与图中线程块中红色的箭头相对应，线程索引是（threadIdx.y，threadIdx.x），而取数据却取的是与之对称的位置的数据，所以后面加的索引值是下面红色部分。

 xindex = blockIdx.y*BLOCK_DIM + threadIdx.x;//观察与上面的不同
 yindex = blockIdx.x*BLOCK_DIM + threadIdx.y;//

感觉表达能力有待提高，总感觉表达不到位，总担心不能给大家解释清

下一篇：使用cublas库遇到的问题（CUDA编程）

https://blog.csdn.net/huwendong666666/article/details/92771292