CUDA 基础编程核心笔记 + C++示例代码（第三章重点）

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

// CUDA核函数：每个线程输出自身的层级ID
__global__ void hello_cuda() {
    // 块内线程ID
    int tid_in_block = threadIdx.x;
    // 网格内块ID
    int bid = blockIdx.x;
    // 块内线程数
    int block_threads = blockDim.x;
    // 全局线程ID（唯一标识整个网格中的线程）
    int global_tid = tid_in_block + bid * block_threads;

    // 输出线程信息（注意：输出顺序不保证，因线程调度无序）
    printf("Grid内块ID：%d, 块内线程ID：%d, 全局线程ID：%d → Hello CUDA!\n",
           bid, tid_in_block, global_tid);
}

int main() {
    std::cout << "CPU端：启动CUDA核函数..." << std::endl;

    // 核函数配置：2个块（gridDim.x=2），每个块5个线程（blockDim.x=5）→ 总线程数=2*5=10
    dim3 grid_dim(2);    // 网格维度：1D，2个块
    dim3 block_dim(5);   // 块维度：1D，5个线程
    hello_cuda<<<grid_dim, block_dim>>>();  // 启动核函数

    // 等待GPU核函数执行完成（同步CPU与GPU）
    cudaDeviceSynchronize();

    // 检查核函数执行错误
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        std::cerr << "CUDA错误：" << cudaGetErrorString(err) << std::endl;
        return 1;
    }

    std::cout << "CPU端：核函数执行完成！" << std::endl;
    return 0;
}

#include <cuda_runtime.h>
#include <iostream>
#include <string>

// 错误检查宏（简化CUDA API调用错误处理）
#define CHECK_CUDA_ERR(err) \
    if (err != cudaSuccess) { \
        std::cerr << "CUDA错误：" << cudaGetErrorString(err) << "（行号：" << __LINE__ << "）" << std::endl; \
        exit(1); \
    }

int main() {
    int device_count;
    cudaError_t err = cudaGetDeviceCount(&device_count);  // 获取GPU设备数量
    CHECK_CUDA_ERR(err);

    std::cout << "系统中GPU设备数量：" << device_count << std::endl;
    std::cout << "------------------------GPU设备属性------------------------" << std::endl;

    for (int dev_id = 0; dev_id < device_count; ++dev_id) {
        cudaDeviceProp device_prop;
        err = cudaGetDeviceProperties(&device_prop, dev_id);  // 获取指定设备属性
        CHECK_CUDA_ERR(err);

        std::cout << "设备" << dev_id << "：" << device_prop.name << std::endl;
        std::cout << "  CUDA算力版本：" << device_prop.major << "." << device_prop.minor << std::endl;
        std::cout << "  每块最大线程数：" << device_prop.maxThreadsPerBlock << std::endl;
        std::cout << "  每块共享内存大小：" << device_prop.sharedMemPerBlock << " 字节" << std::endl;
        std::cout << "  最大网格维度（x/y/z）：" << device_prop.maxGridSize[0] << "/"
                  << device_prop.maxGridSize[1] << "/" << device_prop.maxGridSize[2] << std::endl;
        std::cout << "  全局内存大小：" << (double)device_prop.totalGlobalMem / (1024 * 1024 * 1024) << " GB" << std::endl;
        std::cout << "---------------------------------------------------------" << std::endl;
    }

    return 0;
}

#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <chrono>

using namespace std;
using namespace chrono;

// 错误检查宏
#define CHECK_CUDA_ERR(err) \
    if (err != cudaSuccess) { \
        cerr << "CUDA错误：" << cudaGetErrorString(err) << "（行号：" << __LINE__ << "）" << endl; \
        exit(1); \
    }

// 块大小（适配硬件：通常设为32、64、1024，需≤设备maxThreadsPerBlock）
const int BLOCK_SIZE = 1024;

// CUDA核函数：向量加法（数据并行）
__global__ void vector_add_kernel(const float* d_a, const float* d_b, float* d_c, int size) {
    // 计算全局线程ID（映射到向量索引）
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 避免线程索引超出向量范围（网格大小可能略大于向量所需）
    if (tid < size) {
        d_c[tid] = d_a[tid] + d_b[tid];  // 每个线程处理一个元素（数据并行）
    }
}

// CPU串行向量加法（对比基准）
void cpu_vector_add(const vector<float>& a, const vector<float>& b, vector<float>& c) {
    for (int i = 0; i < a.size(); ++i) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    // 向量大小（1000万元素，体现并行优势）
    const int VECTOR_SIZE = 10'000'000;

    // 1. 主机端内存分配与初始化
    vector<float> h_a(VECTOR_SIZE, 3.14f);   // 向量a：全为3.14
    vector<float> h_b(VECTOR_SIZE, 2.71f);   // 向量b：全为2.71
    vector<float> h_c_cpu(VECTOR_SIZE, 0.0f); // CPU结果
    vector<float> h_c_gpu(VECTOR_SIZE, 0.0f); // GPU结果

    // 2. CPU串行计算
    auto start = high_resolution_clock::now();
    cpu_vector_add(h_a, h_b, h_c_cpu);
    auto cpu_time = duration_cast<milliseconds>(high_resolution_clock::now() - start).count();
    cout << "CPU串行计算耗时：" << cpu_time << " ms" << endl;

    // 3. GPU并行计算
    float *d_a, *d_b, *d_c;
    // 设备端内存分配
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_a, VECTOR_SIZE * sizeof(float)));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_b, VECTOR_SIZE * sizeof(float)));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_c, VECTOR_SIZE * sizeof(float)));

    // 主机→设备数据拷贝
    CHECK_CUDA_ERR(cudaMemcpy(d_a, h_a.data(), VECTOR_SIZE * sizeof(float), cudaMemcpyHostToDevice));
    CHECK_CUDA_ERR(cudaMemcpy(d_b, h_b.data(), VECTOR_SIZE * sizeof(float), cudaMemcpyHostToDevice));

    // 计算网格大小：ceil(VECTOR_SIZE / BLOCK_SIZE)，确保所有元素被覆盖
    dim3 block_dim(BLOCK_SIZE);
    dim3 grid_dim((VECTOR_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE);
    cout << "GPU配置：网格大小=" << grid_dim.x << "，块大小=" << block_dim.x << endl;

    // 核函数执行与计时
    start = high_resolution_clock::now();
    vector_add_kernel<<<grid_dim, block_dim>>>(d_a, d_b, d_c, VECTOR_SIZE);
    CHECK_CUDA_ERR(cudaGetLastError()); // 检查核函数启动错误
    CHECK_CUDA_ERR(cudaDeviceSynchronize()); // 同步GPU
    auto gpu_time = duration_cast<milliseconds>(high_resolution_clock::now() - start).count();

    // 设备→主机数据拷贝（结果回传）
    CHECK_CUDA_ERR(cudaMemcpy(h_c_gpu.data(), d_c, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost));

    // 4. 结果验证
    bool result_valid = true;
    for (int i = 0; i < 10; ++i) { // 验证前10个元素（避免遍历耗时）
        if (abs(h_c_cpu[i] - h_c_gpu[i]) > 1e-5) {
            result_valid = false;
            break;
        }
    }

    // 输出结果
    cout << "GPU并行计算耗时：" << gpu_time << " ms" << endl;
    cout << "结果验证：" << (result_valid ? "正确" : "错误") << endl;
    cout << "GPU加速比：" << (double)cpu_time / gpu_time << "x" << endl;

    // 5. 释放设备内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}