573 lines · plain
1=========================2Compiling CUDA with clang3=========================4 5.. contents::6 :local:7 8Introduction9============10 11This document describes how to compile CUDA code with clang, and gives some12details about LLVM and clang's CUDA implementations.13 14This document assumes a basic familiarity with CUDA. Information about CUDA15programming can be found in the16`CUDA programming guide17<http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html>`_.18 19Compiling CUDA Code20===================21 22Prerequisites23-------------24 25CUDA is supported since llvm 3.9. Clang currently supports CUDA 7.0 through2612.1. If clang detects a newer CUDA version, it will issue a warning and will27attempt to use detected CUDA SDK it as if it were CUDA 12.1.28 29Before you build CUDA code, you'll need to have installed the CUDA SDK. See30`NVIDIA's CUDA installation guide31<https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html>`_ for32details. Note that clang `maynot support33<https://bugs.llvm.org/show_bug.cgi?id=26966>`_ the CUDA toolkit as installed by34some Linux package managers. Clang does attempt to deal with specific details of35CUDA installation on a handful of common Linux distributions, but in general the36most reliable way to make it work is to install CUDA in a single directory from37NVIDIA's `.run` package and specify its location via `--cuda-path=...` argument.38 39CUDA compilation is supported on Linux. Compilation on macOS and Windows may or40may not work and currently have no maintainers.41 42Invoking clang43--------------44 45Invoking clang for CUDA compilation works similarly to compiling regular C++.46You just need to be aware of a few additional flags.47 48You can use `this <https://gist.github.com/855e277884eb6b388cd2f00d956c2fd4>`_49program as a toy example. Save it as ``axpy.cu``. (Clang detects that you're50compiling CUDA code by noticing that your filename ends with ``.cu``.51Alternatively, you can pass ``-x cuda``.)52 53To build and run, run the following commands, filling in the parts in angle54brackets as described below:55 56.. code-block:: console57 58 $ clang++ axpy.cu -o axpy --cuda-gpu-arch=<GPU arch> \59 -L<CUDA install path>/<lib64 or lib> \60 -lcudart_static -ldl -lrt -pthread61 $ ./axpy62 y[0] = 263 y[1] = 464 y[2] = 665 y[3] = 866 67On macOS, replace `-lcudart_static` with `-lcudart`; otherwise, you may get68"CUDA driver version is insufficient for CUDA runtime version" errors when you69run your program.70 71* ``<CUDA install path>`` -- the directory where you installed CUDA SDK.72 Typically, ``/usr/local/cuda``.73 74 Pass e.g. ``-L/usr/local/cuda/lib64`` if compiling in 64-bit mode; otherwise,75 pass e.g. ``-L/usr/local/cuda/lib``. (In CUDA, the device code and host code76 always have the same pointer widths, so if you're compiling 64-bit code for77 the host, you're also compiling 64-bit code for the device.) Note that as of78 v10.0 CUDA SDK `no longer supports compilation of 32-bit79 applications <https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html#deprecated-features>`_.80 81* ``<GPU arch>`` -- the `compute capability82 <https://developer.nvidia.com/cuda-gpus>`_ of your GPU. For example, if you83 want to run your program on a GPU with compute capability of 3.5, specify84 ``--cuda-gpu-arch=sm_35``.85 86 Note: You cannot pass ``compute_XX`` as an argument to ``--cuda-gpu-arch``;87 only ``sm_XX`` is currently supported. However, clang always includes PTX in88 its binaries, so e.g. a binary compiled with ``--cuda-gpu-arch=sm_30`` would be89 forwards-compatible with e.g. ``sm_35`` GPUs.90 91 You can pass ``--cuda-gpu-arch`` multiple times to compile for multiple archs.92 93The `-L` and `-l` flags only need to be passed when linking. When compiling,94you may also need to pass ``--cuda-path=/path/to/cuda`` if you didn't install95the CUDA SDK into ``/usr/local/cuda`` or ``/usr/local/cuda-X.Y``.96 97Flags that control numerical code98---------------------------------99 100If you're using GPUs, you probably care about making numerical code run fast.101GPU hardware allows for more control over numerical operations than most CPUs,102but this results in more compiler options for you to juggle.103 104Flags you may wish to tweak include:105 106* ``-ffp-contract={on,off,fast}`` (defaults to ``fast`` on host and device when107 compiling CUDA) Controls whether the compiler emits fused multiply-add108 operations.109 110 * ``off``: never emit fma operations, and prevent ptxas from fusing multiply111 and add instructions.112 * ``on``: fuse multiplies and adds within a single statement, but never113 across statements (C11 semantics). Prevent ptxas from fusing other114 multiplies and adds.115 * ``fast``: fuse multiplies and adds wherever profitable, even across116 statements. Doesn't prevent ptxas from fusing additional multiplies and117 adds.118 119 Fused multiply-add instructions can be much faster than the unfused120 equivalents, but because the intermediate result in an fma is not rounded,121 this flag can affect numerical code.122 123* ``-fcuda-flush-denormals-to-zero`` (default: off) When this is enabled,124 floating point operations may flush `denormal125 <https://en.wikipedia.org/wiki/Denormal_number>`_ inputs and/or outputs to 0.126 Operations on denormal numbers are often much slower than the same operations127 on normal numbers.128 129* ``-fcuda-approx-transcendentals`` (default: off) When this is enabled, the130 compiler may emit calls to faster, approximate versions of transcendental131 functions, instead of using the slower, fully IEEE-compliant versions. For132 example, this flag allows clang to emit the ptx ``sin.approx.f32``133 instruction.134 135 This is implied by ``-ffast-math``.136 137Standard library support138========================139 140In clang and nvcc, most of the C++ standard library is not supported on the141device side.142 143``<math.h>`` and ``<cmath>``144----------------------------145 146In clang, ``math.h`` and ``cmath`` are available and `pass147<https://github.com/llvm/llvm-test-suite/blob/main/External/CUDA/math_h.cu>`_148`tests149<https://github.com/llvm/llvm-test-suite/blob/main/External/CUDA/cmath.cu>`_150adapted from libc++'s test suite.151 152In nvcc ``math.h`` and ``cmath`` are mostly available. Versions of ``::foof``153in namespace std (e.g. ``std::sinf``) are not available, and where the standard154calls for overloads that take integral arguments, these are usually not155available.156 157.. code-block:: c++158 159 #include <math.h>160 #include <cmath.h>161 162 // clang is OK with everything in this function.163 __device__ void test() {164 std::sin(0.); // nvcc - ok165 std::sin(0); // nvcc - error, because no std::sin(int) override is available.166 sin(0); // nvcc - same as above.167 168 sinf(0.); // nvcc - ok169 std::sinf(0.); // nvcc - no such function170 }171 172``<std::complex>``173------------------174 175nvcc does not officially support ``std::complex``. It's an error to use176``std::complex`` in ``__device__`` code, but it often works in ``__host__177__device__`` code due to nvcc's interpretation of the "wrong-side rule" (see178below). However, we have heard from implementers that it's possible to get179into situations where nvcc will omit a call to an ``std::complex`` function,180especially when compiling without optimizations.181 182As of 2016-11-16, clang supports ``std::complex`` without these caveats. It is183tested with libstdc++ 4.8.5 and newer, but is known to work only with libc++184newer than 2016-11-16.185 186``<algorithm>``187---------------188 189In C++14, many useful functions from ``<algorithm>`` (notably, ``std::min`` and190``std::max``) become constexpr. You can therefore use these in device code,191when compiling with clang.192 193Detecting clang vs NVCC from code194=================================195 196Although clang's CUDA implementation is largely compatible with NVCC's, you may197still want to detect when you're compiling CUDA code specifically with clang.198 199This is tricky, because NVCC may invoke clang as part of its own compilation200process! For example, NVCC uses the host compiler's preprocessor when201compiling for device code, and that host compiler may in fact be clang.202 203When clang is actually compiling CUDA code -- rather than being used as a204subtool of NVCC's -- it defines the ``__CUDA__`` macro. ``__CUDA_ARCH__`` is205defined only in device mode (but will be defined if NVCC is using clang as a206preprocessor). So you can use the following incantations to detect clang CUDA207compilation, in host and device modes:208 209.. code-block:: c++210 211 #if defined(__clang__) && defined(__CUDA__) && !defined(__CUDA_ARCH__)212 // clang compiling CUDA code, host mode.213 #endif214 215 #if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__)216 // clang compiling CUDA code, device mode.217 #endif218 219Both clang and nvcc define ``__CUDACC__`` during CUDA compilation. You can220detect NVCC specifically by looking for ``__NVCC__``.221 222Dialect Differences Between clang and nvcc223==========================================224 225There is no formal CUDA spec, and clang and nvcc speak slightly different226dialects of the language. Below, we describe some of the differences.227 228This section is painful; hopefully you can skip this section and live your life229blissfully unaware.230 231Compilation Models232------------------233 234Most of the differences between clang and nvcc stem from the different235compilation models used by clang and nvcc. nvcc uses *split compilation*,236which works roughly as follows:237 238 * Run a preprocessor over the input ``.cu`` file to split it into two source239 files: ``H``, containing source code for the host, and ``D``, containing240 source code for the device.241 242 * For each GPU architecture ``arch`` that we're compiling for, do:243 244 * Compile ``D`` using nvcc proper. The result of this is a ``ptx`` file for245 ``P_arch``.246 247 * Optionally, invoke ``ptxas``, the PTX assembler, to generate a file,248 ``S_arch``, containing GPU machine code (SASS) for ``arch``.249 250 * Invoke ``fatbin`` to combine all ``P_arch`` and ``S_arch`` files into a251 single "fat binary" file, ``F``.252 253 * Compile ``H`` using an external host compiler (gcc, clang, or whatever you254 like). ``F`` is packaged up into a header file which is force-included into255 ``H``; nvcc generates code that calls into this header to e.g. launch256 kernels.257 258clang uses *merged parsing*. This is similar to split compilation, except all259of the host and device code is present and must be semantically-correct in both260compilation steps.261 262 * For each GPU architecture ``arch`` that we're compiling for, do:263 264 * Compile the input ``.cu`` file for device, using clang. ``__host__`` code265 is parsed and must be semantically correct, even though we're not266 generating code for the host at this time.267 268 The output of this step is a ``ptx`` file ``P_arch``.269 270 * Invoke ``ptxas`` to generate a SASS file, ``S_arch``. Note that, unlike271 nvcc, clang always generates SASS code.272 273 * Invoke ``fatbin`` to combine all ``P_arch`` and ``S_arch`` files into a274 single fat binary file, ``F``.275 276 * Compile ``H`` using clang. ``__device__`` code is parsed and must be277 semantically correct, even though we're not generating code for the device278 at this time.279 280 ``F`` is passed to this compilation, and clang includes it in a special ELF281 section, where it can be found by tools like ``cuobjdump``.282 283(You may ask at this point, why does clang need to parse the input file284multiple times? Why not parse it just once, and then use the AST to generate285code for the host and each device architecture?286 287Unfortunately this can't work because we have to define different macros during288host compilation and during device compilation for each GPU architecture.)289 290clang's approach allows it to be highly robust to C++ edge cases, as it doesn't291need to decide at an early stage which declarations to keep and which to throw292away. But it has some consequences you should be aware of.293 294Overloading Based on ``__host__`` and ``__device__`` Attributes295---------------------------------------------------------------296 297Let "H", "D", and "HD" stand for "``__host__`` functions", "``__device__``298functions", and "``__host__ __device__`` functions", respectively. Functions299with no attributes behave the same as H.300 301nvcc does not allow you to create H and D functions with the same signature:302 303.. code-block:: c++304 305 // nvcc: error - function "foo" has already been defined306 __host__ void foo() {}307 __device__ void foo() {}308 309However, nvcc allows you to "overload" H and D functions with different310signatures:311 312.. code-block:: c++313 314 // nvcc: no error315 __host__ void foo(int) {}316 __device__ void foo() {}317 318In clang, the ``__host__`` and ``__device__`` attributes are part of a319function's signature, and so it's legal to have H and D functions with320(otherwise) the same signature:321 322.. code-block:: c++323 324 // clang: no error325 __host__ void foo() {}326 __device__ void foo() {}327 328HD functions cannot be overloaded by H or D functions with the same signature:329 330.. code-block:: c++331 332 // nvcc: error - function "foo" has already been defined333 // clang: error - redefinition of 'foo'334 __host__ __device__ void foo() {}335 __device__ void foo() {}336 337 // nvcc: no error338 // clang: no error339 __host__ __device__ void bar(int) {}340 __device__ void bar() {}341 342When resolving an overloaded function, clang considers the host/device343attributes of the caller and callee. These are used as a tiebreaker during344overload resolution. See `IdentifyCUDAPreference345<https://clang.llvm.org/doxygen/SemaCUDA_8cpp.html>`_ for the full set of rules,346but at a high level they are:347 348 * D functions prefer to call other Ds. HDs are given lower priority.349 350 * Similarly, H functions prefer to call other Hs, or ``__global__`` functions351 (with equal priority). HDs are given lower priority.352 353 * HD functions prefer to call other HDs.354 355 When compiling for device, HDs will call Ds with lower priority than HD, and356 will call Hs with still lower priority. If it's forced to call an H, the357 program is malformed if we emit code for this HD function. We call this the358 "wrong-side rule", see example below.359 360 The rules are symmetrical when compiling for host.361 362Some examples:363 364.. code-block:: c++365 366 __host__ void foo();367 __device__ void foo();368 369 __host__ void bar();370 __host__ __device__ void bar();371 372 __host__ void test_host() {373 foo(); // calls H overload374 bar(); // calls H overload375 }376 377 __device__ void test_device() {378 foo(); // calls D overload379 bar(); // calls HD overload380 }381 382 __host__ __device__ void test_hd() {383 foo(); // calls H overload when compiling for host, otherwise D overload384 bar(); // always calls HD overload385 }386 387Wrong-side rule example:388 389.. code-block:: c++390 391 __host__ void host_only();392 393 // We don't codegen inline functions unless they're referenced by a394 // non-inline function. inline_hd1() is called only from the host side, so395 // does not generate an error. inline_hd2() is called from the device side,396 // so it generates an error.397 inline __host__ __device__ void inline_hd1() { host_only(); } // no error398 inline __host__ __device__ void inline_hd2() { host_only(); } // error399 400 __host__ void host_fn() { inline_hd1(); }401 __device__ void device_fn() { inline_hd2(); }402 403 // This function is not inline, so it's always codegen'ed on both the host404 // and the device. Therefore, it generates an error.405 __host__ __device__ void not_inline_hd() { host_only(); }406 407For the purposes of the wrong-side rule, templated functions also behave like408``inline`` functions: They aren't codegen'ed unless they're instantiated409(usually as part of the process of invoking them).410 411clang's behavior with respect to the wrong-side rule matches nvcc's, except412nvcc only emits a warning for ``not_inline_hd``; device code is allowed to call413``not_inline_hd``. In its generated code, nvcc may omit ``not_inline_hd``'s414call to ``host_only`` entirely, or it may try to generate code for415``host_only`` on the device. What you get seems to depend on whether or not416the compiler chooses to inline ``host_only``.417 418Member functions, including constructors, may be overloaded using H and D419attributes. However, destructors cannot be overloaded.420 421Clang Warnings for Host and Device Function Declarations422--------------------------------------------------------423 424Clang can emit warnings when it detects that host (H) and device (D) functions are declared or defined with the same signature. These warnings are not enabled by default.425 426To enable these warnings, use the following compiler flag:427 428.. code-block:: console429 430 -Wnvcc-compat431 432Using a Different Class on Host/Device433--------------------------------------434 435Occasionally you may want to have a class with different host/device versions.436 437If all of the class's members are the same on the host and device, you can just438provide overloads for the class's member functions.439 440However, if you want your class to have different members on host/device, you441won't be able to provide working H and D overloads in both classes. In this442case, clang is likely to be unhappy with you.443 444.. code-block:: c++445 446 #ifdef __CUDA_ARCH__447 struct S {448 __device__ void foo() { /* use device_only */ }449 int device_only;450 };451 #else452 struct S {453 __host__ void foo() { /* use host_only */ }454 double host_only;455 };456 457 __device__ void test() {458 S s;459 // clang generates an error here, because during host compilation, we460 // have ifdef'ed away the __device__ overload of S::foo(). The __device__461 // overload must be present *even during host compilation*.462 S.foo();463 }464 #endif465 466We posit that you don't really want to have classes with different members on H467and D. For example, if you were to pass one of these as a parameter to a468kernel, it would have a different layout on H and D, so would not work469properly.470 471To make code like this compatible with clang, we recommend you separate it out472into two classes. If you need to write code that works on both host and473device, consider writing an overloaded wrapper function that returns different474types on host and device.475 476.. code-block:: c++477 478 struct HostS { ... };479 struct DeviceS { ... };480 481 __host__ HostS MakeStruct() { return HostS(); }482 __device__ DeviceS MakeStruct() { return DeviceS(); }483 484 // Now host and device code can call MakeStruct().485 486Unfortunately, this idiom isn't compatible with nvcc, because it doesn't allow487you to overload based on the H/D attributes. Here's an idiom that works with488both clang and nvcc:489 490.. code-block:: c++491 492 struct HostS { ... };493 struct DeviceS { ... };494 495 #ifdef __NVCC__496 #ifndef __CUDA_ARCH__497 __host__ HostS MakeStruct() { return HostS(); }498 #else499 __device__ DeviceS MakeStruct() { return DeviceS(); }500 #endif501 #else502 __host__ HostS MakeStruct() { return HostS(); }503 __device__ DeviceS MakeStruct() { return DeviceS(); }504 #endif505 506 // Now host and device code can call MakeStruct().507 508Hopefully you don't have to do this sort of thing often.509 510Optimizations511=============512 513Modern CPUs and GPUs are architecturally quite different, so code that's fast514on a CPU isn't necessarily fast on a GPU. We've made a number of changes to515LLVM to make it generate good GPU code. Among these changes are:516 517* `Straight-line scalar optimizations <https://docs.google.com/document/d/1momWzKFf4D6h8H3YlfgKQ3qeZy5ayvMRh6yR-Xn2hUE>`_ -- These518 reduce redundancy within straight-line code.519 520* `Aggressive speculative execution521 <https://llvm.org/docs/doxygen/html/SpeculativeExecution_8cpp_source.html>`_522 -- This is mainly for promoting straight-line scalar optimizations, which are523 most effective on code along dominator paths.524 525* `Memory space inference526 <https://llvm.org/doxygen/InferAddressSpaces_8cpp_source.html>`_ --527 In PTX, we can operate on pointers that are in a particular "address space"528 (global, shared, constant, or local), or we can operate on pointers in the529 "generic" address space, which can point to anything. Operations in a530 non-generic address space are faster, but pointers in CUDA are not explicitly531 annotated with their address space, so it's up to LLVM to infer it where532 possible.533 534* `Bypassing 64-bit divides535 <https://llvm.org/docs/doxygen/html/BypassSlowDivision_8cpp_source.html>`_ --536 This was an existing optimization that we enabled for the PTX backend.537 538 64-bit integer divides are much slower than 32-bit ones on NVIDIA GPUs.539 Many of the 64-bit divides in our benchmarks have a divisor and dividend540 which fit in 32-bits at runtime. This optimization provides a fast path for541 this common case.542 543* Aggressive loop unrolling and function inlining -- Loop unrolling and544 function inlining need to be more aggressive for GPUs than for CPUs because545 control flow transfer in GPU is more expensive. More aggressive unrolling and546 inlining also promote other optimizations, such as constant propagation and547 SROA, which sometimes speed up code by over 10x.548 549 (Programmers can force unrolling and inline using clang's `loop unrolling pragmas550 <https://clang.llvm.org/docs/AttributeReference.html#pragma-unroll-pragma-nounroll>`_551 and ``__attribute__((always_inline))``.)552 553Publication554===========555 556The team at Google published a paper in CGO 2016 detailing the optimizations557they'd made to clang/LLVM. Note that "gpucc" is no longer a meaningful name:558The relevant tools are now just vanilla clang/LLVM.559 560| `gpucc: An Open-Source GPGPU Compiler <http://dl.acm.org/citation.cfm?id=2854041>`_561| Jingyue Wu, Artem Belevich, Eli Bendersky, Mark Heffernan, Chris Leary, Jacques Pienaar, Bjarke Roune, Rob Springer, Xuetian Weng, Robert Hundt562| *Proceedings of the 2016 International Symposium on Code Generation and Optimization (CGO 2016)*563|564| `Slides from the CGO talk <http://wujingyue.github.io/docs/gpucc-talk.pdf>`_565|566| `Tutorial given at CGO <http://wujingyue.github.io/docs/gpucc-tutorial.pdf>`_567 568Obtaining Help569==============570 571To obtain help on LLVM in general and its CUDA support, see `the LLVM572community <https://llvm.org/docs/#mailing-lists>`_.573