了解 cuda 中的线程索引：

Question

考虑以下内核：

多线程单块：

   __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 的输出<<<(4000 + 1023)/1024,1024>>>(4000) :

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 之间切换 blockIdx.x 但有时它只是切换回去和在 2 和 3 之间？

By using a loop with stride equal to the grid size, we ensure that all addressing within warps is unit-stride, so we get maximum memory coalescing, just as in the monolithic version.

这是什么意思？

来源：https://developer.nvidia.com/blog/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/

Answer 1

输出没有用，因为无法保证在哪个流式多处理器 (SMP) 的哪个时间点安排哪个块。即使是同一块内的扭曲也不必运行并行。因此，可能块 2 的一个 warp 和块 3 的一个 warp 在不同的 SMP 上并行处理。看来你应该先试着了解一下memory coalescing.

所有这些示例都希望您以 array[tid] 的方式访问全局内存。因此，在不使用块索引的情况下对多个块使用 1. 或 3. 将意味着多个块正在做完全相同的工作。所以 2. 和 4. 是更通用的。对于 2. 你必须小心地为你的问题大小（例如 array 的大小）启动足够的线程，而 4. 中的网格步幅循环确保你会得到正确的结果，即使你启动更少线程。但是如果没有足够的块来填充 GPU，您可能无法获得全部性能。

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

不属于同一个 warp 的任何两个线程可以运行以任何顺序（或并行）。没有任何保证。

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

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

了解 cuda 中的线程索引：

Understanding Thread Indexing in cuda :

cuda