Cuda笔记【1】GPU计算DEMO

声明

• 本文是看小破站某cuda入门教程留下来的笔记，多上PPT上内容，夹杂一点自己的理解，和代码注释
• 教程地址:https://www.bilibili.com/video/av74148375
• git地址（PPT和源码）：https://github.com/huiscliu/tutorials
• 主要目的是为Gstreamer打点基础，不然基本抓瞎

介绍

什么是GPU计算

• CPU基本架构

• GPU架构

  • 核心数远远超过CPU，将核心分成小组SM，一个SM有多个SP
  • 计算的时候数据存在显存中，也叫全局内存

• NVIDIA公司发布CUDA，C是建立在NVIDIA CPUs上的一个通过并行计算，平台和变成模型，基于CUDA编程可以利用GPUs的并行计算引擎来更加高效地解决比较复杂的计算难题
• GPU并不是一个独立运行的计算平台，而需要与CPU协同工作，可以看成是CPU的协处理器，因此我们在说GPU并行计算时，其实是指CPU+GPU的易购计算架构
  • GPU和CPU通过PCIe总线链接在一起协同工作(PCIe延迟稍微高一些)
  • CPU所在的位置成为主机端（host），而GPU所在的位置称为设备端（device）.

为什么要使用GPU计算

• 并行计算引擎强大，可以大幅加快计算速度

CPU与GPU分工与协作

• GPU：
  • 运算核心多，适合数据并行的计算密集型任务：大型矩阵运算
  • 线程是轻量级的
• CPU：
  • 运算核心少，但是可以实现发杂的逻辑运算，一次适合控制密集型任务
  • 线程是重量级的，上下文切换开销大
• 基于二者的计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，而GPU重点处理数据密集型的并行计算程序，从而发挥最大功效。

GPU计算架构

• 尽可能减少拷贝可以加快运行效率
• 最新的架构可以申请一块公共区域，CPU与GPU可以一起访问

程序架构

语言选取

• C/C++
• Python可以调用

编译器

GPU硬件架构综述

一些名词

• SM：

  • Streaming Multiprocessor(32 SP)（流多处理器）
  • Multi-threaded processor core（多线程处理器内核）
  • Fundamental processing unit for CUDA thread block（CUDA螺纹块基本处理单元）

• SP (CUDA Core) Streaming Processor（流处理器）

  • 

• SPA

  • Streaming Processor Array（流处理器阵列）

• TPC/GPC

  • Texture (Graphics) Processor Cluster （流多处理器组成的一个小组）
    • 3 SM + TEX

线程周期

• 线程分为两个级别即Grid&Block，Grid在SPA（流处理器阵列）中启动，线程块block连续分散映射到SM（流多处理器）

• 但是每个block必须是映射到同一个SM中，不能跨SM执行

• block中的thread，每32个组成一个Warps
• Warp: primitive scheduling unit（基本调度单元）
• 每个流多处理器上有个Warp的调度器，根据Warp的状态决定是否调度、先调度哪个。
• 在设计的时候尽量保证每个warp执行相同的计算，否则效率会降低

• 硬件上实现了warp的零开销

• 寄存器是稀有资源，块间共享，在设计时尽量少使用寄存器资源可以保证更多的块处于活跃的状态

CUDA程序执行流程

流程

• 分配host内存，并进行数据初始化
• 分配device内存并从host将数据拷贝到deviceshang
• 调用CUDA的核函数在device上完成指定运算
• 将device上运算结果拷贝到host上（性能）
• 释放device核host的内存

CUDA程序

• kernel是CUDA中的一个重要的概念，kernel是在device上线程中并行执行的函数
• 核函数用**__global__**符号声明,在调用时需要用<<grid,block>>来指定客人呢良药之下的线程数量
• 在CUDA中，每一个县城都要窒息和函数，并且每个线程都会分配一个唯一的县城好thread ID，通过和函数内置变量threadidx获得

// Kernel定义
__global__ void vec_add(double *x, double *y, double *z, int n)
{
   
    int i = get_tid(); // user-defined function
    if (i < n) z[i] = x[i] + y[id];
}
int main()
{
   
    int N = 1000000; // 1M 表示这个函数需要执行多少次~
    //grid 和 block 通过计算得出每个grid和block需要计算的量
    int bs = 256;
    int gs = (N + bs - 1) / bs;
    // kernel, call GPU
    vec_add<<<gs, bs>>>(x, y, z, N);
}

CUDA程序层次结构

• grid 和 block 都是定义为dim3类型的变量
• dim3可以看成是包含三个无符号整数 (x, y, z) 成员的结构体变量，在定 义时，缺省值初始化为1。
• grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构
• 定义的grid和block如下所示，kernel在调用时也必须通过执行配置 <<<grid, block>>>来指定kernel所使用的线程数及结构。
• 不同 GPU 架构, grid 和 block 的维度有限制

dim3 grid(3, 2);//用dim3来定义，3x2=6个grid
dim3 block(5, 3);//5x3=15个block
kernel_fun<<< grid, block >>>(prams.</