nvrtc教程（全篇无废话）

前言

文章使用的cuda版本：12.4
nvrtc（nvidia runtime compilation），是一个用来实现代码运行时动态编译.cu文件的库。使用这个库并不需要配置cuda环境（但仍需要cuda toolkit），因此对于那些死活配置不好cuda环境的（比如我），这可以是另一种使用kernel的办法（但没有直接nvcc来得方便）。
关于nvrtc的使用，如果有能力的话可以参考官方文档

准备工作

首先需要安装cuda toolkit，配置好include目录（默认目录应该是C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\include）和lib目录（默认目录应该是C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.4\lib\x64）以及链接lib（干脆把lib目录下所有lib都链接了吧）。

需要的头文件：

#include<stdio.h>
#include<stdlib.h>
#include<nvrtc.h>
#include<cuda_runtime.h>
#include<cuda.h>

nvrtc.h的函数会返回一个值，可以用来判断函数是否成功运行，如果成功就会返回NVRTC_SUCCESS，返回类型为nvrtcResult，并且可以用nvrtcGetErrorString函数来获取这个错误的字符串形式。写一个简单的宏来实现这个功能

#define NVRTC_SAFE_CALL(x)\
    if(1)\
    {\
        nvrtcResult result = x;\
        if (result != NVRTC_SUCCESS)\
        {\
            printf("\nerror: " #x " failed with error %s\n", nvrtcGetErrorString(result));\
            system("pause");\
            exit(1);\
        }\
    }\

cuda.h和cuda_runtime.h的函数也有类似的返回值CUDA_SUCCESS，同样写一个宏函数：

#define CUDA_SAFE_CALL(x)\
    if(1)\
    {\
        CUresult result = x;\
        if (result != CUDA_SUCCESS)\
        {\
            const char *msg;\
            cuGetErrorName(result, &msg);\
            printf("\nerror: " #x " failed with error %s\n",msg);\
            system("pause");\
            exit(1);\
        }\
    }\

之后每次使用cuda的函数，都会调用这些宏，比如：

NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, source,  0,0,0,0));

此外，在主函数中应初始化cuda：

void CuInit()
{
	//指定设备并创建context
    CUdevice cuDevice;
    CUcontext context;
    CUDA_SAFE_CALL(cuInit(0));
    CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
    CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
}

把.cu文件转化为字符串

nvrtc要求输入的是字符串风格的.cu代码，比如下面这种：

const char *saxpy = "                                           \n\
extern \"C\" __global__                                         \n\
void saxpy(float a, float *x, float *y, float *out, size_t n)   \n\
{                                                               \n\
  size_t tid = blockIdx.x * blockDim.x + threadIdx.x;           \n\
  if (tid < n) {                                                \n\
    out[tid] = a * x[tid] + y[tid];                             \n\
  }                                                             \n\
}  

所以需要把.cu文件转化为字符串，写简单一个函数实现这个功能：

//当size参数没有缺省的话，size返回该字符串的长度
char* LoadFile(const char* path, int* size = 0)
{
    if(!size)
    {
        int _size;
        size = &_size;
    }
	FILE* fp = fopen(path, "rb");
	fseek(fp, 0, 2);
	*size = ftell(fp);
	fseek(fp, 0, 0);
	char* p = (char*)malloc(*size + 1);
	fread(p, 1, *size, fp);
	p[*size] = '\0';
	return p;
}

然后，就可以在主函数里面调用这个函数：

    int size;
    const char* path = "path/to/your/.cu";
    char* source = LoadFile(path, &size);

创建program并编译

获取字符串风格后的.cu文件后，调用nvrtcCreateProgram()创建program并调用nvrtcCompileProgram()编译program：

    nvrtcProgram prog; 
    //nvrtcCreateProgram()参数：
    //第一个为program
    //第二个为字符串风格的.cu
    //第三个是这个文件的文件名（只是一个供cuda内部使用的标识，可以通过cubin找寻这个标识获取这个文件，可以为0）
    //第四，五，六个分别为头文件的数目，include目录，include文件的文件名（具体怎么设置，靠大佬解答了）
    NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, source,  0,0,0,0));
    free(source);
    if(nvrtcCompileProgram(prog, 0,0) != NVRTC_SUCCESS)
    {
        size_t logSize;
        //nvrtcGetProgramLog获取错误日志，nvrtcGetProgramLogSize获取错误日志大小
        NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
        char *log = (char*)malloc(logSize);
        NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
        printf("%s",log);
        free(log);
        system("pause");
        exit(1);
    }

编译为ptx文件

接着从program获取ptx代码:

    size_t ptxSize;
    NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
    char *ptx = (char*)malloc(ptxSize);
    NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));
    NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));

获取kernel

    CUmodule module;
    CUfunction kernel;
    CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));//需要先把ptx加载为module才能获取kernel
    CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, functionName));

