CUDA入门笔记

Posted Lachiven

tags:

篇首语：本文由小常识网(cha138.com)小编为大家整理，主要介绍了CUDA入门笔记相关的知识，希望对你有一定的参考价值。

硬件概念：

一个SM(Streaming Multiprocessor)中的所有SP(Streaming Processor)是分成Warp的，共享同一个Memory和Instruction Unit（指令单元）。

从硬件角度讲，一个GPU由多个SM组成（当然还有其他部分），一个SM包含有多个SP（以及还有寄存器资源，Shared Memory资源，L1cache，Scheduler，SPU，LD/ST单元等等）

SM采用的是SIMT (Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束（warps)，线程束包含32个线程.

线程数尽量为32的倍数，因为线程束(Warp)共享一个内存和指令。假如线程数是1，则Warp会生成一个掩码，当一个指令控制器对一个Warp单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个进程会进入静默状态。

了解硬件的代码：

cu文件:

#include <iostream>
#include "cuda_runtime.h"
int main() 
    int dev = 0;
    cudaDeviceProp devProp;
    cudaGetDeviceProperties(&devProp, dev);
    std::cout << "使用GPU device " << dev << ": " << devProp.name << std::endl;
    std::cout << "SM的数量：" << devProp.multiProcessorCount << std::endl;
    std::cout << "每个线程块的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
    std::cout << "每个线程块的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;
    std::cout << "每个SM的最大 线程    数：" << devProp.maxThreadsPerMultiProcessor << std::endl;
    std::cout << "每个SM的最大 线程束  数——Warp数：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
    return 0;

CMakeLists:

cmake_minimum_required(VERSION 3.8)
project(CUDA_TEST)

find_package(CUDA REQUIRED)

message(STATUS "cuda version: " $CUDA_VERSION_STRING)
include_directories($CUDA_INCLUDE_DIRS)

cuda_add_executable(cuda_test src/test001.cu)
target_link_libraries(cuda_test $CUDA_LIBRARIES)

软件概念

在利用cuda进行编程时，一个grid分为多个block，而一个block分为多个thread

  • 每个 thread 都有自己的一份 register 和 local memory 的空间。
  • 一组thread构成一个 block，这些thread 则共享有一份shared memory。
  • 所有的 thread(包括不同 block 的 thread)都共享一份global memory、constant memory、和 texture memory。
  • 不同的 grid 则有各自的 global memory、constant memory 和 texture memory。
  • 每一个时钟周期内，warp(一个block里面一起运行的thread，其中各个线程对应的数据资源不同(指令相同但是数据不同)包含的thread数量是有限的，现在的规定是32个。一个block中含有16个warp。所以一个block中最多含有512个线程，每次Device（就是显卡）只处理一个grid。
  • 一个sm只会执行一个block里的warp，当该block里warp执行完才会执行其他block里的warp。

CUDA编程CUDA入门笔记

文章目录

1. CUDA结构

显卡内部，有三级结构：网格（grid）、块（block）、线程（thread）。每个显卡只有很少的网格，一个核函数只能运行在一个网格中，而一个网格里有多个块，每个块包含了若干线程。

1对1 1对多 1对多 kernel 网格grid 线程块block 线程thread

kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。线程两层组织结构如下图所示，这是一个gird和block均为2-dim的线程组织。

介绍说明：

  1. Thread：线程，并行的基本单元
  2. Thread Block：线程块，互相合作的线程组。它有以下特点：
  • 允许彼此同步
  • 可以通过共享内存快速交换数据
  • 以1维、2维或3维组织(里面的thread)
  1. Grid：一组线程块
  • 以1维、2维和3维度组织(里面是block)
  • 共享全局内存
  1. Kernel：在GPU上执行的核心程序，这个kernel函数是运行在某个Grid上的。

每一个block和每个thread都有自己的ID，我们通过相应的索引找到相应的线程和线程块。

  • threadIdx，blockIdx
  • Block ID: 1D or 2D
  • Thread ID: 1D, 2D or 3D

2. SP、SM与warp

SP（streaming Process），SM（streaming multiprocessor）是硬件（GPU hardware）概念。而thread，block，grid，warp是软件上的（CUDA）概念。

  • SP(streaming processor)：最基本的处理单元，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。

  • SM(streaming multiprocessor)：多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。

  • warp(线程束)：GPU执行程序时的调度单位，目前cuda的warp的大小为32，同在一个warp的线程，以不同数据资源执行相同的指令,这就是所谓 SIMT。

需要指出，每个SM包含的SP数量依据GPU架构而不同，Fermi架构GF100是32个，GF10X是48个，Kepler架构都是192个，Maxwell都是128个。

简而言之，SP是线程执行的硬件单位，SM中包含多个SP，一个GPU可以有多个SM（比如16个），最终一个GPU可能包含有上千个SP。

下图展示了逻辑层面和硬件层面的对应关系：

  • 每个线程由每个线程处理器（SP）执行
  • 线程块由多核处理器（SM）执行
  • 一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行
  • 关于warp的解释：
  • 一个SM中可以同时有多个warp，这些warp被称为active warp，不同warp处于不同状态，挂起，就绪，执行。但是一个SM上正在被执行的就只有一个warp，这个正在被执行的warp叫做resident warp。 active warp是指已经分配给SM的warp，并且该warp需要的资源（寄存器）也已经分配。
  • 一个SP可以执行一个thread，但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp，warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。一个warp需要占用一个SM运行，多个warps需要轮流进入SM。由SM的硬件warp scheduler负责调度。目前每个warp包含32个threads（Nvidia保留修改数量的权利）。所以，一个GPU上resident thread最多只有 SM*warp个。

block是软件概念，一个block只会由一个sm调度，程序员在开发时，通过设定block的属性，告诉GPU硬件，我有多少个线程，线程怎么组织。而具体怎么调度由sm的warps scheduler负责，block一旦被分配好SM，该block就会一直驻留在该SM中，直到执行结束。一个SM可以同时拥有多个blocks，但需要序列执行。下图显示了GPU内部的硬件架构：

3. CUDA线程索引的方式

CUDA使用多级索引的方式访问线程。

  • 定位Block：第一级索引是(grid.xIdx, grid.yIdy)，通过它我们就能找到了这个线程块的位置。
  • 定位thread：第二级索引(block.xIdx, block.yIdx, block.zIdx)来定位到指定的线程。

grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（水平方向为x轴），定义的grid和block如下所示

dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<< grid, block >>>(prams...);

定义图解如下：

一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置，如图中的Thread (1,1)满足：

threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

一个线程块上的线程是放在同一个流式多处理器（SM)上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数可达1024个。有时候，我们要知道一个线程在blcok中的全局ID，此时就必须还要知道block的组织结构，这是通过线程的内置变量blockDim来获得。它获取线程块各个维度的大小。对于一个2-dim的 b l o c k ( D x , D y ) block(D_x,D_y) block(Dx​,Dy​)，线程 ( x , y ) (x,y) (x,y)的ID值为 ( x + y ∗ D x ) (x+y∗D_x) (x+y∗Dx​)，如果是3-dim的 b l o c k ( D x , D y , D z ) block(D_x,D_y,D_z) block(Dx​,Dy​,Dz​)，线程 ( x , y , z ) (x,y,z) (x,y,z)的ID值为 ( x + y ∗ D x + z ∗ D x ∗ D y ) (x+y∗D_x+z∗D_x∗D_y) (x+y∗Dx​+z∗Dx​∗Dy​)。另外线程还有内置变量gridDim，用于获得网格块各个维度的大小。

