CUDA

合集 - 并行(7)

1.nku并行01-132.并行硬件和并行软件01-133.SIMD01-134.Pthread01-135.OpenMP01-136.MPI01-13

7.CUDA01-13

收起

GPU/CUDA简介#

GPU(Graphics Processing Units)：图形处理单元

GPU——CPU的协处理器（eg：天河一号）

CUDA（Compute Unified Device Architecture）：统一计算 设备架构——建立在GPU基础之上的通用计算开发平台 （Only NVIDIA ！）

CPU vs GPU#

CPU——优化延迟的设计思路

• 大容量cache、复杂的控制逻辑、强大的计算单元
• 低计算密度、低延迟容忍度
  GPU（众核）——优化吞吐率的设计思路
• 小cache、简单的控制、能耗高效的计算单元、高并发
• 高计算密度、高吞吐率、高延迟容忍度

CUDA

• 使用C语言编写GPU 程序，语言层面只 有很小的扩展
• 易于编程
• 属于SIMD模式
• 标准的SPMD模式 （单程序多数据流，所有线程执行相同的Kernel代码）
• 不再需要图形API

CUDA程序设计模型#

• host： 指代CPU端的代码
• device：指代GPU端的代码
• kernel：从host调用，在device端运行的函数

Serial Code (host)
Parallel Kernel (device)
KernelA<<<nBlk,nTid >>>(args);
Serial Code (host)
Parallel Kernel (device)
KernelB<<<nBlk,nTid >>>(args);

• Serial Code (host)：这是主机端的代码，它通常负责管理整个计算任务，包括数据的准备、内存分配等操作。这个部分的代码是在 CPU 上运行的。
• Parallel Kernel (device)：这是在设备端运行的并行内核代码。通常，CUDA 内核（Kernel）是并行执行的，并且是由多个线程在 GPU 上执行的。每个内核函数都有一个特定的线程布局，由 <<<nBlk, nTid>>> 这样的语法指定，其中：
  • nBlk 表示 线程块（block）的数量，
  • nTid 表示每个线程块中的 线程数（threads per block）。
• KernelA<<<nBlk, nTid>>>(args);：这是调用设备端内核 KernelA 的代码。KernelA 是你在 GPU 上运行的函数。<<<nBlk, nTid>>> 语法指定了内核的执行配置，控制有多少线程块（blocks）和每个线程块中有多少个线程（threads）。args 是传递给内核的参数。
• Serial Code (host)：在执行内核后，主机端代码通常会等待内核完成计算，可能还会进行一些后处理操作。
• Parallel Kernel (device)：这是在设备端执行的另一个内核，类似于 KernelA。它在 GPU 上并行地执行。KernelB<<<nBlk, nTid>>>(args); 是调用内核 KernelB，与 KernelA 类似。

#include <iostream> 
#include <cuda_runtime.h> 
using namespace std;

__global__ void demo(void) {

}

int main() {
    demo<<<1,1>>>();
    cout << "Hello, world!"<< endl;
    return 0; 
}

一个kernel对应一个Grid
Grid是三层线程结构

• 线程格Grid
  • 一组线程块Block
    • 一些线程Thread
      • 连续32个线程组成一个warp，在SM（流多处理器）上运行
      • Warp是GPU执行程序时的调度单位，同一个Warp里的线程执行相同的指令
  • Block里支持线程间同步、共享内存
  • Block内一个块至多512个线程、或 1024个线程
  • Block间不能协作
• 同一个Grid上的线程共享相同的全局内存空间

线程层次示例：重点

1 个 Grid 包含 9 个 Blocks ，每个 Block 包含了 16 个 Threads
gridDim：gridDim.x, gridDim.y , gridDim.z 分别表示 grid 各个维度的大小，这里都是 3
blockDim：blockDim.x,blockDim.y , blockDim.z分别表示 block 中各个维度的大小，这里都为 4
blockIdx：blockIdx.x , blockIdx.y , blockIdx.z 分别表示当前 block 所处的Grid的坐标位置
threadIdx threadIdx.x , threadIdx.y , threadIdx.z分别表示当前线程所处的Block的坐标位置
总线程数 :N = gridDim.x * gridDim.y * gridDim.z *blockDim.x * blockDim.y * blockDim.z
当前线程在该 Block 位置 threadId = threadIdx.x +threadIdx.y*blockDim.x + threadIdx.z * blockDim.x * blockDim.y

从0开始编号

Thread工作模式
所有线程执行相同的 kernel 代码 (SPMD)
每个线程使用自己的编号计算不同数据 内存地址 以及执行不同分支

Block 内线程协同计算 共享内存 、 原子操作 、同步机制
Block 间不能协作

连续32个线程组成一个warp，在SM上运行
❑ SM：Stream Multiprocessor,流多处理器
❑ Warp是GPU执行程序时的调度单位，同一个Warp里的线程执行相同的指令
❑ SM动态调度warp

CUDA内存层次简介#

三层内存：

• Global Memory

  • Shared Memory
    • Register

• Device代码可以

  • 读/写每线程独占的寄存器
  • 读/写每kernel共享的全局内存

• Host代码可以

  • 在CPU主存和Device全局内存间传输数据

读写独占的寄存器（~1周期）
读写block内共享的共享内存 （~5周期）
读写全体线程共享的全局内存 （~500周期）
只读全体线程共享的 Host 常量内存（~5周期， 有cache的情况）

CUDA函数声明#

	执行位置	调用位置
__device__ float DeviceFunc()	device	device
__global__ void KernelFunc()	device	host
__host__ float HostFunc()	host	host

CUDA程序优化#

两周优化思路：
掩盖访存
共享内存