NVIDIA CUDA编程实操二：与CPU编程的对比

环境准备

本次使用P100、CUDA10的ubuntu18.04环境：

第一部分：向量加法对比

CPU版本 (vector_add_cpu.c)

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

#define N 1000000

void vector_add_cpu(float *a, float *b, float *c) {
    for(int i = 0; i < N; i++) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    float *a, *b, *c;
    a = (float*)malloc(N * sizeof(float));
    b = (float*)malloc(N * sizeof(float));
    c = (float*)malloc(N * sizeof(float));
    
    // 初始化数据
    for(int i = 0; i < N; i++) {
        a[i] = 1.0f;
        b[i] = 2.0f;
    }
    
    clock_t start = clock();
    vector_add_cpu(a, b, c);
    clock_t end = clock();
    
    printf("Time on CPU: %f ms\n", (double)(end - start) * 1000 / CLOCKS_PER_SEC);
    
    free(a);
    free(b);
    free(c);
    return 0;
}

CUDA版本 (vector_add_gpu.cu)

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

#define N 1000000
#define THREADS_PER_BLOCK 256

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

int main() {
    float *a, *b, *c;
    float *d_a, *d_b, *d_c;
    
    // 分配主机内存
    a = (float*)malloc(N * sizeof(float));
    b = (float*)malloc(N * sizeof(float));
    c = (float*)malloc(N * sizeof(float));
    
    // 初始化数据
    for(int i = 0; i < N; i++) {
        a[i] = 1.0f;
        b[i] = 2.0f;
    }
    
    // 分配设备内存
    cudaMalloc(&d_a, N * sizeof(float));
    cudaMalloc(&d_b, N * sizeof(float));
    cudaMalloc(&d_c, N * sizeof(float));
    
    // 拷贝数据到设备
    cudaMemcpy(d_a, a, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, N * sizeof(float), cudaMemcpyHostToDevice);
    
    // 启动核函数
    int blocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start);
    
    vector_add_gpu<<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b, d_c);
    
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    // 拷贝结果回主机
    cudaMemcpy(c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);
    
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("Time on GPU: %f ms\n", milliseconds);
    
    // 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(a);
    free(b);
    free(c);
    
    return 0;
}

1. 编译并运行两个版本：

gcc vector_add_cpu.c -o cpu_add
./cpu_add

nvcc vector_add_gpu.cu -o gpu_add
./gpu_add

2. 观察执行时间差异

Time on CPU: 6.899000 ms

Time on GPU: 0.038336 ms

#如果计时调整为从分配设备内存开始、到结果拷贝回主机内存结束，则耗时明显增加
Time on GPU: 3.925920 ms

第二部分：核心差异讲解

1. 并行模式：

   • CPU：顺序执行，单线程循环处理每个元素
   • GPU：数千线程并行处理，每个线程处理少量元素

2. 内存模型：

   • CPU：单一内存空间
   • GPU：主机内存和设备内存分离，需要显式传输

3. 线程组织：

   • GPU使用grid-block-thread层次结构：
     • grid：包含多个block
     • block：包含多个thread(最多1024个)
     • thread：最小执行单元

4. 性能考量：

   • GPU有更高的理论计算吞吐量
   • 但数据传输开销显著，适合计算密集型任务

第三部分：代码扩展解读

1. 块(blocks)的计算原理

int blocks = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;

这是一个经典的"向上取整"整数除法技巧，目的是确保所有数据都能被处理。分解来看：

• 问题：当数据总量N不是THREADS_PER_BLOCK的整数倍时，简单除法会向下取整，导致最后一部分数据无法被处理
• 解决方案：通过(N + THREADS_PER_BLOCK - 1)实现向上取整
• 例子：
  • 如果N=1000，THREADS_PER_BLOCK=256
  • 正常除法：1000/256=3.906→取整3（会漏掉最后232个元素）
  • 我们的方法：(1000+256-1)/256=1255/256=4.902→取整4（确保全覆盖）

2. 核函数启动语法

vector_add_gpu<<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b, d_c);

这是CUDA特有的"执行配置"语法，指定了如何组织线程：

• <<< >>>：CUDA核函数特有的执行配置符号
• 第一个参数blocks：grid中的block数量
• 第二个参数THREADS_PER_BLOCK：每个block中的thread数量
• 组合效果：总共启动 blocks × THREADS_PER_BLOCK 个线程

3. 线程索引的完整逻辑

