CUDA学习

超、凢脫俗 2022-08-01 12:15 192阅读 0赞

转自

## [CUDA学习之一][CUDA] ##

CUDA 的核心有三个重要抽象概念：线程组层次结构、共享存储器、屏蔽同步（barrier synchronization），可轻松将其作为 C 语言的最小扩展级公开给程序员。

GPU 专用于解决可表示为数据并行计算的问题——在许多数据元素上并行执行的程序，具有极高的计算密度（数学运算与存储器运算的比率）。由于所有数据元素都执行相同的程序，因此对精密流控制的要求不高；由于在许多数据元素上运行，且具有较高的计算密度，因而可通过计算隐藏存储器访问延迟，而不必使用较大的数据缓存。

最新一代的 NVIDIA GPU 基于 Tesla 架构。在 NVIDIA Tesla 架构中，一个线程块最多可以包含 512 个线程。

可以通过调用 \_syncthreads()\_ 内函数在内核中指定同步点；\_syncthreads()\_ 起到屏障的作用，块中的所有线程都必须在这里等待处理。

int main()

// Kernel invocation

dim3 dimBlock(16, 16); //说明每个块里面有16\*16个线程

dim3 dimGrid((N + dimBlock.x – 1) / dimBlock.x,

(N + dimBlock.y – 1) / dimBlock.y);  //定义二维块。这种方法经常要用

matAdd<<<dimGrid, dimBlock>>>(A, B, C);   //内核函数的调用

一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格。

网格的维度由 <<<…>>> 语法的第一个参数指定。网格内的每个块多可由一个一维或二维索引标识，可通过内置的blockIdx 变量在内核中访问此索引。可以通过内置的 blockDim 变量在内核中访问线程块的维度

存储器层次结构：

CUDA 线程可在执行过程中访问多个存储器空间的数据。每个线程都有一个私有的本地存储器。每个线程块都有一个共享存储器，该存储器对于块内的所有线程都是可见的，并且与块具有相同的生命周期。最终，所有线程都可访问相同的全局存储器。

此外还有两个只读的存储器空间，可由所有线程访问，这两个空间是固定存储器空间和纹理存储器空间。全局、固定和纹理存储器空间经过优化，适于不同的存储器用途。

对于同一个应用程序启动的内核而言，全局、固定和纹理存储器空间都是持久的。

主机和设备编程模式如下图所示：

CUDA软件栈

CUDA 软件栈包含多个层：设备驱动程序、应用程序编程接口（API）及其运行时、两个较高级别的通用数学库，即CUFFT 和 CUBLAS。

支持CUDA的显卡及其计算能力：

<table> 
 <tbody> 
 <tr> 
 <td> &nbsp; </td> 
 <td> 多处理器数量 </td> 
 <td> 计算能力 </td> 
 </tr> 
 <tr> 
 <td> GeForce GTX 280 </td> 
 <td> 30 </td> 
 <td> 1.3 </td> 
 </tr> 
 <tr> 
 <td> GeForce GTX 260 </td> 
 <td> 24 </td> 
 <td> 1.3 </td> 
 </tr> 
 <tr> 
 <td> GeForce 9800 GX2 </td> 
 <td> 2×16 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 9800 GTX </td> 
 <td> 16 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 8800 Ultra, 8800 GTX </td> 
 <td> 16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> GeForce 8800 GT </td> 
 <td> 14 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 9600 GSO, 8800 GS, 8800M GTX </td> 
 <td> 12 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 8800 GTS </td> 
 <td> 12 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> GeForce 9600 GT, 8800M GTS&nbsp; </td> 
 <td> 8 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 9500 GT, 8600 GTS, 8600 GT, 8700M&nbsp;GT, 8600M GT, 8600M GS </td> 
 <td> 4 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 8500 GT, 8400 GS, 8400M GT, 8400M&nbsp;GS </td> 
 <td> 2 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> GeForce 8400M G </td> 
 <td> 1 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> Tesla S1070 </td> 
 <td> 4×30 </td> 
 <td> 1.3 </td> 
 </tr> 
 <tr> 
 <td> Tesla C1060 </td> 
 <td> 30 </td> 
 <td> 1.3 </td> 
 </tr> 
 <tr> 
 <td> Tesla S870 </td> 
 <td> 4×16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Tesla D870 </td> 
 <td> 2×16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Tesla C870 </td> 
 <td> 16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Quadro Plex 1000 Model S4 </td> 
 <td> 4×16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Quadro Plex 1000 Model IV </td> 
 <td> 2×16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Quadro FX 5600 </td> 
 <td> 16 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Quadro FX 3700 </td> 
 <td> 14 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> Quadro FX 3600M </td> 
 <td> 12 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> Quadro FX 4600 </td> 
 <td> 12 </td> 
 <td> 1.0 </td> 
 </tr> 
 <tr> 
 <td> Quadro FX 1700, FX 570, NVS 320M, FX 1600M, FX 570M </td> 
 <td> 4 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> Quadro FX 370, NVS 290, NVS 140M, NVS 135M, FX 360M </td> 
 <td> 2 </td> 
 <td> 1.1 </td> 
 </tr> 
 <tr> 
 <td> Quadro NVS 130M </td> 
 <td> 1 </td> 
 <td> 1.1 </td> 
 </tr> 
 </tbody> 
