CUDA编程之NVRTC——simpleAssert_nvrtc

NVRTC

NVRTC是一个CUDA C++的运行时编译库。它接受以字符串形式表示的CUDA C++源代码,并创建可用于获取PTX的句柄。

由NVRTC生成的PTX字符串可以通过cuModuleLoadData和cuModuleLoadDataEx来加载,并通过CUDA驱动API的cuLinkAddData与其他模块链接。

NVRTC提供了一种在运行时编译和链接CUDA C++代码的方法,这可以带来更好的优化和性能。它是CUDA工具链的一部分,补充了离线编译的功能。

在没有NVRTC(或CUDA中没有运行时编译支持)的情况下,如果用户希望在应用程序或库中实现运行时编译,他们需要在运行时生成一个单独的进程来执行 nvcc。但是这种方法存在以下缺点:

1. 编译开销往往比必要的更高。
2. 因为在linux里设定NVCC编译环境比较麻烦，有坑，而通过上面的对于NVRTC的叙述，我们发现整个编译过程如果利用NVRTC的话会很简单，坑少。

3. 不用我们在编译器里设定NVCC编译环境了。

NVRTC通过提供一个库接口来解决这些问题,消除了生成单独进程、磁盘I/O等相关的开销,同时保持应用程序部署的简单性。

1. 功能:

   • NVRTC提供了在运行时编译和链接CUDA C++代码的功能。
   • 允许开发者动态地编译和使用CUDA内核,而不是在编译时就将它们固定下来。
   • 这种灵活性可以带来更好的优化和性能。

2. 工作流程:

   • 开发者将CUDA C++代码以字符串的形式传递给NVRTC。
   • NVRTC会编译这些代码,并生成可以被CUDA驱动API使用的PTX(并行线程执行)代码。
   • 生成的PTX代码可以通过CUDA驱动API进行加载和链接,从而在运行时执行。

3. 应用场景:

   • 动态生成和执行CUDA内核,以适应不同的输入数据和环境。
   • 在不重新编译整个应用程序的情况下,进行CUDA内核的更新和优化。
   • 在机器学习等领域,根据模型的变化动态编译内核代码。

4. 优势:

   • 提高了代码的灵活性和可扩展性。
   • 可以在运行时针对特定情况进行更好的优化和性能调整。
   • 减少了静态编译所需的时间和资源。

NVRTC为CUDA开发者提供了一种强大的运行时编译工具,可以显著提升应用程序的灵活性和性能。

NVRTC(NVIDIA Runtime Compilation)是CUDA工具包发布的一部分,这些组件在CUDA工具包安装目录中的组织结构如下:

On Windows:

include\nvrtc.h

bin\nvrtc64_Major Release VersionMinor Release Version_0.dll

bin\nvrtc-builtins64_Major Release VersionMinor Release Version.dll

lib\x64\nvrtc.lib

lib\x64\nvrtc_static.lib

lib\x64\nvrtc-builtins_static.lib

doc\pdf\NVRTC_User_Guide.pdf

On Linux:

include/nvrtc.h

lib64/libnvrtc.so

lib64/libnvrtc.so.Major Release Version.Minor Release Version

lib64/libnvrtc.so.Major Release Version.Minor Release Version.<build version>

lib64/libnvrtc-builtins.so

lib64/libnvrtc-builtins.so.Major Release Version.Minor Release Version

lib64/libnvrtc-builtins.so.Major Release Version.Minor Release Version.<build version>

lib64/libnvrtc_static.a

lib64/libnvrtc-builtins_static.a

doc/pdf/NVRTC_User_Guide.pdf

操作系统要求:

• Linux x86_64
• Linux ppc64le
• Linux aarch64
• Windows x86_64

GPU要求:

• SM>=2.0

其他要求:

• 安装CUDA Toolkit
• 安装CUDA driver

1.jpg 2.jpg

CUDA Sample里使用了NVRTC的sample有这些：

3.jpg

这里以simpleAssert_nvrtc为例进行代码讲解。

#include <stdio.h>
#include <cassert>

// Includes CUDA
#include <cuda_runtime.h>
#include "nvrtc_helper.h"

// Utilities and timing functions
#include <helper_functions.h>  // includes cuda.h and cuda_runtime_api.h

