期末速通

并行计算

GPU编程

GPU体系结构

GPU发展简史：

2000年以前
固定渲染流水线
不可编程
2001-2006年
加速渲染
用于OpenGL/DirectX和图形学计算
2006年以后
通用计算

GPU通用计算

基本思路：

  1. 专注于计算。以SIMD/SIMT的方式支持大规模并行计算，并非任何计算过程都可以迁移到GPU。
  2. 不支持复杂的逻辑，不单独使用。仅作为辅助CPU的加速计算设备。

GPU逻辑架构

Turing架构：每个GPU封装6个GPC，每个GPC封装6个TPC，每个TPC封装2个SM。

SM的架构

TU102 V100
分为4个block，96KB共享缓存 分为4个block，128KB共享缓存
每个block有2个FP64、16个FP32、16个INT32、2个Tensor、16384个32位寄存器 每个block有8个FP64、16个FP32、16个INT32、2个Tensor、16384个32位寄存器
1个RT用于Ray Tracing

CUDA程序设计

2006年NVIDIA推出CUDA编程，使用SIMT（单指令多线程）模式。

CPU+GPU异构计算将可用于计算的硬件分为两部分：

  1. Host（主机）：CPU。负责控制、指挥GPU工作。
  2. Device：GPU。负责协处理。

GPU异构计算模型

GPU计算过程：

  1. 将数据从CPU内存拷贝到GPU内存（显存）
  2. 加载GPU程序并执行，使用显存内数据以提升性能。
  3. 将结果数据从GPU内存拷贝回CPU内存。

线程创建与管理 

1
2
3
4
5
6
7
8

// kernel部分
__global__void mykernel(void) {}

// main部分
int main(int argc, char** argv) {
    kernelName<< <gridSize, blockSize> >> (input_params);
    return 0;
}

__global__表示此函数在device上执行，自host调用。mykernel()在此没有进行任何计算。

示例：两个整数求和

kernel部分：

1
2
3

__global__ void add(int* a, int* b, int* c) {
    *c = *a + *b;
}
  1. add()在device运行，所以a，b和c需要指向device内存。
  2. 需要在GPU上分配内存空间。

main()部分：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21

int main(void) {
    int a,b,c;
    int *d_a, *d_b, *d_c;
    // 分配内存
    cudaMalloc((void**) &d_a, size);
    cudaMalloc((void**) &d_b, size);
    cudaMalloc((void**) &d_c, size);
    // 设置输入数据
    a = 2;
    b = 7;
    // 拷贝数据到GPU显存
    cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
    // 执行kernel
    add<< <1,1> >>(d_a, d_b, d_c);
    // 数据拷贝回host内存
    cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);
    // 释放内存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

GPU的优势是大规模并行：

1

add << <N, 1> >>();

这样add()不再只执行一次，而是并行执行了N次。

GPU向量加：并行执行add()可以实现向量加法。定义add()的每个调用为一个block，一组block定义为grid，各个调用通过blockIdx.x标识。

1
2
3

__global__ void add(int* a, int* b, int* c) {
    c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

使用blockIdx.x作为数组的索引，每个block处理各自的数据。在device中，各个block可以并行执行。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

#define N 512
int main(void) {
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = N*sizeof(int);

    cudaMalloc((void**) &d_a, size);
    cudaMalloc((void**) &d_b, size);
    cudaMalloc((void**) &d_c, size);

    a = (int*) malloc (size); random_ints(a, N);
    b = (int*) malloc (size); random_ints(b, N);
    c = (int*) malloc (size);
    
    cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);

    add<< <N,1> >>(d_a, d_b, d_c);

    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

    free(a); free(b); free(c);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

CUDA线程

每个block可以拆分为单独并行的threads。将add()改为使用并行threads，不再使用并行blocks：

1
2
3

__global__ void add(int* a, int* b, int* c) {
    c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}

main()中只需修改：

1
2

- add<< <N,1> >>(d_a, d_b, d_c);
+ add<< <1,N> >>(d_a, d_b, d_c);

Blocks+Threads

同时使用blocks和threads，每个线程需要全局唯一标识

1

int index = threadIdx.x + blockIdx.x * M;

组合使用时的add()：

1
2
3
4

__global__ void add(int* a, int* b, int* c) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    c[index] = a[index] + b[index];
}

修改后的main()：

1
2
3
4