</table>

## [CUDA学习之二][CUDA 1] ##

并发与并行的区别：“并行”是指无论从微观还是宏观，二者都是一起执行的，就好像两个人各拿一把铁锨在挖坑，一小时后，每人一个大坑。

而“并发”在微观上不是同时执行的，只是把时间分成若干段，使多个进程快速交替的执行，从宏观外来看，好像是这些进程都在执行，这就好像两个人用同一把铁锨，轮流挖坑，一小时后，两个人各挖一个小一点的坑，要想挖两个大一点得坑，一定会用两个小时。

从以上本质不难看出，“并发”执行，在多个进程存在资源冲突时，并没有从根本提高执行效率。

Tesla 架构的构建以一个可伸缩的多线程流处理器（SM）阵列为中心。当主机 CPU 上的 CUDA 程序调用内核网格时，网格的块将被枚举并分发到具有可用执行容量的多处理器上。一个线程块的线程在一个多处理器上并发执行。在线程块终止时，将在空闲多处理器上启动新块。

多处理器包含 8 个标量处理器（SP）核心、两个用于先验（transcendental）的特殊函数单元、一个多线程指令单元以及芯片共享存储器。多处理器会在硬件中创建、管理和执行并发线程，而调度开销保持为0。它可通过一条内部指令实现\_syncthreads()\_ 屏障同步。快速的屏障同步与轻量级线程创建和零开销的线程调度相结合，有效地为细粒度并行化提供了支持，举例来说，您可以为各数据元素（如图像中的一个像素、语音中的一个语音元素、基于网格的计算中的一个单元）分配一个线程，从而对问题进行细粒度分解。

为了管理运行各种不同程序的数百个线程，多处理器利用了一种称为 SIMT（单指令、多线程）的新架构。多处理器会将各线程映射到一个标量处理器核心，各标量线程使用自己的指令地址和寄存器状态独立执行。多处理器 SIMT 单元以32 个并行线程为一组来创建、管理、调度和执行线程，这样的线程组称为 warp 块。（此术语源于第一种并行线程技术weaving。半 warp 块可以是一个 warp 块的第一半或第二半。）构成 SIMT warp 块的各个线程在同一个程序地址一起启动，但也可随意分支、独立执行。

为一个多处理器指定了一个或多个要执行的线程块时，它会将其分成 warp 块，并由 SIMT 单元进行调度。将块分割为 warp 块的方法总是相同的，每个 warp 块都包含连续的线程，递增线程 ID，第一个 warp 块中包含线程 0。第 2.1 节介绍了线程 ID 与块中的线程索引之间的关系。

每发出一条指令时，SIMT 单元都会选择一个已准备好执行的 warp 块，并将下一条指令发送到该 warp 块的活动线程。Warp 块每次执行一条通用指令，因此在 warp 块的全部 32 个线程均认可其执行路径时，可达到最高效率。如果一个warp 块的线程通过独立于数据的条件分支而分散，warp 块将连续执行所使用的各分支路径，而禁用未在此路径上的线程，完成所有路径时，线程重新汇聚到同一执行路径下。分支仅在 warp 块内出现，不同的 warp 块总是独立执行的——无论它们执行的是通用的代码路径还是彼此无关的代码路径。

SIMT 架构类似于 SIMD（单指令、多数据）向量组织方法，共同之处是使用单指令来控制多个处理元素。一项主要差别在于 SIMD 向量组织方法会向软件公开 SIMD 宽度，而 SIMT 指令指定单一线程的执行和分支行为。与 SIMD 向量机不同，SIMT 允许程序员为独立、标量线程编写线程级的并行代码，还允许为协同线程编写数据并行代码。为了确保正确性，程序员可忽略 SIMT 行为，但通过维护很少需要使一个 warp 块内的线程分支的代码，即可实现显著的性能提升。在实践中，这与传统代码中的超高速缓冲存储器线作用相似：在以正确性为目标进行设计时，可忽略超高速缓冲存储器线的大小，但如果以峰值性能为目标进行设计，在代码结构中就必须考虑其大小。另一方面，向量架构要求软件将负载并入向量，并手动管理分支。

若要多 GPU 系统上运行的应用程序将多个 GPU 作为 CUDA 设备使用，则这些 GPU 必须具有相同的类型。但如果系统采用的是 SLI 模式，则仅有一个 GPU 可用作 CUDA 设备，因为所有 GPU 都将在驱动程序栈的最低级别融合。要使CUDA 能够将各 GPU 视为独立设备，需要在 CUDA 的控制面板内关闭 SLI 模式。

