1 ===================================
2 Compiling CUDA C/C++ with LLVM
3 ===================================
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>`_.
18 How to Build LLVM with CUDA Support
19 ===================================
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
29 .. code-block:: console
31 $ cd where-you-want-llvm-to-live
32 $ svn co http://llvm.org/svn/llvm-project/llvm/trunk llvm
36 .. code-block:: console
38 $ cd where-you-want-llvm-to-live
40 $ svn co http://llvm.org/svn/llvm-project/cfe/trunk clang
42 #. Apply the temporary patch for CUDA support.
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
48 .. code-block:: console
50 $ cd where-you-want-llvm-to-live
54 Otherwise, go to `its review page <http://reviews.llvm.org/D14452>`_,
55 download the raw diff, and apply it manually using
57 .. code-block:: console
59 $ cd where-you-want-llvm-to-live
61 $ patch -p0 < D14452.diff
63 #. Configure and build LLVM and Clang
65 .. code-block:: console
67 $ cd where-you-want-llvm-to-live
73 How to Compile CUDA C/C++ with LLVM
74 ===================================
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
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).
86 #include <helper_cuda.h> // for checkCudaErrors
90 __global__ void axpy(float a, float* x, float* y) {
91 y[threadIdx.x] = a * x[threadIdx.x];
94 int main(int argc, char* argv[]) {
95 const int kDataLen = 4;
98 float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
99 float host_y[kDataLen];
101 // Copy input data to device.
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));
109 // Launch the kernel.
110 axpy<<<1, kDataLen>>>(a, device_x, device_y);
112 // Copy output data to host.
113 checkCudaErrors(cudaDeviceSynchronize());
114 checkCudaErrors(cudaMemcpy(host_y, device_y, kDataLen * sizeof(float),
115 cudaMemcpyDeviceToHost));
117 // Print the results.
118 for (int i = 0; i < kDataLen; ++i) {
119 std::cout << "y[" << i << "] = " << host_y[i] << "\n";
122 checkCudaErrors(cudaDeviceReset());
126 The command line for compilation is similar to what you would use for C++.
128 .. code-block:: console
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
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``.
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.
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.
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>`_.
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>`_.
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
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.
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.
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.