+ #define THREADS_PER_BLOCK 512

- add<< <N,1> >>(d_a, d_b, d_c);
+ add<< <N/THREADS_PER_BLOCK, THREADS_PER_BLOCK> >>(d_a, d_b, d_c);

向量的长度并不能保证恰好是blockDim.x的整数倍，为了避免数组访问越界，可修改为：

1
2
3
4
5
6

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

此时main()函数修改为：

1

add<< <(N+M-1)/M, M> >>(d_a, d_b, d_c);

线程层次与分组

每个kernel对应一个grid，每个grid对应一组block共享全局内存，每个block包含多个线程。

线程层次

线程分组中，对应的维度必须能被整除，分组的大小不能超过1024

变量存储访问

函数与变量类型

函数类型 特点
__device__ 在GPU中执行，由CPU调用
__global__ 在GPU中执行，由CPU调用
__host__ 在CPU中执行/调用，是默认函数类型
变量类型 特点
__device__ global memory space，grid中所有线程均可访问
__constant__ constant memory space，grid中所有线程均可访问
__shared__ space of a thread block，只能由block中的线程访问
内置向量类型 char，short，int，long，long long，float，double：内置1，2，3，4维向量，结构体数据中含有x,y,z,w；
dim用于表示维度

存储器层次

存储器 特点
Register 线程私有
Local memory 线程私有，KB级
Shared memory 速度快，10KB级；block线程中可访问
Global memory 容量大，速度慢，GB级
Texture memory GPU只读，KB级
Constant memory GPU只读，KB级

存储器层次

内存与显存分配

可分页内存：使用标准C/C++函数操作malloc()/new()。

分页锁定内存：

1
2
3
4
5
6

cudaMallocHost(void** ptr, size_t size)

cudaHostAlloc(void** pHost, size_t size, unsigned int flags)
// cudaHostAllocPortable：多个CPU线程都可以访问
// cudaHostAllocWriteCombined：数据传输快，但只允许host写
// cudaHostAllocMapped：CPU/GPU都可访问

可分页内存转变为分页锁定内存：

1

cudaHostRegister(void** ptr, size_t size, unsigend int flags)

显存分配：

1
2
3
4
5
6

// 一维数组，数据连续存储
cudaMalloc()
// 二维数组，自动对齐，可能会浪费部分内存空间
cudaMallocPitch()
// 三维数组
cudaMalloc3D()

内存映射

ZeroCopy

  1. host内存直接映射到GPU内存空间。
  2. 调用cudaHostAlloc()分配时传入cudaHostAllocMapped标签。
  3. 使用cudaHostRegister()注册时使用cudaHostRegisterMapped标签。
  4. device可直接访问分页锁定内存，硬件在kernel访问显存时通过PCL总线传递，由硬件自动处理数据传输和计算的异步执行。

适用于计算密集型程序：数据访问量小、计算与数据传输重叠、不需要显式内存管理。

统一寻址

host与device使用同一个指针，在host内存和device显存之间根据访问需要自动迁移数据，同时保证host和device均可访问。应用程序并不知道访问时数据所在位置，需要进行显式同步。

统一寻址使得编程简化，数据访问性能提升。

共享与同步

Shared memory中的数据可以由同一个block中的线程访问。在block内部，threads可通过shared memory共享数据。

同步函数：

1
2
3
4
5

void __syncthreads();
int __syncthreads_count(int predicate);
int __syncthreads_and(int predicate);
int __syncthreads_or(int predicate);
void __syncwarp(unsigned mask=0xffffffff);

void __syncthreads();可以同步同一个block之内的全部线程，所有线程都必须运行到barrier。

示例：1D Stencil

对一维数组进行卷积处理：每个输出元素是输入元素一定“半径”内的数据和。

例如，若半径为3，则每个输出元素是7个输入元素的和。

在block内，每个thread处理一个输出元素，则每个block分配blockDim.x个元素。同时，输入数据会被读取多次。

性能优化

维度 优化方向
算法选择 算法并行性是首要因素，时间复杂度低未必合适
并行度 GPU环境中线程切换效率很高，应使用尽可能的线程数
数据传输 次数尽可能少，数据传输量尽可能少，异步传输
© 2025 - 2026 RunicDolphin806
网站总访客数：Loading

使用 Hugo 构建
主题 Stack 由 Jimmy 设计