5.5. Improving CUDA Compilation Time#

5.5.1. Introduction#

This section describes techniques for measuring and diagnosing compile time bottlenecks, compilation flags, and source code refactoring strategies for reducing compile time of CUDA code.

Here’s a recommended workflow for tackling compile time issues:

Recommended workflow for measuring and reducing compile time

There are two ways to measure compile time: compiler-centric and source-centric. In the compiler-centric view, the nvcc/NVRTC “-time” flag reports the breakdown of the compile time across the various compilation phases (e.g. parser, optimizer). In the source-centric view, the user can generate a compile time flame chart that breaks down the compilation time by source elements e.g. header files and template instantiations.

Alternately, the CTAdvisor tool generates a summary of compile time segregated by source elements and suggestions for applicable compiler flags to reduce compile time. Improving compile time is typically an iterative process, successively attacking the next compile time bottleneck until the compile time is satisfactory.

The rest of the document describes how to measure compile time, various applicable compiler flags and source code modifications to reduce compile time.

5.5.2. Measurement#

5.5.2.1. Using the `-time` Flag#

nvcc and NVRTC support the ‘-time <filename>’ flag to measure the compile time for each phase of compilation. The output is appended to the end of the specified file.

Example:

Here’s a sample output generated by NVRTC with the -time flag:

File name, phase name, metric, unit
j4.cu.ptx, CUDA C++ Frontend, 30, ms
j4.cu.ptx, NVVM Optimizer, 1000, ms
j4.cu.ptx, PTX Translation, 14, ms

The per-phase timing breakdown can indicate the potential approaches to improve compile time. In the above case, the NVVM Optimizer phase takes most of the time, so compile time could potentially be improved by tweaking compile time flags that are known to speed up NVVM compilation (e.g –split-compile).

5.5.2.2. Using Compilation Flame Chart (`-fdevice-time-trace`)#

