Separating Only CUDA Kernels into `.cu` Files

In CUDA programs, you can place only the GPU kernels in .cu

files while keeping the rest of your code in .cpp

files.

This organization confines CUDA-specific code to a limited set of files, making the overall program easier to maintain and understand.

This example uses the following two files:

kernel.cu
main.cpp

kernel.cu

: Defines the CUDA kernel.main.cpp

: Launches the kernel and handles error checking.In kernel.cu

, define only the kernel that runs on the GPU.

__global__ void noopKernel()
{
    // A kernel that does nothing
}

A function marked with __global__

is launched from the CPU and executed on the GPU.

The noopKernel

in this example performs no work. It is the smallest possible example for verifying that a kernel can be separated into its own file and launched successfully.

In main.cpp

, declare the kernel and launch it with cudaLaunchKernel

.

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

void noopKernel();

static void checkCuda(cudaError_t error, const char* operation)
{
    if (error != cudaSuccess) {
        std::fprintf(
            stderr,
            "%s failed: %s\n",
            operation,
            cudaGetErrorString(error)
        );
        std::exit(EXIT_FAILURE);
    }
}

int main()
{
    void** kernelArgs = nullptr;

    checkCuda(
        cudaLaunchKernel(
            reinterpret_cast<const void*>(&noopKernel),
            dim3(1, 1, 1),
            dim3(1, 1, 1),
            kernelArgs,
            0,
            nullptr
        ),
        "cudaLaunchKernel"
    );

    checkCuda(cudaGetLastError(), "kernel launch");

    checkCuda(
        cudaDeviceSynchronize(),
        "cudaDeviceSynchronize"
    );

    std::puts("noopKernel completed successfully.");
    return EXIT_SUCCESS;
}
void noopKernel();

Although the implementation of noopKernel

is in kernel.cu

, you still need to declare it so that main.cpp

can reference it.

The function name and parameter list in the declaration must match the definition in kernel.cu

.

cudaLaunchKernel

cudaLaunchKernel(
    reinterpret_cast<const void*>(&noopKernel),
    dim3(1, 1, 1),
    dim3(1, 1, 1),
    kernelArgs,
    0,
    nullptr
)

The main arguments are:

noopKernel

: The kernel to launch.dim3

: Grid dimensions.dim3

: Block dimensions.kernelArgs

: Arguments passed to the kernel.0

: Size of dynamically allocated shared memory.nullptr

: Use the default stream.Since this kernel takes no arguments, kernelArgs

is set to nullptr

.

The checkCuda

function is a helper for checking the return values of CUDA API calls.

static void checkCuda(cudaError_t error, const char* operation)
{
    if (error != cudaSuccess) {
        std::fprintf(
            stderr,
            "%s failed: %s\n",
            operation,
            cudaGetErrorString(error)
        );
        std::exit(EXIT_FAILURE);
    }
}

If an error occurs, it prints the operation name along with the CUDA error message, then terminates the program.

After launching the kernel, errors are checked in two stages:

checkCuda(cudaGetLastError(), "kernel launch");
checkCuda(cudaDeviceSynchronize(), "cudaDeviceSynchronize");

cudaGetLastError

primarily detects errors related to the kernel launch configuration.

cudaDeviceSynchronize

waits until the GPU has finished executing and reports any errors that occurred during kernel execution.

First, compile each source file into an object file.

nvcc -c kernel.cu
g++ -c main.cpp

This produces:

kernel.o
main.o

Finally, link the object files using nvcc

.

nvcc kernel.o main.o -o app

Run the program with:

./app

If execution succeeds, the following message is displayed:

noopKernel completed successfully.

Separating kernels into .cu

files provides several advantages:

g++

or other standard C++ compilers.This organization is well suited for projects where you want to limit the use of .cu

files to only the necessary parts and keep CUDA kernels clearly separated from the rest of the C++ code.

source & further reading

dev.to — original article We built an AI-powered scout for open source. Meet OpenScOut. Automatic Error Recovery in AI Agent Networks AI Agent Answers the Phone and Hands You to a Real Person