天天看点

CUDA学习(九十)

模型:

模块是可动态加载的设备代码和数据包,类似于Windows中的DLL,由nvcc输出(请参见使用NVCC编译)。 所有符号的名称(包括函数,全局变量和纹理或表面引用)都保存在模块范围内,以便独立第三方编写的模块可以在相同的CUDA上下文中进行互操作。

此代码示例加载模块并检索某个内核的句柄:

此代码示例从PTX代码编译和加载新模块并分析编译错误:

此代码示例编译,链接并加载来自多个PTX代码的新模块,并分析链接和编译错误:

内核执行:

cuLaunchKernel()使用给定的执行配置启动一个内核。

参数以指针数组的形式传递(在cuLaunchKernel()的最后一个参数旁边),其中第n个指针对应于第n个参数,并指向参数复制的内存区域,或者作为其中一个额外选项 cuLaunchKernel()的最后一个参数)。

当参数作为附加选项传递时(CU_LAUNCH_PARAM_BUFFER_POINTER选项),它们作为指向单个缓冲区的指针传递,其中通过匹配设备代码中每个参数类型的对齐要求,假定参数相对于彼此适当地偏移。

表3(之前博客中有)列出了内置矢量类型的设备代码中的对齐要求。对于所有其他基本类型,设备代码中的对齐要求与主机代码中的对齐要求相匹配,因此可以使用<code>__alignof()</code>来获取。 唯一的例外是主机编译器在单字边界而不是双字边界(例如,使用gcc的编译标志-mno-align-double)时将双精度和长精度对齐(在64位系统上长对齐) ),因为在设备代码中,这些类型总是在双字边界上对齐。

CUdeviceptr是一个整数,但代表一个指针,所以它的对齐要求是<code>__alignof(void *)</code>。

下面的代码示例使用宏(ALIGN_UP())来调整每个参数的偏移量以满足其对齐要求,并使用另一个宏(ADD_TO_PARAM_BUFFER())将每个参数添加到传递给CU_LAUNCH_PARAM_BUFFER_POINTER选项的参数缓冲区。

结构的对齐要求等于其字段的对齐要求的最大值。 因此,包含内置向量类型CUdeviceptr或非对齐double和long long的结构的对齐要求因此可能在设备代码和主机代码之间有所不同。 这样的结构也可能被填充不同。 例如,以下结构在主机代码中根本没有填充,但是由于对于字段f4的对齐要求为16,所以在字段f之后它填充在具有12个字节的设备代码中。

运行时和驱动程序API之间的互操作性:

应用程序可以将运行时API代码与驱动程序API代码混合。

如果通过驱动程序API创建上下文并使其最新,则后续运行时调用将接收此上下文,而不是创建新上下文。

如果运行时被初始化(隐式地如CUDA C运行时中所述),cuCtxGetCurrent()可用于检索初始化期间创建的上下文。 该上下文可以被后续的驱动程序API调用使用。

可以使用任何API分配和释放设备内存。 CUdeviceptr可以转换为常规指针,反之亦然:

特别是,这意味着使用驱动程序API编写的应用程序可以调用使用运行时API编写的库(例如cuFFT,cuBLAS,...)。

参考手册中设备和版本管理部分的所有功能可以互换使用。

CUDA学习(九十)

继续阅读