[转]Cuda笔记【1】GPU计算DEMO

技术_编程语言 专栏收录该内容
128 篇文章 0 订阅

 

声明

  • 本文是看小破站某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

线程周期

  • 线程分为两个级别即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;
}

 

  • 0
    点赞
  • 0
    评论
  • 1
    收藏
  • 一键三连
    一键三连
  • 扫一扫,分享海报

相关推荐
©️2020 CSDN 皮肤主题: 程序猿惹谁了 设计师:白松林 返回首页
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、C币套餐、付费专栏及课程。

余额充值