## [CUDA学习之三][CUDA 2] ##

CUDA编码

对于函数部分前缀

1.  \_\_device\_\_

使用 \_device\_ 限定符声明的函数具有以下特征：

n         在设备上执行；

n         仅可通过设备调用。

2. \_\_global\_\_

使用 \_global\_ 限定符可将函数声明为内核。此类函数：

n         在设备上执行；

n         仅可通过主机调用。

3. \_\_host\_\_

使用 \_host\_ 限定符声明的函数具有以下特征：

n         在主机上执行；

n         仅可通过主机调用。

仅使用 \_host\_ 限定符声明函数等同于不使用 \_host\_、\_device\_ 或 \_global\_ 限定符声明函数，这两种情况下，函数都将仅为主机进行编译。

函数前缀的一些限制：

\_device\_ 和 \_global\_ 函数不支持递归。

\_device\_ 和 \_global\_ 函数的函数体内无法声明静态变量。

\_device\_ 和 \_global\_ 函数不得有数量可变的参数。

\_device\_ 函数的地址无法获取，但支持 \_global\_ 函数的函数指针。

\_global\_ 和 \_host\_ 限定符无法一起使用。

\_global\_ 函数的返回类型必须为空。

对 \_global\_ 函数的任何调用都必须按规定指定其执行配置。

\_global\_ 函数的调用是异步的，也就是说它会在设备执行完成之前返回。

\_global\_ 函数参数将同时通过共享存储器传递给设备，且限制为 256 字节。

对于变量前缀：

1.\_\_device\_\_

\_device\_ 限定符声明位于设备上的变量。

在接下来的三节中介绍的其他类型限定符中，最多只能有一种可与 \_device\_ 限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符，则变量具有以下特征：

n         位于全局存储器空间中；

n         与应用程序具有相同的生命周期；

可通过网格内的所有线程访问，也可通过运行时库从主机访问。

2.\_\_constant\_\_

\_constant\_ 限定符可选择与 \_device\_ 限定符一起使用，所声明的变量具有以下特征：

n         位于固定存储器空间中；

n         与应用程序具有相同的生命周期；

可通过网格内的所有线程访问，也可通过运行时库从主机访问。

3.\_\_shared\_\_

\_shared\_ 限定符可选择与 \_device\_ 限定符一起使用，所声明的变量具有以下特征：

n         位于线程块的共享存储器空间中；

n         与块具有相同的生命周期；

n         尽可通过块内的所有线程访问。

只有在 \_syncthreads()\_（参见第 4.4.2 节）的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。

限制：

不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。

\_shared\_ 和 \_constant\_ 变量具有隐含的静态存储。

\_device\_、\_shared\_ 和 \_constant\_ 变量无法使用 extern 关键字定义为外部变量。

\_device\_ 和 \_constant\_ 变量仅允许在文件作用域内使用。

不可为设备或从设备指派 \_constant\_ 变量，仅可通过主机运行时函数从主机指派（参见第 4.5.2.3 节和第 4.5.3.6 节）。

\_shared\_ 变量的声明中不可包含初始化。

下面是具体的一个应用：

将共享存储器中的变量声明为外部数组时，例如：

extern \_\_shared\_\_ float shared\[\];

数组的大小将在启动时确定（参见第 4.2.3 节）。所有变量均以这种形式声明，在存储器中的同一地址开始，因此数组中的变量布局必须通过偏移显式管理。例如，如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容：

short array0\[128\];

float array1\[64\];

int array2\[256\];

则应通过以下方法声明和初始化数组：

extern \_\_shared\_\_ char array\[\];

\_\_device\_\_ void func() // \_\_device\_\_ or \_\_global\_\_ function

short\* array0 = (short\*)array;

float\* array1 = (float\*)&array0\[128\];

int\* array2 = (int\*)&array1\[64\];

在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下，编译器可能选择将其置于本地存储器中。

只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间，此类指针即受支持，否则将仅限于指向在全局存储器空间中分配或声明的存储器。

通过获取 \_device\_、\_shared\_ 或 \_constant\_ 变量的地址而获得的地址仅可在设备代码中使用。通过cudaGetSymbolAddress()获取的 \_device\_ 或 \_constant\_ 变量的地址仅可在主机代码中使用。

** **

**对global函数进行配置**

对 \_global\_ 函数的任何调用都必须指定该调用的执行配置。

执行配置定义将用于在该设备上执行函数的网格和块的维度，以及相关的流。可通过在函数名称和括号参数列表之间插入 <<<Dg, Db, Ns, s>>> 形式的表达式来指定，其中：

Dg 的类型为 dim3，指定网格的维度和大小，Dg.x \* Dg.y 等于所启动的块数量，Dg.z 无用；

Db 的类型为 dim3，指定各块的维度和大小，Db.x \* Db.y \* Db.z 等于各块的线程数量；

