CUDA Tutorial

发表于 更新于

内容

  1. CPU体系架构概述
  2. 并行程序设计概述
  3. CUDA开发环境搭建和工具配置
  4. GPU体系架构概述
  5. GPU编程模型
  6. CUDA编程（1）
  7. CUDA编程（2）
  8. CUDA编程（3）
  9. CUDA程序分析和调试工具
  10. CUDA程序基本优化
  11. CUDA程序深入优化
  12. 最新NVIDA GPU 和 CUDA特性

1.CPU体系概述

  1. 桌面级应用以访存 分支操作 数据搬来搬去为主，数值计算占比很低

  2. 取指 译码 执行 访存，流水线 pipeline 指令级并行 减小时钟周期但是增加了延迟和芯片面积。带来的问题：

    • 具有依赖关系的指令 执行顺序
    • 分支怎么处理
    • 流水线长度
      旁路 Bypassing
      停滞 Stalls
      分支 Branches
      分支预测 Branches Prediction
    • +现代预测器准确度>90%
    • -面积增加 延迟增加
      分支断定 Predication GPU中使用了分支断定

  3. 提升IPC(instructions per cycle) 超标量(Superscalar)

    • 寄存器重命名(RegisterRenaming)
    • 乱序执行(Out-of-Order Execution) 重排指令获得最大吞吐率
    • +IPC接近理想状态
    • -面积增加 功耗增加

  4. CPU内部的并行性

    • 指令级并行 Instruction-level Parallelism (ILP)
      • 超标量Superscalar
      • 乱序执行Out-of-Order
    • 数据级并行 Data-level Parallelism (DLP)
      • 矢量计算Vectors
    • 线程级并行 Thread-level Parallelism (TLP)
      • 同步多线程 Simultaneous Mulitithreading (SMT)
      • 多核 Multicore
      • 锁、一致性和同一性 Locks,Coherence and Consisitency
        • 问题：多线程读写同一块数据 解决办法：加锁
        • 问题：谁的数据是正确的？ Coherence 解决办法：缓存一致性协议
        • 问题：什么样的数据是正确的？ Consistency 解决办法：存储器同一性模型

  5. 能量墙/存储墙

  • 结论
  1. CPU为串行程序优化
    • Piplines, Branch Prediction, Superscalar, Out-of-Order(OoO)
    • Reduce execution time with high clock speeds and high utilization
  2. 缓慢的内存带宽(存储器带宽)将会是大问题
  3. 并行处理是方向

2.并行程序设计概述

  • 概念和名词
    • Flynn 矩阵
      • SISD: Single Instruction, Single Data
      • SIMD: Single Instruction, Multiple Data
      • MISD: Multiple Instruction, Single Data
      • MIMD: Multiple Instruction, Multiple Data
    • Task (任务)
    • Parallel Task (并行任务)
    • Serial Execution (串行执行)
    • Parallel Execution (并行执行)
    • Shared Memory (共享存储)
    • Distributed Memory (分布式存储)
    • Communications (通信)
    • Synchronization (同步)
    • Granularity (粒度)
    • Observed Speedup (加速比)
    • Parallel Overhead (并行开销)
    • Scalability (可扩展性)
  • 并行编程模型
    • 共享存储模型 Shared Memory Model
    • 线程模型 Threads Model
    • 消息传递模型 Message Passing Model
    • 数据并行模型 Data Parallel Model
  • 设计并行处理程序和系统
    • 自动和手动并行
    • 理解问题和程序
    • 分块分割 数据分块，任务分割
    • 通信 可扩展性重要影响因素
    • 同步
    • 数据依赖
    • 负载均衡
    • 粒度
    • I/O
    • 成本
    • 性能分析和优化
  • Amdahl’s Law
    • 程序可能的加速比取决于可以被并行化的部分,并行化的可扩展性有极限 取决于可并行部分的比例

3.CUDA开发环境搭建

  • windows cuda zone
  • linux

