CUDA学习（六十四）

API错误和启动失败：
像通常的CUDA运行时一样，任何函数都可能返回一个错误代码。记录返回的最后一个错误代码，并可通过cudaGetLastError（）调用检索。每个线程都会记录错误，以便每个线程都可以识别它生成的最新错误。错误代码类型为cudaError_t。
与主机端启动类似，设备端启动可能由于许多原因（无效参数等）而失败。用户必须调用cudaGetLastError（）来确定启动是否生成错误，但启动后没有错误并不意味着子内核已成功完成。
对于设备端异常（例如，访问无效地址），子网格中的错误将返回给主机，而不是由父节点对cudaDeviceSynchronize（）的调用返回。
启动安装程序API：
内核启动是通过设备运行时库公开的系统级机制，因此可以通过底层的cudaGetParameterBuffer（）和cudaLaunchDevice（）API直接从PTX获取。允许CUDA应用程序自己调用这些API，其要求与PTX相同。在这两种情况下，用户都有责任根据规范以正确的格式正确填充所有必要的数据结构。这些数据结构保证向后兼容性。
与主机端启动一样，设备端运算符<<< >>>映射到底层内核启动API。这样一来，定位PTX的用户就可以实现启动，并且编译器前端可以将<<< >>>转换为这些调用。
CUDA学习（六十四）

这些启动功能的API与CUDA运行时API的API不同，其定义如下：

extern device cudaError_t cudaGetParameterBuffer(void **params);
extern __device__ cudaError_t cudaLaunchDevice(void *kernel,
    void *params, dim3 gridDim,
    dim3 blockDim,
    unsigned int sharedMemSize = 0,
    cudaStream_t stream = 0);

API参考：
设备运行时支持的CUDA Runtime API部分在此处详述。主机和设备运行时API具有相同的语法; 语义是相同的，除非指出。下表提供了相对于主机可用版本的API概述。
CUDA学习（六十四）

设备端从PTX启动：
本部分面向编程语言和编译器实现者，他们针对并行线程执行（PTX）并计划在其语言中支持动态并行。它提供了与在PTX级别支持内核启动相关的低级细节。
设备端内核启动可以使用以下两个可从PTX访问的API实现：cudaLaunchDevice（）和cudaGetParameterBuffer（）。 cudaLaunchDevice（）使用通过调用cudaGetParameterBuffer（）获得的参数缓冲区启动指定的内核，并将参数填充到启动的内核。参数缓冲区可以是NULL，即，如果启动的内核不采用任何参数，则不需要调用cudaGetParameterBuffer（）。
cudaLaunchDevice：
在PTX级别，cudaLaunchDevice（）需要在使用前以如下所示的两种形式之一进行声明。

// PTX-level Declaration of cudaLaunchDevice() when .address_size is 64
.extern.func(.param.b32 func_retval0) cudaLaunchDevice
(
    .param.b64 func,
    .param.b64 parameterBuffer,
    .param.align 4.b8 gridDimension[12],
    .param.align 4.b8 blockDimension[12],
    .param.b32 sharedMemSize,
    .param.b64 stream
)
;
// PTX-level Declaration of cudaLaunchDevice() when .address_size is 32
.extern.func(.param.b32 func_retval0) cudaLaunchDevice
(
    .param.b32 func,
    .param.b32 parameterBuffer,
    .param.align 4.b8 gridDimension[12],
    .param.align 4.b8 blockDimension[12],
    .param.b32 sharedMemSize,
    .param.b32 stream
);

下面的CUDA级别声明映射到上述PTX级别声明之一，并在系统头文件cuda_device_runtime_api.h中找到。该函数在cudadevrt系统库中定义，该库必须与程序链接以便使用设备端内核启动功能。

// CUDA-level declaration of cudaLaunchDevice()
extern "C" __device__
cudaError_t cudaLaunchDevice(void *func, void *parameterBuffer,
    dim3 gridDimension, dim3 blockDimension,
    unsigned int sharedMemSize,
    cudaStream_t stream);

第一个参数是一个指向要启动的内核的指针，第二个参数是参数缓冲区，用于保存启动的内核的实际参数。参数缓冲区的布局在下面的参数缓冲区布局中说明。其他参数指定启动配置，即网格维度，块维度，共享内存大小以及与启动相关的流（有关启动配置的详细说明，请参阅“执行配置”。
cudaGetParameterBuffer：
cudaGetParameterBuffer（）需要在使用前在PTX级别声明。根据地址大小，PTX级声明必须采用以下两种形式之一：

// PTX-level Declaration of cudaGetParameterBuffer() when .address_size is 64
// When .address_size is 64
.extern.func(.param.b64 func_retval0) cudaGetParameterBuffer
(
    .param.b64 alignment,
    .param.b64 size
);
// PTX-level Declaration of cudaGetParameterBuffer() when .address_size is 32
.extern.func(.param.b32 func_retval0) cudaGetParameterBuffer
(
    .param.b32 alignment,
    .param.b32 size
);

以下CUDA级别的cudaGetParameterBuffer（）声明被映射到前述的PTX级声明：

// CUDA-level Declaration of cudaGetParameterBuffer()
extern "C" __device__
void *cudaGetParameterBuffer(size_t alignment, size_t size);

第一个参数指定参数缓冲区的对齐要求，第二个参数指定字节的大小要求。在当前实现中，由cudaGetParameterBuffer（）返回的参数缓冲区始终保证为64字节对齐，并且对齐要求参数将被忽略。但是，建议将正确的对齐要求值（这是要放置在参数缓冲区中的任何参数的最大对齐方向）传递给cudaGetParameterBuffer（）以确保将来可移植。
参数缓冲区布局：
禁止参数缓冲区中的参数重新排序，并且要求对参数缓冲区中放置的每个参数进行对齐。也就是说，每个参数必须放置在参数缓冲区中的第n个字节处，其中n是参数大小的最小倍数，大于前一个参数获取的最后一个字节的偏移量。参数缓冲区的最大大小是4KB。
有关由CUDA编译器生成的PTX代码的更详细说明，请参阅PTX-3.5规范 CUDA学习（六十四）