Ns 的类型为 size\_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用，Ns 是一个可选参数，默认值为 0；

S 的类型为 cudaStream\_t，指定相关流；S 是一个可选参数，默认值为 0。

举例来说，一个函数的声明如下：

\_\_global\_\_ void Func(float\* parameter);

必须通过如下方法来调用此函数：

Func<<< Dg, Db, Ns >>>(parameter);

执行配置的参数将在实际函数参数之前被评估，与函数参数相同，通过共享存储器同时传递给设备。

如果 Dg 或 Db 大于设备允许的最大大小，或 Ns 大于设备上可用的共享存储器最大值，或者小于静态分配、函数参数和执行配置所需的共享存储器数量，则函数将失败。

## [CUDA学习之四][CUDA 3] ##

内置变量

1.gridDim

此变量的类型为 dim3，包含网格的维度。

2. blockIdx

此变量的类型为 uint3，包含网格内的块索引。

3.blockDim

此变量的类型为 dim3，包含块的维度。

4. threadIdx

此变量的类型为 uint3，包含块内的线程索引。

5.warpSize

此变量的类型为 int，包含以线程为单位的 warp 块大小。

限制：

n         不允许接受任何内置变量的地址。

n         不允许为任何内置变量赋值。

默认情况下，\_device\_ 函数总是内嵌的。**\_noinline\_ **函数限定符可用于指示编译器尽可能不要内嵌该函数。函数体必须位于所调用的同一个文件内。

如果函数具有指针参数或者具有较大的参数列表，则编译器不会遵从 \_noinline\_ 限定符。

默认情况下，编译器将展开具有已知行程计数的小循环。**\#pragma unroll** 指令可用于控制任何给定循环的展开操作。它必须紧接于循环之前，而且仅应用于该循环。可选择在其后接一个数字，指定必须展开多少次循环。

例如，在下面的代码示例中：

\#pragma unroll 5

for (int i = 0; i < n; ++i)

循环将展开 5 次。程序员需要负责确保展开操作不会影响程序的正确性（在上面的示例中，如果 n 小于 5，则程序的正确性将受到影响）。

\#pragma unroll 1 将阻止编译器展开一个循环。

如果在 \#pragma unroll 后未指定任何数据，如果其行程计数为常数，则该循环将完全展开，否则将不会展开。

## [CUDA学习之五（通用运行时组件）][CUDA 4] ##

主机和设备函数均可使用通用运行时组件。

**内置向量类型**：char1、uchar1、char2、uchar2、char3、uchar3、char4、uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1、int2、uint2、int3、uint3、int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、long4、ulong4、float1、float2、float3、float4、double2

**dim3 类型**：此类型是一种整形向量类型，基于用于指定维度的 uint3。在定义类型为 dim3 的变量时，未指定的任何组件都将初始化为 1。

**数学函数**：包含了当前支持的 C/C++ 标准库数学函数的完整列表，还分别给出了在设备上执行时的误差范围。

在主机代码中执行时，给定函数将在可用的前提下使用 C 运行时实现。

**计时函数**：

clock\_t clock();

在设备代码中执行时，返回随每一次时钟周期而递增的每个多处理器计数器的值。在内核启动和结束时对此计数器取样，确定两次取样的差别，然后为每个线程记录下结果，这为各线程提供一种度量方法，可度量设备为了完全执行线程而占用的时钟周期数，但不是设备在执行线程指令时而实际使用的时钟周期数。前一个数字要比后一个数字大得多，因为线程是分时的。

**纹理类型**：

CUDA 支持 GPU 用于图形的纹理硬件子集，使之可访问纹理存储器。从纹理存储器而非全局存储器读取数据可带来多方面的性能收益。

内核使用称为纹理获取（texture fetch）的设备函数读取纹理存储器。纹理获取的第一个参数指定称为纹理参考的对象。

纹理参考定义获取哪部分的纹理存储器。必须通过主机运行时函数（将其绑定到存储器的某些区域（即纹理），之后才能供内核使用。多个不同的纹理参考可绑定到同一个纹理，也可绑定到在存储器中存在重叠的纹理。

纹理参考有一些属性。其中之一就是其维度，指定纹理是使用一个纹理坐标（texture coordinate）将纹理作为一维数组寻址、使用两个纹理坐标作为二维数组寻址，还是使用三个纹理坐标作为三维数组寻址。数组的元素称为 texel，即“texture elements（纹理元素）”的简写。

其他属性定义纹理获取的输入和输出数据类型，并指定如何介绍输入坐标、应进行怎样的处理。

纹理参考的部分属性是不变的，在编译时必须为已知，这些属性是在声明纹理参考时指定的。纹理参考在文件作用域内声明，形式为 texture 类型的变量：

texture<Type, Dim, ReadMode> texRef;

其中：

Type 指定获取纹理时所返回的数据类型；Type 仅限于基本整型、单精度浮点类型和第 4.3.1.1 节定义的 1 组件、2 组件和 4 组件向量类型；

