我应该如何让 CMake 也为我的内核创建 PTX 文件

Posted 2023-02-16

【中文标题】我应该如何让 CMake 也为我的内核创建 PTX 文件

【英文标题】：How should I get CMake to also create PTX files for my kernels

【发布时间】：2019-05-05 15:47:51

【问题描述】：

我正在使用具有内在 CUDA 支持的最新 CMake（如果需要，版本 >= 3.8 或更高版本）使用 CUDA 代码构建一个项目。

我如何告诉 CMake（也）为我的各种内核生成 PTX 文件？

我尝试过的东西（？）不起作用：

从 CMake 3.9 开始，我们可以使用 CUDA_PTX_COMPILATION property 定义一个对象库来包含 PTX 而不是各种对象：

add_library(myptx OBJECT a.cu b.cu)
set_property(TARGET myptx PROPERTY CUDA_PTX_COMPILATION ON)

但是，这不是问题的正确解决方案 - 缺少一些东西。假设我们在a.cu:

__global__ void foo()  return; 

在b.cu:

__global__ void bar(int* a)  *a = 5; 

我们运行cmake 和make 并使用以下CMakeLists.txt：

cmake_minimum_required(VERSION 3.9)
add_library(myptx OBJECT a.cu b.cu)
set_property(TARGET myptx PROPERTY CUDA_PTX_COMPILATION ON)

没有生成 PTX 文件，也没有运行 nvcc。不知道为什么。

【问题讨论】：

@Tsyvarev：将半答案移到问题中。

您想要 PTX 做什么？适当编译的二进制有效载荷包含 PTX 和 SASS。这还不够吗？

@talonmies：我不介意 PTX 是否是从二进制后验中提取的。我只想检查 PTX 代码以及构建二进制文件。

【参考方案1】：

您可以将--keep 命令行选项传递给 nvcc，这将使其将所有中间文件保存在用于其他输出的同一目录中。对于较新的 CMake 版本，您可以编写：

target_compile_options(
    some_target
    another_target
    yet_another_target_etc
    PRIVATE 
    "--keep"
)

应该这样做。似乎.ptx 最终出现在二进制文件所在的位置。但是 - 请注意，您还会获得大量其他文件：.cpp4.ii、.cudafe1.c、.cudafe1.cpp、.fatbin 等等。因此，虽然这可行，但它并不是一个很好的解决方案。

【参考方案2】：

尝试开启 CUDA 语言。

cmake_minimum_required(VERSION 3.9)
enable_language(CUDA)
add_library(myptx OBJECT a.cu b.cu)
set_property(TARGET myptx PROPERTY CUDA_PTX_COMPILATION ON)

PTX 位于 $CMAKE_BINARY_DIR/CMakeFiles/myptx.dir