使用 clang 编译 CUDA¶

简介 ¶

本文档描述了如何使用 clang 编译 CUDA 代码，并提供了一些关于 LLVM 和 clang 的 CUDA 实现的细节。

本文档假设读者对 CUDA 有基本的了解。关于 CUDA 编程的信息可以在 CUDA 编程指南中找到。

编译 CUDA 代码 ¶

先决条件 ¶

CUDA 从 llvm 3.9 版本开始支持。Clang 目前支持 CUDA 7.0 到 12.1 版本。如果 clang 检测到更新的 CUDA 版本，它将发出警告，并将尝试使用检测到的 CUDA SDK，就好像它是 CUDA 12.1 一样。

在构建 CUDA 代码之前，您需要安装 CUDA SDK。有关详细信息，请参阅 NVIDIA 的 CUDA 安装指南。请注意，clang 可能不支持某些 Linux 包管理器安装的 CUDA 工具包。Clang 确实尝试处理少数常见 Linux 发行版上 CUDA 安装的特定细节，但总的来说，使其工作的最可靠方法是从 NVIDIA 的 .run 包中将 CUDA 安装到单个目录，并通过 –cuda-path=… 参数指定其位置。

CUDA 编译在 Linux 上受支持。在 MacOS 和 Windows 上的编译可能可以工作，也可能无法工作，目前没有维护人员。

调用 clang ¶

调用 clang 进行 CUDA 编译的工作方式与编译常规 C++ 类似。您只需要注意一些额外的 flag。

您可以使用这个程序作为一个简单的示例。将其保存为 axpy.cu。（Clang 通过注意到您的文件名以 .cu 结尾来检测您正在编译 CUDA 代码。或者，您可以传递 -x cuda。）

要构建和运行，请运行以下命令，并按照下面的描述填写尖括号中的部分

$ clang++ axpy.cu -o axpy --cuda-gpu-arch=<GPU arch> \
    -L<CUDA install path>/<lib64 or lib>             \
    -lcudart_static -ldl -lrt -pthread
$ ./axpy
y[0] = 2
y[1] = 4
y[2] = 6
y[3] = 8

在 MacOS 上，将 -lcudart_static 替换为 -lcudart；否则，在运行程序时，您可能会收到“CUDA driver version is insufficient for CUDA runtime version”错误。

<CUDA 安装路径> – 您安装 CUDA SDK 的目录。通常是 /usr/local/cuda。

如果在 64 位模式下编译，则传递例如 -L/usr/local/cuda/lib64；否则，传递例如 -L/usr/local/cuda/lib。（在 CUDA 中，设备代码和主机代码始终具有相同的指针宽度，因此如果您正在为主机编译 64 位代码，那么您也在为设备编译 64 位代码。）请注意，从 v10.0 CUDA SDK 开始，不再支持 32 位应用程序的编译。
<GPU 架构> – 您的 GPU 的计算能力。例如，如果您想在计算能力为 3.5 的 GPU 上运行您的程序，请指定 --cuda-gpu-arch=sm_35。

注意：您不能将 compute_XX 作为参数传递给 --cuda-gpu-arch；目前仅支持 sm_XX。但是，clang 始终在其二进制文件中包含 PTX，因此，例如，使用 --cuda-gpu-arch=sm_30 编译的二进制文件将向前兼容例如 sm_35 GPU。

您可以多次传递 --cuda-gpu-arch 以针对多个架构进行编译。

-L 和 -l flag 只需要在链接时传递。编译时，如果您没有将 CUDA SDK 安装到 /usr/local/cuda 或 /usr/local/cuda-X.Y 中，您可能还需要传递 --cuda-path=/path/to/cuda。

控制数值代码的 Flag ¶

如果您正在使用 GPU，您可能关心如何使数值代码运行得更快。GPU 硬件比大多数 CPU 允许对数值运算进行更多控制，但这会导致您需要处理更多的编译器选项。

您可能希望调整的 Flag 包括

-ffp-contract={on,off,fast}（在编译 CUDA 时，主机和设备上默认为 fast）控制编译器是否发出 fused multiply-add 操作。
- off：从不发出 fma 操作，并阻止 ptxas 融合乘法和加法指令。
- on：在单个语句中融合乘法和加法，但从不在语句之间融合（C11 语义）。阻止 ptxas 融合其他乘法和加法。
- fast：在任何有利可图的地方融合乘法和加法，甚至跨语句。不阻止 ptxas 融合额外的乘法和加法。
Fused multiply-add 指令可能比未融合的等效指令快得多，但由于 fma 中的中间结果未四舍五入，因此此 flag 可能会影响数值代码。
-fcuda-flush-denormals-to-zero（默认值：off）启用此选项后，浮点运算可能会将非正规化数输入和/或输出刷新为 0。对非正规化数的操作通常比对正规化数的相同操作慢得多。
-fcuda-approx-transcendentals（默认值：off）启用此选项后，编译器可能会发出对更快、近似版本的超越函数的调用，而不是使用较慢、完全符合 IEEE 标准的版本。例如，此 flag 允许 clang 发出 ptx sin.approx.f32 指令。