const char *sampleName = "simpleAssert_nvrtc";


// Auto-Verification Code
bool testResult = true;


// Declaration, forward
void runTest(int argc, char **argv);


// Program main


int main(int argc, char **argv) {
  printf("%s starting...\n", sampleName);

  runTest(argc, argv);

  exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
}

void runTest(int argc, char **argv) {
  int Nblocks = 2;
  int Nthreads = 32;

  // Kernel configuration, where a one-dimensional
  // grid and one-dimensional blocks are configured.

  dim3 dimGrid(Nblocks);
  dim3 dimBlock(Nthreads);

  printf("Launch kernel to generate assertion failures\n");
  char *cubin, *kernel_file;
  size_t cubinSize;

  kernel_file = sdkFindFilePath("simpleAssert_kernel.cu", argv[0]);
  compileFileToCUBIN(kernel_file, argc, argv, &cubin, &cubinSize, 0);

  CUmodule module = loadCUBIN(cubin, argc, argv);
  CUfunction kernel_addr;

  checkCudaErrors(cuModuleGetFunction(&kernel_addr, module, "testKernel"));

  int count = 60;
  void *args[] = {(void *)&count};

  checkCudaErrors(cuLaunchKernel(
      kernel_addr, dimGrid.x, dimGrid.y, dimGrid.z, /* grid dim */
      dimBlock.x, dimBlock.y, dimBlock.z,           /* block dim */
      0, 0,                                         /* shared mem, stream */
      &args[0],                                     /* arguments */
      0));

  // Synchronize (flushes assert output).
  printf("\n-- Begin assert output\n\n");
  CUresult res = cuCtxSynchronize();

  printf("\n-- End assert output\n\n");

  // Check for errors and failed asserts in asynchronous kernel launch.
  if (res == CUDA_ERROR_ASSERT) {
    printf("Device assert failed as expected\n");
  }

  testResult = res == CUDA_ERROR_ASSERT;
}

具体代码解释：

上面这段代码的主要功能和流程:

1. 包含必要的头文件,包括CUDA相关的头文件、辅助函数和计时工具。

2. 定义了一个样本名称常量 sampleName。

3. 定义了一个 testResult 布尔变量,用于记录测试结果。

4. 声明了一个 runTest 函数,用于执行测试。

5. main 函数中主要内容:

   • 打印样本名称,表示程序开始运行。
   • 调用 runTest 函数执行测试。
   • 根据 testResult 的值,退出程序,返回success或failure状态。

6. runTest 函数的实现:

   • 设置 kernel 的网格和块的尺寸。
   • 查找并编译 CUDA kernel 文件到 CUBIN 格式。
   • 加载 CUBIN 文件,并获取 kernel 函数的地址。
   • 设置 kernel 的参数,并启动 kernel 执行。
   • 同步 CUDA 操作,并打印 assert 输出。
   • 检查 kernel 执行过程中是否有 assert 失败,并更新 testResult 变量。

这个程序演示了如何使用 NVRTC 在运行时编译和执行一个 CUDA kernel,该 kernel 会产生一些 assert 失败。它展示了如何检查 CUDA kernel 中的 assert 输出,并验证测试结果。

运行结果：

程序中N的默认值是60，所以有4个test fail

simpleAssert_nvrtc starting...
Launch kernel to generate assertion failures
sdkFindFilePath <simpleAssert_kernel.cu> in ./
> Using CUDA Device [0]: NVIDIA GeForce RTX 4080
> Using CUDA Device [0]: NVIDIA GeForce RTX 4080
> GPU Device has SM 8.9 compute capability

-- Begin assert output

./simpleAssert_kernel.cu:37: void testKernel(int): block: [1,0,0], thread: [28,0,0] Assertion `gtid < N` failed.
./simpleAssert_kernel.cu:37: void testKernel(int): block: [1,0,0], thread: [29,0,0] Assertion `gtid < N` failed.
./simpleAssert_kernel.cu:37: void testKernel(int): block: [1,0,0], thread: [30,0,0] Assertion `gtid < N` failed.
./simpleAssert_kernel.cu:37: void testKernel(int): block: [1,0,0], thread: [31,0,0] Assertion `gtid < N` failed.

-- End assert output

Device assert failed as expected