2.5 KERNEL FUNCTIONS AND THREADING

我们现在准备讨论更多关于CUDA内核功能以及启动这些内核功能的效果。在CUDA中，内核函数指定所有线程在并行阶段执行的代码。由于所有这些线程执行相同的代码，CUDA编程是众所周知的单程序多数据（SPMD）[Ata 1998]并行编程风格的实例，这是一种大规模并行计算系统的流行编程风格。

请注意，SPMD与SIMD（单指令多数据）不同[Flynn 1972]。在SPMD系统中，并行处理单元在数据的多个部分上执行相同的程序。然而，这些处理单元不需要同时执行相同的指令。在SIMD系统中，所有处理单元在任何时候都在执行相同的指令。
“在CUDA 3.0及更高版本中，每个线程块最多可以有1024个线程。一些早期的CUDA版本只允许一个块中最多512个线程。

当程序的主机代码启动内核时，CUDA运行时系统会生成一个线程网格，这些线程被组织成两级层次结构。每个网格都组织成线程块数组，为了简洁，将被称为线程块。网格的所有块大小相同；每个块最多可包含1024个线程。图2.11显示了每个块由256个线程组成的示例。每个线程都由一个卷曲的箭头表示，这个箭头来自一个标有数字的block。启动内核时，每个线程块中的线程总数由主机代码指定。同一内核可以在主机代码的不同部分使用不同数量的线程启动。对于给定的网格，块中的线程数在内置的blockDim变量中可用。

blockDim变量是结构类型，有三个无符号整数字段：x、y和z，这有助于程序员将线程组织成一维、二维或三维数组。对于一维组织，将只使用x字段。对于二维组织，将使用x和y字段。对于三维结构，将使用所有三个字段。组织线程的维度选择通常反映了数据的维度。这是有道理的，因为创建线程是为了并行处理数据。线程的组织反映了数据的组织，这是很自然的。在图2.11中，每个线程块都组织为线程的一维数组，因为数据是一维向量。blockDim.x变量的值指定了每个块中的线程总数，在图2.11中为256。一般来说，线程的数量由于硬件效率的原因，线程块每个维度的线程应该是32的倍数。我们稍后会重温这个问题。

CUDA内核可以访问另外两个内置变量（threadldx，blockldx），这些变量允许线程相互区分，并确定每个线程要处理的数据区域。变量threadldx在块内为每个线程提供唯一的坐标。例如，在图2.11中，由于我们使用的是一维线程组织，因此将只使用threadldx.x。每个线程的sthreadldx.x值显示在图2.11.中每个线程的小阴影框中。每个块中的第一个线程在其threadldx.x变量中具有值 0，第二个线程具有值1，第三个线程具有值2，等。

Blockldx变量为块中的所有线程提供一个公共块坐标。在图2.11中，第一个块中的所有线程在其blockldx.x变量中都有值0，第二个线程块中的值为1，以此等。使用与电话系统的类比，人们可以将threadldx.x视为本地电话号码，将blockldx.x视为区号。两者一起给每条电话线一个全国唯一的电话号码。同样，每个线程可以结合其threadldx和blockIdx值，在整个网格中为自己创建一个唯一的全局索引。

在图2.11中，**一个唯一的全局索引i计算为i = blockldx.x*blockDim.x + threadldx.x。**回想一下，在我们的例子中，blockDim是256。0块中线程的i值从0到255不等。第1块中线程的i值从256到511不等。第2块中线程的i值从512到767不等。也就是说，这三个块中线程的i值形成了从0到767的值的连续覆盖。由于每个线程使用i来访问A、B和C，这些线程涵盖了原始循环的前768次迭代。请注意，我们不会在内核中使用“h_”和“d_”约定，因为没有潜在的混淆。在我们的示例中，我们将无法访问主机内存。通过启动具有更多块的内核，可以处理更大的向量。通过启动具有n个或更多线程的内核，可以处理长度为n的向量。

图2.12显示了向量加法的内核函数。语法是ANSI C，有一些值得注意的扩展。首先，在vecAddKernel函数的声明前面有一个CUDA C特定的关键字“global”。此关键字表示该函数是内核，可以从主机函数调用以在设备上生成线程网格。