这由 -ffast-math 隐含。

标准库支持 ¶

在 clang 和 nvcc 中，大多数 C++ 标准库在设备端不受支持。

`<math.h>` 和 `<cmath>`¶

在 clang 中，math.h 和 cmath 可用，并且通过了从 libc++ 测试套件改编的测试。

在 nvcc 中，math.h 和 cmath 大部分可用。命名空间 std 中的 ::foof 版本（例如 std::sinf）不可用，并且在标准要求重载采用整型参数的情况下，这些通常不可用。

#include <math.h>
#include <cmath.h>

// clang is OK with everything in this function.
__device__ void test() {
  std::sin(0.); // nvcc - ok
  std::sin(0);  // nvcc - error, because no std::sin(int) override is available.
  sin(0);       // nvcc - same as above.

  sinf(0.);       // nvcc - ok
  std::sinf(0.);  // nvcc - no such function
}

`<std::complex>`¶

nvcc 不正式支持 std::complex。在 __device__ 代码中使用 std::complex 是一个错误，但由于 nvcc 对“wrong-side rule”（见下文）的解释，它通常在 __host__ __device__ 代码中有效。但是，我们从实现者那里听说，可能会出现 nvcc 省略对 std::complex 函数的调用的情况，尤其是在没有优化的情况下编译时。

截至 2016-11-16，clang 支持 std::complex，没有这些注意事项。它已经过 libstdc++ 4.8.5 及更高版本的测试，但已知仅适用于 2016-11-16 之后的 libc++ 版本。

`<algorithm>`¶

在 C++14 中，<algorithm> 中的许多有用函数（特别是 std::min 和 std::max）变为 constexpr。因此，当使用 clang 编译时，您可以在设备代码中使用这些函数。

从代码中检测 clang 与 NVCC ¶

尽管 clang 的 CUDA 实现与 NVCC 的实现基本兼容，但您可能仍然想检测何时专门使用 clang 编译 CUDA 代码。

这很棘手，因为 NVCC 可能会调用 clang 作为其自身编译过程的一部分！例如，NVCC 在为设备代码编译时使用主机编译器的预处理器，而该主机编译器实际上可能是 clang。

当 clang 实际编译 CUDA 代码时 – 而不是用作 NVCC 的子工具 – 它定义了 __CUDA__ 宏。__CUDA_ARCH__ 仅在设备模式下定义（但如果 NVCC 使用 clang 作为预处理器，则也会定义）。因此，您可以使用以下代码来检测主机和设备模式下的 clang CUDA 编译

#if defined(__clang__) && defined(__CUDA__) && !defined(__CUDA_ARCH__)
// clang compiling CUDA code, host mode.
#endif

#if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__)
// clang compiling CUDA code, device mode.
#endif

clang 和 nvcc 都在 CUDA 编译期间定义 __CUDACC__。您可以通过查找 __NVCC__ 来专门检测 NVCC。

clang 和 nvcc 之间的方言差异 ¶

没有正式的 CUDA 规范，clang 和 nvcc 使用略有不同的语言方言。下面，我们描述一些差异。

本节内容令人痛苦；希望您可以跳过本节，并幸福地生活在无知的状态中。

编译模型 ¶

clang 和 nvcc 之间的大多数差异源于 clang 和 nvcc 使用的不同编译模型。nvcc 使用分离编译，其工作方式大致如下

在输入 .cu 文件上运行预处理器，将其拆分为两个源文件：H，包含主机源代码，以及 D，包含设备源代码。

对于我们正在编译的每个 GPU 架构 arch，执行以下操作

使用 nvcc proper 编译 D。结果是 P_arch 的 ptx 文件。

可选地，调用 PTX 汇编器 ptxas 以生成文件 S_arch，其中包含 arch 的 GPU 机器代码 (SASS)。

调用 fatbin 将所有 P_arch 和 S_arch 文件组合成单个“fat binary”文件 F。

使用外部主机编译器（gcc、clang 或您喜欢的任何编译器）编译 H。F 被打包到一个头文件中，该头文件被强制包含到 H 中；nvcc 生成调用此头文件的代码，例如启动内核。

clang 使用合并解析。这类似于分离编译，只是所有主机和设备代码都存在，并且在两个编译步骤中都必须在语义上正确。

对于我们正在编译的每个 GPU 架构 arch，执行以下操作

使用 clang 为设备编译输入 .cu 文件。__host__ 代码被解析并且必须在语义上正确，即使我们此时没有为主机生成代码。

