【cuda编程】cuda编程模型

最新推荐文章于 2022-10-19 16:33:06 发布

而与你及

最新推荐文章于 2022-10-19 16:33:06 发布

阅读量378

点赞数

分类专栏： cuda

原文链接：https://www.bilibili.com/video/av74148375?p=1

版权

cuda 专栏收录该内容

1 篇文章 0 订阅

订阅专栏

基础知识

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的整数倍。

参考文献

cuda编程入门

而与你及

关注

0
点赞
踩
2

收藏

觉得还不错? 一键收藏
0
评论
【cuda编程】cuda编程模型

Cuda编程模型基础知识cuda程序执行流程cuda程序cuda程序的层次结构cuda内置变量基础知识cuda编程模型需要cpu和gpu协同工作；host指代cpu及其内存，device指代gpu及其内存；cuda程序既包含host程序又包含device程序，分别在cpu和gpu上运行；host和device之间还要进行通信以便进行数据拷贝。cuda程序执行流程分配host内存...
复制链接

扫一扫