cuda1

CUDA 介绍

CPU ：面向延迟设计

image-20210607010547778(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607010547778.png

• ALU
  • 减少操作延迟
• Cache
  • 将长延迟内存访问转换为短延迟缓存访问
• 控制模块 - 分支预测以减少分支延迟 - 数据转发以减少数据延迟

GPU：面向吞吐量设计

image-20210607010747373(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607010747373.png

• 小缓存
  • 提高内存吞吐量
• 简单控制
  • 无分支预测
  • 无数据转发
• 高效的 ALU
  • 许多、长延迟但大量流水线以实现高吞吐量
• 需要大量线程来容忍延迟
  • 线程逻辑
  • 线程状态

并行代码，GPU 比 CPU 快 10 倍以上，串行代码，CPU 比 GPU 快 10 倍以上

---

CUDA 编程模型

编程模型是底层计算机系统的抽象，它允许表达算法和数据结构。
语言和 API 提供了这些抽象的实现，并允许将算法和数据结构付诸实践——编程模型的存在独立于编程语言和支持 API 的选择。

一些设计目标

• 扩展到 100 个内核、1000 个并行线程

• 让程序员专注于并行算法
  不是并行编程语言的机制。

• 启用异构系统（即 CPU+GPU）
  CPU 和 GPU 是具有独立 DRAM 的独立设备

关键并行抽象

• 并发线程的层次结构
• 轻量级同步原语
• 协作线程的共享内存模型

线程层次结构

• 线程 thread - 由 CUDA 运行时分发
  （由 threadIdx 标识）
• Warp – 最多 32 个线程的调度单元
• 块 block – 用户定义的 1 到 512 个线程组。
  （由 blockIdx 标识）
• 网格 grid – 一组一个或多个块。 为每个 CUDA 核函数创建一个网格

cuda 内存层次结构

• 寄存器

  每个线程内存用于自动变量和寄存器溢出。

• 共享内存

  每块低延迟内存，允许块内数据共享和同步。 线程可以通过这块内存安全地共享数据，并且可以通过 _ _syncthreads() 进行屏障同步

• 全局内存

  可以在块或网格之间共享的设备级内存

硬件

Tesla 架构的主要组件是：
流式多处理器（8800 有 16 个）
标量处理器
内存层次结构
互联网络
主机接口

流式多处理器 Streaming Multiprocessor (SM)

image-20210607012204300(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607012204300.png

• 每个 SM 有 8 个标量处理器 (SP)Each SM has 8 Scalar Processors (SP)
• IEEE 754 32 位浮点支持（不完全支持）- IEEE 754 32-bit floating point support (incomplete support)
• 每个 SP 是一个 1.35 GHz 处理器（32 GFLOPS 峰值）- Each SP is a 1.35 GHz processor (32 GFLOPS peak)
• 支持 32 位和 64 位整数- Supports 32 and 64 bit integers
• 8,192 个动态分区的 32 位寄存器- 8,192 dynamically partitioned 32-bit registers
• 硬件支持 768 个线程（32 个线程的 24 个 SIMT 经线）- Supports 768 threads in hardware (24 SIMT warps of 32 threads)
• 在硬件中完成的线程调度- Thread scheduling done in hardware
• 16KB 低延迟共享内存- 16KB of low-latency shared memory
• 2 个特殊函数单元（平方根倒数、三角函数等）- 2 Special Function Units (reciprocal square root, trig functions, etc)

数据并行 - 向量加法示例

image-20210607155913228(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607155913228.png

// Compute vector sum C = A + B
void vecAdd(float *h_A, float *h_B, float *h_C, int n)
{
    int i;
    for (i = 0; i<n; i++) h_C[i] = h_A[i] + h_B[i];
}

int main()
{
    // Memory allocation for h_A, h_B, and h_C
	  // I/O to read h_A and h_B, N elements
	  …
    vecAdd(h_A, h_B, h_C, N);
}

cudaMalloc()

• 在设备全局内存中分配一个对象
• 两个参数
  • 指向已分配对象的指针的地址
  • 已分配对象的大小（以字节为单位）

cudaFree()

• 从设备全局内存中释放对象
• 一个参数
  • 指向释放对象的指针

cudaMemcpy()

• 内存数据传输
• 需要四个参数
  • 指向目的地的指针
  • 指向源的指针
  • 复制的字节数
  • 转移类型/方向

向量加法主机代码

void vecAdd(float *h_A, float *h_B, float *h_C, int n)
{
    int size = n * sizeof(float); float *d_A, *d_B, *d_C;

    cudaMalloc((void **) &d_A, size);
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
     cudaMalloc((void **) &d_B, size);
     cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
     cudaMalloc((void **) &d_C, size);

     // Kernel invocation code – to be shown later

     cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
     cudaFree(d_A); cudaFree(d_B); cudaFree (d_C);
}

---

cuda 执行模式

异构主机（CPU）+设备（GPU）应用 C 程序

• 主机 C 代码中的串行部分
• 设备 SPMD 内核代码中的并行部分

image-20210607160519981(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607160519981.png

ISA 级别的程序

• 程序是存储在内存中的一组指令，可由硬件读取、解释和执行。
  • CPU 和 GPU 都是基于（不同的）指令集设计的
• 程序指令对存储在存储器和/或寄存器中的数据进行操作。

作为冯诺依曼处理器的线程

线程是“虚拟化的”或“抽象的”

image-20210607160745684(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607160745684.png

并行线程数组

• CUDA 内核由线程网格（数组）执行
  • 网格中的所有线程都运行相同的内核代码（SPMD，Single Program Multiple Data）
  • 每个线程都有用于计算内存地址和做出控制决策的索引

image-20210607160922707(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607160922707.png

线程块：可扩展的合作

image-20210607160955464(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607160955464.png

• 将线程数组分成多个块
  • 块内的线程通过共享内存、原子操作和屏障同步进行协作
  • 不同块中的线程不交互

blockIdx and threadIdx

• 每个线程使用索引来决定要处理的数据

  • blockIdx：1D、2D 或 3D (CUDA 4.0)
  • threadIdx：1D、2D 或 3D

• 处理多维数据时简化内存寻址

  • 图像处理
  • 求解体积上的偏微分方程
  • …

image-20210607161116355(C:\Users\Aerialith\AppData\Roaming\Typora\typora-user-images\image-20210607161116355.png

NVCC 编译器

• NVIDIA 提供了一个 CUDA-C 编译器
  • nvcc
• NVCC 编译设备代码，然后将代码转发到主机编译器（例如 g++）
• 可用于编译和链接 host only 应用程序