CUDA系列：编程实践1

Albert_Liyan
183 阅读7分钟

1 GPU与CUDA结构

认识GPU

处理器关键指标：时延和吞吐

延迟：从发出指令到最终返回结果中间经历的时间间隔。 吞吐量：单位之间内处理的指令的条数。

CPU vs GPU
1 CPU包含多级高速缓存
2 CPU包含很多控制单元：分支预测，流水线前传
3 CPU运算单元core非常强，整型浮点型复杂运算速度快
总结：CPU是延迟导向设计。

1 GPU缓存结构数量少，减少访问缓存次数
2 GPU控制单元简单
3 GPU的运算单元Core非常多，采用长延时流水线以实现高吞吐量。
总结：GPU为吞吐导向设计，GPU核心原则：增加简单指令的吞吐。

二者擅长场景不同：

  • CPU在连续计算部分，延迟优先，CPU比GPU快10倍以上；
  • GPU在并行计算方面，吞吐优先，GPU比CPU单位时间内执行指令数量10倍以上；

GPU适用场景：
1 计算密集：计算比例远大于内存操作；
2 数据并行：大任务可以拆解为执行相同指令的小任务；

CUDA结构

CUDA (Compute Unified Device Architecture)支持GPU通用计算的平台和编程模型，提供C/C++语言扩展和用于编程和管理 GPU的API。

解读CUDA内存模型

CUDA内存模型的最基本单位：SP（线程处理器），每个SP都有自己的register（寄存器）和local mem。不同SP之间相互独立。
多个SP和一块共享内存组成SM（多核处理器）。多核处理器SM中的多个SP是互相并行，互不影响的。每个SM有自己的shared mem（共享内存），shared mem可以被线程块内所有线程访问。
多个SM和一块全局内存，构成GPU。一个GPU的SM共有一块global mem，不同线程块的线程都可适用。

image.png

从内存模型角度，还可表述：每个thread有自己的一份registre和local mem；同一个block中每个thread共享一份shared mem；所有thread（不同block）共享一份global mem。不同grid有各自的global mem。

从软件角度理解：
1 线程处理器SP 对应thread；
2 多核处理器SM 对应线程块 thread block；
3 设备端device 对应线程块组合体grid；

线程块内存模型是软件侧最基本的执行单位，线程块特点：
1 块内线程通过共享内存、原子操作和屏障同步协作；
2 不同块中的线程不能协作；

2 CUDA编程要素

kernel：C++定义的基本函数执行单元。kernel调用时由N个不同CUDA线程并行执行N次；

线程层次

  • Thread：所有线程执行相同核函数，并行执行；32个thread组成一个warp(线程束)，一个warp对应一条指令流。
  • Thread Block：线程块执行在一个SM；同一个Block中线程可以协作；一个硬件SM可执行多个block，一个block只能在一个SM中执行。
  • Thread Grid：多个线程块组成grid。

image.png

变量含义和设置如下：

  • threadIdx.[x y z]，表示执行当前kernel函数的线程在block中索引值, 例如上图中Thread(0,0)
  • blockIdx.[x y z], 表示执行当前kernel函数的线程所在block，在grid中的索引值,例如block(1,1)
  • blockDim.[x y z]表示一个block中包含多少线程
  • gridDim.[x y z] 表示一个gird包含多少block

<<<>>> 运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。

Kernel<<<Dg,Db, Ns, S>>>(param list);
  • 参数Dg定义整个grid的维度和尺寸, 即一个 grid 有多少个 block。为 dim3 类型。Dim3 Dg(Dg.x, Dg.y, 1)
  • 参数Db定义一个 block 的维度和尺寸, Dim3 Db(Db.x, Db.y, Db.z)
  • 参数Ns可选, 设置每个 block 除了静态分配的 shared Memory 以外，最多能动态分配的shared memory 大小，单位为 byte, 默认值是0;
  • 参数S cudaStream_t 类型的可选参数

程序示例如下：

__global__ void add( int *a, int *b, int *c ) {
c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}
add<<<1,4>>>( a, b, c);

理解dim3数据类型:
dim3是一个结构体，有x，y，z三个变量。 使用dim3数据类型指定grid、block大小，x、y、z三个变量分别表示三个维度大小；必须至少指定一个变量x，其他变量不指定默认是1;

计算访问线程的索引的位置

  • 线程块id和线程id定位线程的显存位置 CUDA的核函数需要在设备端执行，执行核函数需要访问每个线程的register和local mem，因此需要确定每个线程在显存的位置 int i= threadIdx.x+ blockDim.x* blockIdx.x

