如何将选项 CU_JIT_LTO 与 CUDA JIT 链接一起使用？
How to use the option CU_JIT_LTO with CUDA JIT linking?
我想知道是否可以使用选项 CU_JIT_LTO 在实时 (JIT) linking 期间改进 link 时间优化 (LTO)。如果是这样，我该如何指定这个选项？
我在一个NVIDIA开发者博客上找到了下面的代码，但是我不明白为什么walltime要给CU_JIT_LTO。博客中未定义 walltime 变量。当我尝试类似的东西时，它对我的内核性能没有影响。
options[0] = CU_JIT_LTO;
values[0] = (void*)&walltime;
...
cuLinkCreate(..., options, values, &linkState);
来源：https://developer.nvidia.com/blog/discovering-new-features-in-cuda-11-4/
我的示例案例使用输入选项 CU_JIT_INPUT_NVVM 到 link 对象，这些对象是使用 LTO 标志（-dlto 或 -code=lto_80）创建的。 linker 似乎已经做了一些 LTO，因为内核比 linking 对象文件没有 LTO 的“表现更好”，但不如 linking 使用 NVCC 的 LTO 好。 （查看示例案例了解详细结果和讨论）
示例案例
为了检查 link 时间优化 (LTO) 的有效性，我使用 4 种不同的方法创建了一个简单的程序，并使用每个线程的寄存器数量作为指标。这在我的系统上给出了以下结果（OS：ubuntu 20.04，CUDA 工具包：11.5.1，NVIDIA 驱动程序：495.44，GPU：NVIDIA RTX 3080）。
method registers/thread
Create program using a single translation file : 30
Link files using NVCC without link time optimization : 44
Link files using NVCC with link time optimization : 30
Link files using NVRTC/JIT with link time optimization : 38
结果解读：
从单个翻译文件创建程序应该会得到最好的结果。编译器可以看到所有的函数实现，并用它来优化内核。这导致 30 registers/thread.
使用 NVCC 与 LTO 链接绝对有效。它使用与从单个 .cu 文件编译的程序相同数量的寄存器 (30)，如果没有 LTO（使用 44 个寄存器）则不同。
使用 NVRTC/JIT 链接文件比 link 使用没有 LTO 的 NVCC 做得“更好”（当我们只关注寄存器使用时），但不如 link ING 与 NVCC 与 LTO。内核使用 38 registers/thread.
注意：我的目标不是减少寄存器的使用，我只是把它作为一个指标。因为来自单个翻译文件的程序使用 30 registers/thread，我假设完全优化的 linked 程序将具有相同的“最终可执行代码”，因此使用相同数量的寄存器。因为事实并非如此，所以我开始查看 JIT 选项。
CU_JIT_LTO 选项：
我试图用 JIT_option CU_JIT_LTO 进一步优化 NVRTC/JIT 中的 linking。但是，我不确定如何使用此选项。我尝试了以下两种方式（有关更多上下文，请参见下面的文件 cuda 代码。linking 的代码从第 41 行开始）：
方法 1：将选项 CU_JIT_LTO 添加到 cuLinkCreate(...)。这似乎没有效果。当 int lto = 0 和 int lto = 1.
时，代码使用相同数量的寄存器
方法 2：在 cuLinkAddFile(...) 和 cuLinkAddData(...) 中添加选项 CU_JIT_LTO。这会立即给出错误 CUDA_ERROR_INVALID_VALUE.
那么现在我的问题是：应该如何使用选项 CU_JIT_LTO？
文件：
下面是两个文件。按照以下步骤 运行 示例（在 linux OS 上）：
- 将cuda代码保存在扩展名为.cu的文件中（例如：code.cu）
- 将 bash 脚本保存在文件中（例如：run.sh）
- 运行 来自终端的命令：
bash run.sh code.cu
cuda代码：
#include <iostream>
#include <stdio.h>
#ifdef RTC
#include <cuda.h>
#include <nvrtc.h>
#define NVRTC_CHECK(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_CHECK(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)
CUmodule compileModule(std::string program)
{
// Compile nvvm from program string ===============
nvrtcProgram prog;
NVRTC_CHECK(nvrtcCreateProgram(&prog, program.c_str(), "programRTC.cu", 0, NULL, NULL));
const char* opts[] = {"-arch=compute_80", "-dlto", "-dc"};
nvrtcResult compileResult = nvrtcCompileProgram(prog, 3, opts);
// Obtain NVVM from the program.
size_t nvvmSize;
NVRTC_CHECK(nvrtcGetNVVMSize(prog, &nvvmSize));
char* nvvm = new char[nvvmSize];
NVRTC_CHECK(nvrtcGetNVVM(prog, nvvm));
// Link files ===============
CUlinkState linker;
// ARE THE OPTIONS SPECIFIED CORRECTLY?
int lto = 1;
CUjit_option options[] = {CU_JIT_LTO};
void* values[] = {(void*)<o};
// METHOD 1: GIVE THE OPTIONS TO 'cuLinkCreate(...)'
// -> HAS NO EFFECT ON THE AMOUNT OF REGISTERS USED
// -------------------------------------------------------------------------------------------
// CUDA_CHECK(cuLinkCreate(0, NULL, NULL, &linker));
CUDA_CHECK(cuLinkCreate(1, options, values, &linker));
// -------------------------------------------------------------------------------------------
// METHOD 2: GIVE THE OPTIONS TO 'cuLinkAddFile(...)' and 'cuLinkAddData(...)'
// -> FUNCTION FAILS WITH ERROR 'CUDA_ERROR_INVALID_VALUE'
// -------------------------------------------------------------------------------------------
CUDA_CHECK(cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", 0, NULL, NULL));
CUDA_CHECK(cuLinkAddData(linker, CU_JIT_INPUT_NVVM, (void*)nvvm, nvvmSize, "programRTC.o", 0,
NULL, NULL));
// CUDA_CHECK(cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", 1, options, values));
// CUDA_CHECK(cuLinkAddData(linker, CU_JIT_INPUT_NVVM, (void*)nvvm, nvvmSize, "programRTC.o", 1,
// options, values));
// -------------------------------------------------------------------------------------------
// Create module ===============
void* cubin;
CUmodule module;
CUDA_CHECK(cuLinkComplete(linker, &cubin, NULL));
CUDA_CHECK(cuModuleLoadDataEx(&module, cubin, 0, NULL, NULL));
// Cleanup
NVRTC_CHECK(nvrtcDestroyProgram(&prog));
CUDA_CHECK(cuLinkDestroy(linker));
return module;
}
#endif // RTC
__device__ double func(double a, double b);
#ifdef FUNC
__device__ double func(double a, double b)
{
return pow(a, b);
}
#endif
#ifdef MAIN
#ifdef RTC
std::string the_program = R"===(
__device__ double func(double a, double b);
extern "C" __global__ void kernel(double* out, double* a, double* b)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid >= 1){
return;
}
a[tid] = 2;
b[tid] = 3;
out[tid] = func(a[tid], b[tid]);
printf("out[%lu] = %f\n", tid, out[tid]);
})===";
#else // RTC
__global__ void kernel(double* out, double* a, double* b)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= 1) {
return;
}
a[tid] = 2;
b[tid] = 3;
out[tid] = func(a[tid], b[tid]);
printf("out[%lu] = %f\n", tid, out[tid]);
}
#endif // RTC
int main()
{
double* a;
double* b;
double* out;
cudaMalloc((void**)&a, sizeof(double));
cudaMalloc((void**)&b, sizeof(double));
cudaMalloc((void**)&out, sizeof(double));
#ifdef RTC
// Create context
CUdevice cuDevice;
CUcontext context;
CUDA_CHECK(cuInit(0));
CUDA_CHECK(cuDeviceGet(&cuDevice, 0));
CUDA_CHECK(cuCtxCreate(&context, 0, cuDevice));
CUmodule module = compileModule(the_program);
CUfunction kernel;
CUDA_CHECK(cuModuleGetFunction(&kernel, module, "kernel"));
size_t n_blocks = 1;
size_t n_threads = 1;
void* args[] = {&out, &a, &b};
CUDA_CHECK(cuLaunchKernel(kernel, n_blocks, 1, 1, // grid dim
n_threads, 1, 1, // block dim
0, NULL, // shared mem and stream
args, 0)); // arguments
CUDA_CHECK(cuCtxSynchronize());
// Cleanup
CUDA_CHECK(cuModuleUnload(module));
CUDA_CHECK(cuCtxDestroy(context));
#else // RTC
kernel<<<1, 1>>>(out, a, b);
cudaDeviceSynchronize();
#endif // RTC
return 0;
}
#endif // MAIN
bash 脚本:
#!/bin/bash
set -e # stop script when an error occurs
SCRIPT=
xCCx=80 # CUDA compute compatibility
# Create program using a single translation file
echo -e "\n---------- main_single ----------\n"
nvcc -DFUNC -DMAIN $SCRIPT -o main_single \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_single # should print 'out[0] = 8.0'
cuobjdump main_single -res-usage | grep kernel -A1
# Link files using NVCC without link time optimization (code=compute_...)
echo -e "\n---------- main_link_nvcc ----------\n"
nvcc -DFUNC $SCRIPT -o func_link_nvcc.o -dc \
-gencode arch=compute_$xCCx,code=compute_$xCCx
nvcc -DMAIN $SCRIPT -o main_link_nvcc.o -dc \
-gencode arch=compute_$xCCx,code=compute_$xCCx
nvcc func_link_nvcc.o main_link_nvcc.o -o main_link_nvcc \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_link_nvcc # should print 'out[0] = 8.0'
cuobjdump main_link_nvcc -res-usage | grep kernel -A1
# Link files using NVCC with link time optimization (code=lto_...)
echo -e "\n---------- main_link_nvcc_lto ----------\n"
nvcc -DFUNC $SCRIPT -o func_link_nvcc_lto.o -dc \
-gencode arch=compute_$xCCx,code=lto_$xCCx
nvcc -DMAIN $SCRIPT -o main_link_nvcc_lto.o -dc \
-gencode arch=compute_$xCCx,code=lto_$xCCx
nvcc func_link_nvcc_lto.o main_link_nvcc_lto.o -o main_link_nvcc_lto -dlto \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_link_nvcc_lto # should print 'out[0] = 8.0'
cuobjdump main_link_nvcc_lto -res-usage | grep kernel -A1
# Link files using NVRTC with link time optimization
echo -e "\n---------- main_link_nvrtc_lto ----------\n"
nvcc -DFUNC $SCRIPT -o func_link_nvrtc_lto.o -dc -ptx \
-gencode arch=compute_$xCCx,code=lto_$xCCx
nvcc -DMAIN -DRTC $SCRIPT -o main_link_nvrtc_lto \
-lnvrtc_static -lnvrtc-builtins_static -lnvptxcompiler_static -lcuda -lpthread \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_link_nvrtc_lto # should print 'out[0] = 8.0'
ncu main_link_nvrtc_lto | grep register/thread
# Registers/thread used on my system with an NVIDIA RTX 3080:
# main_single : 30 registers/thread
# main_link_nvcc : 44 registers/thread
# main_link_nvcc_lto : 30 registers/thread
# main_link_nvrtc_lto : 38 registers/thread
子问题：生成一个 NVVM IR 文件
为了生成与命令 cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", ...) 一起使用的文件 func_link_nvrtc_lto.o，我必须添加 -ptx 标志，如下面的命令所示。我没有在任何文档中找到这个，而是通过尝试和错误找到的。我想知道是否有更好的方法来生成这样的文件。
nvcc -DFUNC $SCRIPT -o func_link_nvrtc_lto.o -dc -ptx \
-gencode arch=compute_$xCCx,code=lto_$xCCx
首先，不幸的是博客 post 中有一个错误 CU_JIT_LTO 值。它应该是：
values[0] = (void*)1;
然而，这并不重要，因为该值被忽略了——它只是使用了 CU_JIT_LTO 的存在。
正如您发现的那样，CU_JIT_LTO 确实应该传递给 cuLinkCreate。
对于您的 sub-question，您的 -ptx 所做的是在生成 nvvm-ir 后停止编译，但这是一个未记录的 side-effect。更简单、更安全的方法是使用：
nvcc -dc -arch=compute_XX,code=lto_XX
创建一个包含 nvvm-ir 的宿主对象。然后将其传递为：
CU_JIT_INPUT_OBJECT to cuLinkAddFile().
我想知道是否可以使用选项 CU_JIT_LTO 在实时 (JIT) linking 期间改进 link 时间优化 (LTO)。如果是这样，我该如何指定这个选项？
我在一个NVIDIA开发者博客上找到了下面的代码，但是我不明白为什么walltime要给CU_JIT_LTO。博客中未定义 walltime 变量。当我尝试类似的东西时，它对我的内核性能没有影响。
options[0] = CU_JIT_LTO;
values[0] = (void*)&walltime;
...
cuLinkCreate(..., options, values, &linkState);
来源：https://developer.nvidia.com/blog/discovering-new-features-in-cuda-11-4/
我的示例案例使用输入选项 CU_JIT_INPUT_NVVM 到 link 对象，这些对象是使用 LTO 标志（-dlto 或 -code=lto_80）创建的。 linker 似乎已经做了一些 LTO，因为内核比 linking 对象文件没有 LTO 的“表现更好”，但不如 linking 使用 NVCC 的 LTO 好。 （查看示例案例了解详细结果和讨论）
示例案例
为了检查 link 时间优化 (LTO) 的有效性，我使用 4 种不同的方法创建了一个简单的程序，并使用每个线程的寄存器数量作为指标。这在我的系统上给出了以下结果（OS：ubuntu 20.04，CUDA 工具包：11.5.1，NVIDIA 驱动程序：495.44，GPU：NVIDIA RTX 3080）。
method registers/thread
Create program using a single translation file : 30
Link files using NVCC without link time optimization : 44
Link files using NVCC with link time optimization : 30
Link files using NVRTC/JIT with link time optimization : 38
结果解读：
从单个翻译文件创建程序应该会得到最好的结果。编译器可以看到所有的函数实现，并用它来优化内核。这导致 30 registers/thread.
使用 NVCC 与 LTO 链接绝对有效。它使用与从单个 .cu 文件编译的程序相同数量的寄存器 (30)，如果没有 LTO（使用 44 个寄存器）则不同。
使用 NVRTC/JIT 链接文件比 link 使用没有 LTO 的 NVCC 做得“更好”（当我们只关注寄存器使用时），但不如 link ING 与 NVCC 与 LTO。内核使用 38 registers/thread.
注意：我的目标不是减少寄存器的使用，我只是把它作为一个指标。因为来自单个翻译文件的程序使用 30 registers/thread，我假设完全优化的 linked 程序将具有相同的“最终可执行代码”，因此使用相同数量的寄存器。因为事实并非如此，所以我开始查看 JIT 选项。
CU_JIT_LTO 选项：
我试图用 JIT_option CU_JIT_LTO 进一步优化 NVRTC/JIT 中的 linking。但是，我不确定如何使用此选项。我尝试了以下两种方式（有关更多上下文，请参见下面的文件 cuda 代码。linking 的代码从第 41 行开始）：
方法 1：将选项 CU_JIT_LTO 添加到 cuLinkCreate(...)。这似乎没有效果。当 int lto = 0 和 int lto = 1.
方法 2：在 cuLinkAddFile(...) 和 cuLinkAddData(...) 中添加选项 CU_JIT_LTO。这会立即给出错误 CUDA_ERROR_INVALID_VALUE.
那么现在我的问题是：应该如何使用选项 CU_JIT_LTO？
文件：
下面是两个文件。按照以下步骤 运行 示例（在 linux OS 上）：
- 将cuda代码保存在扩展名为.cu的文件中（例如：code.cu）
- 将 bash 脚本保存在文件中（例如：run.sh）
- 运行 来自终端的命令：
bash run.sh code.cu
cuda代码：
#include <iostream>
#include <stdio.h>
#ifdef RTC
#include <cuda.h>
#include <nvrtc.h>
#define NVRTC_CHECK(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_CHECK(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)
CUmodule compileModule(std::string program)
{
// Compile nvvm from program string ===============
nvrtcProgram prog;
NVRTC_CHECK(nvrtcCreateProgram(&prog, program.c_str(), "programRTC.cu", 0, NULL, NULL));
const char* opts[] = {"-arch=compute_80", "-dlto", "-dc"};
nvrtcResult compileResult = nvrtcCompileProgram(prog, 3, opts);
// Obtain NVVM from the program.
size_t nvvmSize;
NVRTC_CHECK(nvrtcGetNVVMSize(prog, &nvvmSize));
char* nvvm = new char[nvvmSize];
NVRTC_CHECK(nvrtcGetNVVM(prog, nvvm));
// Link files ===============
CUlinkState linker;
// ARE THE OPTIONS SPECIFIED CORRECTLY?
int lto = 1;
CUjit_option options[] = {CU_JIT_LTO};
void* values[] = {(void*)<o};
// METHOD 1: GIVE THE OPTIONS TO 'cuLinkCreate(...)'
// -> HAS NO EFFECT ON THE AMOUNT OF REGISTERS USED
// -------------------------------------------------------------------------------------------
// CUDA_CHECK(cuLinkCreate(0, NULL, NULL, &linker));
CUDA_CHECK(cuLinkCreate(1, options, values, &linker));
// -------------------------------------------------------------------------------------------
// METHOD 2: GIVE THE OPTIONS TO 'cuLinkAddFile(...)' and 'cuLinkAddData(...)'
// -> FUNCTION FAILS WITH ERROR 'CUDA_ERROR_INVALID_VALUE'
// -------------------------------------------------------------------------------------------
CUDA_CHECK(cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", 0, NULL, NULL));
CUDA_CHECK(cuLinkAddData(linker, CU_JIT_INPUT_NVVM, (void*)nvvm, nvvmSize, "programRTC.o", 0,
NULL, NULL));
// CUDA_CHECK(cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", 1, options, values));
// CUDA_CHECK(cuLinkAddData(linker, CU_JIT_INPUT_NVVM, (void*)nvvm, nvvmSize, "programRTC.o", 1,
// options, values));
// -------------------------------------------------------------------------------------------
// Create module ===============
void* cubin;
CUmodule module;
CUDA_CHECK(cuLinkComplete(linker, &cubin, NULL));
CUDA_CHECK(cuModuleLoadDataEx(&module, cubin, 0, NULL, NULL));
// Cleanup
NVRTC_CHECK(nvrtcDestroyProgram(&prog));
CUDA_CHECK(cuLinkDestroy(linker));
return module;
}
#endif // RTC
__device__ double func(double a, double b);
#ifdef FUNC
__device__ double func(double a, double b)
{
return pow(a, b);
}
#endif
#ifdef MAIN
#ifdef RTC
std::string the_program = R"===(
__device__ double func(double a, double b);
extern "C" __global__ void kernel(double* out, double* a, double* b)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid >= 1){
return;
}
a[tid] = 2;
b[tid] = 3;
out[tid] = func(a[tid], b[tid]);
printf("out[%lu] = %f\n", tid, out[tid]);
})===";
#else // RTC
__global__ void kernel(double* out, double* a, double* b)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= 1) {
return;
}
a[tid] = 2;
b[tid] = 3;
out[tid] = func(a[tid], b[tid]);
printf("out[%lu] = %f\n", tid, out[tid]);
}
#endif // RTC
int main()
{
double* a;
double* b;
double* out;
cudaMalloc((void**)&a, sizeof(double));
cudaMalloc((void**)&b, sizeof(double));
cudaMalloc((void**)&out, sizeof(double));
#ifdef RTC
// Create context
CUdevice cuDevice;
CUcontext context;
CUDA_CHECK(cuInit(0));
CUDA_CHECK(cuDeviceGet(&cuDevice, 0));
CUDA_CHECK(cuCtxCreate(&context, 0, cuDevice));
CUmodule module = compileModule(the_program);
CUfunction kernel;
CUDA_CHECK(cuModuleGetFunction(&kernel, module, "kernel"));
size_t n_blocks = 1;
size_t n_threads = 1;
void* args[] = {&out, &a, &b};
CUDA_CHECK(cuLaunchKernel(kernel, n_blocks, 1, 1, // grid dim
n_threads, 1, 1, // block dim
0, NULL, // shared mem and stream
args, 0)); // arguments
CUDA_CHECK(cuCtxSynchronize());
// Cleanup
CUDA_CHECK(cuModuleUnload(module));
CUDA_CHECK(cuCtxDestroy(context));
#else // RTC
kernel<<<1, 1>>>(out, a, b);
cudaDeviceSynchronize();
#endif // RTC
return 0;
}
#endif // MAIN
bash 脚本:
#!/bin/bash
set -e # stop script when an error occurs
SCRIPT=
xCCx=80 # CUDA compute compatibility
# Create program using a single translation file
echo -e "\n---------- main_single ----------\n"
nvcc -DFUNC -DMAIN $SCRIPT -o main_single \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_single # should print 'out[0] = 8.0'
cuobjdump main_single -res-usage | grep kernel -A1
# Link files using NVCC without link time optimization (code=compute_...)
echo -e "\n---------- main_link_nvcc ----------\n"
nvcc -DFUNC $SCRIPT -o func_link_nvcc.o -dc \
-gencode arch=compute_$xCCx,code=compute_$xCCx
nvcc -DMAIN $SCRIPT -o main_link_nvcc.o -dc \
-gencode arch=compute_$xCCx,code=compute_$xCCx
nvcc func_link_nvcc.o main_link_nvcc.o -o main_link_nvcc \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_link_nvcc # should print 'out[0] = 8.0'
cuobjdump main_link_nvcc -res-usage | grep kernel -A1
# Link files using NVCC with link time optimization (code=lto_...)
echo -e "\n---------- main_link_nvcc_lto ----------\n"
nvcc -DFUNC $SCRIPT -o func_link_nvcc_lto.o -dc \
-gencode arch=compute_$xCCx,code=lto_$xCCx
nvcc -DMAIN $SCRIPT -o main_link_nvcc_lto.o -dc \
-gencode arch=compute_$xCCx,code=lto_$xCCx
nvcc func_link_nvcc_lto.o main_link_nvcc_lto.o -o main_link_nvcc_lto -dlto \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_link_nvcc_lto # should print 'out[0] = 8.0'
cuobjdump main_link_nvcc_lto -res-usage | grep kernel -A1
# Link files using NVRTC with link time optimization
echo -e "\n---------- main_link_nvrtc_lto ----------\n"
nvcc -DFUNC $SCRIPT -o func_link_nvrtc_lto.o -dc -ptx \
-gencode arch=compute_$xCCx,code=lto_$xCCx
nvcc -DMAIN -DRTC $SCRIPT -o main_link_nvrtc_lto \
-lnvrtc_static -lnvrtc-builtins_static -lnvptxcompiler_static -lcuda -lpthread \
-gencode arch=compute_$xCCx,code=sm_$xCCx
./main_link_nvrtc_lto # should print 'out[0] = 8.0'
ncu main_link_nvrtc_lto | grep register/thread
# Registers/thread used on my system with an NVIDIA RTX 3080:
# main_single : 30 registers/thread
# main_link_nvcc : 44 registers/thread
# main_link_nvcc_lto : 30 registers/thread
# main_link_nvrtc_lto : 38 registers/thread
子问题：生成一个 NVVM IR 文件
为了生成与命令 cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", ...) 一起使用的文件 func_link_nvrtc_lto.o，我必须添加 -ptx 标志，如下面的命令所示。我没有在任何文档中找到这个，而是通过尝试和错误找到的。我想知道是否有更好的方法来生成这样的文件。
nvcc -DFUNC $SCRIPT -o func_link_nvrtc_lto.o -dc -ptx \
-gencode arch=compute_$xCCx,code=lto_$xCCx
首先，不幸的是博客 post 中有一个错误 CU_JIT_LTO 值。它应该是：
values[0] = (void*)1;
然而，这并不重要，因为该值被忽略了——它只是使用了 CU_JIT_LTO 的存在。 正如您发现的那样，CU_JIT_LTO 确实应该传递给 cuLinkCreate。
对于您的 sub-question，您的 -ptx 所做的是在生成 nvvm-ir 后停止编译，但这是一个未记录的 side-effect。更简单、更安全的方法是使用：
nvcc -dc -arch=compute_XX,code=lto_XX
创建一个包含 nvvm-ir 的宿主对象。然后将其传递为：
CU_JIT_INPUT_OBJECT to cuLinkAddFile().