History

In C++ projects, it's a common convention to separate declarations and definitions into headers and sourcefiles, respectively. That way, the sourcefiles are compiled independently and linked so that the library or program can access all of the necessary function definitions. This has a number of advantages: it helps avoid ODR violations and ensures that complicated functions are only compiled and optimized once.

However, in the early days of CUDA, the tools did not support device linking. This meant there was no way to put a kernel's definition in a separate translation unit from the __device__ functions it used. As a result, it became a common practice for CUDA programs to put function definitions in headers with an inline attribute (to avoid ODR violations). Performance-wise, putting the function definitions in headers worked great, but it was seen as a faux-pas because of the unnecessary duplication of work to parse, compile and optimize the header's __device__ functions in multiple translation units.

Things changed with the CUDA 5.0 release. The tools added support for device linking which enabled developers to put __device__ function definitions in their own translation units and later link them with kernels. This helped improve compilation times, but it had an important downside: there was a performance penalty. The __device__ functions could be locally optimized and later linked into a kernel, but some optimizations were not possible with "separable compilation". As a result, if you adopted this pattern, the code it produced was noticeably slower than if one had used the header-defined approach. Early adopters of CUDA tended to care deeply about their application's performance, which meant that separable compilation didn't see much use at this time.

CUDA 11 improved separable compilation again by introducing device link-time optimization (dlto). This largely remedied the performance regression associated with separable compilation, but also made the device link step take longer.

So, the question is: as of 2026, should my project define functions in headers, or use separable compilation (and if so, should it do so with or without device linking)? In this document, we'll write up an example project and explore the run time and compile time implications of these different approaches.

Project Outline

Let's consider a simple CUDA C++ application that consists of a small class for doing vector arithmetic:

xxxxxxxxxx
class vec4 {
  public:
    float x;
    float y;
    float z;
    float w;

    __device__ vec4();
    __device__ vec4(float x, float y, float z, float w);

    __device__ vec4 operator+(const vec4& rhs) const;
    __device__ vec4 operator*(float scalar) const;
    __device__ float dot(const vec4& rhs) const;
    __device__ float norm() const;
};

as well as one or more kernels that depend on it:

xxxxxxxxxx
// some arbitrary calculation that uses the vec4 type
__global__ void kernel(float* output, int element_count, int inner_iterations) {

    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < element_count) {
      
        const float seed = (index & 1023U) * (1.0f / 1024.0f);
        vec4 value(seed + 0.25f, seed * 0.5f + 0.5f, 0.75f, 1.0f);
        const vec4 bias(0.013f, -0.021f, 0.034f, -0.008f);

        for (int i = 0; i < inner_iterations; ++i) {
            const float projection = value.dot(bias);
            const float length = value.norm();
            const float scale = 0.25f / (length + 1.0f);
            const vec4 update = bias * (scale + projection * 0.0005f);
            value = value + update;
        }

        output[index] = value.dot(value);
    }

}

The vec4 member functions can be implemented by either:

1. (direct) defining them in the header

2. (separable compilation) defining them in a companion vec4.cu and using -rdc compiler flag

   Note: "relocatable device code" and "separable compilation" are different terms for the same thing

3. (dlto) same as 2, except also using -dlto flag

    

Runtime Performance

If we time the kernel for a problem size of 1048576 elements and 64 inner iterations, we see the following:

xxxxxxxxxx
% ./separable_compilation_test 
GPU: NVIDIA A100-PCIE-40GB (SM 8.0)
Direct header (no RDC):  0.3147 ms  (1.0000x)
Separable compilation:   0.8823 ms  (2.8038x)
Separable + device LTO:  0.3153 ms  (1.0020x)

The runtime of the direct header-based approach is practically the same as separable compilation with device link-time optimization. But, separable compilation without link time optimization is much slower.

But wait-- when separable compilation is enabled, the __device__ functions are still locally optimized. Why is the performance so much worse?

Let's investigate by profiling with NSight compute:

xxxxxxxxxx
ncu -o separable_compilation_test\
    -f\
    --set full\
    --clock-control none\
    --import-source on\
    -c 3\
    ./separable_compilation_test 

Right off the bat, we see an important discrepancy: the direct and dlto version use 22 registers, while the separable compilation version uses 37. Let's inspect the source view in Nsight Compute to track down where those registers are going.

When we look at the SASS associated with a particular line of code from the direct version, we see the instructions we expect: a dot product becomes a bunch of fused multiply-adds. However, if we look at the SASS generated from the same line of CUDA for the separable compilation version we have

No fused multiply adds here. Instead we see only IMAD, MOV and CALL.ABS.NOINC 0x0. What's going on?

CALL.ABS.NOINC indicates calling a function, and the only function associated with that line of source is the dot member function. If we look at the other vec4 member functions, we see a similar pattern: IMAD, MOV and CALL.ABS.NOINC 0x0.

What this means is that device linking from separable compilation alone (i.e. without -dlto) is not able to completely inline the device functions and optimize away the overhead associated with a function call from a separate translation unit. As a result, those function calls emit extra SASS instructions, and end up doing extra work.

However, dlto fixes this issue. In fact, if we compare the SASS code for direct and dlto, we find that they are actually identical.

Note: I'm not saying this is a guarantee that dlto always produces identical SASS.

Compilation Time

This project has an option to generate N kernels in separate sourcefiles that each use the vec4 class. I also include an intentionally long-to-compile member function in the vec4 class so that we can simulate a more expensive abstraction that appears in kernels. Let's take a look at the timings for the different compilation steps, as we vary the number of translation units that use the vec4 class:

The important trends to notice here are:

• direct wins when there is only a single kernel that uses the vec4 class. This makes sense-- putting implementations in headers only becomes problematic when that header ends up being used in many different places.

• separable compilation is the fastest to compile for all but the single-kernel case, but has the aforementioned performance penalty.

• dlto is slightly more expensive to compile than separable compilation, but much cheaper to compile than direct when it comes to multiple translation units.

Summary

For historical reasons, CUDA projects used to put __device__ implementations in headers. Separable compilation's initial release improved compile times, but it also hurt runtime performance. This runtime/compiletime tradeoff soured many developers on the concept. However, as of CUDA 11.2, the dlto option gives us a 3rd option that, in many cases, is the best of both worlds.

If your project uses separable compilation, you might also try turning dlto on and off and measuring the differences in runtime and compile time. The example __device__ functions in this document were small, so the overhead associated with the non-inlined function calls is relatively high.

The sources used to generate the results shown in this document are available here.