docs/CompileCudaWithLLVM.rst

   1 ===================================
   2 Compiling CUDA C/C++ with LLVM
   3 ===================================
   4
   5 .. contents::
   6    :local:
   7
   8 Introduction
   9 ============
  10
  11 This document contains the user guides and the internals of compiling CUDA
  12 C/C++ with LLVM. It is aimed at both users who want to compile CUDA with LLVM
  13 and developers who want to improve LLVM for GPUs. This document assumes a basic
  14 familiarity with CUDA. Information about CUDA programming can be found in the
  15 `CUDA programming guide
  16 <http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html>`_.
  17
  18 How to Build LLVM with CUDA Support
  19 ===================================
  20
  21 The support for CUDA is still in progress and temporarily relies on `this patch
  22 <http://reviews.llvm.org/D14452>`_. Below is a quick summary of downloading and
  23 building LLVM with CUDA support. Consult the `Getting Started
  24 <http://llvm.org/docs/GettingStarted.html>`_ page for more details on setting
  25 up LLVM.
  26
  27 #. Checkout LLVM
  28
  29    .. code-block:: console
  30
  31      $ cd where-you-want-llvm-to-live
  32      $ svn co http://llvm.org/svn/llvm-project/llvm/trunk llvm
  33
  34 #. Checkout Clang
  35
  36    .. code-block:: console
  37
  38      $ cd where-you-want-llvm-to-live
  39      $ cd llvm/tools
  40      $ svn co http://llvm.org/svn/llvm-project/cfe/trunk clang
  41
  42 #. Apply the temporary patch for CUDA support.
  43
  44    If you have installed `Arcanist
  45    <http://llvm.org/docs/Phabricator.html#requesting-a-review-via-the-command-line>`_,
  46    you can apply this patch using
  47
  48    .. code-block:: console
  49
  50      $ cd where-you-want-llvm-to-live
  51      $ cd llvm/tools/clang
  52      $ arc patch D14452
  53
  54    Otherwise, go to `its review page <http://reviews.llvm.org/D14452>`_,
  55    download the raw diff, and apply it manually using
  56
  57    .. code-block:: console
  58
  59      $ cd where-you-want-llvm-to-live
  60      $ cd llvm/tools/clang
  61      $ patch -p0 < D14452.diff
  62
  63 #. Configure and build LLVM and Clang
  64
  65    .. code-block:: console
  66
  67      $ cd where-you-want-llvm-to-live
  68      $ mkdir build
  69      $ cd build
  70      $ cmake [options] ..
  71      $ make
  72
  73 How to Compile CUDA C/C++ with LLVM
  74 ===================================
  75
  76 We assume you have installed the CUDA driver and runtime. Consult the `NVIDIA
  77 CUDA installation Guide
  78 <https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html>`_ if
  79 you have not.
  80
  81 Suppose you want to compile and run the following CUDA program (``axpy.cu``)
  82 which multiplies a ``float`` array by a ``float`` scalar (AXPY).
  83
  84 .. code-block:: c++
  85
  86   #include <helper_cuda.h> // for checkCudaErrors
  87
  88   #include <iostream>
  89
  90   __global__ void axpy(float a, float* x, float* y) {
  91     y[threadIdx.x] = a * x[threadIdx.x];
  92   }
  93
  94   int main(int argc, char* argv[]) {
  95     const int kDataLen = 4;
  96
  97     float a = 2.0f;
  98     float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
  99     float host_y[kDataLen];
 100
 101     // Copy input data to device.
 102     float* device_x;
 103     float* device_y;
 104     checkCudaErrors(cudaMalloc(&device_x, kDataLen * sizeof(float)));
 105     checkCudaErrors(cudaMalloc(&device_y, kDataLen * sizeof(float)));
 106     checkCudaErrors(cudaMemcpy(device_x, host_x, kDataLen * sizeof(float),
 107                                cudaMemcpyHostToDevice));
 108
 109     // Launch the kernel.
 110     axpy<<<1, kDataLen>>>(a, device_x, device_y);
 111
 112     // Copy output data to host.
 113     checkCudaErrors(cudaDeviceSynchronize());
 114     checkCudaErrors(cudaMemcpy(host_y, device_y, kDataLen * sizeof(float),
 115                                cudaMemcpyDeviceToHost));
 116
 117     // Print the results.
 118     for (int i = 0; i < kDataLen; ++i) {
 119       std::cout << "y[" << i << "] = " << host_y[i] << "\n";
 120     }
 121
 122     checkCudaErrors(cudaDeviceReset());
 123     return 0;
 124   }
 125
 126 The command line for compilation is similar to what you would use for C++.
 127
 128 .. code-block:: console
 129
 130   $ clang++ -o axpy -I<CUDA install path>/samples/common/inc -L<CUDA install path>/<lib64 or lib> axpy.cu -lcudart_static -lcuda -ldl -lrt -pthread
 131   $ ./axpy
 132   y[0] = 2
 133   y[1] = 4
 134   y[2] = 6
 135   y[3] = 8
 136
 137 Note that ``helper_cuda.h`` comes from the CUDA samples, so you need the
 138 samples installed for this example. ``<CUDA install path>`` is the root
 139 directory where you installed CUDA SDK, typically ``/usr/local/cuda``.
 140
 141 Optimizations
 142 =============
 143
 144 CPU and GPU have different design philosophies and architectures. For example, a
 145 typical CPU has branch prediction, out-of-order execution, and is superscalar,
 146 whereas a typical GPU has none of these. Due to such differences, an
 147 optimization pipeline well-tuned for CPUs may be not suitable for GPUs.
 148
 149 LLVM performs several general and CUDA-specific optimizations for GPUs. The
 150 list below shows some of the more important optimizations for GPUs. Most of
 151 them have been upstreamed to ``lib/Transforms/Scalar`` and
 152 ``lib/Target/NVPTX``. A few of them have not been upstreamed due to lack of a
 153 customizable target-independent optimization pipeline.
 154
 155 * **Straight-line scalar optimizations**. These optimizations reduce redundancy
 156   in straight-line code. Details can be found in the `design document for
 157   straight-line scalar optimizations <https://goo.gl/4Rb9As>`_.
 158
 159 * **Inferring memory spaces**. `This optimization
 160   <http://www.llvm.org/docs/doxygen/html/NVPTXFavorNonGenericAddrSpaces_8cpp_source.html>`_
 161   infers the memory space of an address so that the backend can emit faster
 162   special loads and stores from it. Details can be found in the `design
 163   document for memory space inference <https://goo.gl/5wH2Ct>`_.
 164
 165 * **Aggressive loop unrooling and function inlining**. Loop unrolling and
 166   function inlining need to be more aggressive for GPUs than for CPUs because
 167   control flow transfer in GPU is more expensive. They also promote other
 168   optimizations such as constant propagation and SROA which sometimes speed up
 169   code by over 10x. An empirical inline threshold for GPUs is 1100. This
 170   configuration has yet to be upstreamed with a target-specific optimization
 171   pipeline. LLVM also provides `loop unrolling pragmas
 172   <http://clang.llvm.org/docs/AttributeReference.html#pragma-unroll-pragma-nounroll>`_
 173   and ``__attribute__((always_inline))`` for programmers to force unrolling and
 174   inling.
 175
 176 * **Aggressive speculative execution**. `This transformation
 177   <http://llvm.org/docs/doxygen/html/SpeculativeExecution_8cpp_source.html>`_ is
 178   mainly for promoting straight-line scalar optimizations which are most
 179   effective on code along dominator paths.
 180
 181 * **Memory-space alias analysis**. `This alias analysis
 182   <http://llvm.org/docs/NVPTXUsage.html>`_ infers that two pointers in different
 183   special memory spaces do not alias. It has yet to be integrated to the new
 184   alias analysis infrastructure; the new infrastructure does not run
 185   target-specific alias analysis.
 186
 187 * **Bypassing 64-bit divides**. `An existing optimization
 188   <http://llvm.org/docs/doxygen/html/BypassSlowDivision_8cpp_source.html>`_
 189   enabled in the NVPTX backend. 64-bit integer divides are much slower than
 190   32-bit ones on NVIDIA GPUs due to lack of a divide unit. Many of the 64-bit
 191   divides in our benchmarks have a divisor and dividend which fit in 32-bits at
 192   runtime. This optimization provides a fast path for this common case.