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.
CUDA support is still in development and works the best in the trunk version of LLVM. Below is a quick summary of downloading and building the trunk version. Consult the Getting Started page for more details on setting up LLVM.
Checkout LLVM
$ cd where-you-want-llvm-to-live
$ svn co http://llvm.org/svn/llvm-project/llvm/trunk llvm
Checkout Clang
$ 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
$ cd where-you-want-llvm-to-live
$ mkdir build
$ cd build
$ cmake [options] ..
$ make
We assume you have installed the CUDA driver and runtime. Consult the NVIDIA CUDA installation guide 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).
#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;
cudaMalloc(&device_x, kDataLen * sizeof(float));
cudaMalloc(&device_y, kDataLen * sizeof(float));
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.
cudaDeviceSynchronize();
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";
}
cudaDeviceReset();
return 0;
}
The command line for compilation is similar to what you would use for C++.
$ clang++ axpy.cu -o axpy --cuda-gpu-arch=<GPU arch> \
-L<CUDA install path>/<lib64 or lib> \
-lcudart_static -ldl -lrt -pthread
$ ./axpy
y[0] = 2
y[1] = 4
y[2] = 6
y[3] = 8
<CUDA install path> is the root directory where you installed CUDA SDK, typically /usr/local/cuda. <GPU arch> is the compute capability of your GPU. For example, if you want to run your program on a GPU with compute capability of 3.5, you should specify --cuda-gpu-arch=sm_35.
Although clang’s CUDA implementation is largely compatible with NVCC’s, you may still want to detect when you’re compiling CUDA code specifically with clang.
This is tricky, because NVCC may invoke clang as part of its own compilation process! For example, NVCC uses the host compiler’s preprocessor when compiling for device code, and that host compiler may in fact be clang.
When clang is actually compiling CUDA code – rather than being used as a subtool of NVCC’s – it defines the __CUDA__ macro. __CUDA_ARCH__ is defined only in device mode (but will be defined if NVCC is using clang as a preprocessor). So you can use the following incantations to detect clang CUDA compilation, in host and device modes:
#if defined(__clang__) && defined(__CUDA__) && !defined(__CUDA_ARCH__)
// clang compiling CUDA code, host mode.
#endif
#if defined(__clang__) && defined(__CUDA__) && defined(__CUDA_ARCH__)
// clang compiling CUDA code, device mode.
#endif
Both clang and nvcc define __CUDACC__ during CUDA compilation. You can detect NVCC specifically by looking for __NVCC__.
If you’re using GPUs, you probably care about making numerical code run fast. GPU hardware allows for more control over numerical operations than most CPUs, but this results in more compiler options for you to juggle.
Flags you may wish to tweak include:
-ffp-contract={on,off,fast} (defaults to fast on host and device when compiling CUDA) Controls whether the compiler emits fused multiply-add operations.
Fused multiply-add instructions can be much faster than the unfused equivalents, but because the intermediate result in an fma is not rounded, this flag can affect numerical code.
-fcuda-flush-denormals-to-zero (default: off) When this is enabled, floating point operations may flush denormal inputs and/or outputs to 0. Operations on denormal numbers are often much slower than the same operations on normal numbers.
-fcuda-approx-transcendentals (default: off) When this is enabled, the compiler may emit calls to faster, approximate versions of transcendental functions, instead of using the slower, fully IEEE-compliant versions. For example, this flag allows clang to emit the ptx sin.approx.f32 instruction.
This is implied by -ffast-math.
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.
To obtain help on LLVM in general and its CUDA support, see the LLVM community.