在核函数内部，通过以下方式计算全局索引：

int i = blockIdx.x * blockDim.x + threadIdx.x;

• blockIdx.x：当前block在grid中的索引(0到blocks-1)
• blockDim.x：每个block的线程数(=THREADS_PER_BLOCK)
• threadIdx.x：当前thread在block中的索引(0到THREADS_PER_BLOCK-1)

4. 为什么需要if(i < N)判断

因为我们的blocks计算是向上取整的，实际启动的线程数可能略多于N：

• 例如N=1000，THREADS_PER_BLOCK=256时：
  • 计算得blocks=4
  • 实际启动线程=4×256=1024
  • 但只需要前1000个线程工作
  • 最后24个线程必须跳过操作(否则会越界访问内存)

5. 为什么这样设计？

这种设计模式提供了：

• 灵活性：适应任意大小的数据集
• 高效性：保持线程块大小最优(如256是常见高效值)
• 可扩展性：自动适应不同GPU(不同GPU有最大grid尺寸限制)

第四部分：硬件扩展解读

1. 硬件配置

Tesla P100内存容量

• Tesla P100配备16GB或12GB HBM2显存
• 对于N=1,000,000的float数组：
  • 单个数组大小：1,000,000 × 4字节 = 3.81MB
  • 三个数组共约11.44MB

Tesla P100核心架构

• 基于Pascal架构
• 包含56或60个SM（Streaming Multiprocessor）
• 每个SM包含：
  • 64个FP32 CUDA核心
  • 32个FP64 CUDA核心
  • 256KB寄存器文件
  • 64KB共享内存/一级缓存（可配置）
  • 4个纹理单元

执行流程

• 当核函数启动时，CUDA runtime会：
  1. 创建所有线程的索引信息
  2. 将线程分组为warp（每组32个）
  3. SM按warp为单位调度执行
  4. 每个线程在运行时通过内置变量计算自己的唯一索引
• 所有线程并行执行相同的指令（SIMT架构）
• 分支处理（如if(i<N)）通过predication机制实现

2. 性能优化提示

1. Block大小选择：

   • Tesla P100每个SM最多2048个线程
   • 建议block大小为128/256/512的倍数
   • 可通过Occupancy Calculator工具计算最优值

2. 内存访问模式：

   • 确保线程对全局内存的访问是合并的(coalesced)
   • 连续的线程访问连续的内存地址

3. 资源限制：

   • 每个block最多1024个线程
   • 每个grid最多2^31-1个block
   • 寄存器使用影响实际并行度

课程总结

• CUDA通过大规模并行获得性能优势
• 理解层次化线程模型
• 注意主机-设备内存分离的特点
• 实际性能受多种因素影响(block大小、数据量等)

附nvidia-smi -q输出

==============NVSMI LOG==============

Timestamp                                 : Wed Jun 11 16:02:25 2025
Driver Version                            : 470.57.02
CUDA Version                              : 11.4