一般来说，CUDA C用三个限定词关键词扩展了C语言，这些关键词可用于函数声明。这些关键字的含义总结在图2.13 中“global”关键字表示正在声明的函数是CUDA C内核函数。请注意，“global”一词的两侧各有两个下划线字符。这种内核函数将在设备上执行，并且只能从主机代码调用，除非在支持动态并行的CUDA系统中，正如我们将在第13章CUDA动态并行性中解释的那样。“device”关键字表示正在声明的函数是CUDA设备函数。设备函数在CUDA设备上执行，只能从内核函数或其他设备函数调用。

我们稍后将解释在不同代CUDA中使用间接函数调用和递归的规则。一般来说，应该避免在其设备函数和内核函数中使用递归和间接函数调用，以实现最大的可移植性。

“host”关键字表示正在声明的函数是CUDA主机函数。主机函数只是一个传统的C函数，在主机上执行，只能从另一个主机函数调用。默认情况下，如果声明中没有任何CUDA关键字，CUDA程序中的所有函数都是主机函数。这是有道理的，因为许多CUDA应用程序是从仅CPU执行环境移植的。程序员将在移植过程中添加内核功能和设备功能。原始功能仍然是主机功能。将所有函数默认为主机函数，使程序员免于更改所有原始函数声明的繁琐工作。

请注意，可以在函数声明中同时使用“host”和“device”。这种组合告诉编译系统为同一函数生成两个版本的对象文件。一个在主机上执行，只能从主机函数调用。另一个在设备上执行，只能从设备或内核函数调用。这支持一个常见的用例，当相同的函数源代码可以重新编译以生成设备版本时。许多用户库功能可能属于这一类。

ANSI C的第二个值得注意的扩展，在图2.12中，是内置变量“threadldx.x”、“blockldx.x”和“blockDim.x”。回想一下，所有线程都执行相同的kernel代码。他们需要一种方法来区分自己，并将每个线程引向数据的特定部分。这些内置变量是线程访问硬件寄存器的手段，这些寄存器为线程提供识别坐标。不同的线程将在其threadldx.X、blockldx.x和blockDim.x变量中看到不同的值。为了简单起见，我们将线程称为$threa{d}_{blockidx.x，threadidx.x}$。请注意，“x”意味着应该有“.y”和“.z”。我们很快就会回到这一点上。

图2.12中有一个自动（局部）变量i。.在CUDA内核函数中，自动变量对每个线程都是私有的。也就是说，将为每个线程生成一个i版本。如果内核以10,000个线程启动，将有10,000个版本的i，每个线程一个。线程分配给其i变量的值对其他线程不可见。我们将在第4章“内存和数据位置”中更详细地讨论这些自动变量。

图2.5和2.12之间的快速比较。揭示了CUDA内核和CUDA内核启动的重要见解。图2.12中的内核函数没有与图2.5.中的循环相对应的循环。读者应该问这个循环去了哪里。答案是，循环现在被线程网格所取代。整个网格形成等价的循环。网格中的每个线程对应于原始循环的一次迭代。这种类型的数据并行性有时也被称为循环并行性，其中原始顺序代码的迭代由线程并行执行。

**请注意，在图2.12.中的addVecKernel中有一个if（i < n）语句。这是因为并非所有向量长度都可以表示为块大小的倍数。**例如，让我们假设矢量长度是100。最小的高效螺纹块尺寸为32。假设我们选择了32个块大小。需要启动四个线程块来处理所有100个矢量元素。然而，这四个线程块将有128个线程。我们需要禁用线程块3中的最后28个线程，使其无法完成原始程序预期之外的工作。由于所有线程都将执行相同的代码，因此所有线程都将针对n（即100）测试其i值。使用if（i <n）语句，前100个线程将执行加法，而最后28个线程不会执行加法。这允许内核处理任意长度的向量。

当主机代码启动内核时，它通过执行配置参数设置网格和线程块尺寸。图2.14.中说明了这一点。配置参数在传统C函数参数之前的“<<<”和“>>>”之间给出。第一个配置参数给出了网格中线程块的数量。第二个指定每个线程块中的线程数。在本例中，每个块中有256个线程。为了确保我们有足够的线程来覆盖所有向量元素，我们将C ceiling 函数应用于n/256.0。使用foating点值256.0可确保我们为除法生成浮动值，以便 ceiling 函数可以正确四舍五入。例如，如果我们有1000个线程，我们将启动ceil（1000/256.0）= 4个线程块。因此，该语句将启动4*256 = 1024线程。使用内核中的if（i < n）语句，如图2.12所示，前1000个线程将对1000个矢量元素进行加法。剩下的24个不会。