4.GPU体系架构概述

  • 为什么需要GPU(Graphic Processing Unit)
    • GPU 是异构 众核 处理器，针对吞吐优化
      • 高效的GPU任务具备的条件
        • 具有成千上万的独立工作
          • 尽量利用大量的ALU单元
          • 大量的片元切换掩藏延迟
        • 可以共享指令流
          • 适用于SIMD处理
        • 最好是计算密集的任务
          • 通信和计算开销比例合适
          • 不要受制于访存带宽
  • 三种方法提升GPU的处理速度
    • 1.Use many “slimmed down cores” to run in parallel
    • 2.Pack cores full of ALUs(by sharing instruction stream across groups of fragments)
      • Option 1: Explicit SIMD vector instructions
      • Option 2: Implicit sharing managed by hardware
    • 3.Avoid latency stalls by interleaving execution of many groups of fragments
  • 实际GPU设计举例
    • NVIDIA GTX 480：Fermi
    • NVIDIA GTX 680: Kepler
  • GPU的存储器设计

5.GPU编程模型

  • 内容

    • CPU和GPU互动模式
    • GPU线程组织模型（不停强化）
    • GPU存储模型
    • 基本的编程问题

  • CPU-GPU交互

    • 各自的物理内存空间
    • 通过PCIE总线互连（8GB/s~16GB/s）
    • 交互开销较大

  • 线程组织架构说明

    • 一个kernel具有大量线程
    • 线程被划分成线程块’Blocks’
      • 一个block内部的线程可以共享’Shared Memory’
      • 可以同步 ‘_syhcthreads()’
    • kernel启动一个’grid’,包含若干线程块
      • 用户设定
    • 线程和线程块具有唯一标识

  • 编程模型

    • 常规意义的GPU用于处理图形图像
    • 操作用于像素，每个像素的操作都类似
    • 可以应用SIMD(single instruction multiple data)
    • Single Instruction Multiple Thread(SIMT)
      • GPU版本的SIMD
      • 大量线程模型获得高度并行
      • 线程切换获得延迟掩藏
      • 多个线程执行相同指令流
      • GPU上大量线程承载和调度

  • CUDA编程模式： Extended C

    • Declspecs (Dclaration Specifier) 声明规范
      • global, device, shared, local, constant
    • 关键词
      • threadIdx, blockIdx
    • Intrinsics
      • __syncthreads
    • 运行期API
      • Memory, symbol, execution, management
    • 函数调用

6.CUDA编程(1)

  • GPU架构概览

    • GPU特别适用于
      • 密集计算，高度可并行计算
      • 图形学
    • 晶体管主要用于：
      • 执行计算
      • 而不是缓存数据，控制指令流

  • GPU计算的历史

    • 2001/2002 研究人员把GPU当作数据并行协处理器
      • GPGPU这个新领域从此诞生
    • 2007 NVIDIA发布CUDA
      • CUDA 全称Compute Uniform Device Architecture 统一计算设备架构
      • GPGPU 发展成GPU Computing
    • 2008 Khronos 发布 OpenCL 规范

  • CUDA的一些信息

    • 层次化线程集合 A hierarchy of thread groups
    • 共享存储 Shared memories
    • 同步 Barrier synchronization

  • CUDA术语

    • Host - 即主机端 通常指CPU
      • 采用ANSI标准C语言编程
    • Device - 即设备端 通常指GPU(数据可并行)
      • 采用ANSI标准C的扩展语言编程
    • Host和Device 拥有各自的存储器
    • CUDA编程
      • 包括主机端和设备端两部分代码
    • Kernel - 数据并行处理函数，类似于OpenCL的shader
    • 通过调用kernel函数在设备端创建轻量级线程
      • 线程由硬件负责创建并调度
    • CUDA核函数(kernels)
      • 在N个不同的CUDA线程上并行执行
    • 线程层次 Thread Hierarchies
      • Grid - 一维或多维线程块(block)
        • 一维或二维
    • Block - 一组线程
      • 一维，二维或三维
        • 例如索引数组，矩阵，体
    • 一个Grid内每个Block的线程数是一样的
    • block内部的每个线程可以
      • 同步 synchronize
      • 访问共享存储器 shared memory
    • 线程块之间彼此独立执行
      • 任意顺序：并行或串行
      • 被任意数量的处理器以任意顺序调度
      • 处理器的数量具有可扩展性
    • Host 可以从device往返传输数据
      • global memory全局存储器
        • cudaMalloc() 在设备端分配global memory
        • cudaFree() 释放存储空间
        • cudaMemcpy() 内存传输
          • Host to host
          • Host to device cudaMemcpyHostToDevice
          • Device to host cudaMemcpyDeviceToHost
          • Device to device
      • Constant memory常量存储器

