7.2 语法

在使用CUDA运行时的时候，内核启动由我们常用的三对尖角括号语法来指定。

Kernel<<gridSize,BlockSize,sharedMem,Stream>>(Parameters...)

其中，Kernel指定待启动的内核名称；

gridSize指定一个dim3结构形式的网格的大小;

blockSize指定了一个dim3结构形式的线程块的维度；

sharedMem指定为每个线程块预留的附加的共享内存 [1]；

Stream指定内核启动所属于的流。

通常，指定网格和块大小的dim3结构包含有3个成员变量（x、y和z）。当使用C++进行编译时，一个带有系统默认参数的构造函数会将x、y和z均默认初始化为1。详见代码清单7-1，摘自NVIDIA SDK中的vector_types.h头文件。

代码清单7-1 dim3结构

struct _device_builtin _dim3
{
    unsigned int x, y, z;
    #if defined(_cplusplus)
        _host _device _dim3(
            unsigned int vx = 1,
            unsigned int vy = 1,
        unsigned int vx = 1) : x(vx), y(vy), z(vz) {}
        _host _device _dim3( uint3 v) : x(v.x), y(v.y), z(v.z) {}
        _host _device _operator uint3(void) {
            uint3 t;
            t.x = x;
            t.y = y;
            t.z = z;
            return t;
}

内核程序可以通过驱动程序API使用cuLaunchKernel()来启动，虽然cuLaunchKernel()将网格和线程块的维度作为独立参数而不是dim3结构。

CUresult cuLaunchKernel ( CUfunction kernel, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void \*\*kernelParams, void \*\*extra   
);

对于三对尖角括号语法，cuLaunchKenel()的参数包括了调用的内核、网格和线程块大小、共享内存数目以及流。其主要的区别在于内核本身的这些参数如何给出：因为ptexas发射出的内核微码包含了描述各个内核参数的元数据[2]，所以kernelParams是一个void*类型的数组，每个元素对应内核的一个参数。又因为参数的类型可以被驱动程序识别，所以参数所占的内存数据（int为4字节、double为8字节，

诸如此类）将会作为用来调用内核的、具有硬件特性的命令的一部分复制到命令缓冲区中。

7.2.1 局限性

所有参与内核启动的C++类都必须是带有以下特性的“简单旧数据”（plain old data, POD）。

·无用户声明的构造函数;
·无用户定义的复制分配操作符；
·无用户定义的析构函数;

无非POD的非静态数据成员;
无私有或保护型的非静态数据;
·无基类;
·无虚函数。

注意，违背了这些规则的类也可以用于CUDA，甚至是CUDA内核程序中。只是，它们绝对不能使用于内核启动中。如此，一个CUDA内核使用的类可以通过使用来自内核启动的POD输入数据来构造。

CUDA内核程序也没有返回值。它们必须通过设备内存（必须显式地复制回CPU）或映射主机内存来传回结果。

7.2.2 高速缓存和一致性

GPU包含了多个高速缓存以便于在发生重用时加速计算。其中，常量缓存（constant cache）被充分利用起来以便于广播式传输到同一个SM的执行单元；而纹理缓存则减少了外部带宽的使用。这两种缓存都不能很好地同GPU的写内存操作保持一致性。例如，没有促使这些缓存和一、二级缓存之间保持一致性的协议，所以无法减少全局内存的延迟和支持聚合带宽。这意味着两件事：

1）当一个内核在运行时，注意不要对那些同时（或者是被一个并发运行的内核）正在通过常量内存和纹理内存进行访问的内存执行写内存的操作。
2）CUDA驱动程序必须在每个内核启动之前使常量缓存和纹理缓存无效。

对于不含TEX指令的内核程序，CUDA驱动程序不需要使常量缓存和纹理缓存无效。因此，未使用纹理的内核程序引发更少的驱动程序开销。

7.2.3 异步与错误处理

内核启动是异步的，因为一旦内核被提交到硬件就会立即与CPU并行执行 [3]。这一异步会使错误处理变得很复杂。如果一个内核遇到一个错误（例如，它读入一个无效内存位置），该错误有时会在内核启动后的一段时间才能传输到驱动程序（和应用程序）中。而检查这种错误最可靠的方法是使用CUDADeviceSynchronize()或cuCtxSynchronize()函数以使其与GPU同步。如果在内核执行中出现一个错误，它们将会返回“unspecified launch failure”的错误代码。

除了CUDADeviceSynchronize()或cuCtxSynchronize()等显式CPU/GPU同步函数外，这个错误代码还可能是来自与CPU隐式同步的函数，如同步的内存复制调用。

无效内核启动

有可能所请求的内核启动无法为硬件执行。例如，指定比硬件能够支持的块内线程数还要多的线程就会出现这种情况。驱动程序会尽可能检测到这些情况并报告错误，而不是试图将该启动提交给硬件。

CUDA运行时和驱动程序API以不同的方式处理这种情况。当一个无效的参数被指定时，驱动程序API的显式API调用（例如

cuLaunchGrid()和cuLaunchKernel()等函数）返回错误代码。但是，当使用CUDA运行时的时候，由于内核是按C/C++一行代码启动的，故而没有API调用来返回错误代码。对应地，那个错误会被“记录”到本地线程槽中，而应用程序则可以通过CUDAGetLastError()函数来查询该错误值。与此相同的错误处理机制还被用在因其他原因（例如，内存访问冲突）导致的无效内核启动上。

7.2.4 超时

GPU在CUDA内核执行期间不能进行上下文切换，所以长时间运行的CUDA内核可能会对使用GPU来与用户交互的系统的交互性产生负面影响。因此，如果GPU运行太久而没有进行上下文切换，许多CUDA系统将实施“超时”以重置它。

在WDDM（Windows Display Driver Model）上，超时是由操作系统执行的。微软已经阐明这个“超时检测和恢复”（Timeout Detection and Recovery，TDR）是如何工作的，详见http://bit.ly/WPPSdQ（其中包括控制TDR行为的注册表项）[4]。虽然特斯拉计算集群（Tesla Compute Cluster，TCC）驱动程序并非适用于所有硬件，但TDR确实可以通过它进行安全禁用。

在Linux上，NVIDIA的驱动程序默认2秒超时。未用于显示的第二块GPU并不实施超时。开发者可以查询一个运行时限制是否在一个给定

GPU上启用，具体通过调用带

CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT标志的

cuDeviceGetAttribute()函数，或检查

CUDADeviceProp::kernelExecTimeoutEnabled属性来实现。

7.2.5 本地内存

由于本地内存是线程私有的，且CUDA中的网格可以包含数千个线程，因此CUDA网格需要的本地内存数目是相当多的。CUDA开发者会用心预先分配资源以降低内核启动等操作因为缺乏资源而失败的可能性。但就本地内存而言，仅一个保守的分配就已经消耗了太多的内存。缘于此，使用大量本地内存的内核需要更长的时间且有可能是同步执行的，因为CUDA驱动程序必须在执行内核启动之前分配好内存。此外，如果内存分配失败，内核也将会因缺乏资源而启动失败。

默认情况下，当CUDA驱动程序必须分配本地内存以运行内核时，它会在内核完成之后释放内存。另外，这种行为还会使得内核启动同步。但这种行为可以通过给cuCtxCreate()指定

CUCTX_LMEM_RESIZE_TO_MAX标志或在创建主上下文之前调用

CUDASetDeviceFlags()函数置CUDADeviceLmemResizeToMax来禁止。它带来的结果是，在一个需要比默认更多本地内存的内核启动之后，增加的可用本地内存不会被释放。

7.2.6 共享内存

共享内存会在内核启动时进行分配，并且它会在内核执行期间而一直保持。除了可以在内核中声明的静态分配方法之外，共享内存也可以被声明为一个未指定大小的extern变量。此时，用于分配大小不定的数组的共享内存量由内核启动的第三个参数或者是cuLaunchKernel()的sharedMemBytes参数指定。

[1] 内核程序可用的共享内存的数目为在此参数值上加上内核中静态申请的共享内存数目。
[2] cuLaunchKernel()不适用于那些没有在CUDA3.2或更高版本上编译的二进制映像文件，因为CUDA3.2是第一个可包含内核参数元数据的版本。
[3] 在大多数平台上，内核会在CPU处理完启动命令之后的数个微秒开始在GPU执行。但在WDDM（Windows Display Driver Model）平台上，它可能需要更长的时间。因为驱动程序必须执行一个内核转换以便将启动提交到硬件中，并且在用户模式下GPU的工作任务会进入队列，平摊用户态到内核态过渡的开销。
[4] 当然，修改注册表只应在测试时使用。

7.2_语法