libnvvm 与 CUDA C 设备库联合链接实战cuda-c-linking 示例全解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplescuda-c-linking 是 cuda-samples 仓库 cpp/7_libNVVM 下的一个标志性示例它演示了一条完整的“无 nvcc 前端编译、纯库级驱动”的 GPU 编程流水线——先用 LLVM C API 在运行时构造一个符合 NVVM IR 规范的模块交由 libnvvm 编译为 PTX再通过 CUDA Driver API 的 JIT linker 将 PTX 与一份由 nvcc 预编译的 CUDA C 设备库链接成 CUBIN最后加载并执行。读完本文你将掌握 NVVM IR 的模块构造要点、libnvvm 编译 API 调用链、cuLink*系列 JIT 链接接口以及 CUDA Driver API 从初始化到内核启动的完整用法。示例要解决的问题常规 CUDA 开发中nvcc同时承担两件事把 CUDA C/C 编译为设备代码以及把设备代码与设备库如数学库、libdevice链接成可加载的镜像。而本示例展示了一条完全脱离 nvcc 设备编译前端的替代路径用 LLVM C API 动态构造一个 LLVM IR 模块该模块符合 NVVM IR 规范即 CUDA 设备代码的 LLVM IR 变体其中包含一个对外部函数的调用调用 libnvvm 把该 IR 模块编译成 PTX借助 CUDA Driver API 中的 JIT linker把生成的 PTX 与一个现成的 CUDA C 设备库链接产出链接后的 CUBIN 镜像在系统中的第一块 CUDA 设备上加载并运行该镜像。这个模式对需要在运行时生成、定制或 JIT 编译 GPU 代码的工具链编译器、代码生成器、自定义算子运行时具有直接的参考价值。源码文件构成原文档列出了三个核心文件实际目录中对应实现如下文件作用cuda-c-linking.cpp主程序LLVM IR 生成、libnvvm 编译、JIT 链接、内核加载与启动math-funcs.cuCUDA C 设备库源码一个简单的 Mandelbrot 设备函数CMakeLists.txt该示例的 CMake 构建脚本负责编译主程序并把设备库打成静态库其中cuda-c-linking.cpp还包含了 common/include/DDSWriter.h 提供的writeDDS()用于把内核计算结果输出为.dds图像文件。构建方法作为 libnvvm 示例的一部分可选构建本示例默认不参与构建需要显式开启。整个 cpp/7_libNVVM/CMakeLists.txt 使用ENABLE_CUDA_C_LINKING_SAMPLE变量控制if (ENABLE_CUDA_C_LINKING_SAMPLE) # Include the LLVM dev package which is required to build cuda-c-linking. find_package(LLVM CONFIG PATHS $ENV{LLVM_HOME}) if (LLVM_FOUND) add_subdirectory(cuda-c-linking) else () message(STATUS Skipping the build of the cuda-c-linking sample: Failed to locate the LLVM package.) endif () else () message(STATUS Skipping the build of the cuda-c-linking sample.) endif ()即满足以下两个条件时才会构建在 CMake 命令行或 CMakeLists.txt 中设置-DENABLE_CUDA_C_LINKING_SAMPLE1系统能找到 LLVM 开发包find_package(LLVM CONFIG PATHS $ENV{LLVM_HOME})。构建时通过以下环境变量定位依赖见 cpp/7_libNVVM/README.mdCUDA_HOMECUDA Toolkit 安装目录例如/usr/local/cuda未指定时 CMake 会尝试从 PATH 中的nvcc推导LIBNVVM_HOMElibnvvm 组件所在目录通常即$CUDA_HOME/nvvmLLVM_HOME仅当使用本地自建 LLVM 时才需要指向 LLVM 安装目录。LLVM 版本约束重要前提由于要用 LLVM C API 构造 IR本示例对 LLVM 版本有硬性要求cuda-c-linking/CMakeLists.txt 开头直接做了版本检查message(STATUS Found LLVM Version ${LLVM_PACKAGE_VERSION}) if (LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL 15 OR LLVM_PACKAGE_VERSION VERSION_LESS 7) message(STATUS The cuda-c-linking sample is expected to build with LLVM development libraries v7 to v14, opaque pointers are not supported in libNVVM for pre-Blackwell architectures.) return() endif ()关键约束与官方说明一致要求LLVM 7 到 14含的开发库头文件与库文件LLVM 15 起默认启用opaque pointers不透明指针而 libnvvm 在 pre-Blackwell 架构上不支持该指针表示Ubuntu 用户安装llvm-dev包即可获得所需头文件与库通常无需显式设置LLVM_HOMEWindows 用户应从 llvm.org 下载 LLVM 14 源码自行构建安装官方 Windows 安装器缺少本示例依赖的部分组件且应使用与 LLVM 相同的构建模式如 LLVM 以 Release 构建则本示例也用-DCMAKE_BUILD_TYPERelease。设备库的预编译nvcc 只负责设备库设备库math-funcs.cu仍由 nvcc 编译成静态库CMakeLists.txt 中注释给出了手动等价命令# nvcc -m64 -archcompute_75 -dc math-funcs.cu -o math-funcs64.o # nvcc -m64 -lib math-funcs64.o -o libmathfuncs64.a-dcdevice only只编译设备代码是关键设备库只产出设备端对象不链接主机代码-lib把对象打包为静态库libmathfuncs64.a目标架构默认使用compute_75若需支持多种 GPU 平台可在同一次 nvcc 调用中多次使用-gencodecompute_XX,codesm_XX。CMake 中对应逻辑为开启CUDA语言、设置CMAKE_CUDA_FLAGS -dc、CMAKE_CUDA_ARCHITECTURES 75并以add_library(mathfuncs64 STATIC math-funcs.cu)生成libmathfuncs64.a。运行环境方面若使用 CUDA Toolkit 13.0 或更新版本配合较新的 user-mode driver580与旧内核驱动550 及更早构建 libnvvm 示例时需在 CMake 调用中加入 stubs 路径-DCMAKE_PREFIX_PATH/usr/local/cuda/lib64/stubs/顶层构建流程在 cpp/7_libNVVM 目录下设置环境变量后Linux 上可执行构建脚本utils/build.sh或直接使用顶层 Makefile/CMake 构建Windows 上可借助 Visual Studio 的 CMake 集成Build All Install libnvvm-samples产物会安装到out/install/arch/bin/或$SAMPLES_INSTALL_DIR。运行方法构建完成后直接运行cuda-c-linking可执行文件即可。原文档给出的调用方式如下Linux$ cd $SAMPLES_INSTALL_DIR $ ./cuda-c-linkingWindows$ cd %SAMPLES_INSTALL_DIR% $ cuda-c-linking.exe程序启动后会打印所用设备信息设备名、计算能力随后生成 1024×512 的 Mandelbrot 图像并保存为mandelbrot.dds。注意运行前提主程序通过cuDeviceGet(device, 0)固定使用0 号设备代码在devMajor 7 devMinor 5时直接报错退出见 cuda-c-linking.cpp即要求设备计算能力不低于sm_75——这与设备库以compute_75预编译一致可执行文件运行时需要在同目录能找到libmathfuncs64.a程序用sys::fs::getMainExecutable定位自身路径并拼接库名见 cuda-c-linking.cpp。命令行选项原文档列出三个用于检查中间产物的选项主程序通过 LLVM 的cl::opt声明cuda-c-linking.cpp选项作用输出文件-save-ptx把 libnvvm 生成的 PTX 内核写入磁盘cuda-c-linking.kernel.ptx-save-ir把生成的 LLVM IR 写入磁盘cuda-c-linking.kernel.ll-save-cubin把链接后的 CUBIN 镜像写入磁盘cuda-c-linking.linked.cubin这些选项非常适合教学与调试-save-ir可以直观看到程序动态构造出的 NVVM IR 是否符合规范-save-ptx可以核对 libnvvm 的编译结果-save-cubin则可验证 JIT 链接是否成功产出最终镜像。源码级原理剖析阶段一用 LLVM C API 构造 NVVM IR 模块generateModule()cuda-c-linking.cpp在纯主机端用 LLVM IRBuilder 构造整个模块关键步骤包括1. 设置 DataLayout 与目标三元组这是 NVVM IR 合规性的基础mod-setDataLayout(e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128- f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32: 64); mod-setTargetTriple(nvptx64-nvidia-cuda);2. 声明外部函数与内核函数通过mod-getOrInsertFunction(mandelbrot, ...)声明一个未定义的外部函数void mandelbrot(float*)它正是要由设备库提供的符号通过Function::Create(..., kernel, *mod)创建入口内核函数kernel(float* ptr)在entry基本块中仅做两件事调用mandelbrot并把参数ptr传给它然后返回。3. 写入 NVVM 必需的元数据MDNode *kernelMD MDNode::get(context, mdVals); NamedMDNode *nvvmAnnot mod-getOrInsertNamedMetadata(nvvm.annotations); nvvmAnnot-addOperand(kernelMD);nvvm.annotations用于把kernel标记为设备内核入口随后写入nvvmir.version元数据声明NVVM IR 版本 2.0{2, 0}供 libnvvm 识别。阶段二libnvvm 把 IR 编译为 PTXgeneratePtx()cuda-c-linking.cpp演示了 libnvvm 的完整调用链nvvmCreateProgram 创建编译单元nvvmProgram nvvmAddModuleToProgram 把 IR 字符串加入编译单元 nvvmCompileProgram 传入 -archcompute_XY 编译为 PTX nvvmGetCompiledResultSize 查询结果长度 nvvmGetCompiledResult 取出 PTX 文本 nvvmDestroyProgram 销毁编译单元其中计算能力参数是运行时动态获取的主程序通过cuDeviceGetAttribute读取 0 号设备的CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR/MINOR拼出-archcompute_XY例如compute_75。编译失败时通过nvvmGetProgramLogSize/nvvmGetProgramLog打印错误日志这也是排查 IR 合规性问题的关键出口。阶段三CUDA Driver API JIT 链接链接阶段使用 Driver API 的cuLink*系列cuda-c-linking.cpp创建链接器并配置日志选项cuLinkCreate(5, linkerOptions, linkerOptionValues, linker)选项数组包含CU_JIT_INFO_LOG_BUFFER、CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES、CU_JIT_ERROR_LOG_BUFFER、CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES与CU_JIT_LOG_VERBOSE用于捕获链接日志加入 PTX 数据cuLinkAddData(linker, CU_JIT_INPUT_PTX, ptx, ptx.size(), compiled-ptx, 0, NULL, NULL)加入设备库文件cuLinkAddFile(linker, CU_JIT_INPUT_LIBRARY, libpath, 0, NULL, NULL)其中libpath即libmathfuncs64.a的路径这里CU_JIT_INPUT_LIBRARY表示按库文件解析其中的设备符号完成链接cuLinkComplete(linker, cubin, cubinSize)得到可加载的 CUBIN 镜像。链接器会解析 PTX 中对mandelbrot的外部引用并从设备库中取出该函数完成符号绑定从而把“动态生成的 IR”与“nvcc 预编译的 CUDA C 库”融为一体——这正是示例名称cuda-c-linking的含义。阶段四加载、启动与输出链接完成后走标准的 Driver API 执行路径cuModuleLoadDataEx(cudaModule, cubin, 0, 0, 0); // 从内存加载 CUBIN cuModuleGetFunction(function, cudaModule, kernel); // 取得内核句柄 cuMemAlloc(devBuffer, sizeof(float) * width * height * 4); cuLaunchKernel(function, gridX, gridY, gridZ, blockX, blockY, blockZ, 0, NULL, params, NULL); cuMemcpyDtoH(data[0], devBuffer, ...); // 回拷结果 writeDDS(mandelbrot.dds, data, width, height); // 写出图像输出尺寸固定为width1024, height512网格按 16×16 线程块组织gridSizeX width/16gridSizeY height/16每个线程负责计算一个像素的 RGBA设备端mandelbrotmath-funcs.cu是一个按 Wikipedia 算法实现的直观版本将像素坐标归一化到复平面区间后做最多 1000 次迭代按迭代次数映射为颜色写入输出缓冲区DDSWriter.h 中的writeDDS()把[0,1]归一化的 float RGBA 数据缩放到[0,255]的 8 位值并写入带标准 DDS 头DDPF_RGB | DDPF_ALPHAPIXELSRGBA 各 8 位的.dds文件。主程序用#define checkCudaErrors(err)包装所有 Driver API 调用出错时通过cuGetErrorName打印 CUDA 错误名与出错文件行号后退出方便快速定位。关键点小结nvcc 的角色被拆分设备库仍由 nvcc 预编译-dc-lib但设备内核本体由运行时生成的 NVVM IR 提供两者通过 JIT linker 汇合NVVM IR 合规性正确设置 DataLayout、target triple、nvvm.annotations与nvvmir.version是 libnvvm 接受模块的前提动态目标架构-archcompute_XY依据实际设备计算能力生成同时程序强制要求 sm_75 及以上可观测性-save-ir、-save-ptx、-save-cubin三个开关让整条流水线的中间产物完全可见是学习 NVVM 工具链的低成本入口版本敏感本示例强依赖 LLVM 7–14 开发库与ENABLE_CUDA_C_LINKING_SAMPLE显式开启构建前务必核对环境。对于希望把“运行时生成 GPU 代码”集成进自有工具链的开发者本示例是理解 NVVM IR libnvvm JIT 链接三者协同工作方式的一份完整、可运行的参考实现。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
libnvvm 与 CUDA C 设备库联合链接实战：cuda-c-linking 示例全解析
libnvvm 与 CUDA C 设备库联合链接实战cuda-c-linking 示例全解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplescuda-c-linking 是 cuda-samples 仓库 cpp/7_libNVVM 下的一个标志性示例它演示了一条完整的“无 nvcc 前端编译、纯库级驱动”的 GPU 编程流水线——先用 LLVM C API 在运行时构造一个符合 NVVM IR 规范的模块交由 libnvvm 编译为 PTX再通过 CUDA Driver API 的 JIT linker 将 PTX 与一份由 nvcc 预编译的 CUDA C 设备库链接成 CUBIN最后加载并执行。读完本文你将掌握 NVVM IR 的模块构造要点、libnvvm 编译 API 调用链、cuLink*系列 JIT 链接接口以及 CUDA Driver API 从初始化到内核启动的完整用法。示例要解决的问题常规 CUDA 开发中nvcc同时承担两件事把 CUDA C/C 编译为设备代码以及把设备代码与设备库如数学库、libdevice链接成可加载的镜像。而本示例展示了一条完全脱离 nvcc 设备编译前端的替代路径用 LLVM C API 动态构造一个 LLVM IR 模块该模块符合 NVVM IR 规范即 CUDA 设备代码的 LLVM IR 变体其中包含一个对外部函数的调用调用 libnvvm 把该 IR 模块编译成 PTX借助 CUDA Driver API 中的 JIT linker把生成的 PTX 与一个现成的 CUDA C 设备库链接产出链接后的 CUBIN 镜像在系统中的第一块 CUDA 设备上加载并运行该镜像。这个模式对需要在运行时生成、定制或 JIT 编译 GPU 代码的工具链编译器、代码生成器、自定义算子运行时具有直接的参考价值。源码文件构成原文档列出了三个核心文件实际目录中对应实现如下文件作用cuda-c-linking.cpp主程序LLVM IR 生成、libnvvm 编译、JIT 链接、内核加载与启动math-funcs.cuCUDA C 设备库源码一个简单的 Mandelbrot 设备函数CMakeLists.txt该示例的 CMake 构建脚本负责编译主程序并把设备库打成静态库其中cuda-c-linking.cpp还包含了 common/include/DDSWriter.h 提供的writeDDS()用于把内核计算结果输出为.dds图像文件。构建方法作为 libnvvm 示例的一部分可选构建本示例默认不参与构建需要显式开启。整个 cpp/7_libNVVM/CMakeLists.txt 使用ENABLE_CUDA_C_LINKING_SAMPLE变量控制if (ENABLE_CUDA_C_LINKING_SAMPLE) # Include the LLVM dev package which is required to build cuda-c-linking. find_package(LLVM CONFIG PATHS $ENV{LLVM_HOME}) if (LLVM_FOUND) add_subdirectory(cuda-c-linking) else () message(STATUS Skipping the build of the cuda-c-linking sample: Failed to locate the LLVM package.) endif () else () message(STATUS Skipping the build of the cuda-c-linking sample.) endif ()即满足以下两个条件时才会构建在 CMake 命令行或 CMakeLists.txt 中设置-DENABLE_CUDA_C_LINKING_SAMPLE1系统能找到 LLVM 开发包find_package(LLVM CONFIG PATHS $ENV{LLVM_HOME})。构建时通过以下环境变量定位依赖见 cpp/7_libNVVM/README.mdCUDA_HOMECUDA Toolkit 安装目录例如/usr/local/cuda未指定时 CMake 会尝试从 PATH 中的nvcc推导LIBNVVM_HOMElibnvvm 组件所在目录通常即$CUDA_HOME/nvvmLLVM_HOME仅当使用本地自建 LLVM 时才需要指向 LLVM 安装目录。LLVM 版本约束重要前提由于要用 LLVM C API 构造 IR本示例对 LLVM 版本有硬性要求cuda-c-linking/CMakeLists.txt 开头直接做了版本检查message(STATUS Found LLVM Version ${LLVM_PACKAGE_VERSION}) if (LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL 15 OR LLVM_PACKAGE_VERSION VERSION_LESS 7) message(STATUS The cuda-c-linking sample is expected to build with LLVM development libraries v7 to v14, opaque pointers are not supported in libNVVM for pre-Blackwell architectures.) return() endif ()关键约束与官方说明一致要求LLVM 7 到 14含的开发库头文件与库文件LLVM 15 起默认启用opaque pointers不透明指针而 libnvvm 在 pre-Blackwell 架构上不支持该指针表示Ubuntu 用户安装llvm-dev包即可获得所需头文件与库通常无需显式设置LLVM_HOMEWindows 用户应从 llvm.org 下载 LLVM 14 源码自行构建安装官方 Windows 安装器缺少本示例依赖的部分组件且应使用与 LLVM 相同的构建模式如 LLVM 以 Release 构建则本示例也用-DCMAKE_BUILD_TYPERelease。设备库的预编译nvcc 只负责设备库设备库math-funcs.cu仍由 nvcc 编译成静态库CMakeLists.txt 中注释给出了手动等价命令# nvcc -m64 -archcompute_75 -dc math-funcs.cu -o math-funcs64.o # nvcc -m64 -lib math-funcs64.o -o libmathfuncs64.a-dcdevice only只编译设备代码是关键设备库只产出设备端对象不链接主机代码-lib把对象打包为静态库libmathfuncs64.a目标架构默认使用compute_75若需支持多种 GPU 平台可在同一次 nvcc 调用中多次使用-gencodecompute_XX,codesm_XX。CMake 中对应逻辑为开启CUDA语言、设置CMAKE_CUDA_FLAGS -dc、CMAKE_CUDA_ARCHITECTURES 75并以add_library(mathfuncs64 STATIC math-funcs.cu)生成libmathfuncs64.a。运行环境方面若使用 CUDA Toolkit 13.0 或更新版本配合较新的 user-mode driver580与旧内核驱动550 及更早构建 libnvvm 示例时需在 CMake 调用中加入 stubs 路径-DCMAKE_PREFIX_PATH/usr/local/cuda/lib64/stubs/顶层构建流程在 cpp/7_libNVVM 目录下设置环境变量后Linux 上可执行构建脚本utils/build.sh或直接使用顶层 Makefile/CMake 构建Windows 上可借助 Visual Studio 的 CMake 集成Build All Install libnvvm-samples产物会安装到out/install/arch/bin/或$SAMPLES_INSTALL_DIR。运行方法构建完成后直接运行cuda-c-linking可执行文件即可。原文档给出的调用方式如下Linux$ cd $SAMPLES_INSTALL_DIR $ ./cuda-c-linkingWindows$ cd %SAMPLES_INSTALL_DIR% $ cuda-c-linking.exe程序启动后会打印所用设备信息设备名、计算能力随后生成 1024×512 的 Mandelbrot 图像并保存为mandelbrot.dds。注意运行前提主程序通过cuDeviceGet(device, 0)固定使用0 号设备代码在devMajor 7 devMinor 5时直接报错退出见 cuda-c-linking.cpp即要求设备计算能力不低于sm_75——这与设备库以compute_75预编译一致可执行文件运行时需要在同目录能找到libmathfuncs64.a程序用sys::fs::getMainExecutable定位自身路径并拼接库名见 cuda-c-linking.cpp。命令行选项原文档列出三个用于检查中间产物的选项主程序通过 LLVM 的cl::opt声明cuda-c-linking.cpp选项作用输出文件-save-ptx把 libnvvm 生成的 PTX 内核写入磁盘cuda-c-linking.kernel.ptx-save-ir把生成的 LLVM IR 写入磁盘cuda-c-linking.kernel.ll-save-cubin把链接后的 CUBIN 镜像写入磁盘cuda-c-linking.linked.cubin这些选项非常适合教学与调试-save-ir可以直观看到程序动态构造出的 NVVM IR 是否符合规范-save-ptx可以核对 libnvvm 的编译结果-save-cubin则可验证 JIT 链接是否成功产出最终镜像。源码级原理剖析阶段一用 LLVM C API 构造 NVVM IR 模块generateModule()cuda-c-linking.cpp在纯主机端用 LLVM IRBuilder 构造整个模块关键步骤包括1. 设置 DataLayout 与目标三元组这是 NVVM IR 合规性的基础mod-setDataLayout(e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128- f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32: 64); mod-setTargetTriple(nvptx64-nvidia-cuda);2. 声明外部函数与内核函数通过mod-getOrInsertFunction(mandelbrot, ...)声明一个未定义的外部函数void mandelbrot(float*)它正是要由设备库提供的符号通过Function::Create(..., kernel, *mod)创建入口内核函数kernel(float* ptr)在entry基本块中仅做两件事调用mandelbrot并把参数ptr传给它然后返回。3. 写入 NVVM 必需的元数据MDNode *kernelMD MDNode::get(context, mdVals); NamedMDNode *nvvmAnnot mod-getOrInsertNamedMetadata(nvvm.annotations); nvvmAnnot-addOperand(kernelMD);nvvm.annotations用于把kernel标记为设备内核入口随后写入nvvmir.version元数据声明NVVM IR 版本 2.0{2, 0}供 libnvvm 识别。阶段二libnvvm 把 IR 编译为 PTXgeneratePtx()cuda-c-linking.cpp演示了 libnvvm 的完整调用链nvvmCreateProgram 创建编译单元nvvmProgram nvvmAddModuleToProgram 把 IR 字符串加入编译单元 nvvmCompileProgram 传入 -archcompute_XY 编译为 PTX nvvmGetCompiledResultSize 查询结果长度 nvvmGetCompiledResult 取出 PTX 文本 nvvmDestroyProgram 销毁编译单元其中计算能力参数是运行时动态获取的主程序通过cuDeviceGetAttribute读取 0 号设备的CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR/MINOR拼出-archcompute_XY例如compute_75。编译失败时通过nvvmGetProgramLogSize/nvvmGetProgramLog打印错误日志这也是排查 IR 合规性问题的关键出口。阶段三CUDA Driver API JIT 链接链接阶段使用 Driver API 的cuLink*系列cuda-c-linking.cpp创建链接器并配置日志选项cuLinkCreate(5, linkerOptions, linkerOptionValues, linker)选项数组包含CU_JIT_INFO_LOG_BUFFER、CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES、CU_JIT_ERROR_LOG_BUFFER、CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES与CU_JIT_LOG_VERBOSE用于捕获链接日志加入 PTX 数据cuLinkAddData(linker, CU_JIT_INPUT_PTX, ptx, ptx.size(), compiled-ptx, 0, NULL, NULL)加入设备库文件cuLinkAddFile(linker, CU_JIT_INPUT_LIBRARY, libpath, 0, NULL, NULL)其中libpath即libmathfuncs64.a的路径这里CU_JIT_INPUT_LIBRARY表示按库文件解析其中的设备符号完成链接cuLinkComplete(linker, cubin, cubinSize)得到可加载的 CUBIN 镜像。链接器会解析 PTX 中对mandelbrot的外部引用并从设备库中取出该函数完成符号绑定从而把“动态生成的 IR”与“nvcc 预编译的 CUDA C 库”融为一体——这正是示例名称cuda-c-linking的含义。阶段四加载、启动与输出链接完成后走标准的 Driver API 执行路径cuModuleLoadDataEx(cudaModule, cubin, 0, 0, 0); // 从内存加载 CUBIN cuModuleGetFunction(function, cudaModule, kernel); // 取得内核句柄 cuMemAlloc(devBuffer, sizeof(float) * width * height * 4); cuLaunchKernel(function, gridX, gridY, gridZ, blockX, blockY, blockZ, 0, NULL, params, NULL); cuMemcpyDtoH(data[0], devBuffer, ...); // 回拷结果 writeDDS(mandelbrot.dds, data, width, height); // 写出图像输出尺寸固定为width1024, height512网格按 16×16 线程块组织gridSizeX width/16gridSizeY height/16每个线程负责计算一个像素的 RGBA设备端mandelbrotmath-funcs.cu是一个按 Wikipedia 算法实现的直观版本将像素坐标归一化到复平面区间后做最多 1000 次迭代按迭代次数映射为颜色写入输出缓冲区DDSWriter.h 中的writeDDS()把[0,1]归一化的 float RGBA 数据缩放到[0,255]的 8 位值并写入带标准 DDS 头DDPF_RGB | DDPF_ALPHAPIXELSRGBA 各 8 位的.dds文件。主程序用#define checkCudaErrors(err)包装所有 Driver API 调用出错时通过cuGetErrorName打印 CUDA 错误名与出错文件行号后退出方便快速定位。关键点小结nvcc 的角色被拆分设备库仍由 nvcc 预编译-dc-lib但设备内核本体由运行时生成的 NVVM IR 提供两者通过 JIT linker 汇合NVVM IR 合规性正确设置 DataLayout、target triple、nvvm.annotations与nvvmir.version是 libnvvm 接受模块的前提动态目标架构-archcompute_XY依据实际设备计算能力生成同时程序强制要求 sm_75 及以上可观测性-save-ir、-save-ptx、-save-cubin三个开关让整条流水线的中间产物完全可见是学习 NVVM 工具链的低成本入口版本敏感本示例强依赖 LLVM 7–14 开发库与ENABLE_CUDA_C_LINKING_SAMPLE显式开启构建前务必核对环境。对于希望把“运行时生成 GPU 代码”集成进自有工具链的开发者本示例是理解 NVVM IR libnvvm JIT 链接三者协同工作方式的一份完整、可运行的参考实现。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
