Introduction

In a previous document, we showed how to write a simple "hello world" program in CUDA. The CPU and GPU code looked almost identical: both called printf to write our "hello world" message out to the terminal.

​x
#include <cstdio>

__global__ void print_hello_gpu() {
    printf("hello world from gpu\n"); 
}

void print_hello_cpu() {
    printf("hello world from cpu\n"); 
}

However, if you try to extend this program to do something marginally more interesting like

xxxxxxxxxx
float important_calculation(float x) {
    return x * x - 3.0f;
}

__global__ void important_calculation_gpu(float x) {
    float answer = important_calculation(x);
    printf("important calculation on gpu: %f\n", answer); 
}

void important_calculation_cpu(float x) {
    float answer = important_calculation(x);
    printf("important calculation on cpu: %f\n", answer); 
}

you'll run into a compilation error to the tune of

xxxxxxxxxx
error: calling a __host__ function("important_calculation(float)") 
from a __global__ function("important_calculation_gpu") is not allowed

error: identifier "important_calculation" is undefined in device code

The CPU code is clearly well-formed, but it looks like something isn't right with the GPU code.

Function Annotations

Well, if we want to be able to invoke a function inside a kernel, it must be marked with the __device__ annotation. Since the important_calculation function has no such annotations, we're not allowed to use it in our kernel.

With that in mind, let's go back and change our code from earlier to insert a __device__ annotation. This should make our program compile, right?

xxxxxxxxxx
// now it's got a __device__ annotation
__device__ float important_calculation(float x) {
    return x * x - 3.0f;
}

__global__ void important_calculation_gpu(float x) {
    float answer = important_calculation(x); // now this line compiles without issue!
    printf("important calculation on gpu: %f\n", answer);
}

void important_calculation_cpu(float x) {
    // error: calling a __device__ function("important_calculation(float)") 
    // from a __host__ function("important_calculation_cpu") is not allowed
    float answer = important_calculation(x);
  
    printf("important calculation on cpu: %f\n", answer); 
}

Unforunately, it fixed the error message inside the kernel, but now the compiler is unhappy with the CPU code. It tells us that __host__ functions can't call __device__ functions, but what is a host function?

__host__ just means it can run on the CPU. By default, "regular" C++ functions are implicitly __host__. However, since we explicitly marked important_calculation as __device__ only, that function as written is not allowed to be used on the CPU side of things.

The fix is simple, since our intention is to use the function in both __host__ and __device__ contexts, we mark it as __host__ __device__.

xxxxxxxxxx
// now it's got both annotations
__host__ __device__ float important_calculation(float x) {
    return x * x - 3.0f;
}

__global__ void important_calculation_gpu(float x) {
    float answer = important_calculation(x); // now this line compiles without issue!
    printf("important calculation on gpu: %f\n", answer);
}

void important_calculation_cpu(float x) {
    float answer = important_calculation(x); // now this line compiles without issue!
    printf("important calculation on cpu: %f\n", answer); 
}

But Why?

The main reason is that kernel code can do things that don't make sense on the CPU, and CPU code can do things that don't make sense on the GPU. For example:

• Code meant to run in a kernel (e.g. __device__ or __global__ functions) can access variables like threadIdx, blockIdx, and blockDim anywhere in the body.

• CPU code can use abstractions and intrinsics that are not defined inside a kernel (e.g. std::thread, or SIMD intrinsics).

One could argue that the compiler could analyze the body of a function to determine whether or not it is suitable to run on CPU or GPU, although this would be challenging. Another important reason is that this aspect of the language dates back to CUDA v1.0. Back then, the hardware and compiler toolchains were much simpler than they were today, so it was likely not practical to try and automatically infer if a function was suitable for CPU/GPU execution, so that decision was left up to the developer.

Software Development Implications

Many regular C++ libraries cannot be used inside CUDA kernels. This includes most of the C++ standard library, Abseil, Boost, etc. Their source code is not appropriately annotated with __device__ and it's not just a simple matter of sprinkling some annotations here and there. Their underlying implementations make use of features not available in kernels (e.g. frequent dynamic memory allocation) so there is no easy way to port those libraries to work in a kernel.

However, all is not lost. At the time of writing, CUDA has been around for two decades and has no shortage of libraries and tools available. NVIDIA provides many useful libraries (most of which are bundled in the CUDA toolkit)

• {CCCL, Thrust, CUB}: Essentially CUDA's version of the C++ Standard library and a lot of other useful tools including containers, algorithms, atomics, etc.

• {cuBLAS, cuFFT, cuSparse, cuDSS, cuOpt}: Optimized libraries for common mathematical operations

And there is also an extensive community of open source libraries.

So, there is no shortage of available libraries for CUDA, but projects do require some refactoring to adopt them.

Anyway, now that we have all that in mind, let's look at some examples of how we can annotate the different kinds of functions to make them usable from inside a kernel.

Examples

Free Functions

