CUDA 编程学习-概念

SPMD 概念

一个Grid执行同一个kernel函数。

线程

  • threadIdx

可以获取多维线程id,也可以获取一维ID

1
2
3
4
5
6
7
8
9
10
// 三维
int tx = threadIdx.x;
int ty = threadIdx.y;
int tz = threadIdx.z

// 二维
int tx = threadIdx.x;
int ty = threadIdx.y;
// 一维
int threadID = threadIdx;
  • blockIdx

计算公式

1
threadID = blockIdx.x * blockDim.x + threadIdx
  • dim3 数据结构

struct 含有 x, y, z

1
2
3
4
5
6
7
8
gridDim.x-线程网络X维度上线程块的数量
gridDim.y-线程网络Y维度上线程块的数量
blockDim.x-一个线程块X维度上的线程数量
blockDim.y-一个线程块Y维度上的线程数量
blockIdx.x-线程网络X维度上的线程块索引
blockIdx.y-线程网络Y维度上的线程块索引
threadIdx.x-线程块X维度上的线程索引
threadIdx.y-线程块Y维度上的线程索引

函数调用

__device__ deviceFunc() 在设备上调用并执行

__global__ globalFunc() 在主机调用,在设备执行,多用于核函数

__host__ hostFunc() 在主机调用,在主机执行。

同步

__syncthreads()线程会停止,直到所有线程均到达该位置。

可以通过判断对不同的块或者线程进行同步。

线程调度

wrap 32个线程 为一个调度单位。

块划分为wrap

  • 第一个wrap包含线程ID 0 - 31
  • 第二个wrap包含线程ID 32 - 63

wrap的概念说明,做并行化算法时,线程分配的最小粒度要大于等于32,最好32的整数倍

存储器

  • 设备代码:

    1. 读写每个线程寄存器
    2. 读写每个线程局部存储器
    3. 读写每个块的共享存储器
    4. 读写网格的全局存储器
    5. 读写网格的常数存储器
  • 主机代码:

在每个网格的全局存储器和常数存储器间传输数据

变量类型与存储器

变量 存储器 作用域
自动变量(数组以外) 寄存器 thread
自动数组变量 局部存储器 thread
__device__,__shared__ 共享存储器 block
__device__ 全局存储器 grid
__device__, __constant__ 常数存储器 grid
  • __constant__ 或者 __device__ __constant__存储在常数存储区。
  • 如果一个变量前只有__device__ 则变量存放在全局存储区
  • __device__ __shared____shared__相同,存储在共享存储区。

提高并行效率: 减少全局存储器访问的次数与数据流量。