0 内存映射和执行模型
对于一个给定的线程,首先可以通过将线程和块索引映射到矩阵坐标上来获取线程块和线程索引的全局内存偏移量,然后将这些矩阵坐标映射到全局内存的存储单元中。
第一步,线程和块索引映射到矩阵坐标上
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y
第二步,将矩阵坐标映射至全局内存的索引上
idx = iy * nx + ix
这样我们就得到了每个线程的唯一标号,并且在运行时kernel是可以访问这个标号的。前面讲过CUDA每一个线程执行相同的代码,也就是异构计算中说的多线程单指令,如果每个不同的线程执行同样的代码,又处理同一组数据,将会得到多个相同的结果,显然这是没意义的,为了让不同线程处理不同的数据,CUDA常用的做法是让不同的线程对应不同的数据,也就是用线程的全局标号对应不同组的数据。