Attached GPUs                             : 1
GPU 00000000:00:0D.0
    Product Name                          : Tesla P100-PCIE-16GB
    Product Brand                         : Tesla
    Display Mode                          : Enabled
    Display Active                        : Disabled
    Persistence Mode                      : Enabled
    MIG Mode
        Current                           : N/A
        Pending                           : N/A
    Accounting Mode                       : Disabled
    Accounting Mode Buffer Size           : 4000
    Driver Model
        Current                           : N/A
        Pending                           : N/A
    Serial Number                         : 0320418056407
    GPU UUID                              : GPU-c046a1d5-7146-be85-e430-d08976b3825f
    Minor Number                          : 0
    VBIOS Version                         : 86.00.4D.00.01
    MultiGPU Board                        : No
    Board ID                              : 0xd
    GPU Part Number                       : 900-2H400-0000-000
    Module ID                             : 0
    Inforom Version
        Image Version                     : H400.0201.00.08
        OEM Object                        : 1.1
        ECC Object                        : 4.1
        Power Management Object           : N/A
    GPU Operation Mode
        Current                           : N/A
        Pending                           : N/A
    GSP Firmware Version                  : N/A
    GPU Virtualization Mode
        Virtualization Mode               : Pass-Through
        Host VGPU Mode                    : N/A
    IBMNPU
        Relaxed Ordering Mode             : N/A
    PCI
        Bus                               : 0x00
        Device                            : 0x0D
        Domain                            : 0x0000
        Device Id                         : 0x15F810DE
        Bus Id                            : 00000000:00:0D.0
        Sub System Id                     : 0x118F10DE
        GPU Link Info
            PCIe Generation
                Max                       : 3
                Current                   : 3
            Link Width
                Max                       : 16x
                Current                   : 16x
        Bridge Chip
            Type                          : N/A
            Firmware                      : N/A
        Replays Since Reset               : 0
        Replay Number Rollovers           : 0
        Tx Throughput                     : 0 KB/s
        Rx Throughput                     : 0 KB/s
    Fan Speed                             : N/A
    Performance State                     : P0
    Clocks Throttle Reasons
        Idle                              : Active
        Applications Clocks Setting       : Not Active
        SW Power Cap                      : Not Active
        HW Slowdown                       : Not Active
            HW Thermal Slowdown           : Not Active
            HW Power Brake Slowdown       : Not Active
        Sync Boost                        : Not Active
        SW Thermal Slowdown               : Not Active
        Display Clock Setting             : Not Active
    FB Memory Usage
        Total                             : 16280 MiB
        Used                              : 0 MiB
        Free                              : 16280 MiB
    BAR1 Memory Usage
        Total                             : 16384 MiB
        Used                              : 2 MiB
        Free                              : 16382 MiB
    Compute Mode                          : Default
    Utilization
        Gpu                               : 0 %
        Memory                            : 0 %
        Encoder                           : 0 %
        Decoder                           : 0 %
    Encoder Stats
        Active Sessions                   : 0
        Average FPS                       : 0
        Average Latency                   : 0
    FBC Stats
        Active Sessions                   : 0
        Average FPS                       : 0
        Average Latency                   : 0
    Ecc Mode
        Current                           : Enabled
        Pending                           : Enabled
    ECC Errors
        Volatile
            Single Bit            
                Device Memory             : 0
                Register File             : 0
                L1 Cache                  : N/A
                L2 Cache                  : 0
                Texture Memory            : 0
                Texture Shared            : 0
                CBU                       : N/A
                Total                     : 0
            Double Bit            
                Device Memory             : 0
                Register File             : 0
                L1 Cache                  : N/A
                L2 Cache                  : 0
                Texture Memory            : 0
                Texture Shared            : 0
                CBU                       : N/A
                Total                     : 0
        Aggregate
            Single Bit            
                Device Memory             : 0
                Register File             : 0
                L1 Cache                  : N/A
                L2 Cache                  : 0
                Texture Memory            : 0
                Texture Shared            : 0
                CBU                       : N/A
                Total                     : 0
            Double Bit            
                Device Memory             : 0
                Register File             : 0
                L1 Cache                  : N/A
                L2 Cache                  : 0
                Texture Memory            : 0
                Texture Shared            : 0
                CBU                       : N/A
                Total                     : 0
    Retired Pages
        Single Bit ECC                    : 0
        Double Bit ECC                    : 0
        Pending Page Blacklist            : No
    Remapped Rows                         : N/A
    Temperature
        GPU Current Temp                  : 34 C
        GPU Shutdown Temp                 : 85 C
        GPU Slowdown Temp                 : 82 C
        GPU Max Operating Temp            : N/A
        GPU Target Temperature            : N/A
        Memory Current Temp               : N/A
        Memory Max Operating Temp         : N/A
    Power Readings
        Power Management                  : Supported
        Power Draw                        : 26.23 W
        Power Limit                       : 250.00 W
        Default Power Limit               : 250.00 W
        Enforced Power Limit              : 250.00 W
        Min Power Limit                   : 125.00 W
        Max Power Limit                   : 250.00 W
    Clocks
        Graphics                          : 405 MHz
        SM                                : 405 MHz
        Memory                            : 715 MHz
        Video                             : 835 MHz
    Applications Clocks
        Graphics                          : 1189 MHz
        Memory                            : 715 MHz
    Default Applications Clocks
        Graphics                          : 1189 MHz
        Memory                            : 715 MHz
    Max Clocks
        Graphics                          : 1328 MHz
        SM                                : 1328 MHz
        Memory                            : 715 MHz
        Video                             : 1328 MHz
    Max Customer Boost Clocks
        Graphics                          : 1328 MHz
    Clock Policy
        Auto Boost                        : N/A
        Auto Boost Default                : N/A
    Voltage
        Graphics                          : N/A
    Processes                             : None