Building reusable CUDA libraries with CMake and Python bindings

Prerequisites

This article assumes you have read Device functions, host functions, and CUDA function qualifiers, including the __global__/__device__/__host__ qualifier rules and how host code launches kernels. Everything so far in this series has lived in a single .cu file compiled directly with nvcc. This article is about the next step: packaging kernels and their host wrappers into a library that other CUDA projects, C++ programs, and Python code can consume without ever seeing your kernel source.

Designing the API surface, not just the kernel

The hardest part of a reusable CUDA library is rarely the kernel itself; it is the boundary between your library and its caller. Four questions need explicit answers before you write a header:

• Who owns device memory? Does the library allocate and return device pointers the caller must free, or does the caller always allocate and merely pass pointers in? Mixing these conventions within one API is the single most common source of leaks and double-frees in CUDA libraries.
• Which stream does an operation run on? A library function that always uses the default stream, or worse, calls cudaDeviceSynchronize() internally, silently breaks any caller trying to overlap your library’s work with its own kernels or transfers. Every asynchronous entry point should accept a cudaStream_t parameter explicitly.
• How are dimensions and shapes communicated? Raw pointers plus separate int size parameters are simplest and match how the rest of this series has looked, but you must decide (and document) what happens for zero-length inputs, mismatched dimensions between arguments, or sizes that overflow a 32-bit int.
• How are errors reported? CUDA C++ has no exceptions that cross a C ABI boundary safely. A library’s public functions should return an explicit status code (either cudaError_t directly or a library-specific error enum) rather than throwing, aborting, or printing to stderr and continuing.

None of these are CUDA-specific software engineering concerns in the abstract, but CUDA raises the stakes: an ownership mistake corrupts device memory instead of host memory, a hidden synchronization silently serializes a caller’s carefully overlapped stream pipeline, and a swallowed error can mean a kernel silently never ran at all.

Anatomy of the library: header, implementation, kernels

A minimal but complete reusable CUDA library separates three concerns into three kinds of files:

graph TD
  H["vecops.h
public C API
extern C, no CUDA types leak"] --> IMPL[".cu implementation
host wrapper functions"]
  IMPL --> K["internal __global__ kernels
not exported, no external linkage"]
  IMPL --> ERR["error translation
cudaError_t -> library status"]
  CONSUMER["Consumer C++ program"] -->|"#include vecops.h
link against libvecops"| H
  style H fill:#4a90d9,color:#fff
  style IMPL fill:#e8744f,color:#fff
  style K fill:#9b59b6,color:#fff

The public header never mentions __global__, <<<...>>>, or any CUDA-specific type beyond cudaStream_t (which is just an opaque pointer type, safe to expose in a C header). Kernels live entirely inside the .cu implementation file and are never declared in the header at all, which is what keeps them out of the library’s exported symbol table.

The public header (vecops.h)

#ifndef VECOPS_H
#define VECOPS_H

#include <cuda_runtime.h>   /* only for cudaStream_t; safe in a C header */
#include <stddef.h>

#ifdef __cplusplus
extern "C" {
#endif

/* Visibility macro: exported on the build, hidden by default elsewhere.
 * VECOPS_BUILDING is defined only when compiling the library itself. */
#if defined(_WIN32)
#  ifdef VECOPS_BUILDING
#    define VECOPS_API __declspec(dllexport)
#  else
#    define VECOPS_API __declspec(dllimport)
#  endif
#else
#  ifdef VECOPS_BUILDING
#    define VECOPS_API __attribute__((visibility("default")))
#  else
#    define VECOPS_API
#  endif
#endif

typedef enum {
    VECOPS_SUCCESS = 0,
    VECOPS_ERROR_INVALID_ARGUMENT = 1,
    VECOPS_ERROR_INVALID_DEVICE = 2,
    VECOPS_ERROR_CUDA_FAILURE = 3
} vecops_status_t;

/* y[i] = alpha * x[i] + y[i], the classic axpy pattern.
 * Ownership: caller allocates and frees d_x and d_y (device pointers).
 * Stream: operation is enqueued on `stream`; caller is responsible for
 *         synchronizing before reading d_y on the host.
 * Returns VECOPS_SUCCESS or a specific error code; never throws, never
 * calls cudaDeviceSynchronize internally. */
VECOPS_API vecops_status_t vecops_axpy(
    float alpha,
    const float* d_x,
    float* d_y,
    size_t n,
    cudaStream_t stream);

/* Human-readable description of the last error code, analogous to
 * cudaGetErrorString. Never returns NULL. */
VECOPS_API const char* vecops_status_string(vecops_status_t status);

#ifdef __cplusplus
}
#endif

#endif /* VECOPS_H */

The implementation (vecops.cu)

#define VECOPS_BUILDING
#include "vecops.h"
#include <cstdio>

namespace {  // internal linkage: this kernel is not part of the library's ABI

__global__ void axpy_kernel(float alpha, const float* x, float* y, size_t n) {
    size_t i = static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = alpha * x[i] + y[i];
    }
}

}  // namespace

extern "C" vecops_status_t vecops_axpy(
    float alpha, const float* d_x, float* d_y, size_t n, cudaStream_t stream)
{
    // Validate inputs before touching the device. Null pointers and a
    // zero-length request that isn't explicitly meaningful are the two
    // most common misuse cases to catch here.
    if (d_x == nullptr || d_y == nullptr) {
        return VECOPS_ERROR_INVALID_ARGUMENT;
    }
    if (n == 0) {
        return VECOPS_SUCCESS;  // no-op is a valid, well-defined outcome
    }

    int device_count = 0;
    cudaError_t err = cudaGetDeviceCount(&device_count);
    if (err != cudaSuccess || device_count == 0) {
        return VECOPS_ERROR_INVALID_DEVICE;
    }

    const int threads_per_block = 256;
    const size_t blocks = (n + threads_per_block - 1) / threads_per_block;

    axpy_kernel<<<static_cast<unsigned int>(blocks), threads_per_block, 0, stream>>>(
        alpha, d_x, d_y, n);

    // Catch launch-time errors (bad configuration, no kernel image for this
    // architecture) without forcing a synchronization the caller did not
    // ask for. This does NOT catch errors that occur during kernel
    // execution; those surface on a later synchronizing call the caller
    // makes, exactly as with any other asynchronous CUDA API.
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        return VECOPS_ERROR_CUDA_FAILURE;
    }

    return VECOPS_SUCCESS;
}

extern "C" const char* vecops_status_string(vecops_status_t status) {
    switch (status) {
        case VECOPS_SUCCESS:                 return "success";
        case VECOPS_ERROR_INVALID_ARGUMENT:  return "invalid argument";
        case VECOPS_ERROR_INVALID_DEVICE:    return "no CUDA device available";
        case VECOPS_ERROR_CUDA_FAILURE:      return "CUDA runtime error";
        default:                             return "unknown status";
    }
}

Notice what the wrapper deliberately does not do: it never calls cudaDeviceSynchronize(), and it never blocks waiting for the kernel to finish. cudaGetLastError() after the launch only surfaces launch-time failures (invalid configuration, no compiled binary for the running GPU’s architecture); execution-time failures (an illegal memory access inside the kernel) will only be visible the next time the caller synchronizes that stream. This is intentional: a library that hides synchronization decisions from its caller breaks exactly the kind of overlapped, multi-stream pipelines covered in streams and asynchronous execution and advanced stream patterns.

Static vs. shared libraries and the C ABI

A CUDA library can be built as either a static library (.a/.lib, linked directly into the consumer’s binary at build time) or a shared library (.so/.dll, loaded at runtime and shared across processes). The extern "C" linkage and the visibility macros above matter for both, but for different reasons:

• Static libraries avoid symbol-visibility and ABI-versioning concerns almost entirely, because the consumer recompiles/relinks against your exact object code every time. The main cost is a larger consumer binary and a full rebuild of everything whenever the library changes.
• Shared libraries need a stable, C-only ABI at the boundary, because C++ name mangling and standard-library ABI are not guaranteed stable across compiler versions. extern "C" disables name mangling for the exported functions; hiding all other symbols by default (visibility("default") applied only to the functions you intend to export, with a project-wide -fvisibility=hidden default) prevents accidental exposure of internal helpers, C++ classes, or STL types whose ABI you cannot guarantee across builds.

A library targeting both C and C++ consumers, and eventually a Python binding, should default to the C ABI shown above regardless of which build type you choose, because it is the lowest common denominator every consumer (C, C++, Python via ctypes/pybind11, Rust via FFI) can call without extra glue code.

Separate compilation, RDC, and device linking

By default, nvcc compiles each .cu file’s device code independently and does not allow device code in one translation unit to call a __device__ function defined in another. This is fine for a single-file library like the one above, but breaks down the moment your library spans multiple .cu files with device functions that call each other across file boundaries.

Relocatable Device Code (RDC) lifts this restriction. Enabling it (-rdc=true on the nvcc command line, or the CUDA_SEPARABLE_COMPILATION target property in CMake, shown below) makes nvcc emit relocatable device object code that a separate device link step (an internal nvcc -dlink invocation) resolves, exactly analogous to how the host linker resolves host-side symbols across object files.

// device_math.cu -- a __device__ function defined in one translation unit
__device__ float square_device(float x) { return x * x; }

// kernels.cu -- calls it from a different translation unit; only works
// with RDC enabled, because the compiler must defer resolving this
// cross-file device call to a separate device-link step
extern __device__ float square_device(float x);

__global__ void square_kernel(float* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) data[i] = square_device(data[i]);
}