7.CUDA编程(2)

  • 目录

    • 内置类型和函数 Built-ins and functions

      • 函数的声明
        • global void KernelFunc(),返回值必须是void. Executed on the:device Only callable from the:host
        • device float DeviceFunc(),曾经默认内联，现在有些变化. Executed on the:device Only callable from the:device
        • host float HostFunc() Executed on the:host Only callable from the:host
      • Global和device函数
        • 尽量少用递归（不鼓励）
        • 不要用静态变量
        • 少用malloc（现在允许但不鼓励）
        • 小心通过指针实现的函数调用
      • 向量数据类型
        • type name
          • 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>构造
          • 通过.x, .y, .z, .w 访问
      • 数学函数
        • Intrinsic function 内建函数
          • 仅面向 Device设备端
          • 更快但精度降低
          • 以__为前缀，例如：__exp, __log,__pow,…

    • 线程同步 Synchronizing threads

      • 块内线程可以同步
        • 调用__syncthreads 创建一个barrier栅栏
        • 每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续指令
      • 要求线程执行时间尽量接近 -> 防止块内大部分 线程等待时间超长，降低效率
      • 为什么只在一个块内同步 -> 全局同步开销大
      • __syncthreads()会导致暂停 死锁

    • 线程调度 Scheduling threads

      • 术语 Streaming Processor(SP) Streaming Multi-Processor(SM)
      • G80架构
        • 16个SMs
        • 每个含8个SPs,总共128个SPs
        • 每个SM驻扎多达768个线程
        • 总共同时执行12,288个线程
      • GT200架构
        • 30个SMs
        • 每个含8个SPs,总共含240个SPs
        • 每个SM驻扎多达8个block,或1024个线程
        • 同时执行，多达240个block，或30,720个线程
      • Warp -块内的一组线程
        • G80/GT200 -32个线程
        • 运行于同一个SM
        • 线程调度的基本单位
        • threadIdx值连续
        • 一个实现细节 -理论上从硬件上保证每个warp内的线程执行到相同位置
      • SM implements zero-overhead warp scheduling
        • At any time,only one of the warps is executed by SM
        • Warps whose next instruction has its operands ready for consumption are eligible for execution
        • All threads in a warp execute the same instruction when selected

    • 存储模型 Memory model

      • Device code can:

      • R/W per-thread register

      • R/W per-thread local memory

      • R/W per-block shared memory

      • R/W per-grid global memory

      • Read Only per-grid constant memory

      • Host code can

      • R/W per-grid global and constant memory

      • 寄存器Registers

        • 每个线程专用
        • 快速，片上，可读写

      • 局部存储器Local Memory

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

      • 共享存储器Shared Memory

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

      • 全局存储器Global Memory

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

      • 常量存储器Constant Memory

        • 短延时，高带宽，当所有线程访问同一位置时只读
        • 存储区global memory 但是有缓存
        • Host主机端可读写
      • 变量声明 存储器 作用域 生命期
        必须是单独的自动变量而不能是数组 register thread kernel
        自动变量数组 local thread kernel
        __shared__int sharedVar; shared block kernel
        __device__int globalVar; global grid application
        __constant__int constantVar; constant grid application

    • 重访 Matrix multiply

    • 原子函数 Atomic functions

8.CUDA编程(3)