了解cuda中的线程索引：

Question

考虑以下内核:

1. 多线程单块：

   __global__ Kernel(){
             int tid = threadIdx.x;
          }

    __global__ Kernel(){
         int tid = threadIdx.x + blockIdx.x* blockDim.x;
        }

    __global__ Kernel(int n){
           for(int tid = threadIdx.x; tid < n; tid += blockDim.x){
           }
        }

    __global__ Kernel(int n){
           for(int tid = threadIdx.x + blockIdx.x*blockDim.x; tid < n; tid += blockDim.x * grdiDim.x){
           }
        }

现在1将从单个块启动请求的线程数，但最大线程数限制为1024（我的硬件）。只需生成另一个块。

2从多个块中启动请求数量的线程。现在，可以启动的最大线程数增加到了65535 * 1024。

3是块步幅循环，4是网格步幅循环。

我不明白的是3和4是如何迭代的？我读过的几篇文章中，状态网格步幅循环在批中迭代，这意味着什么？

调用时从3输出，作为内核执行

Local Thread Id : 56  Block ID : 2 Global Thread Id : 3128
Local Thread Id : 57  Block ID : 2 Global Thread Id : 3129
Local Thread Id : 58  Block ID : 2 Global Thread Id : 3130
Local Thread Id : 59  Block ID : 2 Global Thread Id : 3131
Local Thread Id : 60  Block ID : 2 Global Thread Id : 3132
Local Thread Id : 61  Block ID : 2 Global Thread Id : 3133
Local Thread Id : 62  Block ID : 2 Global Thread Id : 3134
Local Thread Id : 63  Block ID : 2 Global Thread Id : 3135
Local Thread Id : 448  Block ID : 3 Global Thread Id : 3520
Local Thread Id : 449  Block ID : 3 Global Thread Id : 3521
Local Thread Id : 450  Block ID : 3 Global Thread Id : 3522
Local Thread Id : 451  Block ID : 3 Global Thread Id : 3523
Local Thread Id : 452  Block ID : 3 Global Thread Id : 3524
Local Thread Id : 453  Block ID : 3 Global Thread Id : 3525
Local Thread Id : 454  Block ID : 3 Global Thread Id : 3526
Local Thread Id : 455  Block ID : 3 Global Thread Id : 3527
Local Thread Id : 456  Block ID : 3 Global Thread Id : 3528

有时切换发生在0,1,2,3块Idx. x之间，但有时它只是在2和3之间来回切换？

通过使用步幅等于网格大小的循环，我们确保warps中的所有寻址都是单位步幅，因此我们获得了最大的内存合并，就像在单片版本中一样。

这是什么意思？

来源：html" target="_blank">https://developer.nvidia.com/blog/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/

Answer 1

输出没有用，因为无法保证哪个块将在哪个时间点调度在哪个流多处理器 （SMP） 上。即使是同一块内的经线也不必并行运行。因此，块 2 的一个翘曲和块 3 的一个翘曲可能在不同的 SMP 上并行处理。看来你应该首先尝试理解记忆合并。

所有这些示例都希望您以数组[tid]的形式访问全局内存。所以使用 1.或 3.使用多个块而不使用块索引将意味着多个块正在执行完全相同的工作。所以2。和 4.是更通用的。与 2.您必须小心地启动足够的线程来满足您的问题大小（例如数组的大小），而网格在 4 中跨步循环。确保您将获得正确的结果，即使您启动的线程较少。但是，如果没有足够的块来填充 GPU，则可能无法获得全部性能。

需要理解的一点是，每个块都由扭曲（到目前为止，其大小为32个通道或“线程”）组成，其工作方式类似于在向量寄存器上工作的SIMD向量通道。扭曲的线程通常是同步工作的，这样，只要这些元素不大于8字节，协作加载连续的内存区域（通道0加载元素0，通道1加载元素1，…通道31加载元素31）就会更快。

不属于同一扭曲的任何两条线都可以按任何顺序（或并行）运行。没有任何保证。

它是如何迭代的？for循环应该很清楚。块DIM. x*grdiDIM. x是线程的总数。如果您使用3。只有一个块，块DIM. x是线程的总数。因此，不像在CPU上那样，每个线程都在连续的块上工作（用于缓存局部性），线程以交错的方式工作（用于内存合并）。

有不同的方法来迭代并保持内存合并（例如，每个扭曲或每个块处理连续的内存块），但这种方法最容易写下来和理解。此外，假设您使用的线程不会超过一次可以调度的线程，网格步幅循环在某种程度上比其他版本具有更多的“缓存局部性”，至少在查看由SMP共享的L2缓存时是这样。