Cuda笔记【1】GPU计算DEMO

椰子奶糖

已于 2023-06-29 14:56:31 修改

阅读量2.8k

点赞数 2

分类专栏： # CUDA编程文章标签：笔记

于 2020-03-15 17:41:35 首次发布

本文链接：https://blog.csdn.net/symuamua/article/details/104881741

版权

CUDA编程专栏收录该内容

6 篇文章 1 订阅

订阅专栏

声明

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

椰子奶糖

关注

2
点赞
踩
2

收藏

觉得还不错? 一键收藏
打赏
0
评论
Cuda笔记【1】GPU计算DEMO

声明本文是看小破站某cuda入门教程留下来的笔记，多上PPT上内容，夹杂一点自己的理解，和代码注释教程地址:https://www.bilibili.com/video/av74148375git地址（PPT和源码）：https://github.com/huiscliu/tutorials主要目的是为Gstreamer打点基础，不然基本抓瞎文章目录声明介绍什么是GPU计算为什么要使...
复制链接

扫一扫