一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些Threads的组织方式也可以是一维，二维或者三维的。

CUDA中每一个线程都有一个唯一的标识ID—ThreadIdx，这个ID随着Grid和Block的划分方式的不同而变化，这里给出Grid和Block不同划分方式下线程索引ID的计算公式。

1、 grid划分成1维，block划分为1维

    int threadId = blockIdx.x *blockDim.x + threadIdx.x;

2、 grid划分成1维，block划分为2维

int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;

3、 grid划分成1维，block划分为3维

    int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z  
                       + threadIdx.z * blockDim.y * blockDim.x  
                       + threadIdx.y * blockDim.x + threadIdx.x;

4、 grid划分成2维，block划分为1维

    int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
    int threadId = blockId * blockDim.x + threadIdx.x;

5、 grid划分成2维，block划分为2维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;

6、 grid划分成2维，block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
                       + (threadIdx.z * (blockDim.x * blockDim.y))  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;

7、 grid划分成3维，block划分为1维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * blockDim.x + threadIdx.x;

8、 grid划分成3维，block划分为2维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;

9、 grid划分成3维，block划分为3维

    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
                     + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
                       + (threadIdx.z * (blockDim.x * blockDim.y))  
                       + (threadIdx.y * blockDim.x) + threadIdx.x;

4 CUDA的内存

CUDA中的内存模型分为以下几个层次：

  • 每个线程都用自己的registers（寄存器）和local memory（局部内存）
  • 每个线程块(block)内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
  • 每个grid都有自己的global memory（全局内存），constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用。

线程访问这几类存储器的速度是：register > local memory >shared memory > global memory。

下面这幅图表示就是这些内存在计算机架构中的所在层次。

5 编程模型

在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。

典型的CUDA程序的执行流程如下：

  1. 分配host内存，并进行数据初始化；
  2. 分配device内存，并从host将数据拷贝到device上；
  3. 调用CUDA的核函数在device上完成指定的运算；
  4. 将device上的运算结果拷贝到host上；
  5. 释放device和host上分配的内存。

核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。

（1）怎么确定是在CPU还是GPU上跑？

由于GPU实际上是异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词开区别host和device上的函数，主要的三个函数类型限定词如下：

  • __global__：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
  • __device__：在device上执行，仅可以从device中调用，不可以和__global__同时用。
  • __host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。

通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑。

函数执行位置调用位置
__device__ float DeviceFunc()devicedevice
__global void KernelFunc()devicehost
__host__ float HostFunc()hosthost

（2）CPU与GPU的数据传输

首先介绍在GPU内存分配回收内存的函数接口：

  • cudaMalloc(): 在设备端分配global memory
  • cudaFree(): 释放存储空间

CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：

cudaMemcpy(void *dst, void *src, size_t nbytes, enum cudaMemcpyKind direction);

enum cudaMemcpyKind的类型：

  • cudaMemcpyHostToDevice（CPU到GPU）
  • cudaMemcpyDeviceToHost（GPU到CPU）
  • cudaMemcpyDeviceToDevice（GPU到GPU）

（3）怎么用代码表示线程组织模型

我们可以用dim3类来表示网格和线程块的组织方式，网格grid可以表示为一维和二维格式，线程块block可以表示为一维、二维和三维的数据格式。

dim3 DimGrid(100, 50);  //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8);  //每个线层块内包含256个线程，线程块内的维度是4*8*8

参考：

https://blog.csdn.net/xiaohu2022/article/details/79599947

以上是关于CUDA入门笔记的主要内容，如果未能解决你的问题，请参考以下文章

BBuf的CUDA笔记三，reduce优化入门学习笔记

BBuf的CUDA笔记一，解析OneFlow Element-Wise 算子实现

BBuf的CUDA笔记一，解析OneFlow Element-Wise 算子实现

关于ChatGPT的一切；CUDA入门之矩阵乘；PyTorch 2.0发布｜AI系统前沿动态

CUDA入门

CUDA编程之快速入门