GPU编程(二): GPU架构了解一下!

目录

  • 前言
  • GPU架构
  • GPU处理单元
  • 概念GPU
  • GPU线程与SM
  • GPU线程
  • SM
  • 加法
  • 统一内存
  • 乘法
  • 最后

前言

在实际CUDA编程之前, 先来了解下GPU的结构. 和CPU相比显得粗暴又强大(手动滑稽).


GPU架构

GPU处理单元

从这张GPU概念内核图开始讲起, 会发现和CPU内核是不同的, 少了三级缓存, 分支预测等等. 但是增加了ALU的数量, 扩大了上下文存储池(Pool of context storge).
可以看到, 上下文存储池分成4份, 也就是说, 可以执行4条指令流, 比方说指令1阻塞, 立马切换指令2, 指令2阻塞切换指令3, 这就起到了隐藏延迟的效果. 当然数量到底是多少是很讲究的, 不是越多越好.
总的来看, 内核含8个ALU, 4组执行环境(Execution context), 每组有8个Ctx. 这样, 一个这样的内核可以并发(concurrent but interleaved)执行4条指令流(instruction streams), 32个并发程序片元(fragment).


概念GPU

复制16个上述的处理单元, 得到一个GPU. 实际肯定没有这么简单的, 所以说是概念GPU.

这个GPU含16个处理单元, 128个ALU, 64组执行环境(Execution context), 512个并发程序片元(fragment).
祭出n多年前的卡皇GTX 480, 有480个CUDA核(也就是ALU), 内存带宽177.4GB/s. 而GTX 980 Ti有2816个CUDA核, 内存带宽336.5GB/s.
但是带宽依旧是瓶颈, 虽然比CPU带宽高了一个数量级, 但是可以看到, GTX 980 Ti的带宽也就是多年前GTX 480的两倍左右.


GPU线程与SM

由于目前还没有完全依靠GPU运行得机器, 一般来说, 都是异构的, CPU+GPU. 这一点是要特别注意的, 也就是Host与Device.


GPU线程

在CUDA架构下, 显示芯片执行时的最小单位是thread. 数个thread可以组成一个block. 一个block中的thread能存取同一块共享的内存, 而且可以快速进行同步的动作. 不同block中的thread无法存取同一个共享的内存, 因此无法直接互通或进行同步. 因此, 不同block中的thread能合作的程度是比较低的. 上图:

然后依据thread, block和grid, 有着不同的存储. 核心就是thread. 可以结合下图进行理解:

  • 每个处理器上有一组本地32位寄存器(Registers);
  • 并行数据缓存或共享存储器(Shared Memory), 由所有标量处理器核心共享, 共享存储器空间就位于此处;
  • 只读固定缓存(Constant Cache), 由所有标量处理器核心共享, 可加速从固定存储器空间进行的读取操作(这是设备存储器的一个只读区域);
  • 一个只读纹理缓存(Texture Cache), 由所有标量处理器核心共享, 加速从纹理存储器空间进行的读取操作(这是设备存储器的一个只读区域), 每个多处理器都会通过实现不同寻址模型和数据过滤的纹理单元访问纹理缓存.

SM

看到上图, GPU硬件的一个核心组件是SM, SM是英文名是Streaming Multiprocessor, 翻译过来就是流式多处理器. SM的核心组件包括CUDA核心, (其实就是ALU, 如上图绿色小块就是一个CUDA核心)共享内存, 寄存器等, SM可以并发地执行数百个线程, 并发能力就取决于SM所拥有的资源数. 当一个kernel被执行时, 它的gird中的线程块被分配到SM上, 一个线程块只能在一个SM上被调度. SM一般可以调度多个线程块, 这要看SM本身的能力. 那么有可能一个kernel的各个线程块被分配多个SM, 所以grid只是逻辑层, 而SM才是执行的物理层.
下图是我GT 750M的显卡信息:

SM采用的是SIMT(Single-Instruction, Multiple-Thread, 单指令多线程)架构, 基本的执行单元是线程束(wraps), 线程束包含32个线程, 这些线程同时执行相同的指令, 但是每个线程都包含自己的指令地址计数器和寄存器状态,也有自己独立的执行路径.


加法

试着用CUDA编程做一个矩阵加法:

#include <stdio.h>__global__ void add(float * x, float *y, float * z, int n){        int index = threadIdx.x + blockIdx.x * blockDim.x;        int stride = blockDim.x * gridDim.x;                for (int i = index; i < n; i += stride){                z[i] = x[i] + y[i];        }}int main(){        int N = 1 << 20;        int nBytes = N * sizeof (float);        float *x, *y, *z;        x = (float*)malloc(nBytes);        y = (float*)malloc(nBytes);        z = (float*)malloc(nBytes);        for (int i = 0; i < N; i++){                x[i] = 10.0;                y[i] = 20.0;        }        float *d_x, *d_y, *d_z;        cudaMalloc((void**)&d_x, nBytes);        cudaMalloc((void**)&d_y, nBytes);        cudaMalloc((void**)&d_z, nBytes);        cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);        cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);                dim3 blockSize(256);        // 4096        dim3 gridSize((N + blockSize.x - 1) / blockSize.x);                add << < gridSize, blockSize >> >(d_x, d_y, d_z, N);        cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);        float maxError = 0.0;        for (int i = 0; i < N; i++){                maxError = fmax(maxError, (float)(fabs(z[i] - 30.0)));        }        printf ("max default: %.4f\n", maxError);        cudaFree(d_x);        cudaFree(d_y);        cudaFree(d_z);        free(x);        free(y);        free(z);        return 0;}

由于我是用mac ssh访问linux主机的, 所以看中文都是乱码的, 就没打注释. 简单说下.

  • 逻辑部分:
    申请1M的float, 放入10.0. 申请1M的float, 放入20.0, 然后加起来. 但是我们不存在直接看结果的. 就循环计算误差值, 输出最大的那个误差值. 最后看到是0就代表全部计算正确了.
  • CUDA部分:
    cudaMalloc((void**)&d_x, nBytes);是很抢眼的, 意思也很简单, 在GPU中申请空间, 而不是CPU.
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);将CPU中的数据放入到GPU, 注意第二个是源数据, 第三个是方向.
    dim3 blockSize(256);猜猜也知道了, 就是申请256个block. dim3 gridSize()同理.
    最后cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);从GPU中把结果拷贝回CPU, 注意第三个参数和之前的不同.
    记得释放申请的空间.

统一内存

时不我待. 在高版本的CUDA中, 升级了申请空间的操作. CUDA 6.x引入统一内存(Unified Memory). 具体内容建议查阅我给出的链接, 说的非常细致. 简单来说, 就是申请一次就好, 不用先CPU后GPU, 再拷贝来拷贝去, 太傻.

#include <stdio.h>__global__ void add(float * x, float *y, float * z, int n){        // 同之前, 略}int main(){    int N = 1 << 20;    int nBytes = N * sizeof(float);    float *x, *y, *z;    cudaMallocManaged((void**)&x, nBytes);    cudaMallocManaged((void**)&y, nBytes);    cudaMallocManaged((void**)&z, nBytes);    for (int i = 0; i < N; ++i)    {        x[i] = 10.0;        y[i] = 20.0;    }    dim3 blockSize(256);    // 4096    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);        add << < gridSize, blockSize >> >(x, y, z, N);    cudaDeviceSynchronize();        float maxError = 0.0;    for (int i = 0; i < N; i++){        maxError = fmax(maxError, (float)(fabs(z[i] - 30.0)));    }    printf ("max default: %.4f\n", maxError);        cudaFree(x);    cudaFree(y);    cudaFree(z);    return 0;}

之前看起来有多蠢, 现在就有多简洁(手动滑稽). 注意到cudaMallocManaged((void**)&x, nBytes);, 这样的申请就无需再拷贝操作了.
cudaDeviceSynchronize();同步device, 保证结果能正确访问. 具体原理请参看之前链接.


乘法

学过线性代数的都知道矩阵乘法是那种不难单头疼的工作, 原来可能用MATLAB偷懒, 现在可以考虑CUDA了(手动滑稽).
先定义一个Matrix:

struct Matrix{    int width;    int height;    float *elements;};

定义矩阵操作函数和乘法函数:

__device__ float getElement(Matrix *A, int row, int col){        return A->elements[row * A->width + col];}__device__ void setElement(Matrix *A, int row, int col, float value){        A->elements[row * A->width + col] = value;}__global__ void matMulKernel(Matrix *A, Matrix *B, Matrix *C){        float Cvalue = 0.0;        int row = threadIdx.y + blockIdx.y * blockDim.y;        int col = threadIdx.x + blockIdx.x * blockDim.x;                for (int i = 0; i < A->width; ++i)        {                Cvalue += getElement(A, row, i) * getElement(B, i, col);        }        setElement(C, row, col, Cvalue);}

主函数并不难写, 主要是理解调度原理:

