在这里简单总结一下driver API的大致特点,主要是解释其中几个重要的概念。

下面是driver API中几个常见的术语:

cuda dirver API总结

dirver API的核心为定义了一系列object,通过handle来显式控制object。

context
context类比CPU process,resource以及其他action被封装在context中,每个CPU thread每个时刻只能有一个current context。context可通过cuCtxCreate()来创建,cuCtxCreate()会将创建的context push到该host thread的context stack的top,使其成为当前host thread的current context。同时,每个context都有一个counter,记录其被多少个thread引用。counter在cuCtxCreate()时初始化为1,之后每被一个thread通过cuCtxAttach()引用一次,则usage count加1,每被cuCtxDetach()一次,count减1,当减为0时该context会被销毁。
与cuCtxCreate()功能相反的一个函数为: cuCtxDestroy();
通过cuCtxPopCurrent()可以将当前host thread的current context解耦成为float context,被解耦的context之后可被其他thread设置为current current(通过cuCtxPushCurrent())。

如下图所示:
cuda dirver API总结

module
module是一种可以动态加载的package,类似dll,计算需要的function(kernel)以及global variables等,因此,在发起一个kernel前,需要首先将包含该kernel的module加载到context中,然后再从module中获取该kernel的句柄(通过函数名)。
如下图所示:

cuda dirver API总结

注意:module一般为已经编译好的ptx文件或者cubin文件,因此需要先将源代码通过nvcc编译得到ptx或者cubin文件。
ptx相比cubin的优点在于:可移植。二进制代码是architecture-specific的,而ptx代码则是在加载时通过driver实时编译的,因此可移植。

kernel launch
kernel launch对应的函数为:cuLaunchKernel();
其中function对应的参数可以通过在指定kernel维度等信息后通过一系列执行memory的pointer指定,也可通过kernel的最后一个参数extra指定。通过extra指定时需要考虑对齐问题。
例如:
cuda dirver API总结

Interoperability between Runtime and Driver APIs
dirver API和runtime API可以混合使用。比如:通过driver API创建的context可被之后的runtime API使用而不会重新创建一个新的context,同样,runtime API初始化后的context之后可被driver API通过调用cuCtxGetCurrent()使用。
memory分配和释放可以采用两种API中的任意一种,由driver API创建的memory handle CUdeviceptr可被映射为regular runtime pointer,反过来,runtime创建的pointer也可被映射为一个driver handle。如下所示:
cuda dirver API总结

driver API和runtime API的混合使用的意义在于:dirver API编写的程序可以方便地调用runtime library,如cuBLAS等。

相关文章: