CUDA 并行编程进阶笔记 + C++示例代码（第四章重点）

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

using namespace std;
using namespace chrono;

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

// CPU串行素数测试
bool checkPrimeCPU(long long num) {
    if (num <= 1) return false;
    if (num = 2) return true;
    if (num % 2 = 0) return false;
    for (long long i = 3; i * i <= num; i += 2) {
        if (num % i = 0) return false;
    }
    return true;
}

// GPU核函数：素数测试（每个线程处理一个奇数）
__global__ void checkPrimeGPU(long long start, long long end, bool* d_results) {
    // 计算全局线程ID
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 映射到待测试数字（仅处理奇数，跳过偶数）
    long long num = start + (tid * 2);
    // 超出范围则退出
    if (num > end) return;

    bool isPrime = true;
    if (num <= 1) {
        isPrime = false;
    } else if (num = 2) {
        isPrime = true;
    } else if (num % 2 = 0) {
        isPrime = false;
    } else {
        for (long long i = 3; i * i <= num; i += 2) {
            if (num % i = 0) {
                isPrime = false;
                break;
            }
        }
    }

    // 存储结果到设备内存（线程ID对应结果索引）
    d_results[tid] = isPrime;
}

int main() {
    // 测试范围（100,001 ~ 190,001，仅奇数）
    const long long start = 100001;
    const long long end = 190001;
    const int total_nums = (end - start + 1) / 2;  // 奇数个数

    // 1. CPU串行测试
    auto cpu_start = high_resolution_clock::now();
    int cpu_prime_count = 0;
    for (long long num = start; num <= end; num += 2) {
        if (checkPrimeCPU(num)) cpu_prime_count++;
    }
    auto cpu_time = duration_cast<microseconds>(high_resolution_clock::now() - cpu_start).count() / 1000.0;
    cout << "CPU测试结果：" << cpu_prime_count << " 个素数，耗时：" << cpu_time << " ms" << endl;

    // 2. GPU并行测试
    // 线程配置
    const int threadsPerBlock = 256;
    const int blocksPerGrid = (total_nums + threadsPerBlock - 1) / threadsPerBlock;

    // 分配设备内存（存储结果）
    bool* d_results;
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_results, total_nums * sizeof(bool)));

    // 分配主机内存（接收结果）
    bool* h_results = new bool[total_nums];

    // GPU计时
    cudaEvent_t gpu_start, gpu_stop;
    CHECK_CUDA_ERR(cudaEventCreate(&gpu_start));
    CHECK_CUDA_ERR(cudaEventCreate(&gpu_stop));
    CHECK_CUDA_ERR(cudaEventRecord(gpu_start, 0));

    // 启动核函数（预热执行，不计入时间）
    checkPrimeGPU<<<blocksPerGrid, threadsPerBlock>>>(start, end, d_results);
    CHECK_CUDA_ERR(cudaDeviceSynchronize());

    // 正式执行并计时
    CHECK_CUDA_ERR(cudaEventRecord(gpu_start, 0));
    checkPrimeGPU<<<blocksPerGrid, threadsPerBlock>>>(start, end, d_results);
    CHECK_CUDA_ERR(cudaGetLastError());  // 检查核函数启动错误
    CHECK_CUDA_ERR(cudaEventRecord(gpu_stop, 0));
    CHECK_CUDA_ERR(cudaEventSynchronize(gpu_stop));

    // 计算GPU耗时
    float gpu_time;
    CHECK_CUDA_ERR(cudaEventElapsedTime(&gpu_time, gpu_start, gpu_stop));

    // 设备→主机拷贝结果
    CHECK_CUDA_ERR(cudaMemcpy(h_results, d_results, total_nums * sizeof(bool), cudaMemcpyDeviceToHost));

    // 统计GPU素数个数
    int gpu_prime_count = 0;
    for (int i = 0; i < total_nums; ++i) {
        if (h_results[i]) gpu_prime_count++;
    }

    cout << "GPU测试结果：" << gpu_prime_count << " 个素数，耗时：" << gpu_time << " ms" << endl;
    cout << "GPU相对CPU加速比：" << cpu_time / gpu_time << "x" << endl;

    // 释放资源
    delete[] h_results;
    CHECK_CUDA_ERR(cudaFree(d_results));
    CHECK_CUDA_ERR(cudaEventDestroy(gpu_start));
    CHECK_CUDA_ERR(cudaEventDestroy(gpu_stop));

    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); \
    }

// 块大小（适配GPU硬件，≤1024）
const int THREADS_PER_BLOCK = 256;

// CPU串行向量加法
void vectorAddCPU(const float* h_A, const float* h_B, float* h_C, int N) {
    for (int i = 0; i < N; ++i) {
        h_C[i] = h_A[i] + h_B[i];
    }
}

// GPU核函数：向量加法（完全并行）
__global__ void vectorAddGPU(const float* d_A, const float* d_B, float* d_C, int N) {
    // 全局线程ID（映射到向量索引）
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    // Guard Clause：避免索引越界
    if (i < N) {
        d_C[i] = d_A[i] + d_B[i];
    }
}

int main() {
    // 向量大小（可修改为10000/100000/10000000/100000000）
    const int N = 10'000'000;
    const size_t data_size = N * sizeof(float);

    // 1. 主机内存分配与初始化
    vector<float> h_A(N, 1.23f);  // 向量A：全1.23
    vector<float> h_B(N, 4.56f);  // 向量B：全4.56
    vector<float> h_C_CPU(N, 0.0f);  // CPU结果
    vector<float> h_C_GPU(N, 0.0f);  // GPU结果

    // 2. CPU串行计算
    auto cpu_start = high_resolution_clock::now();
    vectorAddCPU(h_A.data(), h_B.data(), h_C_CPU.data(), N);
    auto cpu_time = duration_cast<milliseconds>(high_resolution_clock::now() - cpu_start).count();
    cout << "CPU串行计算耗时：" << cpu_time << " ms" << endl;

    // 3. GPU并行计算（含内存传输）
    float *d_A, *d_B, *d_C;
    // 分配设备内存
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_A, data_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_B, data_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_C, data_size));

    // 计时：总耗时（传输+计算+回传）
    auto gpu_total_start = high_resolution_clock::now();

    // 主机→设备：拷贝输入向量
    CHECK_CUDA_ERR(cudaMemcpy(d_A, h_A.data(), data_size, cudaMemcpyHostToDevice));
    CHECK_CUDA_ERR(cudaMemcpy(d_B, h_B.data(), data_size, cudaMemcpyHostToDevice));

    // 线程配置
    const int blocksPerGrid = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;

    // GPU计算计时（仅计算）
    cudaEvent_t compute_start, compute_stop;
    CHECK_CUDA_ERR(cudaEventCreate(&compute_start));
    CHECK_CUDA_ERR(cudaEventCreate(&compute_stop));
    CHECK_CUDA_ERR(cudaEventRecord(compute_start, 0));

    // 启动核函数（预热）
    vectorAddGPU<<<blocksPerGrid, THREADS_PER_BLOCK>>>(d_A, d_B, d_C, N);
    CHECK_CUDA_ERR(cudaDeviceSynchronize());

    // 正式计算
    vectorAddGPU<<<blocksPerGrid, THREADS_PER_BLOCK>>>(d_A, d_B, d_C, N);
    CHECK_CUDA_ERR(cudaGetLastError());
    CHECK_CUDA_ERR(cudaEventRecord(compute_stop, 0));
    CHECK_CUDA_ERR(cudaEventSynchronize(compute_stop));

    // 计算仅计算耗时
    float compute_time;
    CHECK_CUDA_ERR(cudaEventElapsedTime(&compute_time, compute_start, compute_stop));

    // 设备→主机：拷贝结果
    CHECK_CUDA_ERR(cudaMemcpy(h_C_GPU.data(), d_C, data_size, cudaMemcpyDeviceToHost));

    // 总耗时结束
    auto gpu_total_time = duration_cast<milliseconds>(high_resolution_clock::now() - gpu_total_start).count();

    // 4. 结果验证（前10个元素）
    bool valid = true;
    for (int i = 0; i < 10; ++i) {
        if (abs(h_C_CPU[i] - h_C_GPU[i]) > 1e-5) {
            valid = false;
            break;
        }
    }

    // 输出结果
    cout << "GPU仅计算耗时：" << compute_time << " ms" << endl;
    cout << "GPU总耗时（含传输）：" << gpu_total_time << " ms" << endl;
    cout << "结果验证：" << (valid ? "正确" : "错误") << endl;
    cout << "GPU计算加速比：" << (double)cpu_time / compute_time << "x" << endl;
    cout << "GPU总耗时加速比：" << (double)cpu_time / gpu_total_time << "x" << endl;

    // 5. 释放资源
    CHECK_CUDA_ERR(cudaFree(d_A));
    CHECK_CUDA_ERR(cudaFree(d_B));
    CHECK_CUDA_ERR(cudaFree(d_C));
    CHECK_CUDA_ERR(cudaEventDestroy(compute_start));
    CHECK_CUDA_ERR(cudaEventDestroy(compute_stop));

    return 0;
}

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

using namespace std;
using namespace chrono;

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

// 3D点结构体
struct Point {
    float x, y, z;
};

// GPU核函数：计算两组点的欧氏距离
__global__ void calculateEuclideanDistanceGPU(Point* d_lineA, Point* d_lineB, float* d_distances, int numPoints) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < numPoints) {
        float dx = d_lineA[idx].x - d_lineB[idx].x;
        float dy = d_lineA[idx].y - d_lineB[idx].y;
        float dz = d_lineA[idx].z - d_lineB[idx].z;
        // 用sqrtf（GPU优化的单精度平方根函数）
        d_distances[idx] = sqrtf(dx*dx + dy*dy + dz*dz);
    }
}

int main() {
    const int numPoints = 1'000'000;  // 100万个点
    const size_t point_size = numPoints * sizeof(Point);
    const size_t dist_size = numPoints * sizeof(float);

    // 1. 主机内存初始化
    vector<Point> h_lineA(numPoints);
    vector<Point> h_lineB(numPoints);
    vector<float> h_distances(numPoints, 0.0f);

    // 随机生成点坐标（0~100）
    srand(time(0));
    for (int i = 0; i < numPoints; ++i) {
        h_lineA[i] = {rand()%100, rand()%100, rand()%100};
        h_lineB[i] = {rand()%100, rand()%100, rand()%100};
    }

    // 2. 设备内存分配
    Point *d_lineA, *d_lineB;
    float *d_distances;
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_lineA, point_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_lineB, point_size));
    CHECK_CUDA_ERR(cudaMalloc((void**)&d_distances, dist_size));

    // 3. 数据拷贝与计算
    CHECK_CUDA_ERR(cudaMemcpy(d_lineA, h_lineA.data(), point_size, cudaMemcpyHostToDevice));
    CHECK_CUDA_ERR(cudaMemcpy(d_lineB, h_lineB.data(), point_size, cudaMemcpyHostToDevice));

    // 线程配置
    const int threadsPerBlock = 256;
    const int blocksPerGrid = (numPoints + threadsPerBlock - 1) / threadsPerBlock;

    // GPU计时
    cudaEvent_t start, stop;
    CHECK_CUDA_ERR(cudaEventCreate(&start));
    CHECK_CUDA_ERR(cudaEventCreate(&stop));
    CHECK_CUDA_ERR(cudaEventRecord(start, 0));

    calculateEuclideanDistanceGPU<<<blocksPerGrid, threadsPerBlock>>>(d_lineA, d_lineB, d_distances, numPoints);
    CHECK_CUDA_ERR(cudaGetLastError());
    CHECK_CUDA_ERR(cudaEventRecord(stop, 0));
    CHECK_CUDA_ERR(cudaEventSynchronize(stop));

    // 4. 结果回传与耗时计算
    float gpu_time;
    CHECK_CUDA_ERR(cudaEventElapsedTime(&gpu_time, start, stop));
    CHECK_CUDA_ERR(cudaMemcpy(h_distances.data(), d_distances, dist_size, cudaMemcpyDeviceToHost));

    cout << "GPU计算100万个3D点欧氏距离耗时：" << gpu_time << " ms" << endl;
    cout << "前5个距离结果：";
    for (int i = 0; i < 5; ++i) {
        cout << h_distances[i] << " ";
    }
    cout << endl;

    // 5. 性能分析（Docker环境）
    // 提示：在Docker中启动时添加--cap-add=SYS_ADMIN，执行：ncu -o distance_report ./euclidean_distance
    cout << "性能分析命令（Docker）：ncu -o distance_report ./euclidean_distance" << endl;

    // 6. 释放资源
    CHECK_CUDA_ERR(cudaFree(d_lineA));
    CHECK_CUDA_ERR(cudaFree(d_lineB));
    CHECK_CUDA_ERR(cudaFree(d_distances));
    CHECK_CUDA_ERR(cudaEventDestroy(start));
    CHECK_CUDA_ERR(cudaEventDestroy(stop));

    return 0;
}