CUDA中step的计算方式及设计原理是什么？

源代码（来自《Programming in Parallel with CUDA》第25-26页） #

15.1     __global__ void gpu_sin(float *sums, int steps, int terms, float step_size)
15.2     {
15.3         int step = blockIdx.x*blockDim.x+threadIdx.x;
15.4         if(step<steps){
15.5            float x = step_size*step;
15.6            sums[step] = sinsum(x,terms); // store values
15.7         }
15.8     }

书中原文翻译 #

第15.3行声明了变量step，它的作用等价于示例1.1第24行中同名的for循环索引变量。该变量由内置变量blockDim.x、blockIdx.x和threadIdx.x定义，这些变量的值取决于主机调用内核时使用的启动参数：

• blockDim.x被设置为threads，即内核使用的线程块大小。
• blockIdx.x被设置为当前线程所属线程块的序号，范围为[0, blocks-1]。
• threadIdx.x被设置为当前线程在其线程块内的序号，范围为[0, threads-1]。
• step = blockDim.x * blockIdx.x + threadIdx.x 的取值范围为[0, threads × blocks - 1]。

用户问题 #

我无法理解最后一点内容，请问该step表达式是如何计算的？其设计原理是什么？

一、step表达式的计算过程 #

用实际数值举例最直观：
假设我们设置线程块大小blockDim.x = 256（每个块包含256个线程），总共启动4个线程块（即blocks=4）。

• 第0个线程块（blockIdx.x=0）：
  块内线程的threadIdx.x从0到255，计算step = 256*0 + threadIdx.x，得到的step范围是0~255。
• 第1个线程块（blockIdx.x=1）：
  step = 256*1 + threadIdx.x，范围是256~511。
• 第2个线程块（blockIdx.x=2）：
  step范围是512~767。
• 第3个线程块（blockIdx.x=3）：
  step范围是768~1023。

把所有块的step范围连起来，就是0~(256*4-1)=1023，刚好覆盖所有线程对应的唯一索引，无重复无遗漏。

本质上，blockDim.x * blockIdx.x是当前线程块的起始全局索引，加上threadIdx.x就是当前线程在全局所有线程中的唯一序号，也就是step的值。

二、设计原理 #

这个表达式的核心目的是给每个线程分配一个全局唯一的索引，用来对应串行代码中for循环的迭代变量。

在串行代码里，我们用for循环逐个处理step从0到steps-1的任务；而CUDA并行中，这些任务被拆分给不同线程，每个线程处理一个step对应的计算。

通过blockIdx.x和threadIdx.x组合计算的优势：

1. 利用线程块的批量特性，blockDim.x * blockIdx.x快速定位当前块的任务起始位置，避免每个线程重复计算全局偏移。
2. 块内的threadIdx.x负责块内的任务细分，确保块内每个线程拿到不同的子索引。
3. 最终的step能覆盖所有启动的线程，实现和串行for循环完全对应的任务分配逻辑，让并行代码无缝匹配串行的迭代逻辑。

内容的提问来源于stack exchange，提问作者user366312