Dim 指定纹理参考的维度，其值为 1、2 或 3；Dim 是一个可选的参数，默认值为 1；

ReadMode 等于 cudaReadModeNormalizedFloat 或 cudaReadModeElementType；如果是cudaReadModeNormalizedFloat，且Type 为 16 位或 8 位整型类型，则值将作为浮点类型返回，对于所有整型数据而言，无符号整型将映射为 \[0.0, 1.0\]，有符号整型将映射为 \[-1.0, 1.0\]，例如，一个值为 0xff 的无符号 8 位纹理元素将被读取为 1；如果是cudaReadModeElementType，则不执行任何转换操作；ReadMode 是一个可选的参数，默认值为cudaReadModeElementType。

纹理参考的其他属性是可变的，可通过主机运行时在运行时更改。它们指定纹理坐标是否为规范化的，以及寻址模式和纹理过滤，下面将介绍相关内容。

默认情况下，使用 \[0, N） 范围内的浮点坐标引用纹理，其中的 N 是纹理在对应于坐标的维度中的大小。例如，有一个大小为 64x32 的纹理，在 x 和 y 维度引用此纹理时坐标分别处于 \[0, 63\] 和 \[0, 31\] 范围内。规范化的纹理坐标将在 \[0.0, 1.0)的范围内指定，而非 \[0, N)，因此在规范化的坐标内，同一 64x32 纹理的寻址范围在 x 和 y 维度均为 \[0, 1)。一般情况下，纹理坐标与纹理大小无关，规范化的纹理坐标通常足以满足一些应用程序的需求。

寻址模式定义在纹理坐标超出范围时将出现怎样的情况。在使用非规范化纹理坐标时，超出 \[0, N) 范围的纹理坐标将被调整：小于 0 的值被设置为 0，大于或等于 N 的值被设置为 N-1。在使用规范化纹理坐标时，默认寻址模式也是调整坐标：小于 0.0 或大于 1.0 的值将被调整到范围 \[0.0, 1.0) 内。对于规范化坐标，“warp 块”的寻址模式也可指定。Warp 块寻址往往在纹理包含周期信号时使用。它仅使用纹理坐标的一部分，例如，1.25 被视为 0.25，-1.25 被视为 0.75.

线性纹理过滤只能对配置为返回浮点数据的纹理进行。这将在相邻 texel 间执行低精度插值。在启用时，位于纹理获取位置周围的 texel 将被读取，纹理获取的返回值将根据纹理坐标在 texel 间的位置进行插值。对于一维纹理执行简单的线性插值，而对于二维纹理则执行双线性插值。

纹理可以是线性存储器或 CUDA 数组的任意区域。

在线性存储器内分配的纹理：

维度仅能为 1；

不支持纹理过滤；

仅可使用非规范化整型纹理坐标寻址；

不支持多种寻址模式：超出范围的纹理访问将返回零。

硬件会对纹理基址实施对齐要求。为了抽象这种来自程序员的对齐要求，绑定设备存储器上的纹理参考的函数将传回一个字节偏移，必须将其应用到纹理获取，之后才能读取所需的存储器。CUDA 分配例程返回的基址指针符合这种对齐限制，因此应用程序可通过向 cudaBindTexture()/cuTexRefSetAddress() 传递所分配的指针来完全避免偏移。

## [CUDA学习之六（设备运行时组件）][CUDA 5] ##

设备运行时组件仅可用于设备函数。

1.数学函数

设备运行时组件中存在准确性略低而速度更快的版本；其名称相同，但带有一个\_前缀（如\_sinf(x)）。编译器有一个 (-use\_fast\_math) 选项，用于强制要求所有函数编译其准确性略低的版本（如果存在）。

2.同步函数

void \_\_syncthreads();

同步块中的所有线程。一旦所有线程均达到此同步点，执行将正常恢复。

\_syncthreads() 用于调整同一个块的线程之间的通信。在一个块内的某些线程访问共享或全局存储器中的相同地址时，部分访问操作可能存在写入后读取、读取后写入或写入后写入之类的风险。可通过在这些访问操作间同步线程来避免这些数据风险。

\_syncthreads() 允许在条件代码中使用，但仅当条件估值在整个线程块中都相同时才允许使用，否则代码执行将有可能挂起，或者出现意料之外的副作用。

3.纹理函数

来自线性存储器的纹理：

对于来自线性存储器的纹理，通过 tex1Dfetch() 系列函数访问纹理，示例如下：

template<class Type>

Type tex1Dfetch(

texture<Type, 1, cudaReadModeElementType> texRef,

int x);

float tex1Dfetch(

texture<unsigned char, 1, cudaReadModeNormalizedFloat> texRef,

int x);

float tex1Dfetch(

texture<signed char, 1, cudaReadModeNormalizedFloat> texRef,

int x);

float tex1Dfetch(

texture<unsigned short, 1, cudaReadModeNormalizedFloat> texRef,

int x);

float tex1Dfetch(

texture<signed short, 1, cudaReadModeNormalizedFloat> texRef,

int x);

这些函数会使用纹理坐标 x 获取绑定到纹理参考 texRef 的线性存储器区域。不支持纹理过滤和寻址模式。对于整型来说，这些函数可选择将整型转变为单精度浮点类型。

除了上述函数以外，还支持 2 元组和 4 元组，示例如下：

float4 tex1Dfetch(

texture<uchar4, 1, cudaReadModeNormalizedFloat> texRef,

int x);

以上示例将使用纹理坐标 x 获取绑定到纹理参考 texRef 的线性存储器。

来自 CUDA 数组的纹理：

对于来自 CUDA 数组的纹理，可通过 tex1D()、tex2D()、tex3D() 访问纹理：

template<class Type, enum cudaTextureReadMode readMode>

Type tex1D(texture<Type, 1, readMode> texRef,

float x);

template<class Type, enum cudaTextureReadMode readMode>

Type tex2D(texture<Type, 2, readMode> texRef,

float x, float y);

template<class Type, enum cudaTextureReadMode readMode>

Type tex3D(texture<Type, 3, readMode> texRef,

float x, float y, float z);

这些函数将使用纹理坐标 x、y 和 z 获取绑定到纹理参考 texRef 的 CUDA 数组。纹理参考的不变（编译时）和可变（运行时）属性相互结合，共同确定坐标的解释方式、在纹理获取过程中发生的处理以及纹理获取所提供的返回值。

4.原子函数

原子函数对位于全局或共享存储器内的一个 32 位或 64 位字执行读取-修改-写入原子操作。例如，atomicAdd() 将在全局或共享存储器内的某个地址读取 32 位字，将其与一个整型相加，并将结果写回同一地址。之所以说这样的操作是原子的，是因为它可在不干扰其他线程的前提下执行。换句话说，在操作完成中，其他任何线程都无法访问此地址。

原子操作仅适用于有符号和无符号整型（但 atomicExch() 是一个例外情况，它支持单精度浮点数字）。

## [CUDA学习之七（主机运行时组件）][CUDA 6] ##

只有主机函数才能使用主机运行时组件。

它提供了具有以下功能的函数：

n         设备管理；

n         上下文管理；

n         存储器管理；

n         代码模块管理；

n         执行控制；

n         纹理参考管理；

n         与 OpenGL 和 Direct3D 的互操作性。

它包含两个 API：

n         一个称为 CUDA 驱动程序 API 的低级 API；

n         一个称为 CUDA 运行时 API 的高级 API，它是在 CUDA 驱动程序 API 的基础之上实现的。

这些 API 是互斥的：一个应用程序仅能使用其中之一。

CUDA 驱动程序 API 是通过 nvcuda 动态库提供的，其所有入口点都带有 cu 前缀。

CUDA 运行时 API 是通过 cudart 动态库提供的，其所有入口点都带有 cuda 前缀。

运行时API

**cudaGetDeviceCount() **和 **cudaGetDeviceProperties()** 提供了一种方法，用于枚举这些设备并检索其属性：

int deviceCount;

cudaGetDeviceCount(&deviceCount);

int device;

for (device = 0; device < deviceCount; ++device) \{

cudaDeviceProp deviceProp;

cudaGetDeviceProperties(&deviceProp, device);

cudaSetDevice() 用于选择与主机线程相关的设备：

cudaSetDevice(device);

必须首先选择设备，之后才能调用 \_global\_ 函数或任何来自运行时 API 的函数。如果未通过显式调用 cudaSetDevice() 完成此任务，将自动选中设备 0，随后对 cudaSetDevice() 的任何显式调用都将无效。

存储器管理：

线性存储器是使用 cudaMalloc() 或 cudaMallocPitch() 分配的，使用 cudaFree() 释放。

以下示例代码将在线性存储器中分配一个包含 256 个浮点元素的数组：

float\* devPtr;

cudaMalloc((void\*\*)&devPtr, 256 \* sizeof(float));

建议在分配二维数组时使用 cudaMallocPitch()，因为它能确保合理填充已分配的存储器，满足第 5.1.2.1 节介绍的对齐要求，从而确保访问行地址或执行二维数组与设备存储器的其他区域之间的复制（使用 cudaMemcpy2D()）时获得最优性能。所返回的间距（或步幅）必须用于访问数组元素。以下代码示例将分配一个 widthxheight 的二维浮点值数组，并显示如何在设备代码中循环遍历数组元素：

// host code

float\* devPtr;

int pitch;

cudaMallocPitch((void\*\*)&devPtr, &pitch,

width \* sizeof(float), height);

myKernel<<<100, 512>>>(devPtr, pitch);

// device code

\_\_global\_\_ void myKernel(float\* devPtr, int pitch)

for (int r = 0; r < height; ++r) \{

float\* row = (float\*)((char\*)devPtr + r \* pitch);

for (int c = 0; c < width; ++c) \{

float element = row\[c\];

CUDA 数组是使用 cudaMallocArray() 分配的，使用 cudaFreeArray() 释放。cudaMallocArray() 需要使用cudaCreateChannelDesc() 创建的格式描述。

以下代码示例分配了一个 widthxheight 的 CUDA 数组，包含一个 32 位的浮点组件：

cudaChannelFormatDesc channelDesc =

cudaCreateChannelDesc<float>();

cudaArray\* cuArray;

cudaMallocArray(&cuArray, &channelDesc, width, height);

cudaGetSymbolAddress() 用于检索指向为全局存储器空间中声明的变量分配的存储器的地址。所分配存储器的大小是通过cudaGetSymbolSize() 获取的。

参考手册列举了用于在 cudaMalloc() 分配的线性存储器、cudaMallocPitch() 分配的线性存储器、CUDA 数组和为全局或固定存储器空间中声明的变量分配的存储器之间复制存储器的所有函数。

下面的代码示例将二维数组复制到之前代码示例中分配的 CUDA 数组中：

cudaMemcpy2DToArray(cuArray, 0, 0, devPtr, pitch,

width \* sizeof(float), height,

cudaMemcpyDeviceToDevice);

下面的代码示例将一些主机存储器数组复制到设备存储器中：

float data\[256\];

int size = sizeof(data);

float\* devPtr;

cudaMalloc((void\*\*)&devPtr, size);

cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice);

下面的代码示例将一些主机存储器数组复制到固定存储器中：

\_\_constant\_\_ float constData\[256\];

float data\[256\];

cudaMemcpyToSymbol(constData, data, sizeof(data));

流管理：

以下代码示例创建两个流：

cudaStream\_t stream\[2\];

for (int i = 0; i < 2; ++i)

cudaStreamCreate(&stream\[i\]);

这些流均通过以下代码示例定义为一个序列，包括一次从主机到设备的存储器复制、一次内核启动、一次从设备到主机的存储器复制：

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(inputDevPtr + i \* size, hostPtr + i \* size,

size, cudaMemcpyHostToDevice, stream\[i\]);

for (int i = 0; i < 2; ++i)

myKernel<<<100, 512, 0, stream\[i\]>>>

(outputDevPtr + i \* size, inputDevPtr + i \* size, size);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(hostPtr + i \* size, outputDevPtr + i \* size,

size, cudaMemcpyDeviceToHost, stream\[i\]);

cudaThreadSynchronize();

两个流均会将其输入数组 hostPtr 的一部分复制到设备存储器的 inputDevPtr 数组中，通过调用 myKernel() 处理设备上的inputDevPtr，并将结果 outputDevPtr 复制回 hostPtr 的相同部分。使用两个流处理 hostPtr 允许一个流的存储器复制与另外一个流的内核执行相互重叠。hostPtr 必须指向分页锁定的主机存储器，这样才能出现重叠：

float\* hostPtr;

cudaMallocHost((void\*\*)&hostPtr, 2 \* size);

最后调用了 cudaThreadSynchronize()，目的是在进一步处理之前确定所有流均已完成。cudaStreamSynchronize() 可用于同步主机与特定流，允许其他流继续在该设备上执行。通过调用 cudaStreamDestroy() 可释放流。

时间管理：

下面的代码示例创建了两个事件：

cudaEvent\_t start, stop;

cudaEventCreate(&start);

cudaEventCreate(&stop);

这些事件可用于为上一节的代码示例计时，方法如下：

cudaEventRecord(start, 0);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(inputDev + i \* size, inputHost + i \* size,

size, cudaMemcpyHostToDevice, stream\[i\]);

for (int i = 0; i < 2; ++i)

myKernel<<<100, 512, 0, stream\[i\]>>>

(outputDev + i \* size, inputDev + i \* size, size);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(outputHost + i \* size, outputDev + i \* size,

size, cudaMemcpyDeviceToHost, stream\[i\]);

cudaEventRecord(stop, 0);

cudaEventSynchronize(stop);

float elapsedTime;

cudaEventElapsedTime(&elapsedTime, start, stop);

cudaEventDestroy(start);

cudaEventDestroy(stop);

纹理参考管理：

在内核使用纹理参考从纹理存储器中读取之前，必须使用 **cuTexRefSetAddress() **或 **cuTexRefSetArray()**将纹理参考绑定到纹理。

如果模块 cuModule 包含定义如下的纹理参考 texRef：

texture<float, 2, cudaReadModeElementType> texRef;

则下面的代码示例将检索 texRef 的句柄：

CUtexref cuTexRef;

cuModuleGetTexRef(&cuTexRef, cuModule, “texRef”);

下面的代码示例将 texRef 绑定到 devPtr 指向的线性存储器：

cuTexRefSetAddress(NULL, cuTexRef, devPtr, size);

下面的代码示例将 texRef 绑定到 CUDA 数组 cuArray：

cuTexRefSetArray(cuTexRef, cuArray, CU\_TRSA\_OVERRIDE\_FORMAT);

参考手册列举了用于设置寻址模式、过滤模式和其他针对纹理参考的标记的各种函数。在将纹理绑定到纹理参考时所指定的格式必须与声明纹理参考时指定的参数相匹配；否则纹理获取的结果将无法确定。

Open GL互操作性

必须使用 cuGLInit() 初始化与 OpenGL 的互操作性。

首先必须将一个缓冲对象注册到 CUDA，之后才能进行映射。可通过 **cuGLRegisterBufferObject() **完成：

GLuint bufferObj;

cuGLRegisterBufferObject(bufferObj);

注册完成后，内核即可使用 **cuGLMapBufferObject()** 返回的设备存储器地址读取或写入缓冲对象：

GLuint bufferObj;

CUdeviceptr devPtr;

int size;

cuGLMapBufferObject(&devPtr, &size, bufferObj);

解除映射是通过 **cuGLUnmapBufferObject()** 完成的，可使用 **cuGLUnregisterBufferObject()** 取消注册。

Direct3D 互操作性要求在创建 CUDA 上下文时指定 Direct3D 设备。通过使用 **cuD3D9CtxCreate() **而非**cuCtxCreate()** 创建 CUDA 上下文即可实现此目标。。

随后即可使用 **cuD3D9RegisterResource()** 将 Direct3D 资源注册到 CUDA：

LPDIRECT3DVERTEXBUFFER9 buffer;

cuD3D9RegisterResource(buffer, CU\_D3D9\_REGISTER\_FLAGS\_NONE);

LPDIRECT3DSURFACE9 surface;

cuD3D9RegisterResource(surface, CU\_D3D9\_REGISTER\_FLAGS\_NONE);

**cuD3D9RegisterResource() **可能具有较高的开销，通常仅为每个资源调用一次。使用**cuD3D9UnregisterVertexBuffer()** 可取消注册。

将资源注册到 CUDA 之后，即可在需要时分别使用 **cuD3D9MapResources() **和 **cuD3D9UnmapResources()**任意多次地映射和解除映射。内核可使用 **cuD3D9ResourceGetMappedPointer() **返回的设备存储器地址和**cuD3D9ResourceGetMappedSize()**、**cuD3D9ResourceGetMappedPitch()** 及**cuD3D9ResourceGetMappedPitchSlice() **返回的大小和间距信息来读取和写入已映射的资源。通过Direct3D 访问已映射的资源将导致不确定的结果。

下面的代码示例使用 0 填充了一个缓冲区：

CUdeviceptr devPtr;

cuD3D9ResourceGetMappedPointer(&devPtr, buffer);

size\_t size;

cuD3D9ResourceGetMappedSize(&size, buffer);

cuMemset(devPtr, 0, size);

在下面的代码示例中，每个线程都访问大小为 **(width, height) **的二维表面的一个像素，像素格式为 float4：

// host code

CUdeviceptr devPtr;

cuD3D9ResourceGetMappedPointer(&devPtr, surface);

size\_t pitch;

cuD3D9ResourceGetMappedPitch(&pitch, surface);

cuModuleGetFunction(&cuFunction, cuModule, “myKernel”);

cuFuncSetBlockShape(cuFunction, 16, 16, 1);

int offset = 0;

cuParamSeti(cuFunction, offset, devPtr);

offset += sizeof(devPtr);

cuParamSeti(cuFunction, 0, width);

offset += sizeof(width);

cuParamSeti(cuFunction, 0, height);

offset += sizeof(height);

cuParamSeti(cuFunction, 0, pitch);

offset += sizeof(pitch);

cuParamSetSize(cuFunction, offset);

cuLaunchGrid(cuFunction,

(width+Db.x–1)/Db.x, (height+Db.y–1)/Db.y);

// device code

\_\_global\_\_ void myKernel(unsigned char\* surface,

int width, int height, size\_t pitch)

int x = blockIdx.x \* blockDim.x + threadIdx.x;

int y = blockIdx.y \* blockDim.y + threadIdx.y;

if (x >= width || y >= height) return;

float\* pixel = (float\*)(surface + y \* pitch) + 4 \* x;

[CUDA]: http://blog.csdn.net/yanghangjun/article/details/5586240
[CUDA 1]: http://blog.csdn.net/yanghangjun/article/details/5586378
[CUDA 2]: http://blog.csdn.net/yanghangjun/article/details/5586883
[CUDA 3]: http://blog.csdn.net/yanghangjun/article/details/5586969
[CUDA 4]: http://blog.csdn.net/yanghangjun/article/details/5587269
[CUDA 5]: http://blog.csdn.net/yanghangjun/article/details/5587376
[CUDA 6]: http://blog.csdn.net/yanghangjun/article/details/5588284