GPU概述
GPU是异构、众核处理器，针对吞吐优化。
GPU架构是围绕一个流式多处理器(SM)的可拓展阵列搭建的。可以通过复制这种架构的构建块来实现GPU的硬件并行。
GPU中的每一个SM都能支持数百个进程并行执行，每个GPU通常有多个SM，所以在一个GPU上可以并行执行数千个进程。
当启动一个内核网格时，它的线程块被分布在了可用的SM上执行。线程块一旦被调度到一个SM上，其中的线程只会在那个指定的SM上 并发执行。多个线程块可能会被分配到同一个SM上。同一线程中的指令利用指令级并行性进行流水线化。
Fermi架构
以NVIDIA GeForce GTX 480为例：
- 480个流处理器(stream processors, 也即是CUDA核心)
- SIMT execution(单指令多线程，single instruction multi thread)
- 15个核心
- 每个核心拥有2组，每组16个SIMD Functional units（每个组简称SM，每个SM含有32个cuda core）
Kepler架构
相比Fermi架构更复杂，但基本概念相似。
GPU内存和线程关系
- Local Memory 线程私有
- Shared Memory 每个线程块共有
- Global Memory 每个SM中的kernel共有
CPU-GPU交互
- 拥有各自独立的物理内存空间
- 通过PCIE总线互连(8GB/s~16GB/s)
- 交互开销大
CUDA函数声明
| 声明 | 执行位置 | 调用位置 |
|---|---|---|
| __device__ | device | device |
| __global__ | device | host |
| __host__ | host | host |
| __device__and__host__ | deviceAndHost | deviceAndHost |
值得注意的是_global__的返回值必须是void。
- 在设备上尽量少用递归
- 不要使用静态变量
- 少用malloc
- 小心通过指针实现的函数调用
这是数据并行处理函数。通过调用kernel函数在设备端创建轻量级线程。线程由硬件负责创建并调度。
内存传输
gpu内存访问关系见上。
cudaMalloc()分配的是global memory。
线程同步
线程同步只发生在线程块内，全局的线程没有同步机制。
线程块内的所有线程同步。可以保证数据的一致性。
调用__syncthread()创建一个同步点，每个线程块内的所有线程都执行到同步点之后才能够继续向下执行。
mds[i] = md[j];
__syncthread();
func(mds[i], mds[i+1]);
如在上例中，如果没有线程同步，可能会造成数据不可重复度。
要求
- 线程的执行时间尽量接近。否则会造成时间大量浪费
- 线程同步仅在线程块内进行。因为更大范围的线程同步需要繁杂的同步机制，降低效能。
潜在问题
- 同步一定会使得线程停止。
- 可能造成死锁。如下例：
if(someFunc())
{
__syncthread();
}
else
{
__syncthread();
}