mirror of
https://github.com/RPCS3/llvm-mirror.git
synced 2024-11-25 20:23:11 +01:00
170 lines
6.2 KiB
ReStructuredText
170 lines
6.2 KiB
ReStructuredText
===================================
|
|
Compiling CUDA C/C++ with LLVM
|
|
===================================
|
|
|
|
.. contents::
|
|
:local:
|
|
|
|
Introduction
|
|
============
|
|
|
|
This document contains the user guides and the internals of compiling CUDA
|
|
C/C++ with LLVM. It is aimed at both users who want to compile CUDA with LLVM
|
|
and developers who want to improve LLVM for GPUs. This document assumes a basic
|
|
familiarity with CUDA. Information about CUDA programming can be found in the
|
|
`CUDA programming guide
|
|
<http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html>`_.
|
|
|
|
How to Build LLVM with CUDA Support
|
|
===================================
|
|
|
|
Below is a quick summary of downloading and building LLVM. Consult the `Getting
|
|
Started <http://llvm.org/docs/GettingStarted.html>`_ page for more details on
|
|
setting up LLVM.
|
|
|
|
#. Checkout LLVM
|
|
|
|
.. code-block:: console
|
|
|
|
$ cd where-you-want-llvm-to-live
|
|
$ svn co http://llvm.org/svn/llvm-project/llvm/trunk llvm
|
|
|
|
#. Checkout Clang
|
|
|
|
.. code-block:: console
|
|
|
|
$ cd where-you-want-llvm-to-live
|
|
$ cd llvm/tools
|
|
$ svn co http://llvm.org/svn/llvm-project/cfe/trunk clang
|
|
|
|
#. Configure and build LLVM and Clang
|
|
|
|
.. code-block:: console
|
|
|
|
$ cd where-you-want-llvm-to-live
|
|
$ mkdir build
|
|
$ cd build
|
|
$ cmake [options] ..
|
|
$ make
|
|
|
|
How to Compile CUDA C/C++ with LLVM
|
|
===================================
|
|
|
|
We assume you have installed the CUDA driver and runtime. Consult the `NVIDIA
|
|
CUDA installation Guide
|
|
<https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html>`_ if
|
|
you have not.
|
|
|
|
Suppose you want to compile and run the following CUDA program (``axpy.cu``)
|
|
which multiplies a ``float`` array by a ``float`` scalar (AXPY).
|
|
|
|
.. code-block:: c++
|
|
|
|
#include <helper_cuda.h> // for checkCudaErrors
|
|
|
|
#include <iostream>
|
|
|
|
__global__ void axpy(float a, float* x, float* y) {
|
|
y[threadIdx.x] = a * x[threadIdx.x];
|
|
}
|
|
|
|
int main(int argc, char* argv[]) {
|
|
const int kDataLen = 4;
|
|
|
|
float a = 2.0f;
|
|
float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
|
|
float host_y[kDataLen];
|
|
|
|
// Copy input data to device.
|
|
float* device_x;
|
|
float* device_y;
|
|
checkCudaErrors(cudaMalloc(&device_x, kDataLen * sizeof(float)));
|
|
checkCudaErrors(cudaMalloc(&device_y, kDataLen * sizeof(float)));
|
|
checkCudaErrors(cudaMemcpy(device_x, host_x, kDataLen * sizeof(float),
|
|
cudaMemcpyHostToDevice));
|
|
|
|
// Launch the kernel.
|
|
axpy<<<1, kDataLen>>>(a, device_x, device_y);
|
|
|
|
// Copy output data to host.
|
|
checkCudaErrors(cudaDeviceSynchronize());
|
|
checkCudaErrors(cudaMemcpy(host_y, device_y, kDataLen * sizeof(float),
|
|
cudaMemcpyDeviceToHost));
|
|
|
|
// Print the results.
|
|
for (int i = 0; i < kDataLen; ++i) {
|
|
std::cout << "y[" << i << "] = " << host_y[i] << "\n";
|
|
}
|
|
|
|
checkCudaErrors(cudaDeviceReset());
|
|
return 0;
|
|
}
|
|
|
|
The command line for compilation is similar to what you would use for C++.
|
|
|
|
.. code-block:: console
|
|
|
|
$ 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
|
|
$ ./axpy
|
|
y[0] = 2
|
|
y[1] = 4
|
|
y[2] = 6
|
|
y[3] = 8
|
|
|
|
Note that ``helper_cuda.h`` comes from the CUDA samples, so you need the
|
|
samples installed for this example. ``<CUDA install path>`` is the root
|
|
directory where you installed CUDA SDK, typically ``/usr/local/cuda``.
|
|
|
|
Optimizations
|
|
=============
|
|
|
|
CPU and GPU have different design philosophies and architectures. For example, a
|
|
typical CPU has branch prediction, out-of-order execution, and is superscalar,
|
|
whereas a typical GPU has none of these. Due to such differences, an
|
|
optimization pipeline well-tuned for CPUs may be not suitable for GPUs.
|
|
|
|
LLVM performs several general and CUDA-specific optimizations for GPUs. The
|
|
list below shows some of the more important optimizations for GPUs. Most of
|
|
them have been upstreamed to ``lib/Transforms/Scalar`` and
|
|
``lib/Target/NVPTX``. A few of them have not been upstreamed due to lack of a
|
|
customizable target-independent optimization pipeline.
|
|
|
|
* **Straight-line scalar optimizations**. These optimizations reduce redundancy
|
|
in straight-line code. Details can be found in the `design document for
|
|
straight-line scalar optimizations <https://goo.gl/4Rb9As>`_.
|
|
|
|
* **Inferring memory spaces**. `This optimization
|
|
<http://www.llvm.org/docs/doxygen/html/NVPTXFavorNonGenericAddrSpaces_8cpp_source.html>`_
|
|
infers the memory space of an address so that the backend can emit faster
|
|
special loads and stores from it. Details can be found in the `design
|
|
document for memory space inference <https://goo.gl/5wH2Ct>`_.
|
|
|
|
* **Aggressive loop unrooling and function inlining**. Loop unrolling and
|
|
function inlining need to be more aggressive for GPUs than for CPUs because
|
|
control flow transfer in GPU is more expensive. They also promote other
|
|
optimizations such as constant propagation and SROA which sometimes speed up
|
|
code by over 10x. An empirical inline threshold for GPUs is 1100. This
|
|
configuration has yet to be upstreamed with a target-specific optimization
|
|
pipeline. LLVM also provides `loop unrolling pragmas
|
|
<http://clang.llvm.org/docs/AttributeReference.html#pragma-unroll-pragma-nounroll>`_
|
|
and ``__attribute__((always_inline))`` for programmers to force unrolling and
|
|
inling.
|
|
|
|
* **Aggressive speculative execution**. `This transformation
|
|
<http://llvm.org/docs/doxygen/html/SpeculativeExecution_8cpp_source.html>`_ is
|
|
mainly for promoting straight-line scalar optimizations which are most
|
|
effective on code along dominator paths.
|
|
|
|
* **Memory-space alias analysis**. `This alias analysis
|
|
<http://reviews.llvm.org/D12414>`_ infers that two pointers in different
|
|
special memory spaces do not alias. It has yet to be integrated to the new
|
|
alias analysis infrastructure; the new infrastructure does not run
|
|
target-specific alias analysis.
|
|
|
|
* **Bypassing 64-bit divides**. `An existing optimization
|
|
<http://llvm.org/docs/doxygen/html/BypassSlowDivision_8cpp_source.html>`_
|
|
enabled in the NVPTX backend. 64-bit integer divides are much slower than
|
|
32-bit ones on NVIDIA GPUs due to lack of a divide unit. Many of the 64-bit
|
|
divides in our benchmarks have a divisor and dividend which fit in 32-bits at
|
|
runtime. This optimization provides a fast path for this common case.
|