线程索引和线程ID的对应关系：

  • 一维块，线程的threadID=threadIdx.x；
  • 二维块，（blockDim.x, blockDim.y），threadID=threadIdx.x+threadIdx.y *　blockDim.x;
  • 三维块，(blockDim.x, blockDim.y, blockDim.z），threadID=threadIdx.x +threadIdx.y * blockDim.x+threadIdx.z * blockDim.x * blockDim.y

CUDA执行流程

1 加载核函数
2 将Grid分配到一个Device
3 根据<<<…>>>内的执行设置的第一个参数，Giga threads engine将block分配到SM中。一个Block内的线程一定会在同一个SM内，一个SM可以有很多个Block。
4 根据<<<…>>>内的执行设置的第二个参数，Warp调度器会调用线程。
5 Warp调度器为了提高运行效率，会将每32个线程分为一组，称作一个warp。 6 在某个时刻，每个Warp指令会被SM内部的某些单元执行(Warp具体分配给多少个SP是不确定的，可以根据deviceQuery查询，例如计算能力7.5可能给1组（16个）SP，连续2个周期执行)

3 pytorch自定义算子

参考github.com/ifromeast/c… 构建自定义算子：两个n*n tensor相加，block和grid都是二维结构（dim3 block(16, 16), dim3 grid(n/block.x, n/block.y)）。

代码结构

include存放cuda算子的头文件 kernel存放cuda算子的具体实现(.cu)和cpp torch接口封装(.cpp)


├── include  
│   └── add2.h # cuda算子的头文件
├── kernel
│   ├── add2_kernel.cu # cuda算子的具体实现
│   └── add2.cpp # cuda算子的cpp torch封装
├── CMakeLists.txt
├── setup.py
└── time.py # 比较cuda算子和torch实现的时间差异

Torch 使用CUDA算子 主要分为三个步骤：
1 编写CUDA算子和对应的调用函数
2 编写torch cpp函数建立PyTorch和CUDA之间的联系，使用pybind11封装或TORCH_LIBRARY封装
3 用PyTorch的cpp扩展库进行编译和调用

编译cpp和cuda

JIT编译

JIT即时编译，就是说在python代码运行的时候再去编译cpp和cuda文件

编译流程：先加载需要即时编译的文件，然后调用接口函数 extra_include_paths表示包含的头文件目录，sources表示需要编译的代码

from torch.utils.cpp_extension import load
cuda_module = load(name="add2",
                    extra_include_paths=["include"],
                    sources=["kernel/add2_ops.cpp", "kernel/add2_kernel.cu"],
                    verbose=True)
cuda_module.torch_launch_add2(cuda_c, a, b, n)  //调用接口函数

执行编译命令：

python run_time.py --compiler jit
Setuptools编译

编译流程：编写setup.py，首先调用Torch CUDAExtension模块注册算子add2，同时include_dirs加上头文件目录，ext_modules 加上算子及封装函数。

from setuptools import setup
from torch.utils.cpp_extension import BuildExtension, CUDAExtension

setup(
    name="add2",
    include_dirs=["include"],
    ext_modules=[
        CUDAExtension(
            "add2",
            ["kernel/add2_ops.cpp", "kernel/add2_kernel.cu"],
        )
    ],
    cmdclass={
        "build_ext": BuildExtension
    }
)

编译执行：

python setup.py install

在python端调用：

import torch
import add2
add2.torch_launch_add2(c, a, b, n)

执行命令：

python run_time.py --compiler setup
CMAKE编译调用

编写一个CMakeLists.txt文件，需要关注：依赖库的匹配、编译过程及软连接的建立。 cpp用TORCH_LIBRARY封装

TORCH_LIBRARY(add2, m) {
    m.def("torch_launch_add2", torch_launch_add2);
}

编译命令：

mkdir build
cd build
cmake ..
make

编译过程中报错Policy CMP0104 is not set: CMAKE_CUDA_ARCHITECTURES now detected for NVCC， set(CMAKE_CUDA_ARCHITECTURES 61) //61表示显卡版本（T4），其他报错根据日志检查配置即可。

执行命令：

python run_time.py --compiler cmake

在python端调用方式：

import torch
torch.ops.load_library("build/libadd2.so")
torch.ops.add2.torch_launch_add2(c, a, b, n)

执行结果

Running cuda...
Cuda time:  106.382us
Running torch...
Torch time:  106.716us
Kernel test passed.

参考:
mp.weixin.qq.com/s/kxYSw_fR4… zhuanlan.zhihu.com/p/645330027 blog.csdn.net/lansebingxu…
cppdebug.com/archives/55…

avatar
文章
阅读
粉丝