如何在CUDA中使用驱动程序API

网友投稿 378 2022-10-23

如何在CUDA中使用驱动程序API

它是一个基于句柄的命令式 API:大多数对象都由不透明的句柄引用,这些句柄可以指定给函数来操作对象。

在调用驱动程序 API 的任何函数之前,必须使用cuInit()初始化驱动程序 API。 然后必须创建一个附加到特定设备的 CUDA 上下文,并使其成为当前调用主机线程,如上下文中所述。

任何想要在未来设备架构上运行的应用程序都必须加载 PTX,而不是二进制代码。 这是因为二进制代码是特定于体系结构的,因此与未来的体系结构不兼容,而 PTX 代码在加载时由设备驱动程序编译为二进制代码。

以下是使用驱动程序 API 编写的内核示例的主机代码:

完整的代码可以在 vectorAddDrv CUDA 示例中找到。

L.1. Context

主机线程一次可能只有一个设备上下文当前。当使用 cuCtxCreate() 创建上下文时,它对调用主机线程是当前的。如果有效上下文不是线程当前的,则在上下文中操作的 CUDA 函数(大多数不涉及设备枚举或上下文管理的函数)将返回 CUDA_ERROR_INVALID_CONTEXT。

每个主机线程都有一堆当前上下文。 cuCtxCreate() 将新上下文推送到堆栈顶部。可以调用 cuCtxPopCurrent() 将上下文与主机线程分离。然后上下文是“浮动的”,并且可以作为任何主机线程的当前上下文推送。 cuCtxPopCurrent() 还会恢复先前的当前上下文(如果有)。

还为每个上下文维护使用计数。 cuCtxCreate() 创建使用计数为 1 的上下文。cuCtxAttach() 增加使用计数,而 cuCtxDetach() 减少使用计数。当调用 cuCtxDetach() 或 cuCtxDestroy() 时使用计数变为 0,上下文将被销毁。

驱动程序 API 可与运行时互操作,并且可以通过 cuDevicePrimaryCtxRetain() 从驱动程序 API 访问由运行时管理的主上下文(参见初始化)。

使用计数有助于在相同上下文中运行的第三方编写的代码之间的互操作性。例如,如果加载三个库以使用相同的上下文,则每个库将调用 cuCtxAttach() 来增加使用计数,并在库使用上下文完成时调用 cuCtxDetach() 来减少使用计数。对于大多数库,预计应用程序会在加载或初始化库之前创建上下文;这样,应用程序可以使用自己的启发式方法创建上下文,并且库只需对传递给它的上下文进行操作。希望创建自己的上下文的库(可能会或可能没有创建自己的上下文的 API 客户端不知道)将使用 cuCtxPushCurrent() 和 cuCtxPopCurrent(),如下图所示。

L.2. Module

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

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

CUmodule cuModule;cuModuleLoad(&cuModule, "myModule.ptx");CUfunction myKernel;cuModuleGetFunction(&myKernel, cuModule, "MyKernel");

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

#define BUFFER_SIZE 8192CUmodule cuModule;CUjit_option options[3];void* values[3];char* PTXCode = "some PTX code";char error_log[BUFFER_SIZE];int err;options[0] = CU_JIT_ERROR_LOG_BUFFER;values[0] = (void*)error_log;options[1] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;values[1] = (void*)BUFFER_SIZE;options[2] = CU_JIT_TARGET_FROM_CUCONTEXT;values[2] = 0;err = cuModuleLoadDataEx(&cuModule, PTXCode, 3, options, values);if (err != CUDA_SUCCESS) printf("Link error:\n%s\n", error_log);

此代码示例从多个 PTX 代码编译、链接和加载新模块,并解析链接和编译错误:

#define BUFFER_SIZE 8192CUmodule cuModule;CUjit_option options[6];void* values[6];float walltime;char error_log[BUFFER_SIZE], info_log[BUFFER_SIZE];char* PTXCode0 = "some PTX code";char* PTXCode1 = "some other PTX code";CUlinkState linkState;int err;void* cubin;size_t cubinSize;options[0] = CU_JIT_WALL_TIME;values[0] = (void*)&walltime;options[1] = CU_JIT_INFO_LOG_BUFFER;values[1] = (void*)info_log;options[2] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;values[2] = (void*)BUFFER_SIZE;options[3] = CU_JIT_ERROR_LOG_BUFFER;values[3] = (void*)error_log;options[4] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;values[4] = (void*)BUFFER_SIZE;options[5] = CU_JIT_LOG_VERBOSE;values[5] = (void*)1;cuLinkCreate(6, options, values, &linkState);err = cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void*)PTXCode0, strlen(PTXCode0) + 1, 0, 0, 0, 0);if (err != CUDA_SUCCESS) printf("Link error:\n%s\n", error_log);err = cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void*)PTXCode1, strlen(PTXCode1) + 1, 0, 0, 0, 0);if (err != CUDA_SUCCESS) printf("Link error:\n%s\n", error_log);cuLinkComplete(linkState, &cubin, &cubinSize);printf("Link completed in %fms. Linker Output:\n%s\n", walltime, info_log);cuModuleLoadData(cuModule, cubin);cuLinkDestroy(linkState);

完整的代码可以在 ptxjit CUDA 示例中找到。

L.3. Kernel Execution

cuLaunchKernel() 启动具有给定执行配置的内核。

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

表 4 列出了内置向量类型的设备代码中的对齐要求。对于所有其他基本类型,设备代码中的对齐要求与主机代码中的对齐要求相匹配,因此可以使用 __alignof() 获得。唯一的例外是当宿主编译器在一个字边界而不是两个字边界上对齐 double 和 long long(在 64 位系统上为 long)(例如,使用 gcc 的编译标志 -mno-align-double ) 因为在设备代码中,这些类型总是在两个字的边界上对齐。

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

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

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

typedef struct { float f; float4 f4;} myStruct;

版权声明:本文内容由网络用户投稿,版权归原作者所有,本站不拥有其著作权,亦不承担相应法律责任。如果您发现本站中有涉嫌抄袭或描述失实的内容,请联系我们jiasou666@gmail.com 处理,核实后本网站将在24小时内删除侵权内容。

上一篇:谷歌经纬度查询(谷歌经纬度查询位置工具)
下一篇:云端技能包 | 百亿级日志之云原生实时流实战(1)
相关文章

 发表评论

暂时没有评论,来抢沙发吧~