CUDA Vector Addition Optimization Guide

Introduction

Vector addition is the "Hello World" of CUDA programming. While seemingly simple, it teaches fundamental concepts that apply to all GPU kernels: memory coalescing, thread organization, and bandwidth optimization. A well-optimized vector add achieves near-theoretical memory bandwidth.

Common Performance Issues

• Not achieving peak memory bandwidth due to poor coalescing
• Incorrect grid/block dimensions causing incomplete processing
• Missing bounds checking for non-power-of-2 sizes
• Excessive kernel launch overhead for small vectors

Optimization Techniques

__global__ void vectorAdd(float* a, float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = idx; i < n; i += stride) {
        c[i] = a[i] + b[i];
    }
}

__global__ void vectorAdd4(float4* a, float4* b, float4* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float4 va = a[idx];
        float4 vb = b[idx];
        c[idx] = make_float4(va.x+vb.x, va.y+vb.y, va.z+vb.z, va.w+vb.w);
    }
}

Implementation Comparison

Before (Naive Implementation)

Simple one-element-per-thread approach.

__global__ void vectorAddNaive(float* a, float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// Launch
int threads = 256;
int blocks = (n + threads - 1) / threads;
vectorAddNaive<<<blocks, threads>>>(d_a, d_b, d_c, n);

After (Optimized Implementation)

Combines vectorized loads with grid-stride loop.

__global__ void vectorAddOpt(float4* a, float4* b, float4* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;

    for (int i = idx; i < n; i += stride) {
        float4 va = a[i];
        float4 vb = b[i];
        c[i] = make_float4(va.x+vb.x, va.y+vb.y, va.z+vb.z, va.w+vb.w);
    }
}

// Launch with limited blocks for grid-stride
int threads = 256;
int blocks = min((n/4 + threads - 1) / threads, 256);
vectorAddOpt<<<blocks, threads>>>((float4*)d_a, (float4*)d_b, (float4*)d_c, n/4);

Performance Results

Frequently Asked Questions

Related Guides

Ready to optimize your CUDA code? Download RightNow AI and get real-time performance analysis for your kernels.