RDC has a real cost: it disables some cross-function inlining and optimization opportunities the compiler would otherwise take when everything is visible in one translation unit, and it adds a device-link step to your build. Prefer keeping tightly coupled __device__ helpers and the kernels that call them in the same .cu file when practical, and reach for RDC only when your library’s kernel logic is genuinely spread across multiple files by design (for example, a plugin-style architecture where different .cu files implement different operator kernels sharing common device-side utilities).

Modern CMake for a CUDA library

CMake treats CUDA as a first-class language since CMake 3.18+ style native CUDA support (enable_language(CUDA) or listing CUDA in project(... LANGUAGES ...)), which is preferred over the older FindCUDA module approach.

cmake_minimum_required(VERSION 3.24)
project(vecops LANGUAGES CXX CUDA)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)

# Target architectures: list the compute capabilities you actually support,
# rather than relying on a default that may not match your deployment
# hardware. "native" (CMake 3.24+) builds only for the architecture of the
# GPU present on the build machine, useful for local development only.
set(CMAKE_CUDA_ARCHITECTURES 70 75 80 86 90)

add_library(vecops SHARED
    src/vecops.cu
)

set_target_properties(vecops PROPERTIES
    CUDA_SEPARABLE_COMPILATION ON      # enable RDC for this target
    POSITION_INDEPENDENT_CODE ON       # required for shared libraries
    CXX_VISIBILITY_PRESET hidden       # hide symbols by default
    CUDA_VISIBILITY_PRESET hidden
    PUBLIC_HEADER include/vecops.h
)

target_include_directories(vecops
    PUBLIC
        $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>
        $<INSTALL_INTERFACE:include>
)

target_compile_definitions(vecops PRIVATE VECOPS_BUILDING)

# Install the library, its public header, and export a CMake package so
# downstream projects can `find_package(vecops)` and `target_link_libraries`
# against an imported target instead of hardcoding paths.
include(GNUInstallDirs)
install(TARGETS vecops
    EXPORT vecopsTargets
    LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}
    ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}
    RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}
    PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
)
install(EXPORT vecopsTargets
    FILE vecopsTargets.cmake
    NAMESPACE vecops::
    DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/vecops
)

The three properties worth calling out specifically:

