使用 NVIDIA CUDA Toolkit 12.4 編譯器創建運行時 Fatbin

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計劃如何以及為什么利用該功能的預覽。

CUDA C++ can be passed to either NVRTC or NVCC, each of which can produce PTX, CUBINs, or LTOIR from the CUDA C++ code. These PTX, CUBINs, and LTOIR can be passed into either nvFatbin or fatbinary, both of which can produce fatbins from the inputs. — *圖 1。與使用 nvFatbin 庫相比，現有的 fatbinary 命令行工具*

如何使運行時 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。

使用 NVIDIA CUDA Toolkit 12.4 編譯器創建運行時 Fatbin

新庫提供了運行時 fatbin 創建支持

如何使運行時 fatbin 創建正常工作

使用 NVCC 離線生成 fatbins

在運行時生成 fatbins

CUDA 次版本兼容性

向后兼容性

更大的圖景

nvPTX 編譯器

NVRTC

nvJitLink

相關資源

標簽

關于作者

使用 NVIDIA CUDA Toolkit 12.4 編譯器創建運行時 Fatbin

新庫提供了運行時 fatbin 創建支持

如何使運行時 fatbin 創建正常工作

使用 NVCC 離線生成 fatbins

在運行時生成 fatbins

CUDA 次版本兼容性

向后兼容性

更大的圖景

nvPTX 編譯器

NVRTC

nvJitLink

相關資源

標簽

關于作者

相關文章

CUDA 運行時中的動態加載機制

借助 NVIDIA TensorRT 輕量級引擎，為 AI 應用提供卓越性能和最小占用空間

相關文章

借助 NVIDIA DriveOS LLM SDK 簡化自動駕駛汽車應用的 LLM 部署

衡量 AI Guardrails 在生成式 AI 應用中的有效性和性能

聚焦：NAVER Place 利用 NVIDIA TensorRT-LLM 優化 SLM 基礎的垂直服務

使用 DeepSeek-R1 NIM 構建具有專家推理功能的 AI 智能體

圖像和視頻理解的視覺語言模型提示工程實踐指南