我想知道是否可以使用选项 CU_JIT_LTO 在即时 (JIT) 链接期间改进链接时间优化 (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来链接使用 LTO 标志（-dlto或-code=lto_80）创建的对象。似乎链接器已经做了一些 LTO，因为内核“执行得更好”，而不是链接没有 LTO 的目标文件，但不如使用 NVCC 链接 LTO。（有关详细结果和讨论，请参见示例案例）
示例案例
为了检查链接时间优化 (LTO) 的有效性，我使用 4 种不同的方法创建了一个简单的程序，并使用每个线程的寄存器数量作为指标。这在我的系统上给出了以下结果（操作系统：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 个寄存器/线程。
使用 NVCC 与 LTO 链接绝对有效。它使用与从单个 .cu 文件编译的程序相同数量的寄存器 (30)，而没有 LTO 的情况则不同（使用 44 个寄存器）。
使用 NVRTC/JIT 链接文件比在没有 LTO 的情况下链接 NVCC 做得“更好”（当我们只关注寄存器使用时），但不如使用 LTO 链接 NVCC。内核使用 38 个寄存器/线程。
注意：我的目标不是减少寄存器的使用，我只是把它作为一个指标。因为来自单个翻译文件的程序使用 30 个寄存器/线程，所以我假设一个完全优化的链接程序将具有相同的“最终可执行代码”，因此使用相同数量的寄存器。因为情况并非如此，所以我开始研究 JIT 选项。
CU_JIT_LTO 选项：
我尝试使用 JIT_option CU_JIT_LTO 进一步优化 NVRTC/JIT 案例中的链接。但是，我不确定如何使用此选项。我尝试了以下两种方式（有关更多上下文，请参见下面的文件cuda 代码。链接代码从第 41 行开始）：
方法 1：将选项 CU_JIT_LTO 添加到cuLinkCreate(...). 这似乎没有任何效果。int lto = 0该代码在和时使用相同数量的寄存器int lto = 1。
方法 2：将选项 CU_JIT_LTO 添加到cuLinkAddFile(...)和cuLinkAddData(...)。这会立即给出错误 CUDA_ERROR_INVALID_VALUE。
所以现在我的问题是：应该如何使用选项 CU_JIT_LTO？
文件：
下面是两个文件。按照以下步骤运行示例（在 linux 操作系统上）：
- 将 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=$1
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 文件
要生成func_link_nvrtc_lto.o与命令一起使用的文件cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", ...)，我必须添加-ptx标志，如下面的命令所示。我没有在任何文档中找到这一点，而是通过尝试和错误。我想知道是否有更好的方法来生成这样的文件。
nvcc -DFUNC $SCRIPT -o func_link_nvrtc_lto.o -dc -ptx \
-gencode arch=compute_$xCCx,code=lto_$xCCx