此步骤的输出是 ptx 文件 P_arch。

调用 ptxas 以生成 SASS 文件 S_arch。请注意，与 nvcc 不同，clang 始终生成 SASS 代码。

调用 fatbin 将所有 P_arch 和 S_arch 文件组合成单个 fat binary 文件 F。

使用 clang 编译 H。__device__ 代码被解析并且必须在语义上正确，即使我们此时没有为设备生成代码。

F 传递给此编译，clang 将其包含在特殊的 ELF section 中，工具（如 cuobjdump）可以在其中找到它。

（您此时可能会问，为什么 clang 需要多次解析输入文件？为什么不只解析一次，然后使用 AST 为主机和每个设备架构生成代码？

不幸的是，这行不通，因为我们必须在主机编译期间和为每个 GPU 架构的设备编译期间定义不同的宏。）

clang 的方法使其对于 C++ 边缘情况非常健壮，因为它不需要在早期阶段决定保留哪些声明以及丢弃哪些声明。但它有一些您应该注意的后果。

基于 `host` 和 `device` 属性的重载 ¶

让 “H”、“D” 和 “HD” 分别代表 “__host__ 函数”、“__device__ 函数” 和 “__host__ __device__ 函数”。没有属性的函数行为与 H 相同。

nvcc 不允许您创建具有相同签名的 H 和 D 函数

// nvcc: error - function "foo" has already been defined
__host__ void foo() {}
__device__ void foo() {}

但是，nvcc 允许您使用不同的签名“重载” H 和 D 函数

// nvcc: no error
__host__ void foo(int) {}
__device__ void foo() {}

在 clang 中，__host__ 和 __device__ 属性是函数签名的一部分，因此拥有（其他方面）相同签名的 H 和 D 函数是合法的

// clang: no error
__host__ void foo() {}
__device__ void foo() {}

HD 函数不能被具有相同签名的 H 或 D 函数重载

// nvcc: error - function "foo" has already been defined
// clang: error - redefinition of 'foo'
__host__ __device__ void foo() {}
__device__ void foo() {}

// nvcc: no error
// clang: no error
__host__ __device__ void bar(int) {}
__device__ void bar() {}

在解析重载函数时，clang 会考虑调用者和被调用者的主机/设备属性。这些在重载解析期间用作决胜因素。有关完整规则集，请参阅 IdentifyCUDAPreference，但从高层次来看，它们是

D 函数优先调用其他 D 函数。HD 函数的优先级较低。

同样，H 函数优先调用其他 H 函数或 __global__ 函数（具有相同的优先级）。HD 函数的优先级较低。

HD 函数优先调用其他 HD 函数。

在为设备编译时，HD 函数调用 D 函数的优先级低于 HD 函数，调用 H 函数的优先级更低。如果它被迫调用 H 函数，则如果我们为此 HD 函数发出代码，则程序格式错误。我们称之为“wrong-side rule”，请参见下面的示例。

在为主机编译时，规则是对称的。

一些示例

__host__ void foo();
__device__ void foo();

__host__ void bar();
__host__ __device__ void bar();

__host__ void test_host() {
  foo();  // calls H overload
  bar();  // calls H overload
}

__device__ void test_device() {
  foo();  // calls D overload
  bar();  // calls HD overload
}

__host__ __device__ void test_hd() {
  foo();  // calls H overload when compiling for host, otherwise D overload
  bar();  // always calls HD overload
}

Wrong-side rule 示例

__host__ void host_only();

// We don't codegen inline functions unless they're referenced by a
// non-inline function.  inline_hd1() is called only from the host side, so
// does not generate an error.  inline_hd2() is called from the device side,
// so it generates an error.
inline __host__ __device__ void inline_hd1() { host_only(); }  // no error
inline __host__ __device__ void inline_hd2() { host_only(); }  // error

__host__ void host_fn() { inline_hd1(); }
__device__ void device_fn() { inline_hd2(); }

// This function is not inline, so it's always codegen'ed on both the host
// and the device.  Therefore, it generates an error.
__host__ __device__ void not_inline_hd() { host_only(); }

出于 wrong-side rule 的目的，模板函数也像 inline 函数一样表现：除非它们被实例化（通常是作为调用它们的过程的一部分），否则它们不会被 codegen。

clang 在 wrong-side rule 方面的行为与 nvcc 的行为相匹配，只是 nvcc 仅对 not_inline_hd 发出警告；设备代码允许调用 not_inline_hd。在其生成的代码中，nvcc 可能会完全省略 not_inline_hd 对 host_only 的调用，或者它可能会尝试在设备上为 host_only 生成代码。您得到的结果似乎取决于编译器是否选择内联 host_only。

成员函数（包括构造函数）可以使用 H 和 D 属性重载。但是，析构函数不能重载。

