Mastering CUDA Function Type Annotations: A Complete Guide

Overview

CUDA provides several function type annotations that define where a function can be called and where it runs. Understanding these annotations is essential for effective CUDA programming because they determine execution space and call constraints.

Function Type Annotations

1. __global__ – Kernel function

Purpose: Define a function that runs on the GPU and is called from the host (CPU).

Characteristics:

Executes on the device (GPU)

Called from the host (CPU)

Must return void Cannot be called from other device functions

Executes asynchronously unless synchronized

Syntax:

__global__ void KernelFunction(parameters) {
    // GPU code
}

2. __device__ – Device function

Purpose: Define a function that runs on the GPU and can be called from other GPU functions.

Characteristics:

Executes on the device (GPU)

Can only be called from __global__ or other __device__ functions

Can return any type

Cannot be called from host code

Inlined by default for performance

Syntax:

__device__ ReturnType DeviceFunction(parameters) {
    // GPU code
    return value;
}

3. __host__ – Host function

Purpose: Define a function that runs on the CPU (default behavior).

Characteristics:

Executes on the host (CPU)

Can only be called from host code

Annotation can be omitted

Cannot be called from device code

Syntax:

__host__ ReturnType HostFunction(parameters) {
    // CPU code
    return value;
}
// Equivalent to:
ReturnType HostFunction(parameters) {
    // CPU code
    return value;
}

4. Combined annotation __host__ __device__

Functions marked with both annotations can be compiled for both host and device execution.

Characteristics:

Compiled for CPU and GPU

Can be called from host and device code

Useful for utility functions

Code must be compatible with both architectures

Syntax:

__host__ __device__ ReturnType DualFunction(parameters) {
    // Code that works on both CPU and GPU
    return value;
}

Memory Space Annotations

__shared__ – Shared memory

Purpose: Declare variables in shared memory within a thread block.

Characteristics:

Shared among all threads in a block

Much faster than global memory

Size limited (typically 48 KB–96 KB per block)

Lifetime matches block execution

__constant__ – Constant memory

Purpose: Declare read‑only variables in constant memory.

Characteristics:

Read‑only from device code

Cached for faster access

Total size limited to 64 KB

Initialized from host code

__managed__ – Unified memory

Purpose: Declare variables that can be accessed from both CPU and GPU and are migrated automatically.

Characteristics:

Automatic migration between CPU and GPU

Simplifies memory management

May impact performance

Requires compute capability 3.0+

Practical Examples

Example 1: Basic kernel calling a device function

// Device function for distance calculation
__device__ float computeDistance(float x1, float y1, float x2, float y2) {
    float dx = x2 - x1;
    float dy = y2 - y1;
    return sqrtf(dx * dx + dy * dy);
}

// Global kernel that uses the device function
__global__ void calculateDistances(float* x1, float* y1, float* x2, float* y2,
                                 float* distances, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        distances[idx] = computeDistance(x1[idx], y1[idx], x2[idx], y2[idx]);
    }
}

Example 2: Host‑device dual function

// Function usable on both CPU and GPU
__host__ __device__ float lerp(float a, float b, float t) {
    return a + t * (b - a);
}

// Device code that uses the dual function
__device__ float interpolateValue(float* array, float index, int size) {
    int lower = (int)index;
    int upper = min(lower + 1, size - 1);
    float t = index - lower;
    return lerp(array[lower], array[upper], t);
}

// Host code that also uses the dual function
__host__ void preprocessData(float* data, int size) {
    for (int i = 0; i < size - 1; i++) {
        data[i] = lerp(data[i], data[i + 1], 0.5f);
    }
}

Example 3: Memory‑space annotations

// Constant memory declaration
__constant__ float convolution_kernel[9];

// Global kernel using shared and constant memory
__global__ void convolutionWithShared(float* input, float* output,
                                    int width, int height) {
    // Tile in shared memory (16×16 threads + 2‑pixel halo)
    __shared__ float tile[18][18];

    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int gx = blockIdx.x * blockDim.x + tx;
    int gy = blockIdx.y * blockDim.y + ty;

    // Load data into shared memory (details omitted)
    __syncthreads();

    // Convolution using constant memory kernel
    if (gx < width && gy < height) {
        float result = 0.0f;
        for (int i = 0; i < 3; i++) {
            for (int j = 0; j < 3; j++) {
                result += tile[ty + i][tx + j] * convolution_kernel[i * 3 + j];
            }
        }
        output[gy * width + gx] = result;
    }
}

Best Practices and Guidelines

1. Function design principles

Keep device functions simple: avoid complex control flow inside __device__ functions.

Minimize parameter passing: use references or pointers where possible.

Consider inlining: small __device__ functions are automatically inlined.

2. Memory management

Use __shared__ memory for data shared within a block.

Use __constant__ memory for read‑only data accessed by all threads.

Use global memory for large data sets.

3. Performance considerations

Avoid branch divergence; minimize if‑else statements in kernels.

Optimize memory access patterns; ensure coalesced accesses.

Use __host__ __device__ functions judiciously to promote code reuse.

4. Debugging and development

Start with simple __global__ kernels and gradually add complexity.

Test incrementally; verify each function type works as expected.

Always check CUDA error codes.

Common Pitfalls and Solutions

1. Call restrictions

Problem: Attempting to call a __device__ function from host code.

// Incorrect: leads to compilation error
__device__ int deviceFunc() { return 42; }

int main() {
    int result = deviceFunc(); // Error!
    return 0;
}

Solution: Call the device function from a kernel.

__device__ int deviceFunc() { return 42; }

__global__ void kernel(int* result) {
    *result = deviceFunc(); // Correct
}

2. Return‑type limitation

Problem: __global__ functions returning non‑void.

// Incorrect: global functions must return void
__global__ int badKernel() {
    return 42; // Error!
}

Solution: Use output parameters.

// Correct: use an output parameter
__global__ void goodKernel(int* output) {
    *output = 42;
}

3. Memory‑space confusion

Problem: Incorrectly accessing different memory spaces.

__shared__ float sharedData[256];

__global__ void kernel() {
    // Wrong: trying to copy shared memory to host
    cudaMemcpy(hostPtr, sharedData, sizeof(float) * 256, cudaMemcpyDeviceToHost);
}

Solution: Copy via global memory.

__global__ void kernel(float* globalOutput) {
    __shared__ float sharedData[256];
    // Process data in shared memory (omitted)

    // Copy to global memory
    if (threadIdx.x < 256) {
        globalOutput[threadIdx.x] = sharedData[threadIdx.x];
    }
}

Advanced Topics

1. Dynamic parallelism

CUDA supports launching kernels from device code (compute capability 3.5+).

__global__ void parentKernel() {
    // Launch child kernel from device
    childKernel<<<1, 1>>>();
    cudaDeviceSynchronize(); // Synchronize child kernel
}

__global__ void childKernel() {
    printf("Greeting from child kernel!
");
}

2. Cooperative groups

Modern CUDA programming with cooperative groups:

#include <cooperative_groups.h>
using namespace cooperative_groups;

__global__ void cooperativeKernel() {
    thread_block block = this_thread_block();

    // Synchronize all threads in the block
    block.sync();

    // Cooperative group operation
    int sum = reduce(block, threadIdx.x, plus<int>());
}

Conclusion

Understanding CUDA function type annotations is the foundation of effective GPU programming. These annotations control:

Execution location (CPU vs GPU)

Call context (where a function can be invoked)

Memory access (which memory spaces a function can use)

Performance characteristics (how to optimise functions)

By mastering these concepts and following the best practices, you can write efficient, maintainable CUDA code that fully exploits the power of GPU computation.