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 變量必須在函數外聲明