CUDA Toolkit 12.4 引入了一个新的 nvFatbin 库，用于在运行时创建 fatbins。fatbins，也称为 NVIDIA 设备代码fat 二进制文件sm_61和sm_90。
到目前为止，要生成 fatbin，必须依赖命令行工具fatbinary，这不适合动态代码生成。这使得动态生成 fatbins 变得困难，因为您需要将生成的代码放入一个文件中，然后使用exec或类似命令调用fatbinary，并处理输出，这显著增加了动态生成 fatbins 的难度，并导致多次尝试通过各种容器模仿 fatbins。
CUDA Toolkit 12.4 引入了 nvFatbin，这是一个新的库，能够通过编程创建 fatbin，从而大大简化了这项任务，不再需要写入文件、调用exec、解析命令行输出和从目录中获取输出文件。
新库提供了运行时 fatbin 创建支持
使用 nvFatbin 库类似于任何其他熟悉的库，如NVRTC、nvPTXCompiler 和 nvJitLink。nvFatbin 库有静态和动态版本，适用于所有平台，这些平台都随 nvrtc 提供。
经过适当考虑，通过 nvFatbin 库创建的 Fatbin 符合 CUDA 兼容性保证。本文主要涵盖通过 nvFatbin 库的运行时 fatbin 创建，并在适当的时候强调与现有命令行 fatbinary 的差异。我们将通过代码示例、兼容性保证和优点深入了解该功能的细节。作为额外的奖励，我们还提供了NVIDIA TensorRT计划如何以及为什么利用该功能的预览。
如何使运行时 fatbin 创建正常工作
创建稍后要引用的句柄，以便将相关的设备代码插入到 fatbinary 中。
nvFatbinCreate(&handle, numOptions, options);
使用取决于输入类型的函数，添加要放入 fatbin 的设备代码。
nvFatbinAddCubin(handle, data, size, arch, name);
nvFatbinAddPTX(handle, data, size, arch, name, ptxOptions);
nvFatbinAddLTOIR(handle, data, size, arch, name, ltoirOptions);
对于 PTX 和LTO-IR（一种用于 JIT LTO 的中间表示形式），请在此处指定在 JIT 编译期间使用的其他选项。
检索得到的 fatbin。为此，显式分配一个缓冲区。执行此操作时，请确保查询生成的 fatbin 的大小，以确保分配了足够的空间。
nvFatbinSize(linker, &fatbinSize);
void* fatbin = malloc(fatbinSize);
nvFatbinGet(handle, fatbin);
清理把手。
nvFatbinDestroy(&handle);
使用 NVCC 离线生成 fatbins
要使用 NVCC 离线生成一个 fatbin，请添加选项-fatbin。例如，给定文件loader.cu，以下命令将生成一个 fatbin，其中包含一个用于sm_90的条目，该条目包含代码的 LTO-IR 版本，名为loader.fatbin。
nvcc -arch lto_90 -fatbin loader.cu
如果指定 -arch=sm_90，nvcc 将创建一个 fatbin，该 fatbin 同时包含 PTX 和 CUBIN（SASS）。该对象包含特定于sm_90的 SASS 指令和 PTX，以后可以对任何架构>=90 进行 JIT。
nvcc -arch sm_90 -fatbin loader.cu
要创建具有多个条目的 fatbin，请使用指定多个体系结构-gencode:
nvcc -gencode arch=compute_80,code=sm_80 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_52,code=compute_52
这将创建一个包含sm_80 ELF, sm_90 ELF和compute_52 PTX。您可以使用cuobjdump查看 fatbin 的内容。
在运行时生成 fatbins
除了前面描述的离线编译和运行时 fatbin 创建模型（图 1）外，还可以在运行时完全构建 fatbin，方法是使用 NVRTC 生成对象代码，然后使用 nvFatbin API 将它们添加到 fatbin。以下代码示例对使用 nvFatbin API 进行了相关修改。
#include <nvrtc.h>
#include <cuda.h>
#include <nvFatbin.h>
#include <nvrtc.h>
#include <iostream>
#define NUM_THREADS 128
#define NUM_BLOCKS 32
#define NVRTC_SAFE_CALL(x) \
do { \
nvrtcResult result = x; \
if (result != NVRTC_SUCCESS) { \
std::cerr << "\nerror: " #x " failed with error " \
<< nvrtcGetErrorString(result) << '\n'; \
exit(1); \
} \
} while(0)
#define CUDA_SAFE_CALL(x) \
do { \
CUresult result = x; \
if (result != CUDA_SUCCESS) { \
const char *msg; \
cuGetErrorName(result, &msg); \
std::cerr << "\nerror: " #x " failed with error " \
<< msg << '\n'; \
exit(1); \
} \
} while(0)
#define NVFATBIN_SAFE_CALL(x) \
do \
{ \
nvFatbinResult result = x; \
if (result != NVFATBIN_SUCCESS) \
{ \
std::cerr << "\nerror: " #x " failed with error " \
<< nvFatbinGetErrorString(result) << '\n';\
exit(1); \
} \
} while (0)
const char *fatbin_saxpy = " \n\
__device__ float compute(float a, float x, float y) { \n\
return a * x + y; \n\
} \n\
\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] = compute(a, x[tid], y[tid]); \n\
} \n\
} \n";
size_t process(const void* input, const char* input_name, void** output, const char* arch)
{
// Create an instance of nvrtcProgram with the code string.
nvrtcProgram prog;
NVRTC_SAFE_CALL(
nvrtcCreateProgram(&prog, // prog
(const char*) input, // buffer
input_name, // name
0, // numHeaders
NULL, // headers
NULL)); // includeNames
// specify that LTO IR should be generated for LTO operation
const char *opts[1];
opts[0] = arch;
nvrtcResult compileResult = nvrtcCompileProgram(prog, // prog
1, // numOptions
opts); // options
// Obtain compilation log from the program.
size_t logSize;
NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
char *log = new char[logSize];
NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
std::cout << log << '\n';
delete[] log;
if (compileResult != NVRTC_SUCCESS) {
exit(1);
}
// Obtain generated CUBIN from the program.
size_t CUBINSize;
NVRTC_SAFE_CALL(nvrtcGetCUBINSize(prog, &CUBINSize));
char *CUBIN = new char[CUBINSize];
NVRTC_SAFE_CALL(nvrtcGetCUBIN(prog, CUBIN));
// Destroy the program.
NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
*output = (void*) CUBIN;
return CUBINSize;
}
int main(int argc, char *argv[])
{
void* known = NULL;
size_t known_size = process(fatbin_saxpy, "fatbin_saxpy.cu", &known, "-arch=sm_52");
CUdevice cuDevice;
CUcontext context;
CUmodule module;
CUfunction kernel;
CUDA_SAFE_CALL(cuInit(0));
CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
// Dynamically determine the arch to make one of the entries of the fatbin with
int major = 0;
int minor = 0;
CUDA_SAFE_CALL(cuDeviceGetAttribute(&major,
CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevice));
CUDA_SAFE_CALL(cuDeviceGetAttribute(&minor,
CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevice));
int arch = major*10 + minor;
char smbuf[16];
sprintf(smbuf, "-arch=sm_%d", arch);
void* dynamic = NULL;
size_t dynamic_size = process(fatbin_saxpy, "fatbin_saxpy.cu", &dynamic, smbuf);
sprintf(smbuf, "%d", arch);
// Load the dynamic CUBIN and the statically known arch CUBIN
// and put them in a fatbin together.
nvFatbinHandle handle;
const char* fatbin_options[] = {"-cuda"};
NVFATBIN_SAFE_CALL(nvFatbinCreate(&handle, fatbin_options, 1));
NVFATBIN_SAFE_CALL(nvFatbinAddCubin(handle,
(void *)dynamic, dynamic_size, smbuf, "dynamic"));
NVFATBIN_SAFE_CALL(nvFatbinAddCubin(handle,
(void *)known, known_size, "52", "known"));
size_t fatbinSize;
NVFATBIN_SAFE_CALL(nvFatbinSize(handle, &fatbinSize));
void *fatbin = malloc(fatbinSize);
NVFATBIN_SAFE_CALL(nvFatbinGet(handle, fatbin));
NVFATBIN_SAFE_CALL(nvFatbinDestroy(&handle));
CUDA_SAFE_CALL(cuModuleLoadData(&module, fatbin));
CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "saxpy"));
// Generate input for execution, and create output buffers.
#define NUM_THREADS 128
#define NUM_BLOCKS 32
size_t n = NUM_THREADS * NUM_BLOCKS;
size_t bufferSize = n * sizeof(float);
float a = 5.1f;
float *hX = new float[n], *hY = new float[n], *hOut = new float[n];
for (size_t i = 0; i < n; ++i) {
hX[i] = static_cast<float>(i);
hY[i] = static_cast<float>(i * 2);
}
CUdeviceptr dX, dY, dOut;
CUDA_SAFE_CALL(cuMemAlloc(&dX, bufferSize));
CUDA_SAFE_CALL(cuMemAlloc(&dY, bufferSize));
CUDA_SAFE_CALL(cuMemAlloc(&dOut, bufferSize));
CUDA_SAFE_CALL(cuMemcpyHtoD(dX, hX, bufferSize));
CUDA_SAFE_CALL(cuMemcpyHtoD(dY, hY, bufferSize));
// Execute SAXPY.
void *args[] = { &a, &dX, &dY, &dOut, &n };
CUDA_SAFE_CALL(
cuLaunchKernel(kernel,
NUM_BLOCKS, 1, 1, // grid dim
NUM_THREADS, 1, 1, // block dim
0, NULL, // shared mem and stream
args, 0)); // arguments
CUDA_SAFE_CALL(cuCtxSynchronize());
// Retrieve and print output.
CUDA_SAFE_CALL(cuMemcpyDtoH(hOut, dOut, bufferSize));
for (size_t i = 0; i < n; ++i) {
std::cout << a << " * " << hX[i] << " + " << hY[i]
<< " = " << hOut[i] << '\n';
}
// Release resources.
CUDA_SAFE_CALL(cuMemFree(dX));
CUDA_SAFE_CALL(cuMemFree(dY));
CUDA_SAFE_CALL(cuMemFree(dOut));
CUDA_SAFE_CALL(cuModuleUnload(module));
CUDA_SAFE_CALL(cuCtxDestroy(context));
delete[] hX;
delete[] hY;
delete[] hOut;
// Release resources.
free(fatbin);
delete[] ((char*)known);
delete[] ((char*)dynamic);
return 0;
}
请参见nvFatbin以查看完整的示例。
nvFatbin 库直接从输入文件创建 fatbin，不进行任何链接或编译，也不依赖 CUDA 驱动程序，可以在没有 GPU 的系统上运行。
处理输入的是 nvFatbin 库的工具包版本，重要的是已编译输入的工具包版本。
nvFatbin 库保留了对旧输入的支持，无论版本如何。这并不取代驾驶员在装载所述版本时施加的任何限制，这些限制独立于使用 fatbin 作为容器格式。生成的输出 fatbin 仅与 nvFatbin 库的主版本相同或更高版本的 CUDA 驱动程序兼容。
此外，nvFatbin 可以处理来自较新 NVCC 或 NVRTC 的输入，只要它们在同一主要版本中。因此，目标系统上的 nvFatbin 库版本必须至少与用于生成任何输入的工具包的最新版本相同或更新。
例如，12.4 附带的 nvFatbin 可以支持任何 CUDA Toolkit 12.X 或更早版本生成的代码，但不能保证与 CUDA Toolkit13.0 或更高版本生成的任何代码一起使用。
离线工具 fatbinary 和 nvFatbin 都产生相同的输出文件类型，使用相同的输入类型，因此在线和离线工具在某些情况下可以互换使用。例如，NVCC 编译的 CUBIN 可以在运行时由 nvFatbin 放入 fatbin，而 NVRTC 编译的 CUBIN 可以由离线工具 fatbinary 离线放入 fatbin。这两个 fatbin 创建工具也遵循相同的兼容性规则。
NVIDIA 仅保证 nvFatbin 与使用相同或更低主要版本的 CUDA Toolkit 代码创建的输入兼容。如果您试图使用 nvFatbin 从 12.4 创建一个 fatbin，而 PTX 是在未来 CUDA Toolkit 13 版本中创建的，您可能会看到失败。然而，它应该支持与较旧的 CUDA 工具包，如 11.8 的输入兼容。
CUDA 次版本兼容性
如前所述，nvFatbin 库将与来自同一 CUDA 工具包主要版本的所有输入兼容，无论次要版本如何。这意味着 nvFatbin 的 12.4 版本将与 12.5 版本的输入兼容
一些新引入的功能将不适用于以前的版本，例如 fatbin 条目中添加了新类型。但是，任何格式只要在版本中已经被接受，就将继续被接受。
向后兼容性
nvFatbin 库支持来自 CUDA 工具包早期版本的输入。
更大的图景
既然有了所有主要编译器组件的运行时等价物，它们是如何相互作用的？
nvPTX 编译器
运行时 PTX 编译器 nvPTXCompiler 既是一个独立的工具，也集成到 NVRTC 和 nvJitLink 中以方便使用。它可以与 nvFatbin 一起使用，创建用于放入 fatbin 的 CUBIN。
NVRTC
运行时编译器 NVRTC 可用于编译 CUDA 程序，它支持 PTX 和 LTO-IR，以及 CUBIN，这是通过集成 nvPTXCompiler 实现的，尽管您可以手动生成 PTX，然后使用 nvPTXCompiler 来生成 CUBIN。所有这些结果格式都可以通过 nvFatbin 中放入一个 fatbin。
nvJitLink
运行时链接器 nvJitLink 可与 NVRTC 一起用于在运行时编译和链接 CUDA 程序。结果可以直接通过驱动程序 API 运行，也可以通过 nvFatbin 放入 fatbin 中。
随着 nvFatbin 的引入，动态生成灵活的库比以往任何时候都更容易。
TensorRT 希望存储用于现有体系结构的 CUBIN，以及用于未来体系结构的 PTX，这样，在保持兼容的同时，尽可能使用代码的优化版本。虽然对于未来的体系结构来说可能不是最佳的，但它确保了现有体系结构的最佳代码，并且仍将与未来的体系架构保持兼容。
在引入 nvFatbin 之前，您必须想出一种替代方法来处理这一问题，以避免不必要地将相关数据写入文件，从而导致重复工作来制作类似于 fatbin 的在线格式。
现在有了 nvFatbin，您和 TensorRT 背后的团队可以使用库来处理该操作，防止不必要的 I/O 操作，并避免使用自定义格式来存储带有 PTX 的 CUBIN。
