Nvidia CUDA初级教程7 CUDA编程二

视频：https://www.bilibili.com/video/BV1kx411m7Fk?p=8
讲师：周斌

本节内容：

• 内置类型和函数 Built-ins and functions
• 线程同步 Synchronizing
• 线程调度 Scheduling threads
• 存储模型 Memory model
• 重访 Matrix multiply
• 原子函数 Atomic functions

函数的声明

	执行	调用
__global__ void KernelFunc()	device	host
__device__ float DeviceFunc()	device	device
__host__ float Host	host	host

• __device__ 和 __host__ 可以同时修饰一个函数
• __global__ 的返回值必须是 void
• __device__ 曾经默认内联，现在有些变化
• 对于 global 和 device：
  • 尽量少用递归（不鼓励）
  • 不要用静态变量
  • 少用 malloc（现在允许但不鼓励）
  • 小心通过指针实现函数调用

向量数据类型

• char[1-4], uchar[1-4]
• short[1-4], ushort[1-4]
• int[1-4], uint[1-4]
• long[1-4], ulong[1-4]
• longlong[1-4], ulonglong[1-4]
• float[1-4]
• double1, double2

---

• 同时适用于 host 和 device 代码

• 通过函数 make_<type name> 构造

  int2 i2 = make_int2(1, 2);
  float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f);
  

• 通过 .x, .y, .z, ,w 访问

  int x = i2.x;
  int y = i2.y;
  

数学函数

• 部分函数列表

  • sqrt, rsqrt
  • exp, log
  • sin, cos, tan, sincos
  • asin, acos, atan2
  • trunc, ceil, floor

• Intrinsic function 内建函数

  • 仅面向 device 设备端

  • 更快，但是精度降低

  • 以 __ 为前缀，例如：

    __exp, __log, __sin, __pow, …

线程层次回顾

线程同步

• 块内的线程可以同步
  • 调用 __syncthreads 创建一个 barrier
  • 每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续指令

Mds[i] = Md[j];
__syncthreads();
func(Mds[i], Mds[i+1]);

• 要求线程的执行时间尽量接近

• 只在一个块内进行同步

• 线程同步可能会导致死锁

  if (someFunc()) {__syncthreads();
  }
  else {__syncthreads();	// 注意这两个barrier不是同一个
  }
  

  线程调度

• 多线程切换，达到延迟掩藏的效果。

• warp - 块内的一组线程

  • 运行于同一个SM

  • 线程调度的基本单位

  • 一个warp内是天然同步的（硬件保证）

  • warp 调度是零开销的

  • 一个SM上某个时刻只会有一个warp再执行

  • threadIdx 值连续

  • 一个实现细节 - 理论上

    • warpSize

  • warp内执行不同的分支的情况：divergent warp

    其他的分支需要等待该分支进行

举例：

• 如果一个 SM 分配了 3 个 block，其中每个 block 含 256 个线程，总共有多少个 warp（warp大小为32）？

  一个 block 内有 256/32 = 8个 warp，一个 SM 内共有 8 * 3 = 24个

• GT200 的一个 SM 最多可以驻扎 1024 个线程，那相当于多少个 warp？

  1024 / 32 = 32

每个 warp 含 32 个小牛橙，但是每个 SM 只有 8 个 SPs，如何分配？

当一个 SM 调度一个 warp 时：

• 指令已经预备
• 在第一个周期 8 个线程进入 SPs
• 在第二三四个周期也分别进入 8 个线程
• 因此，分发一个 warp 需要4个周期

另一个问题：

一个 kernel 包含：

• 1 次对 global memory 的读操作（200 cycles）
• 4 次独立的 multiples/adds 操作

需要多少个 warp 才可以隐藏内存延迟？

解：

每个 warp 含 4 个 multiple/adds 操作需要16 个周期，我们需要覆盖 200 个周期，200 / 16 = 12.5 ，ceil(12.5)=13，需要 13 个 warps。

内存模型回顾

…

内存模型

寄存器 registers - G80

• 每个 SM，多达 768 个 threads，8K 个寄存器，即每个线程可以分到 8K / 768 = 10 个寄存器

• 超出限制后，线程数将因为 block 的减少而减少

  因为同一个 block 必须在同一个 SM 内

  例如，每个线程用到 11 个寄存器，而由于每个 block 含 256 个线程，则：

  • 一个 SM 可以驻扎多少个线程？512（两个block）
  • 一个 SM 可以驻扎多少个 warp？ 16
  • warp 数少了意味着什么？效率降低

local memory

• 存储于 global memory，作用域是每个 thread
• 用于存储自动变量数组，通过常量索引访问

shared memory

• 每个块
• 快速，片上，可读写
• 全速随机访问

global memory

• 长延迟（100个周期）
• 片外，可读写
• 随机访问影响性能
• host 主机端可读写

constant memory

• 短延时，高带宽，当所有线程访问同一位置时只读
• 存储于 global memory，但是有缓存
• host 主机端可读写
• 容量：64KB

变量声明

变量声明	存储器	作用域	生命期
必须是单独的自动变量而不能是数组	register	thread	kernel
自动变量数组	local	thread	kernel
__shared__ int sharedVar;	shared	block	kernel
__device__ int globalVar;	global	grid	application
__constant__ int constantVar	constant	grid	application

关于 global and constant 变量

• Host 可以通过以下函数访问：
  • cudaGetSymbolAddress()
  • cudaGetSymbolSize()
  • cudaMemcpyToSymbol()
  • cudaMemcpyFromSymbol()
• constants 变量必须在函数外声明