• CMAKE_CUDA_ARCHITECTURES controls which compute capabilities nvcc generates code for. Listing several (rather than one) produces a “fat binary” containing a cubin per listed architecture plus PTX for forward compatibility, at the cost of longer compile times and a larger binary.
• POSITION_INDEPENDENT_CODE ON is required for any code that ends up in a shared library; without it, some platforms will fail to link or will silently produce a non-relocatable, unsafe shared object.
• install(EXPORT ...) generates a CMake package config that lets a consumer write find_package(vecops REQUIRED) and target_link_libraries(myapp PRIVATE vecops::vecops) instead of manually specifying include paths and library files, which is what makes the library genuinely reusable across projects rather than just reusable within one repository.

A minimal C++ consumer

# consumer/CMakeLists.txt
cmake_minimum_required(VERSION 3.24)
project(vecops_consumer LANGUAGES CXX)

find_package(vecops REQUIRED)

add_executable(consumer main.cpp)
target_link_libraries(consumer PRIVATE vecops::vecops)

// consumer/main.cpp
#include <vecops.h>
#include <cuda_runtime.h>
#include <vector>
#include <cstdio>

int main() {
    const size_t n = 1 << 20;
    std::vector<float> h_x(n, 1.0f), h_y(n, 2.0f);

    float *d_x = nullptr, *d_y = nullptr;
    cudaMalloc(&d_x, n * sizeof(float));
    cudaMalloc(&d_y, n * sizeof(float));
    cudaMemcpy(d_x, h_x.data(), n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y.data(), n * sizeof(float), cudaMemcpyHostToDevice);

    vecops_status_t status = vecops_axpy(2.0f, d_x, d_y, n, /*stream=*/0);
    if (status != VECOPS_SUCCESS) {
        fprintf(stderr, "vecops_axpy failed: %s\n", vecops_status_string(status));
        return 1;
    }

    cudaStreamSynchronize(0);  // the consumer, not the library, decides when to sync
    cudaMemcpy(h_y.data(), d_y, n * sizeof(float), cudaMemcpyDeviceToHost);

    printf("y[0] = %f (expected 4.0)\n", h_y[0]);  // 2*1 + 2 = 4

    cudaFree(d_x);
    cudaFree(d_y);
    return 0;
}

Every name used here (vecops_axpy, vecops_status_t, VECOPS_SUCCESS, the vecops::vecops target) matches exactly what the header and CMake files above declare, which is worth double-checking any time you copy a multi-file example like this one: a mismatch between the header’s function signature and the .cu file’s definition is a link-time or (worse, with a mismatched signature that still links) a runtime error that is easy to introduce when editing files separately.

Python bindings

Once the C ABI exists, exposing it to Python is a separate, optional layer, not a reason to change anything about the library itself.

pybind11 is the most common approach for a from-scratch binding: write a thin C++ wrapper that accepts NumPy or CuPy device pointers and calls into your library.

// bindings.cpp -- conceptual sketch, not a full build; pybind11 itself is
// not added as a dependency here, only described.
#include <pybind11/pybind11.h>
#include <vecops.h>

namespace py = pybind11;

// Accept raw device pointers as Python integers (as CuPy's
// `.data.ptr` or a `__cuda_array_interface__` pointer would provide),
// so this binding does not depend on any specific array library.
void py_axpy(float alpha, uintptr_t d_x, uintptr_t d_y, size_t n, uintptr_t stream) {
    vecops_status_t status = vecops_axpy(
        alpha,
        reinterpret_cast<const float*>(d_x),
        reinterpret_cast<float*>(d_y),
        n,
        reinterpret_cast<cudaStream_t>(stream));
    if (status != VECOPS_SUCCESS) {
        throw std::runtime_error(vecops_status_string(status));
    }
}

PYBIND11_MODULE(vecops_py, m) {
    m.def("axpy", &py_axpy, "y = alpha * x + y on the GPU");
}

This is where the C API’s explicit stream parameter and status-code error contract pay off directly: the pybind11 wrapper can convert a CuPy or PyTorch stream handle straight into a cudaStream_t and convert a vecops_status_t failure into a Python exception at the boundary, without the underlying library needing any Python-specific code at all.