xxxxxxxxxx
__device__ double free_function(double x) {
    return x * x;
}

Member Functions / Functors

xxxxxxxxxx
struct stateless_function_object {
    __device__ double operator()(double x) {
        return x * x;
    }
};

xxxxxxxxxx
struct stateful_function_object {
    __device__ double operator()(double x) {
        return scale * x;
    }

    double scale;
};

Lambda Functions

xxxxxxxxxx
__global__ void my_kernel(double * out, double * in) {
    // if the lambda is created and used in the same kernel, 
    // it doesn't need a __device__ annotation
    auto f = [](double x){ return x * x; };
    *out = f(*in);
}

xxxxxxxxxx
template < typename callable >
__global__ void my_kernel(double * out, double * in, callable f) {
    *out = f(*in);
}

int main() {

    double * d_in = ... ;
    double * d_out = ... ;
  
    // if the lambda is created on the host, but is used 
    // inside a kernel, it needs a __device__ annotation
    // and the compilation flag "--extended-lambda"
    auto f = [] __device__ (double x){ return x * x; };
    my_kernel<<<1,1>>>(d_out, d_in, f);

}

Virtual Functions

Virtual functions are problematic inside kernels.

Take a look the following code, and guess what happens when we run the program.

xxxxxxxxxx
struct base_class {
    __device__ virtual double operator()(double x) = 0;
};

struct derived_class : public base_class {
    __device__ double operator() (double x) final {
        return x * x;
    };
};

__global__ void my_kernel(double * out, double * in, base_class * f) {
    *out = (*f)(*in);
}

...

double * d_out = ...;
double * d_in = ...;
derived_class f;
my_kernel<<< 1, 1 >>>(d_out, d_in, &f);

When we check for error codes after launching my_kernel, we get

CUDA error: an illegal memory access was encountered

so something is not right, but what is going on?

First of all, f lives in CPU memory, so trying to dereference a host pointer from a kernel leads to an invalid memory access. Okay, so let's actually make a copy of the object in device memory and pass that to the kernel:

xxxxxxxxxx
derived_class * d_f;
cudaMalloc(&d_f, sizeof(derived_class));
cudaMemcpy(&d_f, &f, sizeof(derived_class), cudaMemcpyDeviceToHost);
my_kernel<<< 1, 1 >>>(d_out, d_in, d_f);

Nope, that didn't fix it, we still get an error when invoking my_kernel. This time the error code is:

CUDA error: misaligned address

Debugging with compute-sanitizer (more info on this tool later) reveals more info about what is going wrong, saying that this line of the kernel

xxxxxxxxxx
    *out = (*f)(*in);

has an Invalid __global__ read of size 8 bytes. However, all 3 variables (out, in and f) are in device memory and accessed properly, so what is the problem?

It turns out the vtable is the issue. We made sure that the object itself and the operands are available in GPU memory, but that's not enough. Dynamic dispatch uses another pointer to track the virtual functions of a base class, but that vtable still lives on the host, not the device. So, when the program tried to look up the appropriate function pointer for the derived class' virtual function, it triggers an invalid memory access!

In reality, virtual functions inside kernels are both dangerous and detrimental to performance. As a result, the CUDA programming guide makes the following, unambiguous statement:

"It is not allowed to pass as an argument to a __global__ function an object of a class derived from virtual base classes."

So: think twice before porting virtual function-based interfaces to the GPU!

Constexpr Functions

constexpr functions already satisfy many of the same requirements as __device__ functions:

• they can only call other constexpr functions

• (pre C++20)

  • they couldn't dynamically allocate memory

  • they can't call virtual virtual functions

As a result, nvcc has a flag to make constexpr functions implicitly work inside kernels:

xxxxxxxxxx
// requires --expt-relaxed-constexpr flag
constexpr double constexpr_function(double x) {
    return x * x;
}

__global__ void my_kernel(double * out, double * in) {
    *out = constexpr_function(*in);
}

Summary

• For code to run on a GPU, it must be in a __global__-annotated function (i.e. a kernel)

• Not just any code can be executed inside a kernel

  • For example: "regular" C++ functions are not callable

  • Functions must be annotated with __device__ to be callable

    • As a result, __device__ functions can't call __host__ functions!

  • Many "regular" C++ libraries will not work out-of-the-box inside a kernel!

    • However, there are many libraries that do work inside kernels.

• Polymorphism:

  • Many C++ projects rely heavily on virtual functions

    • dynamic dispatch presents a number of problems on GPUs:

      • What are the resource requirements (registers, shared memory) of the virtual function?

      • Different execution paths for threads within a warp (thread divergence)

      • Where does the vtable live in memory?

    • Inside a kernel, this pattern often performs poorly or plainly doesn't work

    • I strongly recommended avoiding using virtual functions inside a kernel

  • Static Polymorphism is more prevalent in CUDA kernels

    • Compile time information about resource usage

    • Supports inlining and other optimizations

    • no vTable to worry about