声明
- 本文是看小破站某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
- Texture (Graphics) Processor Cluster (流多处理器组成的一个小组)
线程周期
- 线程分为两个级别即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...);//如上,一共15x6=90个
dim3 grid(128);//或者直接指定个数
dim3 block(256);
kernel_fun<<< grid, block >>>(prams...);
dim3 grid(100, 120);
dim3 block(16,16,1);//16x16x1=256个
kernel_fun<<< grid, block >>>(prams...);
- 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都编译。
CUDA 内置变量
- 一个线程需要两个内置的坐标变量(blockIdx,threadIdx)来唯一标识,它们都是 dim3类型变量,其中blockIdx指明线程所在grid中的位置,而threaIdx指明线程所在 block中的位置:
- threadIdx包含三个值: threadIdx.x, threadIdx.y, threadIdx.z
- blockIdx同样包含三个值: blockIdx.x, blockIdx.y, blockIdx.z
- 一个线程块上的线程是放在同一个流式多处理器(SM)上的
- 单个SM的资源有限,这导致线程块中的线程数是有限制的,现代GPUs的线程块可 支持的线程数可达1024个。
- 有时候,我们要知道一个线程在blcok中的全局ID,此时就必须还要知道block的组 织结构,这是通过线程的内置变量blockDim来获得。它获取线程块各个维度的大小。
- 对于一个2-dim的block ,线程 的ID值为 ,如果是3-dim的block ,线程 的ID值为 。 另外线程还有内置变量gridDim,用于获得网格块各个维度的大小。
/* get thread id: 1D block and 2D grid */
//定义宏,下同
#define get_tid() (blockDim.x * (blockIdx.x + blockIdx.y * gridDim.x) + threadIdx.x)
/* get block id: 2D grid */
#define get_bid() (blockIdx.x + blockIdx.y * gridDim.x)
__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];
}
完整代码
- 以及少量注释
#include <stdio.h>
#include <cuda.h>
typedef float FLOAT;
#define USE_UNIX 1
/* get thread id: 1D block and 2D grid */
#define get_tid() (blockDim.x * (blockIdx.x + blockIdx.y * gridDim.x) + threadIdx.x)
/* get block id: 2D grid */
#define get_bid() (blockIdx.x + blockIdx.y * gridDim.x)
/* warm up, start GPU, optional */
void warmup();
/* get time stamp */
double get_time(void);
/* host, add */
void vec_add_host(FLOAT *x, FLOAT *y, FLOAT *z, int N);
/* device function */
__global__ void vec_add(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
/* 1D block */
int idx = get_tid();
if (idx < N) z[idx] = z[idx] + y[idx] + x[idx];
}
void vec_add_host(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
int i;
for (i = 0; i < N; i++) z[i] = z[i] + y[i] + x[i];
}
/* a little system programming */
#if USE_UNIX
#include <sys/time.h>
#include <time.h>
//linux 获得时间戳
double get_time(void)
{
struct timeval tv;
double t;
gettimeofday(&tv, (struct timezone *)0);
t = tv.tv_sec + (double)tv.tv_usec * 1e-6;
return t;
}
#else
#include <windows.h>
//windows 获得时间戳
double get_time(void)
{
LARGE_INTEGER timer;
static LARGE_INTEGER fre;
static int init = 0;
double t;
if (init != 1) {
QueryPerformanceFrequency(&fre);
init = 1;
}
QueryPerformanceCounter(&timer);
t = timer.QuadPart * 1. / fre.QuadPart;
return t;
}
#endif
/* warm up GPU */
__global__ void warmup_knl()
{
int i, j;
i = 1;
j = 2;
i = i + j;
}
void warmup()
{
int i;
for (i = 0; i < 8; i++) {
warmup_knl<<<1, 256>>>();
}
}
int main()
{
int N = 20000000;
int nbytes = N * sizeof(FLOAT);
/* 1D block */
int bs = 256;
/* 2D grid */
int s = ceil(sqrt((N + bs - 1.) / bs));
dim3 grid = dim3(s, s);
FLOAT *dx = NULL, *hx = NULL;
FLOAT *dy = NULL, *hy = NULL;
FLOAT *dz = NULL, *hz = NULL;
int itr = 30;
int i;
double th, td;
/* warm up GPU 预热 */
warmup();
/* allocate GPU mem 申请GPU内存*/
cudaMalloc((void **)&dx, nbytes);
cudaMalloc((void **)&dy, nbytes);
cudaMalloc((void **)&dz, nbytes);
if (dx == NULL || dy == NULL || dz == NULL) {
printf("couldn't allocate GPU memory\n");
return -1;
}
printf("allocated %.2f MB on GPU\n", nbytes / (1024.f * 1024.f));
/* alllocate CPU mem 申请CPU内存*/
hx = (FLOAT *) malloc(nbytes);
hy = (FLOAT *) malloc(nbytes);
hz = (FLOAT *) malloc(nbytes);
if (hx == NULL || hy == NULL || hz == NULL) {
printf("couldn't allocate CPU memory\n");
return -2;
}
printf("allocated %.2f MB on CPU\n", nbytes / (1024.f * 1024.f));
/* init */
for (i = 0; i < N; i++) {
hx[i] = 1;
hy[i] = 1;
hz[i] = 1;
}
/* copy data to GPU 数据cpu->gpu*/
//参数解释(dx为GPU内存地址,hx为CPU地址,nbytes表示一共拷贝多少字节,cudaMemcpyHostToDevice:表示方向,为cpu->gpu)
cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dz, hz, nbytes, cudaMemcpyHostToDevice);
/* warm up */
warmup();
/* call GPU */
//cudaThreadSynchronize()强行加锁,这样可以得到GPU运行时间,否则异步调用不会等待
cudaThreadSynchronize();
td = get_time();
for (i = 0; i < itr; i++) vec_add<<<grid, bs>>>(dx, dy, dz, N);
cudaThreadSynchronize();
td = get_time() - td;
/* CPU */
th = get_time();
for (i = 0; i < itr; i++) vec_add_host(hx, hy, hz, N);
th = get_time() - th;
printf("GPU time: %e, CPU time: %e, speedup: %g\n", td, th, th / td);
cudaFree(dx);
cudaFree(dy);
cudaFree(dz);
free(hx);
free(hy);
free(hz);
return 0;
}