运行函数

	//参数：
	//第一个为kernel
	//之后六个参数为grid和block的size
	//接着两个参数为共享内存的大小和工作流的指定
	//args即为传递的参数，等会讲；最后的0我也不知道
	CUDA_SAFE_CALL(cuLaunchKernel(kernel, grid[0],grid[1],grid[2], block[0],block[1],block[2], 0,0, args,0));

至此，我们可以把上述封装成头函数了：

#include<stdio.h>
#include<stdlib.h>
#include<nvrtc.h>
#include<cuda_runtime.h>
#include<cuda.h>

#define NVRTC_SAFE_CALL(x)\
    if(1)\
    {\
        nvrtcResult result = x;\
        if (result != NVRTC_SUCCESS)\
        {\
            printf("\nerror: " #x " failed with error %s\n", nvrtcGetErrorString(result));\
            system("pause");\
            exit(1);\
        }\
    }\


#define CUDA_SAFE_CALL(x)\
    if(1)\
    {\
        CUresult result = x;\
        if (result != CUDA_SUCCESS)\
        {\
            const char *msg;\
            cuGetErrorName(result, &msg);\
            printf("\nerror: " #x " failed with error %s\n",msg);\
            system("pause");\
            exit(1);\
        }\
    }\
   
char* LoadFile(const char* path, int* size = 0)
{
    if(!size)
    {
        int _size;
        size = &_size;
    }
	FILE* fp = fopen(path, "rb");
	fseek(fp, 0, 2);
	*size = ftell(fp);
	fseek(fp, 0, 0);
	char* p = (char*)malloc(*size + 1);
	fread(p, 1, *size, fp);
	p[*size] = '\0';
	return p;
}

void CuInit()
{
    CUdevice cuDevice;
    CUcontext context;
    CUDA_SAFE_CALL(cuInit(0));
    CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
    CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
}

CUfunction GetKernel(const char* path, const char* functionName)
{
    int size;
    char* source = LoadFile(path, &size);
    nvrtcProgram prog; 
    NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, source,  0,0,0,0));
    free(source);
    if(nvrtcCompileProgram(prog, 0,0) != NVRTC_SUCCESS)
    {
        size_t logSize;
        NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
        char *log = (char*)malloc(logSize);
        NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
        printf("%s",log);
        free(log);
        system("pause");
        exit(1);
    }
    size_t ptxSize;
    NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
    char *ptx = (char*)malloc(ptxSize);
    NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));
    NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
    CUmodule module;
    CUfunction kernel;
    CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
    CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, functionName));
    return kernel;
}

void UseKernel(const char* path, const char* functionName, void** args, int* grid, int* block)
{
    CUfunction kernel = GetKernel("1.cu", "test");
    CUDA_SAFE_CALL(cuLaunchKernel(kernel, grid[0],grid[1],grid[2], block[0],block[1],block[2], 0,0, args,0));
}

args设置

cuLaunchKernel()要求的args是一个数组，包含了需要传递参数的指针，比如像这样的：

void* args[] = {&d_out, &a,&x,&b,&y};

如果要传递常量，那么只需在主存创建该变量，然后传递指向在主存上该变量的指针；如果要传递数组，那么首先要把在主存上的数组数据拷贝到显存上，接着传递指向在显存上该数组的指针。

此外，kernel返回类型必须void，这意味着返回值需要在参数列表里面，然后把结果从显存拷贝回内存。

举个例子，我要为下面的kernel传递参数（kernel前面记得加 extern “C”）：

extern "C" __global__
void test(int* out, int a, int b, int* coord)
{
    out[0] = a*coord[0]+b*coord[1];
}                                                                        

那么，我在我的主函数里面要这么调用：

#include<stdio.h>
#include<stdlib.h>
#include<LoadCuda.h>
#include<cuda.h>

int grid[3] = {2,2,1};
int block[3] = {2,2,2};

int main()
{   
    CuInit();
    int a=1,b=2;
    int coord[] = {3,4};
    CUdeviceptr d_coord;//创建一个显存指针
    cuMemAlloc(&d_coord, 2*sizeof(int));//分配一块显存并让d_coord指向该显存
    cuMemcpyHtoD(d_coord, coord, 2*sizeof(int));//把主存上coord指向的数据传给d_coord指向的区域
    CUdeviceptr d_out;
    cuMemAlloc(&d_out, sizeof(int));//返回值也要分配显存
    void* args[] = {&d_out, &a,&b,&d_coord};
    UseKernel("1.cu", "test", args, grid, block);
    int out;
    cuMemcpyDtoH(&out, d_out, 4);//把数据传回主存
    printf("%d\n",out);
    system("pause");
    return 0;
}

至此，你应该学会如何基本使用nvrtc.h来加载kernel了