向量加法CUDA实现

/**
 * 向量加法 - CUDA基础实现
 * C = A + B (element-wise)
 * 
 * 编译: nvcc -o vecadd vecadd.cu
 * 运行: ./vecadd
 */

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

// ==================== 错误检查宏 ====================
// 这个宏会检查每个CUDA API调用是否成功
// 如果失败，会打印错误信息和位置，然后退出
#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA Error: %s at %s:%d\n", \
                cudaGetErrorString(err), __FILE__, __LINE__); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

// ==================== CUDA Kernel ====================
// __global__ 修饰的函数在GPU上执行
// 从CPU调用（host function）
__global__ void vec_add_kernel(const float* A, const float* B, float* C, int n) {
    // 计算全局线程索引
    // threadIdx.x: 线程在Block内的局部索引 (0 ~ blockDim.x-1)
    // blockIdx.x: Block在Grid中的索引 (0 ~ gridDim.x-1)
    // blockDim.x: 每个Block的线程数
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    // 边界检查：确保不访问越界内存
    if (idx < n) {
        C[idx] = A[idx] + B[idx];
    }
}

// ==================== CPU参考实现 ====================
void vec_add_cpu(const float* A, const float* B, float* C, int n) {
    for (int i = 0; i < n; i++) {
        C[i] = A[i] + B[i];
    }
}

// ==================== 结果验证 ====================
int verify(const float* cpu_result, const float* gpu_result, int n) {
    for (int i = 0; i < n; i++) {
        if (fabs(cpu_result[i] - gpu_result[i]) > 1e-5) {
            printf("❌ Mismatch at index %d: CPU=%f, GPU=%f\n", i, cpu_result[i], gpu_result[i]);
            return 0;
        }
    }
    printf("✅ Results match!\n");
    return 1;
}

// ==================== 主函数 ====================
int main() {
    const int N = 1 << 20;  // 1M elements
    const int size = N * sizeof(float);
    
    printf("向量加法: N = %d (%.2f MB)\n", N, size / 1024.0 / 1024.0);
    
    // ==================== 1. 分配Host内存 ====================
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);
    float *h_C_gpu = (float*)malloc(size);
    
    // 初始化输入数据
    for (int i = 0; i < N; i++) {
        h_A[i] = float(i) * 0.01f;
        h_B[i] = float(i) * 0.02f;
    }
    
    // ==================== 2. 分配Device内存 ====================
    float *d_A, *d_B, *d_C;
    CUDA_CHECK(cudaMalloc(&d_A, size));
    CUDA_CHECK(cudaMalloc(&d_B, size));
    CUDA_CHECK(cudaMalloc(&d_C, size));
    
    // ==================== 3. 数据传输 Host → Device ====================
    CUDA_CHECK(cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice));
    
    // ==================== 4. 配置并启动Kernel ====================
    int threads_per_block = 256;  // 每个Block 256个线程
    int blocks = (N + threads_per_block - 1) / threads_per_block;  // 向上取整
    
    printf("启动配置: Grid=%d, Block=%d, Total=%d threads\n", 
           blocks, threads_per_block, blocks * threads_per_block);
    
    // 启动Kernel
    // <<>> 是CUDA的启动语法
    vec_add_kernel<<<blocks, threads_per_block>>>(d_A, d_B, d_C, N);
    
    // 检查Kernel启动错误
    CUDA_CHECK(cudaGetLastError());
    // 等待Kernel执行完成
    CUDA_CHECK(cudaDeviceSynchronize());
    
    // ==================== 5. 数据传输 Device → Host ====================
    CUDA_CHECK(cudaMemcpy(h_C_gpu, d_C, size, cudaMemcpyDeviceToHost));
    
    // ==================== 6. 验证结果 ====================
    vec_add_cpu(h_A, h_B, h_C, N);
    verify(h_C, h_C_gpu, N);
    
    // ==================== 7. 清理资源 ====================
    CUDA_CHECK(cudaFree(d_A));
    CUDA_CHECK(cudaFree(d_B));
    CUDA_CHECK(cudaFree(d_C));
    free(h_A);
    free(h_B);
    free(h_C);
    free(h_C_gpu);
    
    printf("✅ 完成!\n");
    return 0;
}