首页手记 CUDA编程(三): GPU架构了解一下!

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

标签：

机器学习人工智能

前言

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

GPU架构

GPU处理单元

从这张GPU概念内核图开始讲起, 会发现和CPU内核是不同的, 少了三级缓存, 分支预测等等. 但是增加了ALU的数量, 扩大了上下文存储池(Pool of context storge). 其实这里的ALU就是实际的CUDA核, 上下文会对应实际的warp.
可以看到, 上下文存储池分成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的两倍左右. 而且, 336.5GB/s看起来很爽吧, 可是你在水平不行的情况下, 可能连1%的性能都发挥不出来, 也就是说, 不懂CUDA编程, 直接移植c代码, 还不如用CPU直接跑.

GPU线程与SM

由于目前还没有完全依靠GPU运行得机器, 一般来说, 都是异构的, CPU+GPU. 这一点是要特别注意的, 也就是Host与Device. 而通常，我们将在 CPU上执行的代码称为主机代码, 而将在GPU上运行的代码称为设备代码.

GPU线程

在CUDA架构下, 显示芯片执行时的最小单位是thread. 数个thread可以组成一个block. 一个block中的thread能存取同一块共享的内存(shared memory), 而且可以快速进行同步的动作, 特别要注意, 这是块(block)同步. 不同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才是执行的物理层.

通常, 当调用要在GPU上运行的函数时, 我们将此种函数称为已启动的核(kernel)函数要补充的是, 核函数启动方式为异步: CPU代码将继续执行而无需等待核函数完成启动. 调用CUDA运行时提供的函数cudaDeviceSynchronize将导致主机(CPU)代码暂作等待, 直至设备 (GPU)代码执行完成, 才能在CPU上恢复执行. 否则CPU完事了, GPU还在算, 到哪里去找计算返回的结果?

下图是我GT 750M的显卡信息:

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

这个warp以后还会说到.

加法

试着用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;
}

说下代码内容:

逻辑部分:
申请1M的float, 放入10.0. 申请1M的float, 放入20.0, 然后加起来. 但是我们不存在直接看结果的, 数量太多了, 可以考虑头打印5个值, 尾打印5个值. 这里就循环计算误差值, 输出最大的那个误差值. 最后看到是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, 注意第三个参数和之前的不同.
记得释放申请的空间.
然后看到核函数, _global_ 关键字表明以下函数将在GPU上运行并可全局调用, 而在此种情况下, 则指由CPU或GPU调用.
int index = threadIdx.x + blockIdx.x * blockDim.x;是获取线程位置. 其实不管是多少维度的矩阵, 在计算机内存中都是线性存储的, 所以要来一个一维展开. 依据之前说的结构, CUDA核函数可以访问能够识别如下两种索引的特殊变量: 正在执行核函数的线程(位于线程块内)索引和线程所在的线程块(位于网格内)索引. 这两种变量分别为threadIdx.x和blockIdx.x. CUDA核函数可以访问给出块中线程数的特殊变量: blockDim.x.
int stride = blockDim.x * gridDim.x;是计算网格中的总线程数，即网格中的线程块数乘以每个线程块中的线程数：gridDim.x * blockDim.x. 这里的操作也是优化手段, 以后会再说到.

统一内存

时不我待. CUDA的最新版本(版本6和更高版本)已能轻松分配可用于CPU主机和任意数量GPU设备的内存. CUDA 6.x引入统一内存(Unified Memory). 具体内容建议查阅我给出的链接, 说的非常细致. 简单来说, 就是申请一次就好, 不用先CPU后GPU, 再拷贝来拷贝去, 太傻. 但是注意, 之后还是会说把Host数据往Device迁移的操作, 为了提升性能.

#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();之前也说过了, 主要是由于核函数是异步启动, 需要CPU等待GPU代码执行完成.

数据迁移

在主机到设备和设备到主机的内存传输过程中, 可以使用一种技术来减少页错误和按需内存迁移成本, 此技术称为异步内存预取. 这在后边的性能提升部分会说, 这里先简单演示下用法:

int main()
{
    // 同前, 略
    int deviceId;
    cudaGetDevice(&deviceId);                                        

    cudaMemPrefetchAsync(x, nBytes, deviceId); 
    cudaMemPrefetchAsync(y, nBytes, deviceId);
    cudaMemPrefetchAsync(z, nBytes, deviceId);
                                                                                                                              
    dim3 blockSize(256);
    // 同前, 略 
}