Changing the source code instead of adding compiler flags may lead to significant compile time improvement for many programs. To understand what source code changes are needed, we first need information on source code elements that are incurring the most compilation time. A flame chart provides a graphical view of the compilation time tagged to individual source elements (e.g. templates). Invoke nvcc / NVRTC to generate a device compilation flame chart using the -fdevice-time-trace <filename> flag. This produces a .json file that follows the “Trace Event” format, a widely recognized format used for profiling. The .json trace file can be opened and using the Perfetto trace viewer (https://ui.perfetto.dev/).

The flame chart enables us to visualize the end-to-end workflow of the compilation process, providing a comprehensive timeline of all major stages:

Preprocessing
Host and device code compilation
Device linking
Binary generation

This holistic view is particularly valuable for understanding how different phases interact and contribute to the total compile time. By examining this visualization, you can identify stages that dominate the workflow, determine whether delays stem from specific code constructs or the compiler pipeline itself, and prioritize areas for optimization.

Example:

Compilation flame chart for multiple target architectures

The flame chart shows compilation for multiple architectures (compute_90, compute_100), with significant time being spent on the frontend. The compilation for multiple architectures could be overlapped by allowing nvcc to use multiple threads to compile (-threads N)

The flame chart can also be used to find header files and templates that incur significant compile time.

Example:

Flame chart showing recursive template instantiation

The ‘recursiveTree’ class has a deep nested template instantiation. Restructuring the code to avoid excessive instantiations will improve the overall compile time in this case.

Please refer to this blog post for more details on using compile time flame charts.

5.5.2.3. Using Compile Time Advisor#

Compile Time Advisor is a command line tool that is part of the CUDA Toolkit. It processes .json file generated by the nvcc/NVRTC –fdevice-time-trace flag, and summarizes the compile time diagnostics and advice in a generated text file.

Currently, Compile Time Advisor provides five different types of advice:

Advice	Description
Expensive Templates	Identifies template functions/classes that took the longest to instantiate and number of instantiations.
Expensive Headers	Identifies header files that took the longest to include and breaks down the time spent on processing each sub-header file.
Parallelizable Compilation	Checks if source files were compiled for multiple architectures (multiple -gencode targets) and the –threads flag was unused (using –threads will overlap per-architecture compilation and reduce overall compilation time).
Split Compilation	Identifies source files with high optimization time that could benefit from the –split-compile flag to run optimizations in parallel.
Compile Time Breakdown	Breaks down the compile time of each translation unit into various compilation steps, lists source files that took the largest compilation time, and helps identify which stages of compilation are a bottleneck.

For example, here’s a sample output of Compile Time Advisor, truncated to only the “Expensive Templates” and “Expensive Headers” advice:

//-------------------------------------------------------------
Trace file loading is complete
***********************************************
Start expensive template advisor
Top template functions/classes by instantiation time:
  [0.00s]: (2 times in 2 files) std::__detail::__hyperg
    [0.00s]: (1 times in 1 files) std::__detail::__hyperg<float>(float, float, float, float)
      [0.00s]: (1x) saxpy.cu, recursively instantiates 12 templates with max depth 5
…
***********************************************
Start expensive headers advisor
Top header files by include processing time:
 [0.05s]: (1 times in 1 files) /usr/local/cuda-13.0/bin/../targets/x86_64-linux/include/cuda_runtime.h
    Processing this header itself takes 0.00s (1.65%)
    includes /usr/local/cuda-13.0/bin/../targets/x86_64-linux/include/crt/common_functions.h (1x) taking 0.03s (58.73%)
…
//-------------------------------------------------------------

Please see Compile Time Advisor documentation for further details on using the tool.

5.5.3. Compiler Flags to Improve Compile Time#

Using the approaches described in the previous section, we can identify the compiler phases and program characteristics impacting compile time. We can then experiment with flags described in this section to reduce compile time.

The following table summarizes the available flags:

Flag	Description	Tool	Typical Usecase	Perf impact
–threads <N>	Number of threads for compilation.	nvcc	Multi-arch compilation with –gencode.	None
–split-compile <N>	Number of threads used for concurrent optimization.	nvcc, nvrtc, nvlink, nvJitLink, ptxas	Translation unit with multiple kernels or non-inlined device functions.	Minimal
–split-compile-extended <N>	Number of threads used for concurrent optimization. More aggressive form of spit-compile. LTO-only.	nvcc, nvrtc, nvlink, nvJitLink	LTO (link time optimization) processing, where non-trivial run time slow down is acceptable.	Program-dependent
–Ofast-compile <min/mid/max>	Device optimization levels	nvcc, nvrtc, nvlink, nvJitLink, ptxas	Tradeoff modest to significant increase in run time for significant savings in compile time.	Program-dependent
–gen-opt-lto	Run the NVVM optimizer before saving LTO-IR	nvcc, nvrtc	LTO-IR fragments are mostly self-contained; move optimization cost to build time instead of run time.	Program-dependent
–pch –use-pch	Precompiled Header support. Automatic (-pch) and explicit (-use-pch).	nvrtc	Multiple programs being compiled that use a common header prefix, and frontend (parsing) time is a bottleneck. Note: forces frontend serialization, so may affect scaling with concurrent NVRTC invocations.	None
–minimal	Disable certain rarely used language features	nvrtc	Frontend time is the bottleneck. Fixed cost savings on the order of tens of milliseconds on modern machines.	None
–opt-level <N>	Optimization level for ptxas	ptxas	Tradeoff significant run time increase for faster compile time.	Program-dependent
–allow-expensive-optimizations=false	Disallow aggressive optimizations	ptxas	Tradeoff potentially significant run time increase for faster compile time.	Program-dependent
–register-usage-level <N>	Control aggressiveness of optimizations affecting register usage	ptxas	Tradeoff potentially significant run time increase for faster compile time.	Program-dependent
–jobserver	Enable GNU jobserver support	nvcc, ptxas	Avoid oversubscribing the number of CPU cores	None

A detailed description of the options follows, grouped by category.

5.5.3.1. Concurrent Compilation Options#

5.5.3.1.1. Threads#

--threads <N> (-t)

Description

Specify the maximum number of threads to be created in parallel when compiling for multiple architectures. If <N> is 1 or if compiling for one architecture, this option is ignored. If <N> is 0, the number of threads will be the number of CPUs on the machine.

Supported By

nvcc

When to Use

The CUDA file is being compiled for multiple target architectures using the “-gencode” flag. In this case, the multiple target compilations can be overlapped using different threads, reducing overall compile time.

When using GNU make or Ninja build systems, can pair with --jobserver flag, to avoid oversubscribing the number of CPU cores on the system.

Run Time Impact

None – this option does not affect the quality of the generated code.

5.5.3.1.2. Split Compile#

--split-compile <N> (-split-compile)

Description

Specify the maximum amount of concurrent threads to be utilized when running compiler optimizations. If <N> is 1, this option is ignored. If <N> is 0, the number of threads will be the number of CPUs on the machine.

Supported By

nvcc, nvrtc, nvlink, nvJitLink, ptxas

When to Use

The NVVM or PTXAS compilation steps are compile time bottlenecks.
Code contains multiple kernel or non-inlined device functions.

The compiler optimizer will run per-function optimizations concurrently using the specified number of threads. When using GNU make or Ninja build systems, can pair with --jobserver flag, to avoid oversubscribing the number of CPU cores on the system.

Run Time Impact

Minimal – In most cases, this should not affect the quality of the generated device machine code.

5.5.3.1.3. Extended Split Compile#

--extended-split-compile <N> (-split-compile)

Description

Specify the maximum amount of concurrent threads to be utilized when running compiler optimizations in LTO mode. If <N> is 1, this option is ignored. If <N> is 0, the number of threads will be the number of CPUs on the machine.

This option is a more aggressive form of split compilation, and can potentially impact performance of the compiled binary. It is only available in LTO mode.

Supported By

nvcc, nvrtc, nvlink, nvJitLink

When to Use

LTO optimization is turned on.
NVVM processing is a bottleneck.
Potential increase in program run time is tolerable.

When using GNU make or Ninja build systems, can pair with --jobserver flag, to avoid oversubscribing the number of CPU cores on the system.

Run Time Impact

Program-dependent – for kernels that contain mostly inlined code, impact may be minimal.

5.5.3.1.4. Jobserver#

--jobserver (-jobserver)

Description

Enable GNU Make Jobserver support. GNU jobserver protocol is supported by GNU Make and Ninja build systems. The option prevents CPU oversubscription when using concurrent compilation options (–threads/–split-compile/–extended-split-compile) and concurrent build commands (e.g. GNU make’s “-j” flag) at the same time.

Supported By

nvcc, ptxas

When to Use

Build uses GNU make or ninja build system and concurrent builds (e.g. GNU make’s “-j” flag).
Concurrent compilation options are active(–threads/–split-compile/–extended-split-compile).

Run Time Impact

None – this option does not affect the quality of the generated code.

5.5.3.2. NVRTC-specific Options#

5.5.3.2.1. Precompiled headers#

--pch (-pch)

--use-pch <filename> (-use-pch)

Description

Enable automatic precompiled header (PCH) mode(-pch)or explicit PCH mode (--use-pch). Please see NVRTC documentation for details and example code.

Supported By

nvrtc

When to Use

There are multiple NVRTC programs to be compiled, and they use a common prefix of headers. E.g.:

#include “foo.h”
#include “bar.h”

Frontend compilation time is significant.

Note: Using precompiled headers serializes frontend processing, negatively impacting scaling if multiple concurrent NVRTC invocations from different host threads are in progress. This impact becomes more significant as the number of concurrent NVRTC invocations increases; in this scenario, experiment with turning off PCH to remove scaling bottleneck.

Run Time Impact

None – this option does not affect the quality of the generated code.

5.5.3.2.2. Minimal#

--minimal (-minimal)

Description

Omit certain language features to reduce compile time for small programs. In particular, the following are omitted:

Texture and surface functions and associated types, e.g., cudaTextureObject_t.
CUDA Runtime Functions that are provided by the cudadevrt device code library, typically named with prefix “cuda”, e.g., cudaMalloc.
Kernel launch from device code.
Types and macros associated with CUDA Runtime and Driver APIs, provided by cuda/tools/cudart/driver_types.h, typically named with prefix “cuda”, e.g., cudaError_t.

Supported By

nvrtc

When to Use

Small source programs, where frontend processing time is significant.
Reduces some fixed costs in the compiler, typically saving on the order of tens of microseconds.

Run Time Impact

None – this option does not affect the quality of the generated code.

5.5.3.3. PTXAS-specific Options#

--opt-level <N> (-O)

--allow-expensive-optimizations=false (-allow-expensive-optimizations)

--register-usage-level <0…10> (-regUsageLevel)

Description

These options control the PTXAS optimizer.

--opt-level sets the overall optimization level (default: 3).
--allow-expensive-optimizations=false turns off compile time expensive optimizations (default: automatically turned on with optimization level >= 2).
--register-usage-level controls the aggressiveness of the optimizations that target register usage. Higher values aggressively optimize the source program, trading off additional register usage for potential improvements in the generated code. (Default: 5).

Supported By

ptxas

When to Use

PTX compilation is a compile time bottleneck.
Measured impact on runtime is within acceptable limit.

Run Time Impact

Program-dependent – experiment with different flag values and measure compile time speedup versus run time increases to find an acceptable tradeoff.

5.5.3.4. Miscellaneous Options#

5.5.3.4.1. Fast Compile#

--Ofast-compile <level> (-Ofc)

Description

Specify fast-compile level for device code, which controls the tradeoff between compilation speed and runtime performance by disabling certain optimizations at varying levels.

Levels:

‘max’: Focus only on the fastest compilation speed, disabling many optimizations.
‘mid’: Balance compile time and runtime, disabling expensive optimizations.
‘min’: More minimal impact on both compile time and runtime, minimizing some expensive optimizations.
‘0’: Disables fast-compile.

The option is disabled by default.

Supported By

nvcc, nvrtc, nvlink, nvJitLink, ptxas

When to Use

Compile time is more important than peak run time (e.g. a fast initial compilation to reduce user facing latency, with a slower background compilation with full optimization).

Run Time Impact

Program-dependent – experiment with different levels and measure compile time speedup versus run time increases to find an acceptable tradeoff.

5.5.3.4.2. Optimized LTO-IR#

--gen-opt-lto (-gen-opt-lto)

Description

By default, nvcc and NVRTC will generate unoptimized LTO-IR code, and the NVVM optimizer will run only after the LTO-IR fragments are linked together. However, for workflows that generate machine code online using the nvJitLink library to link LTO-IR fragments, the build time overhead of running the NVVM optimizer may be significant.

As an alternative, use this flag to direct nvcc and NVRTC to run the NVVM optimizer passes on the LTO-IR before the LTO-IR is emitted. This allows moving some optimizations from the online phase to the LTO-IR generation phase, with the aim of reducing the overall online processing time.

Note: This flag causes the optimizer to run twice (once during LTO-IR generation, and again after the LTO-IR fragments are linked), so the generated machine code may be different than with the default LTO-IR processing mode, where all optimization are deferred until the LTO-IR fragments are linked.

Supported By

nvcc, nvrtc

When to Use

Workflow uses online LTO-IR linking with nvJitLink library.
LTO-IR fragments being linked are more-or-less self-contained, and have only coarse-grained interaction e.g. , different stages of a FFT pipeline. This allows the offline optimizer to optimize the code more effectively, reducing the work needed to be done by the NVVM optimizer invocation that occurs online after the LTO-IR fragments are linked.
On the other hand, if the LTO fragments have fine-grained interaction, then the offline optimizer that operates on each LTO-IR fragment in isolation may not be able to effectively optimize the code, leaving most of the optimization work to be still done by the online optimizer invocation.
As noted earlier, running the optimizer twice may change the generated machine code compared to the default LTO workflow, so ensure that performance of the generated code has not degraded unexpectedly.

Run Time Impact

Program-dependent – experiment to see if run time impact is acceptable.

5.5.4. User Code and Workflow Changes#

In this section, we discuss changes to user code and compiler invocation to speed up compile time.

5.5.4.1. Remove Unused Headers#

In large software projects, there is typically a common “prefix” of headers that is included by every translation unit. For example, some machine learning implementations generate NVRTC code that unconditionally includes the CUDA C++ header files that define the __half and __bfloat16 types, even if the types are not actually used in the generated NVRTC program:

#include <cuda_fp16.h>
#include <cuda_bf16.h>
/* other includes */

Unfortunately, these headers contain a large number of inline __device__ functions, which increases parsing time.

Action:

Look at the per-phase compilation time break up reported by the -time flag. Alternately, generate the compile time trace using -fdevice-time-trace and run the Compile Time Advisor, and check the report for “Expensive Headers Advisor” to find expensive headers, and “Compile Time Breakdown Advisor” to check if frontend processing time is significant.

Remove unused headers and boilerplate code e.g. type definitions, static/inline helper functions.

5.5.4.2. Avoid Excessive `forceinline` Functions#

In general, the CUDA compiler inlines aggressively because it enables important optimizations such as promotion of local variables to registers. In addition, marking a function with __forceinline__ directs the compiler to ignore the inliner heuristics and try very hard to inline the target function.

However, indiscriminate inlining can cause a code explosion, causing the NVVM and PTX translation phases to slow down.

Action:

Avoid __forceinline__ annotation where possible, and let the compiler heuristics decide whether to inline the function or not.
Apply ___noinline__ annotation to functions that are not expected to be invoked frequently.

Example:

__device__ int  large_func(int in) { ... } // a large function
__device__ void parent() {
  #pragma unroll
  for (int i = 0; i < 100; ++i) {
    if (unlikely_condition) {
          result += large_func(i); // <- Code explosion due
                                    // to inlining “large_func”
        }
      }
    }

In the example, the call to large_func() seldom occurs, however the enclosing loop is marked with #pragma unroll , so there is a code explosion when the loop is unrolled and the large_func is inlined.
Fix the compile time problem by marking large_func as __noinline__.

5.5.4.3. Avoid excessive loop unrolling#

In general, fully unrolling loops with known bounds will generate more optimized code. However, unrolling aggressively increases the amount of code that needs to be processed by the NVVM and PTXAS optimizers, causing compile time slowdowns.

Action:

Measure runtime impact of disabling loop unrolling with “#pragma unroll 1” on outermost loops. In many cases, the boost in compile time might be an acceptable tradeoff for the measured runtime degradation.
In this context, look for outermost loops with array index expressions that are input or thread-id dependent; such loops may be good candidates to disable unrolling (since these arrays cannot be promoted to registers even with the loop fully unrolled).

Example:

int arr[32]{};

#pragma unroll    // problematic unroll for compile time
for(int i = 0; i < 32; ++i) {
  arr[i%threadIdx.x] = large_func(i); // runtime dependent
                                      // index
}
...
#pragma unroll
for (int i = 0; i < 32; ++i) {
  result[i] = arr[i];
}

In this example, unrolling the first loop results in a code explosion because large_func() is inlined multiple times in the unrolled loop body. However, since the arr index expression is thread-dependent (i%threadIdx.x), unrolling the loop body does not allow the elements of arr to be promoted to registers; the array stays in local memory.
Therefore, disable the loop unrolling with #pragma unroll 1 will boost compile time without likely degrading run time appreciably.

5.5.4.4. Use Concurrent NVRTC and nvJitLink Invocation#

In some contexts, user code may have a large number of CUDA C++ programs that need to be compiled with NVRTC and nvJitLink e.g., when processing a graph corresponding to Deep Learning model, where each node corresponds to a separate CUDA C++ kernel.

Action:

Use multiple host threads to invoke NVRTC (or nvJitLink) concurrently for different programs, to reduce overall compilation time.

5.5.4.5. Be Careful When Using `llvm::SplitModule()` For Partitioning#

In order to achieve good throughput, some clients may choose to split LLVM code into separate LLVM modules, and then process each module concurrently with separate threads to generate PTX and GPU machine code (e.g. openxla/xla).
If you encounter poor speedup with concurrent compilation, check that the partitioned modules are evenly partitioned, i.e. there should not be a few modules that are significantly larger than other modules, causing the corresponding compilation threads to take a lot longer to finish. For example, this may occur if using llvm::SplitModule() for partitioning, and there are calls to a common device function from multiple kernel functions (in this case, both kernels and the device function will end up in the same partition).

Action:

One potential solution is to inline aggressively before invoking llvm::SplitModule() - this ends up creating functions (kernels) that do not call other functions, allowing llvm::SplitModule() to split the functions into separate modules that can be more efficiently processed concurrently.

Example:

__device__  static int get() { /* some code */ }
 __global__ void kernel1(int *ptr) { *ptr = get(); }
 __global__ void kernel2(int *ptr) { *ptr = get()+1; }

llvm::SplitModule() will likely put all the functions in the same module, and no concurrent compilation will be possible. After inlining, before invoking llvm::SplitModule(), get() is inlined into kernel1() and kernel2(), and then kernel1() and kernel2() can be in different partitions that can be concurrently compiled.

5.5. Improving CUDA Compilation Time#

5.5.1. Introduction#

5.5.2. Measurement#

5.5.2.1. Using the -time Flag#

5.5.2.2. Using Compilation Flame Chart (-fdevice-time-trace)#

5.5.2.3. Using Compile Time Advisor#

5.5.3. Compiler Flags to Improve Compile Time#

5.5.3.1. Concurrent Compilation Options#

5.5.3.1.1. Threads#

5.5.3.1.2. Split Compile#

5.5.3.1.3. Extended Split Compile#

5.5.3.1.4. Jobserver#

5.5.3.2. NVRTC-specific Options#

5.5.3.2.1. Precompiled headers#

5.5.3.2.2. Minimal#

5.5.3.3. PTXAS-specific Options#

5.5.3.4. Miscellaneous Options#

5.5.3.4.1. Fast Compile#

5.5.3.4.2. Optimized LTO-IR#

5.5.4. User Code and Workflow Changes#

5.5.4.1. Remove Unused Headers#

5.5.4.2. Avoid Excessive __forceinline__ Functions#

5.5.4.3. Avoid excessive loop unrolling#

5.5.4.4. Use Concurrent NVRTC and nvJitLink Invocation#

5.5.4.5. Be Careful When Using llvm::SplitModule() For Partitioning#

5.5.2.1. Using the `-time` Flag#

5.5.2.2. Using Compilation Flame Chart (`-fdevice-time-trace`)#

5.5.4.2. Avoid Excessive `forceinline` Functions#

5.5.4.5. Be Careful When Using `llvm::SplitModule()` For Partitioning#