DAY86：阅读Kernel Execution

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第85天，我们正在讲解Driver API，希望在接下来的15天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

本文共计442字，阅读时间15分钟

I.3. Kernel Execution

cuLaunchKernel() launches a kernel with a given execution configuration.

Parameters are passed either as an array of pointers (next to last parameter of cuLaunchKernel()) where the nth pointer corresponds to the nth parameter and points to a region of memory from which the parameter is copied, or as one of the extra options (last parameter of cuLaunchKernel()).

When parameters are passed as an extra option (the CU_LAUNCH_PARAM_BUFFER_POINTER option), they are passed as a pointer to a single buffer where parameters are assumed to be properly offset with respect to each other by matching the alignment requirement for each parameter type in device code.

Alignment requirements in device code for the built-in vector types are listed in Table 3. For all other basic types, the alignment requirement in device code matches the alignment requirement in host code and can therefore be obtained using __alignof(). The only exception is when the host compiler aligns double and long long (and long on a 64-bit system) on a one-word boundary instead of a two-word boundary (for example, using gcc's compilation flag -mno-align-double) since in device code these types are always aligned on a two-word boundary.

CUdeviceptr is an integer, but represents a pointer, so its alignment requirement is __alignof(void*).

The following code sample uses a macro (ALIGN_UP()) to adjust the offset of each parameter to meet its alignment requirement and another macro (ADD_TO_PARAM_BUFFER()) to add each parameter to the parameter buffer passed to the CU_LAUNCH_PARAM_BUFFER_POINTER option.

The alignment requirement of a structure is equal to the maximum of the alignment requirements of its fields. The alignment requirement of a structure that contains built-in vector types,CUdeviceptr, or non-aligned double and long long, might therefore differ between device code and host code. Such a structure might also be padded differently. The following structure, for example, is not padded at all in host code, but it is padded in device code with 12 bytes after field f since the alignment requirement for field f4 is 16.

本文备注/经验分享：

本章节主要说明了cuLaunchKernel()函数， 如同你在<<<启动形状配置>>>(参数)（见之前的“启动配置”章节），本函数一样可以允许你指定启动形状配置，和kernel所需要的参数。因为历史的原因，如何设定kernel的参数有2个版本，其中今天的driver api用户，可以通过一个简化的方式来指定参数： 每个kernel的参数，被提前保存在某个缓冲区中。然后每个参数需要一个指向他们的指针，所有的参数被打包成一个指针的列表，最终你只需要提供该指针列表给cuLaunchKernel就可以了。这种方式比较灵活，甚至你可以考虑当参数比较零散的情况下，分布在内存中的各处，最终的指针列表可能指向的内容在内存中七零八落的。 而另外一种方式则是比较传统的方式，需要考虑每个kernel参数的大小，对齐方式等信息，然后将它们*必须*放置在一个缓冲区中, 这样在该缓冲区中，除了参数之间的因为对齐等因素，所引入的参数之间的位置空缺外，都是连续排列的， 这种方式是以前的唯一的一种方式，比较复杂，但是可能有更好的性能（微乎其微） 如果你采用了第一种的简化版本，driver api实际上是会采用第二种方式自动为你打包形成一个统一的整体的，现在的情况下应当不再使用第二种方式了。但是用户需要知道他的存在： （1）直接打包好的第二种方式，可能有更好的性能。 （2）在kernel的内部，实际上总是看到参数是连续排列的（除了之间可能的paddings之类的空白）。 实际上，这个函数版本已经简化了很多了，这是当年的CUDA 3.2 (v2 api)所引入的新功能之一： 之前的3.2之前的版本，存在两个重要的问题：一个是这种手工准备参数缓冲区的方式，必须需要用户知道设备上和host上的参数对齐，参数大小等信息才能手工处理。所以你会看到大量的类似ALIGN_UP之类的宏，非常麻烦。 第二个问题则是，设定kernel的启动形状(<<<>>>内部的），和参数（在小括号（）内部的），不是一步完成的： 这样存在多个线程设定同一个context下的kernel启动时候，所需要的在Host上用类似临界区之类的东西，保护这个过程的麻烦。CUDA 3.2起引入了v2版本的API，使得首先能一步到位（启动形状，shared memory, stream等等信息）， 其次刚才的段落中的第一种的方式的简化版本（只需要N个指针指向N个参数），方便了很多。随着现在的CPU性能的提升，几个简单的指针指向的内容的整理工作，已经不会对CPU或者kernel启动造成任何性能上的负担了，（除非你启动非常非常小的kernel，us的那种，然后大量的要求启动，同时一个慢速的CPU，但是这种应用场合，应当考虑是否可以将频繁的大量启动放到GPU本身上来完成，通过动态并行）。 你看到的本章节的后面的大幅段落，都是当年的老driver api的残留遗迹，注意这里的ALIGN_UP宏，还是写的很巧妙的，当年也是折腾死一帮人。 注意因为曾经我们在前几天的时候，说过OpenCL实际上来源于CUDA Driver API，NV在早期版本的CUDA中，刚刚有OpenCL的时候，还提供过一张对照表，里面有OpenCL的各个API，以及用法概念；和CUDA Driver API的详细对比，描述了它们的异同。不过现在NV已经将OpenCL的相关内容，完全从CUDA开发包中去掉了，用户现在已经看不到了他。现在的用户如果需要在N卡上开发OpenCL，请使用AMD或者Intel的开发包，不妨碍的。但是本函数的部分遗迹，还依然保留在现在的OpenCL中，并形成了OpenCL中唯一的，不是在Host CPU上，线程安全的地方，即OpenCL的kernel的参数设定（一个或者多个），在最终要求OpenCL启动Kernel（EnqueueNDRangeKernel）之间，必须使用Host上的锁，或者临界区之类的东西保护起来，这点和今天章节中的你看到的这个新版API（能同时设定形状和参数，无论参数你是用的新简化方式，还是老繁琐方式设定的）形成了鲜明的对比。CUDA已经在3.2的时代，专门的一次性解决了这个问题。而至今OpenCL还在五花八门的提供方案（包括拆成或者克隆成多个kernel对象，每个CPU线程一个，规避多个CPU线程同时启动同一个kernel），还是传统的需要锁去保护这个参数设定到启动的过程。这点实际上是之前我们论坛的OpenCL板块，大部分人能遭遇的唯一的OpenCL中的著名的坑。你看到CUDA已经在8年前就为你解决了这个问题，哪怕是你在使用Driver API。算是非常良心了。 所以这是为何我们一开头，就说，如果需要将三种著名的GPU上的开发方式进行比较的话，难度是OpenCL > CUDA Driver API > CUDA Runtime API的，这里多少可见一斑。 然后本章节的后面部分一般用户可以不看， 但是如果是在维护老代码，则最好还是要看一下的（3.2之前的老代码，现在还有么？） 一般情况下，后面的部分用不到的。 然后需要说明的是： （1）两种参数设定方式你只能设定一种，同时指定两种会导致kernel启动失败。 （2）如果一个kernel每次启动需要很长的参数列表（例如几十个），但是参数列表中每次只有轻微的参数的值发生变化，则可以考虑老方式设定，然后只需要改动这个缓冲区中的一小部分值即可。 这种方式也非常方便。 很多OpenCL用户，无师自通的发明了这点，虽然OpenCL里面只有分步的设定和启动，没有参数缓冲区的概念，但他们在脑海中构建了这么一个东西，在kernel的大量参数中，将不需要变动的参数抽取出来，提前设定好，然后每次启动kernel，只设定部分变化的参数，然后启动kernel。这样因为在实际的应用中，提前设定好的参数（一般在初始化过程中）和改变的参数相差较远，代码阅读者或者维护者会看到一种，只设定了不全的参数就开始要求启动kernel的奇特现象。 我个人是不建议这种方式的，每次我看到这种设定方式，都将残缺的参数从初始化的过程中抽取出来，单独的和kernel启动放置在一起。这种方式提升了很高的代码可维护性。同时几乎没有任何CPU上的性能损失。类似的，CUDA也不建议每次只更改那些变化的参数，而总是完全的指定。这样也是会提升代码可读性很多。 你可以试想一样，runtime api的<<<>>>()语法里面，参数列表可是齐全的，并没有参数。 再考虑一下，从当年引入这个新启动函数（2010）到今天都8年了，CPU的性能又提升了多少！ 所以尽量使用本章节前面的方式。至少你在出先问题的时候，能减少一点此方面的问题，除了这几天的Context，Module, Kernel启动，其实Driver API和Runtime API的区别并不大，除了新的Runtime没有的概念外（例如Context，Module等），更多的东西只是Runtime API中的，在Driver API中的另外一种用法而已。例如CUdeviceptr实际上只是一个普通的void *指针，还是比较容易入门的。但是这里需要说明一个重要的东西，特别是老代码（3.2）的维护软件， 3.2之前（忘记是否包含3.2了），是运行混合模式代码的，什么是混合模式，就是64-bit的CPU代码 + 32-bit的GPU代码，这种。 所以当年的CUdeviceptr并不完全的和Host上的指针(void *)等价,但是从当年的3.2左右起，为了简化用户的使用，取消了这个模式，32-bit的CPU代码必须配合32-bit的GPU代码使用，64-bit的CPU必须和64-bit的GPU代码使用，试图利用本章节和前几天章节的知识，构造一个用cuModuleLoad, 在64-bit的CPU进程中，载入32-bit的GPU代码的trick，已经无法成功。这其实当年引发了一定的争议，因为我们的GPU是RISC，并不像CPU那样有专门的快速便捷的地址计算。所有的地址计算都需要普通的SP，进行普通加减乘逻辑运算等。32-bit的地址变换可能有时候只需要1条指令，但是64-bit的地址变换，可能需要2条（包括从Maxwell引入的辅助地址计算的LEA指令也是如此，一次只能算一半）。这导致了很著名的64-bit代码损失性能的问题。这种损失直到今天也有，主要是两点： (1)每个kernel里面的每个指针使用，原本只需要占用1个寄存器的，现在需要占用2个连续的寄存器（因为GPU的通用寄存器是32-bit的） （2）刚才说的增加了运算量。 但是好处也有两点： （1）使用Driver API之类的用户，不再需要考虑32-bit/64-bit的CPU/Host不等的混合问题。 （2）现在的高性能通用计算往往需要大量的数据，任何超过4GB的数据总量（例如显存使用）以前32-bit下都很为难，现在可以随意的用了。甚至12GB，24GB大显存用满的kernel，也很常见。 各有利弊吧。

有不明白的地方，请在本文后留言