Cuda编程模型
基础知识
- cuda编程模型需要cpu和gpu协同工作;
- host指代cpu及其内存,device指代gpu及其内存;
- cuda程序既包含host程序又包含device程序,分别在cpu和gpu上运行;
- host和device之间还要进行通信以便进行数据拷贝。
cuda程序执行流程
- 分配host内存,并进行数据初始化;
- 分配device内存,并将host上的数据拷贝到device上;
- 调用cuda的核函数在device上完成指定的运算;
- 将device上的运算结果拷贝回host;
- 释放host和device上分配的内存。
cuda程序
cuda核函数kernel是在device上的线程中并行执行的函数。核函数用__global__符号声明,在调用时要用<<<grid,block>>>指定kernel要执行的线程的数量。
在cuda中,每一个线程都要执行核函数,每一个线程会分配一个唯一的线程号thread ID,这个ID值可以通过核函数的内置变量threadIdx来获得。
GPU代码片段:
/* get thread id: 1D block + 2D grid */
# define get_id() (blockDim.x * (blockIdx.x + blockIdx.y * gridDim.x) + threadIdx.x)
/* get block id: 2D grid */
# define get_bid() (gridDim.x * blockIdx.y + blockIdx.x)
// Kernel的定义
__gloabal__ void vec_add(double *x, double *y, double *z, int n)
{
int i = get_tid() // 用户自定义函数,用来得到线程的ID
if(i < n) z[i] = x[i] + y[i]
}
int main()
{
int N = 1000000;
int bs = 256;
int gs = (N + bs - 1) / bs;
// kernel call gpu
vec_add<<<gs, bs>>> (x, y, z, N)
}
cuda程序的层次结构
一个kernel所启动的所有线程称为一个网格(grid)。同一个网格上的线程共享相同的全局内存空间,grid是线程结构的第一个层次。网格又可以分为很多的线程块block,一个线程块里包含很多的线程,这是第二个层次。第三个层次是warp:32个线程为一组。
grid和block都是定义为dim3类型的变量,也可以灵活的定义为1-dim,2-dim,3-dim。dim3可以看成包含三个无符号整数(x, y, z)成员的结构体变量,在定义时,缺省值为1。
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_func<<<grid, block>>>(params...)
在cuda中通过函数类型限定词来区别host和device上的函数:
- global:在device上执行,从host中调用,返回类型必须是void,不支持可变参数;注意用__global__定义的kernel是异步的,即前一个kernel执行完才能执行后一个kernel,也就是串行,并且host不会等待kernel执行完就开始下一步;
- device:在device上执行,仅可以从device中调用,不可和__global__同时使用;
- host:在host上执行,仅可以在host上调用,一般省略不写,不可以和__global__同时使用,但可以和__device__同时使用,此时函数会在device和host都编译。
cuda内置变量
一个线程需要有两个内置的坐标变量(blockIdx,threadIdx)来唯一辨识,它们都是dim3类型变量,其中blockIdx指明线程所在grid中的位置,而threadIdx指明线程所在block中的位置。有时候,我们要知道一个线程在block中的全局ID,此时就还必须要知道block的结构,这是通过线程的内置变量blockDim来获得。它获取线程块的各个维度的大小。类似的内置变量还有gridDim,用于获取网格块的各个维度的大小。
blockIdx:blockIdx.x,blockIdx.y,blockIdx.z
threadIdx:threadIdx.x,threadIdx.y,threadIdx.z
一个线程块上的线程是放在同一个流式多处理器(SM)上的。
向量加法程序实例
#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>
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>
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 */
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 */
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 */
cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dz, hz, nbytes, cudaMemcpyHostToDevice);
/* warm up */
warmup();
/* call 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;
}
GPU内存
GPU 内存类型
- 每个线程都有自己的私有本地内存(Local Memory);
- 每个线程块都有共享内存(Shared Memory)可被线程块中的线程共享,生命周期与线程块一致;
- 所有线程都可以访问全局内存(Global Memory);
- 只读内存块:常量内存和纹理内存;
- L1 cache、L2 cache。
可编程内存:可以用代码控制的内存;
不可编程内存:不对用户开放的内存。
在CPU、GPU内存结构中,一级二级缓存都是不可编程内存的存储设备。
内存作用域&生命周期
- cuda中每个线程都有自己的私有本地内存和寄存器;
- 所有线程都可以访问和读取常量内存和纹理内存,但不能写;
- 全局内存和常量内存和纹理内存有相同的生命周期。
寄存器
- 速度最快的内存空间,与CPU不同的是GPU的寄存器储量要多一些;
- 当我们在核函数中不加修饰的声明一个变量时,此变量就存储在寄存器中;并且在核函数中定义常数长度的数组也是在寄存器中分配地址的;
- 寄存器对于每个线程来说是私有的,寄存器通常保存被频繁使用的私有变量。寄存器的生命周期和核函数一致,从开始运行到运行结束,执行完毕后,寄存器就不能访问了;
- 如果一个线程里面的变量太多,寄存器不够用的话,本地内存就过来帮忙存储多出来的变量,但是效率会降低。
本地内存
本地内存实质上是和全局内存一样在同一块存储区域当中的,特点是高延迟、低带宽。
共享内存
- 在核函数中使用如下修饰符的内存称为共享内存:shared;
- 共享内存跟主存比速度要快很多,也就是延迟低、带宽高,类似一级缓存,但是可以被编程;
- 使用共享内存的时候一定要注意不要因为过度使用共享内存,而导致SM上活跃的线程束减少,也就是说,一个线程块使用的共享内存过多,导致其他线程块没法启动;
- 共享内存在核函数中声明,生命周期和线程块一致,线程块运行开始,此块共享内存被分配,此块结束,则共享内存被释放;
- 共享内存是块内线程可见的,所有就有竞争问题存在,也可以通过共享内存进行通信,为了避免内存竞争,可以使用同步语句:void __syncthreads();语句相当于在线程块执行时各个线程的一个障碍点,当块内所有线程都执行到本障碍点时才能进行下一步的计算
共享内存的访问冲突
- 共享内存分成大小相同的内存块bank,实现高速并行访问;
- 为了实现内存高带宽的同时访问,共享内存被划分成了可以同时访问的等大小的内存块banks。因此,内存读写n个地址的行为则可以以b个独立的bank同时操作的方式进行,这样有效带宽就提高到了一个bank的b倍;
- 如果多个线程请求的内存地址被映射到了同一个bank上,那么这些请求就变成了串行。硬件把这些请求分成x个没有冲突的请求序列,带宽就降成了原来的x分之一。但是如果一个warp内的所有线程都访问同一个内存地址的话,会产生一次广播,这些请求会一次完成。
全局内存
也就是GPU的显存。全局内存访问是对齐的,也就是一次要读取指定大小(32,64,128)整数倍字节的内存。
内存管理
GPU全局内存的分配与释放
- 内存分配:cudaMalloc(void **devPtr, size_t size);
- 内存释放:cudaFree(void *devPtr);
Host内存分配与释放
Host内存属于CPU内存,传输速度比CPU内存快很多。
- 内存分配:cudaMallocHost(void **devPtr, size_t size);
- 内存释放:cudaFreeHost(void *devPtr);
统一(Unified)内存分配与释放
Unified内存可以被CPU和GPU访问。
- 内存分配:cudaMallocManaged(void **devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal);
flags = cudaMemAttachGlobal:内存可以被任意处理器访问(CPU、GPU);
flags = cudaMemAttachHost:内存只能被CPU访问; - 内存释放:cudaFreeHost(void *devPtr);
CPU与GPU内存同步拷贝
cudaMemcpy(void *dst, const void *src, size_t size, cudaMemcpyKind kind)
代码实例
#include <stdio.h>
#include <cuda.h>
typedef double FLOAT;
__global__ void sum(FLOAT *x)
{
int tid = threadIdx.x;
x[tid] += 1;
}
int main()
{
int N = 32;
int nbytes = N * sizeof(FLOAT);
FLOAT *dx = NULL, *hx = NULL;
int i;
/* allocate GPU mem */
cudaMalloc((void **)&dx, nbytes);
if (dx == NULL) {
printf("couldn't allocate GPU memory\n");
return -1;
}
/* alllocate CPU host mem: memory copy is faster than malloc */
hx = (FLOAT *)malloc(nbytes);
if (hx == NULL) {
printf("couldn't allocate CPU memory\n");
return -2;
}
/* init */
printf("hx original: \n");
for (i = 0; i < N; i++) {
hx[i] = i;
printf("%g\n", hx[i]);
}
/* copy data to GPU */
cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
/* call GPU */
sum<<<1, N>>>(dx);
/* let GPU finish */
cudaThreadSynchronize();
/* copy data from GPU */
cudaMemcpy(hx, dx, nbytes, cudaMemcpyDeviceToHost);
printf("\nhx from GPU: \n");
for (i = 0; i < N; i++) {
printf("%g\n", hx[i]);
}
cudaFree(dx);
free(hx);
return 0;
}
cuda程序执行与硬件映射
上图表示一一映射关系。网格与GPU相对应,线程块与SM相对应,线程与cuda核心相对应。
GPU流式多处理器
- GPU硬件的一个核心组件是SM,也就是流失多处理器;
- SM的核心组件包括cuda核心、共享内存、寄存器等,SM可以并发的执行数百个线程,并发的能力取决于SM拥有的资源数;
- 当一个kernel被执行时,它的grid中的线程块被分配到SM上,一个线程块只能在一个SM上调度;
- SM一般可以调度多个线程块。那么有可能一个kernel的各个线程块被分配到多个SM上,所以grid只是逻辑层,而SM才是执行的物理层。
warp技术细节
- SM采用的是SIMT,即单指令多线程架构,基本执行单元是线程束warp,一个warp包含32个线程。一个warp内的线程执行的指令是一样的;
- 线程同时执行相同的指令,但是每个线程都有自己的指令地址计数器,寄存器状态等,也有自己独立的执行路径;
- 当线程块被分配到某个SM上时,它将进一步划分为多个线程束warp,它才是SM的基本执行单元;
- 因为资源限制,一个SM同时并发的线程束是有限的,SM要为每个线程块分配共享内存,而且也要为每个线程束中的线程分配独立的寄存器。所以SM的配置会影响并发的效率;
- kernel的grid和block的配置不同,性能也会出现差异。由于SM的基本执行单元是包含32个线程的线程束,所以block的大小一般要设置为32的整数倍。