RidgeRun CUDA Optimisation Guide/Optimisation Recipes/Fine optimisations/Inlining

RidgeRun CUDA Optimisation Guide
GPU Architecture
Execution process Memory hierarchy Communication & Concurrency
Optimisation Workflow
Tools CUDA Memcheck CUDA Profiler Computational Budget Tool
Optimisation Recipes
Types of optimisations Coarse optimisations Workload offloading Problem size Communication overlapping Correct memory access patterns Inter-thread communication Memory management Fine optimisations Increase arithmetic intensity Function approximation Condition and loops replacement Inlining
Common pitfalls when optimising
Examples
Empirical Experiments
Simple bounding test Multi-threaded bounding test
Contact Us

GPU bound: inlining

Inlining is another optimisation technique. Take it into account when the kernel calls another function so many times. Often, the function invocations involve context switching, memory jumping, and other phenomena that cause performance degradation.

However, take into account that inlining sometimes causes performance degradation in some cases. Review this thread.

In general, for inlining, you can provide a hint to the compiler by using:

• The __noinline__ function qualifier can be used as a hint for the compiler not to inline the function if possible.
• The __forceinline__ function qualifier can be used to force the compiler to inline the function.
• The __noinline__ and __forceinline__ function qualifiers cannot be used together, and neither function qualifier can be applied to an inline function.

By default, the compiler inlines any __device__ function when deemed appropriate.