GPU

GPU加速是通过大量线程并行实现的

GPU 和CPU通信 (GPU与CPU通过PCIe总线连接)

硬件资源：

Streaming multiprocessor: GPU大核

流多处理器，多个SP+其他资源(warp scheduler, register, shared memory等)，

Streaming processor/ CUDA core/核心

流处理器, 基本计算单元，可以并行运算的单元，被分组为warp

Wrap

一个warp包含了多个32个整数倍的sp, Wrap是GPU调度单位，用一个warp中线程执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径

软件资源 cuda program：

Grid: 多个block构成一个Grid

Block 多个threads会被组成一个block，block内的线程共享内存通信

Thread 一个cuda的并行程序会被很多个threads执行

一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识

kernel的这种线程组织结构天然适合vector,matrix等运算，如我们将利用上图2-dim结构实现两个矩阵的加法，每个线程负责处理每个位置的两个元素相加，代码如下所示。线程块大小为(16, 16)，然后将N*N大小的矩阵均分为不同的线程块来执行加法运算。

CUDA

定义

A general purpose parallel computing platform and programming model that leverages the parallel compute engine in NVIDIA GPUs to solve many complex computational problems in a more efficient way than on a CPU.

CUDA是NVIDIA推出的用于自家GPU的并行计算框架，也就是说CUDA只能在NVIDIA的GPU上运行，而且只有当要解决的计算问题是可以大量并行计算的时候才能发挥CUDA的作用。

CUDA programming

在 CUDA 的架构下，一个程序分为两个部份：host 端和 device 端。Host 端是指在 CPU 上执行的部份，而 device 端则是在显示芯片上执行的部份。Device 端的程序又称为 “kernel”。通常 host 端程序会将数据准备好后，复制到显卡的内存中，再由显示芯片执行 device 端程序，完成后再由 host 端程序将结果从显卡的内存中取回。执行流程:

1. 分配host内存，并进行数据初始化；
2. 分配device内存，并从host将数据拷贝到device上；
3. 调用CUDA的核函数在device上完成指定的运算；
4. 将device上的运算结果拷贝到host上；
5. 释放device和host上分配的内存。

内存模型

每个线程有自己的私有本地内存（Local Memory），而每个线程块有包含共享内存（Shared Memory）,可以被线程块中所有线程共享，其生命周期与线程块一致。此外，所有的线程都可以访问全局内存（Global Memory）。还可以访问一些只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）。

重要的APIs

// 在device上申请一定字节大小的显存
cudaError_t cudaMalloc(void** devPtr, size_t size);
// 释放分配的内存使用cudaFree函数
cudaError_t cudaFree(void** devPtr);
//负责host和device之间数据通信的cudaMemcpy函数
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)
// CUDA unified memory
cudaError_t cudaMallocManaged(void **devPtr, size_t size, 
  

例子

int main()
{
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    // 申请host内存
    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;
    }

    // 申请device内存
    float *d_x, *d_y, *d_z;
    cudaMalloc((void**)&d_x, nBytes);
    cudaMalloc((void**)&d_y, nBytes);
    cudaMalloc((void**)&d_z, nBytes);

    // 将host数据拷贝到device
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
    // 定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    // 执行kernel
    add << < gridSize, blockSize >> >(d_x, d_y, d_z, N);

    // 将device得到的结果拷贝到host
    cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);

    // 检查执行结果
    float maxError = 0.0;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    std::cout << "最大误差: " << maxError << std::endl;

    // 释放device内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    // 释放host内存
    free(x);
    free(y);
    free(z);

    return 0;
}

单独在host和device上进行内存分配，并且要进行数据拷贝，这是很容易出错的。好在CUDA 6.0引入统一内存（Unified Memory）来避免这种麻烦，简单来说就是统一内存使用一个托管内存来共同管理host和device中的内存，并且自动在host和device中进行数据传输。CUDA中使用 cudaMallocManaged 函数分配托管内存， 利用统一内存，可以将上面的程序简化如下：

unsigned int flag=0);
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;
    }

    // 定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    // 执行kernel
    add << < gridSize, blockSize >> >(x, y, z, N);

    // 同步device 保证结果能正确访问
    cudaDeviceSynchronize();
    // 检查执行结果
    float maxError = 0.0;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    std::cout << "最大误差: " << maxError << std::endl;

    // 释放内存
    cudaFree(x);
    cudaFree(y);
    cudaFree(z);

    return 0;
}

相比之前的代码，使用统一内存更简洁了，值得注意的是kernel执行是与host异步的，由于托管内存自动进行数据传输，这里要用cudaDeviceSynchronize()函数保证device和host同步，这样后面才可以正确访问kernel计算的结果。

CUDA initialization

作用:

其中一个就是创建 cuda context。即调用这些函数的时候，需要已经有context 存在了。cuda context 非常重要，它作为一个容器，管理了所有对象的生命周期，大多数的CUDA函数调用需要contex.

一个device 对应一个context，所有线程都可以使用。

创建cuda context

隐式调用 （cuda runtime 软件层的库， 是隐式调用）

cuda runtime创建的context 是针对所有线程的，即一个device 对应一个context，所有线程都可以使用。

cuda runtime 不提供API直接创建CUDA context，而是通过lazy initialization。在调用每一个CUDART库函数时，它会检查当前是否有context存在，假如需要context，那么才自动创建。

cuda runtime将context和device的概念合并了，即在一个gpu上操作可看成在一个context下。因而cuda runtime提供的函数形式类似cudaDeviceSynchronize()而不是与driver API 对应的cuCtxSynchronize()。

显式调用 （cuda driver API，驱动层的库，显式调用）

cuda driver API 创建的context是针对一个线程的，即一个device，对应多个context，每个context对应多个线程，线程之间的context可以转移。

在driver API中，每一个cpu线程必须要创建 context，或者从其他cpu线程转移context。如果没有context，就会报错。怎样才回到导致报错呢？即如果没有创建context，就直接调用 driver api创建上面那些对象，就会报错。因为上面的那些对象在runtime 和driver api 中都有函数可以创建。因此，注意注意！！！

每个cpu线程都有一个current context的栈，新建新的context就入栈。针对每一个线程只能有一个出栈变成可使用的current context，而这个游离的context可以转移到另一个cpu线程，通过函数cuCtxPushCurrent/cuCtxPopCurrent来实现。

当context被销毁，里面分配的资源也都被销毁，一个context内分配的资源其他的context不能使用。

注意：

1、隐式调用的context是primary context； 显示调用的context是standard context

2、每次cuda初始化比较费时间，其中一个工作可能就是使用runtime 进行了隐式调用context。因此，如果要避免这部分，有一个方法就是使用cudasetdevice 或者 cudaFree(0) 提前创建context。

The canonical way to force runtime API context establishment is to call cudaFree(0). If you have multiple devices, call cudaSetDevice() with the ID of the device you want to establish a context on, then cudaFree(0) to establish the context.

CUDNN

是NVIDIA打造的针对深度神经网络的加速库，是一个用于深层神经网络的GPU加速库。如果你要用GPU训练模型，cuDNN不是必须的，但是一般会采用这个加速库

• ← Previous Post
• Next Post →