CUDA学习（九十）

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

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

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

#define BUFFER_SIZE 8192
CUmodule 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 8192
CUmodule 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);

内核执行：
cuLaunchKernel（）使用给定的执行配置启动一个内核。
参数以指针数组的形式传递（在cuLaunchKernel（）的最后一个参数旁边），其中第n个指针对应于第n个参数，并指向参数复制的内存区域，或者作为其中一个额外选项 cuLaunchKernel（）的最后一个参数）。
当参数作为附加选项传递时（CU_LAUNCH_PARAM_BUFFER_POINTER选项），它们作为指向单个缓冲区的指针传递，其中通过匹配设备代码中每个参数类型的对齐要求，假定参数相对于彼此适当地偏移。
表3（之前博客中有）列出了内置矢量类型的设备代码中的对齐要求。对于所有其他基本类型，设备代码中的对齐要求与主机代码中的对齐要求相匹配，因此可以使用__alignof（）来获取。 唯一的例外是主机编译器在单字边界而不是双字边界（例如，使用gcc的编译标志-mno-align-double）时将双精度和长精度对齐（在64位系统上长对齐） ），因为在设备代码中，这些类型总是在双字边界上对齐。
CUdeviceptr是一个整数，但代表一个指针，所以它的对齐要求是__alignof（void *）。
下面的代码示例使用宏（ALIGN_UP（））来调整每个参数的偏移量以满足其对齐要求，并使用另一个宏（ADD_TO_PARAM_BUFFER（））将每个参数添加到传递给CU_LAUNCH_PARAM_BUFFER_POINTER选项的参数缓冲区。

#define ALIGN_UP(offset, alignment) \
(offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)
char paramBuffer[1024];
size_t paramBufferSize = 0;
#define ADD_TO_PARAM_BUFFER(value, alignment) \
do { \
paramBufferSize = ALIGN_UP(paramBufferSize, alignment); \
memcpy(paramBuffer + paramBufferSize, \
&(value), sizeof(value)); \
paramBufferSize += sizeof(value); \
} while (0)
int i;
ADD_TO_PARAM_BUFFER(i, __alignof(i));
float4 f4;
ADD_TO_PARAM_BUFFER(f4, 16); // float4's alignment is 16
char c;
ADD_TO_PARAM_BUFFER(c, __alignof(c));
float f;
ADD_TO_PARAM_BUFFER(f, __alignof(f));
CUdeviceptr devPtr;
ADD_TO_PARAM_BUFFER(devPtr, __alignof(devPtr));
float2 f2;
ADD_TO_PARAM_BUFFER(f2, 8); // float2's alignment is 8
void* extra[] = {
    CU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
    CU_LAUNCH_PARAM_BUFFER_SIZE, &paramBufferSize,
    CU_LAUNCH_PARAM_END
};
cuLaunchKernel(cuFunction,
    blockWidth, blockHeight, blockDepth,
    gridWidth, gridHeight, gridDepth,
    0, 0, 0, extra);

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

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

运行时和驱动程序API之间的互操作性:
应用程序可以将运行时API代码与驱动程序API代码混合。
如果通过驱动程序API创建上下文并使其最新，则后续运行时调用将接收此上下文，而不是创建新上下文。
如果运行时被初始化（隐式地如CUDA C运行时中所述），cuCtxGetCurrent（）可用于检索初始化期间创建的上下文。 该上下文可以被后续的驱动程序API调用使用。
可以使用任何API分配和释放设备内存。 CUdeviceptr可以转换为常规指针，反之亦然：

CUdeviceptr devPtr;
float* d_data;
// Allocation using driver API
cuMemAlloc(&devPtr, size);
d_data = (float*)devPtr;
// Allocation using runtime API
cudaMalloc(&d_data, size);
devPtr = (CUdeviceptr)d_data;

特别是，这意味着使用驱动程序API编写的应用程序可以调用使用运行时API编写的库（例如cuFFT，cuBLAS，...）。
参考手册中设备和版本管理部分的所有功能可以互换使用。