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