主机和设备函数声明的 Clang 警告 ¶

当 clang 检测到主机 (H) 和设备 (D) 函数使用相同的签名声明或定义时，它可以发出警告。默认情况下，这些警告未启用。

要启用这些警告，请使用以下编译器 flag

-Wnvcc-compat

在主机/设备上使用不同的类 ¶

有时您可能希望拥有具有不同主机/设备版本的类。

如果类的所有成员在主机和设备上都相同，您只需为类的成员函数提供重载。

但是，如果您希望您的类在主机/设备上具有不同的成员，您将无法在两个类中都提供有效的 H 和 D 重载。在这种情况下，clang 可能会对您不满意。

#ifdef __CUDA_ARCH__
struct S {
  __device__ void foo() { /* use device_only */ }
  int device_only;
};
#else
struct S {
  __host__ void foo() { /* use host_only */ }
  double host_only;
};

__device__ void test() {
  S s;
  // clang generates an error here, because during host compilation, we
  // have ifdef'ed away the __device__ overload of S::foo().  The __device__
  // overload must be present *even during host compilation*.
  S.foo();
}
#endif

我们认为您实际上不希望在 H 和 D 上拥有具有不同成员的类。例如，如果您将其中一个作为参数传递给内核，它在 H 和 D 上将具有不同的布局，因此将无法正常工作。

为了使这样的代码与 clang 兼容，我们建议您将其分离为两个类。如果您需要编写在主机和设备上都能工作的代码，请考虑编写一个重载的包装函数，该函数在主机和设备上返回不同的类型。

struct HostS { ... };
struct DeviceS { ... };

__host__ HostS MakeStruct() { return HostS(); }
__device__ DeviceS MakeStruct() { return DeviceS(); }

// Now host and device code can call MakeStruct().

不幸的是，这种习惯用法与 nvcc 不兼容，因为它不允许您基于 H/D 属性进行重载。这是一个与 clang 和 nvcc 都兼容的习惯用法

struct HostS { ... };
struct DeviceS { ... };

#ifdef __NVCC__
  #ifndef __CUDA_ARCH__
    __host__ HostS MakeStruct() { return HostS(); }
  #else
    __device__ DeviceS MakeStruct() { return DeviceS(); }
  #endif
#else
  __host__ HostS MakeStruct() { return HostS(); }
  __device__ DeviceS MakeStruct() { return DeviceS(); }
#endif

// Now host and device code can call MakeStruct().

希望您不必经常做这种事情。

优化 ¶

现代 CPU 和 GPU 在架构上差异很大，因此在 CPU 上运行速度快的代码不一定在 GPU 上运行速度快。我们对 LLVM 进行了一些更改，以使其生成良好的 GPU 代码。这些更改包括

直线标量优化 – 这些优化减少了直线代码中的冗余。
激进的推测执行 – 这主要是为了促进直线标量优化，这些优化在沿支配路径的代码上最有效。
内存空间推断 – 在 PTX 中，我们可以操作位于特定“地址空间”（global、shared、constant 或 local）中的指针，或者我们可以操作位于“generic”地址空间中的指针，该地址空间可以指向任何内容。在非 generic 地址空间中的操作速度更快，但 CUDA 中的指针没有显式注释其地址空间，因此在可能的情况下，由 LLVM 推断它。
绕过 64 位除法 – 这是我们为 PTX 后端启用的一项现有优化。

在 NVIDIA GPU 上，64 位整数除法比 32 位整数除法慢得多。我们基准测试中的许多 64 位除法的除数和被除数在运行时都适合 32 位。此优化为此常见情况提供了快速路径。
激进的循环展开和函数内联 – 对于 GPU 而言，循环展开和函数内联需要比 CPU 更激进，因为 GPU 中的控制流转移成本更高。更激进的展开和内联也促进了其他优化，例如常量传播和 SROA，有时可以将代码加速 10 倍以上。

（程序员可以使用 clang 的循环展开 pragma 和 __attribute__((always_inline)) 强制展开和内联。）

出版物 ¶

Google 团队在 CGO 2016 上发表了一篇论文，详细介绍了他们对 clang/LLVM 所做的优化。请注意，“gpucc” 不再是一个有意义的名称：相关的工具现在只是普通的 clang/LLVM。

gpucc：一个开源 GPGPU 编译器
Jingyue Wu, Artem Belevich, Eli Bendersky, Mark Heffernan, Chris Leary, Jacques Pienaar, Bjarke Roune, Rob Springer, Xuetian Weng, Robert Hundt
2016 年代码生成和优化国际研讨会 (CGO 2016) 会议记录

CGO 演讲的幻灯片

CGO 上的教程

获取帮助 ¶

要获得关于 LLVM 及其 CUDA 支持的帮助，请参阅 LLVM 社区。

文档

参与贡献

附加链接

本页