PyTorch custom operators are a related but distinct path worth knowing about conceptually: PyTorch’s extension mechanism lets a C++/CUDA operator receive torch::Tensor arguments directly (with PyTorch handling device placement, dtype, and autograd registration), which removes the need for the manual pointer marshaling shown above, at the cost of tying the binding to PyTorch’s tensor and build system. Which of the two you reach for is a project-level dependency decision (pybind11 for a framework-agnostic binding, a PyTorch extension for a PyTorch-only operator); this article is not adding either as a required dependency, only describing the shape of both options so you can choose deliberately.

Testing invalid inputs

A library’s error-handling contract is only real if you test it. At minimum, verify:

// Null pointer: must return VECOPS_ERROR_INVALID_ARGUMENT, not crash.
assert(vecops_axpy(1.0f, nullptr, d_y, 1024, 0) == VECOPS_ERROR_INVALID_ARGUMENT);

// Zero-length: must be a defined no-op, not undefined behavior.
assert(vecops_axpy(1.0f, d_x, d_y, 0, 0) == VECOPS_SUCCESS);

// Invalid/destroyed stream: the underlying CUDA call fails; the library
// must translate that into VECOPS_ERROR_CUDA_FAILURE, not crash or hang.
cudaStream_t bad_stream;
cudaStreamCreate(&bad_stream);
cudaStreamDestroy(bad_stream);   // now invalid
assert(vecops_axpy(1.0f, d_x, d_y, 1024, bad_stream) == VECOPS_ERROR_CUDA_FAILURE);

// No CUDA device present (simulate with CUDA_VISIBLE_DEVICES="" in the
// test environment, or mock cudaGetDeviceCount in a unit test harness):
// must return VECOPS_ERROR_INVALID_DEVICE, not segfault.

// Unsupported architecture: running the library's compiled binary on a
// GPU whose compute capability was not in CMAKE_CUDA_ARCHITECTURES
// produces cudaErrorInvalidDeviceFunction at kernel launch time. Confirm
// this surfaces as VECOPS_ERROR_CUDA_FAILURE via cudaGetLastError(),
// rather than an unhandled crash, by deliberately building a test binary
// without the current test GPU's architecture listed.

That last case is easy to overlook: if CMAKE_CUDA_ARCHITECTURES does not include the compute capability of the GPU actually running your code, and no PTX fallback was embedded either, nvcc produced no usable kernel image for that GPU at all. The failure only appears at launch time as cudaErrorInvalidDeviceFunction, not at build time, which is exactly why the wrapper’s explicit cudaGetLastError() check after every launch, shown earlier, is not optional boilerplate.

In practice

Design the ownership, stream, and error contract before writing the first kernel; retrofitting an explicit stream parameter or a status-code return type onto an API that started as “just call cudaDeviceSynchronize and print errors” breaks every existing caller. Keep kernels and their __device__ helpers in internal linkage (anonymous namespaces or static) so they never appear in your exported symbol table, and expose only a small, C-linkage, pointer-and-length surface at the header. Reach for RDC only when your library’s kernels are genuinely split across files by design; it costs both optimization opportunities and build complexity. Use modern CMake’s native CUDA language support and install(EXPORT ...) so find_package works for downstream consumers instead of requiring them to hand-write include and link paths. Add a Python binding as a thin, separate layer once the C ABI is stable, and test the failure paths (null pointers, invalid streams, missing devices, unsupported architectures) as deliberately as you test the success path, because CUDA’s failure modes are easy to get wrong silently.

What comes next

You now know how to turn a working kernel into a distributable, testable library with an explicit ownership and error contract, rather than a single file compiled ad hoc with nvcc. The next article, CUDA synchronization and atomics, returns to kernel internals: __syncthreads(), memory fences, and atomic operations for safe concurrent access within a block and across the device, all of which your library’s kernels will eventually need once they do more than one independent operation per thread.