int main(){    int width = 1 << 10;    int height = 1 << 10;        Matrix *A, *B, *C;        cudaMallocManaged((void**)&A, sizeof(Matrix));    cudaMallocManaged((void**)&B, sizeof(Matrix));    cudaMallocManaged((void**)&C, sizeof(Matrix));        int nBytes = width * height * sizeof(float);        cudaMallocManaged((void**)&A->elements, nBytes);    cudaMallocManaged((void**)&B->elements, nBytes);    cudaMallocManaged((void**)&C->elements, nBytes);    A->height = height;    A->width = width;    B->height = height;    B->width = width;    C->height = height;    C->width = width;        for (int i = 0; i < width * height; ++i)    {        A->elements[i] = 1.0;        B->elements[i] = 2.0;    }    dim3 blockSize(32, 32);    dim3 gridSize((width + blockSize.x - 1) / blockSize.x,        (height + blockSize.y - 1) / blockSize.y);    matMulKernel << < gridSize, blockSize >> >(A, B, C);    cudaDeviceSynchronize();        float maxError = 0.0;    for (int i = 0; i < width * height; ++i)        maxError = fmax(maxError, fabs(C->elements[i] - 2 * width));    printf ("max fault: %f\n", maxError);    return 0;}

和加法差不多, 就是多些矩阵操作. 不多说了~
然后矩阵计算代码参考这篇文章.


最后

CUDA近期就先更到这里, 下次一更新应该是内核必须懂部分的. 喜欢记得点赞哦, 有意见或者建议评论区见~


(0)

相关推荐

  • long、int转为byte[]

    很多数据类型的基本数据结构都是byte[],如何快速初始化byte[]的length位为0. 1. 一位一位的赋值 byte[] data = new byte[10000]; for (int i ...

  • CUDA学习2-编程部分

    CUDA编程 函数声明 host:主机端,通常指CPU device:设备端,通常指GPU(数据可并行) kernel:数据并行处理函数,在主机端调用kernel可以在设备端创建大量轻量级线程 hos ...

  • CUDA-GDB安装+环境配置

    在GPU上开发大规模并行应用程序时,需要一个调试器,GDB调试器能够处理系统中每个GPU上同时运行的数千个线程.CUDA-GDB提供了无缝的调试体验,可以同时调试应用程序的CPU和GPU部分. 就像G ...

  • 英特尔独显GPU将采用全新架构、独特外观,2020年问世

    相比处理器市场,全球GPU市场的总价值并不高,高性能GPU一年的营收也就100多亿美元而已,主要掌握在AMD.NVIDIA两家公司中,Q3季度NVIDIA出货的独显份额达到了74%,AMD大约是25. ...

  • Magic Leap将在夏天正式发布产品,GPU还采用帕斯卡架构

    可能你不知道,但是增强现实公司Magic Leap公布他们的产品形态已经是去年十二月底的事情,遥远吧?换句话说,在这过去的超过半年的时间,Magic Leap有做过哪些事情呢?首先他们在三月份的时候表 ...

  • 面向对象编程(二)

    使用__slots__ 正常情况下,当我们定义了一个class,创建了一个class的实例后,我们可以给该实例绑定任何属性和方法,这就是动态语言的灵活性.先定义class: class Student ...

  • 核心银行系统 之二 企业架构与核心建设

    关键字:银行核心系统 企业架构 核心建设 字数:8000字/建议阅读时间:20分钟 作者:CS1026 这篇文章来自前辈CS1026的分享,经过笔者同意后首发于微信公众号平台:小代嘚吧嘚,专注于分享银 ...

  • VS2010 Windows API 串口编程 (二)

    转载请注明出自:blog.csdn.net/mingojiang 目录 一串口通信基础 1.1串口通信原理与特点 1.2串口通信的传输方式 1.3串口通信的同步技术 1.4串行接口标准 二 API函数 ...

  • 深度揭秘:"万物GPU"(The GPU of Everything)

    几个月前,Imagination在其"Inspire年度技术大会"上发布了全新的第十代PowerVR图形处理架构IMG A系列(IMG A-Series).作为Imaginatio ...

  • 三星全新旗舰芯片曝光,用上AMD架构GPU,小米12或将首发

    没有人会怀疑三星在手机研发创新领域的实力,虽然三星手机被各种黑,但三星自身供应链的强大是有目共睹的.在这个世界上三星是唯一能自己掌控并制造屏幕,内存芯片,存储芯片,相机镜头,处理器,电池的厂商.这种能 ...

  • 前端:浏览器、GPU 工作原理简要及动画编程启示

    " 最近作者在 VIPKID 企业内部做了一次关于'动画增强技术方案'的分享,在原分享的基础之上,加入了对浏览器工作原理的考察,并补充动画编码启示若干,烩成此篇,欢迎讨论雅正.本文大约 33 ...

  • AMD确认CDNA2架构计算加速卡将搭载双GPU核心

    6月12日,据外媒报道,Linux系统最近的更新代码中显示,AMD代号为"Aldebaran"的计算加速卡将配备两个GPU核心,印证了此前业界对于CDNA2架构计算加速卡的部分猜测 ...