Learn CUDA Programming(1)

Access the Power of GPU

  • 一般有三种方法来实现加速:Libraries or OpenACC or Programming Languages
  • 函数库:

例如:

cuBLAS: 线性代数库

cuSPARSE: 稀疏矩阵的线性代数库

cuDNN: 深度神经网络

GPU Architecture

异构计算(Heterogeneous Computing)

Terminology:

  • Host: The CPU and tis memory(host memory) 主机端(CPU)
  • Device: The GPU and its memory(device memory) 设备端(GPU)

CPU与GPU的简易架构图

  • CPU有很多逻辑控制单元,而GPU有很多处理核心
  • CPU是缓存优化的处理器,缓存特别大,而且有很多控制单元
  • GPU是一个并行吞吐优化的处理器,计算核心特别多
  • 因此我们知道,GPU的单精度处理能力比CPU快很多的原因其实就是GPU将更多的晶体管用于了计算而不是缓存

程序的开发中如何选择CPU和GPU

  • 从图中可以看到CPU的缓存延迟非常短,这也对应它的架构里有很多复杂的逻辑控制
  • GPU则不同,图中的W是warp,1 warp = 32 Threads,GPU是以warp为单位进行切换的。我们可以通过海量的线程的切换来对缓存延迟做隐藏,如果你的程序是计算密集型的,而且有很好的并行特性,那么gpu是更好的选择。
  • GPU的架构有两个部分组成,第一部分是计算单元,第二部分是内存单元
  • 图中的绿色部分就是一个CUDA Core
  • 从图中可以看到,GPU和CPU是通过PCle进行链接的
  • 一个GPU有很多SM(流多处理器)组成,每一个SM有很多核,因此有越多的SM,那么GPU就能在同一时间内处理更多的任务。

SM的架构

  • Kepler架构中的SM称为SMX,表示处理器更强劲,其逻辑控制单元是一个整体,控制192个CUDA Cores
  • Maxwell架构中的SM称为SMM,其逻辑控制单元有4个,每个控制32个CUDA Cores
  • 可以简单认为,一个SMM中包含了4个SMX,这样的话就避免了因为逻辑控制单元过少而造成核心的冗余,效率更高。从架构上来看,Maxwell一个CUDA核心相比于Kepler提升了35%

Memory Hierarchy

因为经常要跟内存打交道,所以需要对GPU的内存架构非常熟悉,GPU的架构分为三个层次,如图所示,和CPU类似,可以看做成一个金字塔,从下往上越来越快,首先是Global Memory,有很大的内存空间,它的空间很大,但是缓存延迟很高,因此我们需要加入一个缓存,称之为L2缓存,GPU的缓存要像Global Memory一样读取数据的话,首先我们要看这些数据是否在L2中已经有缓存,有的话称之为缓存命中,这样就不用再从Global Memory里面读数据,这样就加快了读写。L2缓存是多个缓存所共享的。再上面有很多寄存器,L1缓存等。对于不同的内存有不同的使用和优化方案。

GPU in Computer System

CUDA Programming Basics

配置CUDA开发环境

  • 注意如果使用C/C++进行开发的话,我们使用的是NVCC的编译器

Heterogeneous Computing 异构编程

  • 一般我们把GPU运行的程序构造成一个函数,在我们需要的时候进行调用,这个函数就是在GPU上运行

  • 这个函数的之前和之后都是在CPU上单线程运行的,因此我们把整个程序分为:顺序执行,并行执行,顺序执行

  • 顺序执行就是利用CPU的一个线程来执行

  • 并行执行就是用GPU的多线程来进行计算

  • 因此如果一部分功能是计算密集型的,我们就可以把它写成一个函数,对它进行并行,并不是把所有的代码都在GPU上运行,而是把需要的部分放到GPU上

CUDA Kernels

这里我们讲解一个概念,我们要将一部分代码在gpu上运行,这部分代码是一个函数,我们称为一个kernel(核函数),CPU(Host)执行functions,GPU(Device)执行kernels。

Hello World

//hello_world.c
#include <stdio.h>

void hello_world_kernel()
{
    printf("Hello World\n");
}

int main()
{
    hello_world_kernel();
}

Compile & Run:
gcc hello_world.c
./a.out

在GPU上执行:

//hello_world.cu
#include <stdio.h>

_global_ hello_world_kernel()
{
    printf("Hello World\n");
}

int main()
{
    hello_world_kernel<<<1.1>>>();
}

Compile & Run:
nvcc hello_world.cu
./a.out

不同之处

  • 文件后缀名改为.cu
  • _global_表示了该函数为一个核函数,标识这个函数会在gpu上运行
  • “<<<...,...>>>”具体含义是GPU中的一个配置
  • 编译器改为nvcc

GPU memory management

hello world程序非常简单明了,但是在编程过程中我们经常要和内存打交道,那么如何进行GPU内存管理呢?

我们很熟悉对CPU的内存管理,我们会使用

malloc();//动态分配内存
memset();//对空间进行初始化
free();//释放一段内存

对应于,在GPU中也有类似的函数,它们分别是

cudaMalloc(void** pointer, sie_t nbytes);//pointer指定空间,大小是nbytes
cudaMemset(void* pointer, int value, size_t count);
cudaFree(void* pointer);

例如:

int nbytes = 1024*sizeof(int);
int* d_a = 0;
cudaMalloc((void**)&d_a, nbytes);
cudaMemset(d_a, 0, nbytes);
cudaFree(d_a);

因此一定要分清楚,哪部分数据在cpu上,哪部分数据在gpu上

Data Copies

__host__ cudaMemcpy(void* dst, void* src, size_t nbytes, cudaMemcpyKind direction);
  • __global__是在gpu上进行的,__host__就是在cpu上执行
  • 四个形参:拷贝目的指针,拷贝原指针,拷贝大小,拷贝方向
  • 线程阻塞的,拷贝不完成,cpu不会执行下面的代码
  • 常见拷贝方向有:cudaMemcpyHostToDevice(CPU到GPU)、cudaMemcpyDeviceToHost(GPU到CPU)、cudaMemcpyDeviceToDevice(GPU到GPU)
  • 设定拷贝方向时,需要注意src和dst的指针是CPU的还是GPU的,要和拷贝方向一致
  • 上述函数为线程阻塞的,同样还有一个异步的函数:cudaMemcpyAsync();

三步流程

第一步:通过调用cudaMemcpy函数将CPU中的数据拷贝到GPU中,通过PCI Bus,方向就是HostToDevice

第二步:通过CPU启动kernel核函数,然后开始一个并行计算

Kernel<<<grid, block>>>();

第三步:计算完以后,再调用cudaMemcpy将数据从GPU端拷贝到CPU端,方向是DeviceToHost

例子:

代码分为五个步骤:

  • 在CPU端分配n个整型变量的空间
  • 在GPU端分配n个整型变量的空间
  • 初始化GPU内存为0
  • 将数据从GPU拷贝到CPU
  • 打印变量
#include<stdio.h>

int main()
{
    int dimx = 16;
    int num_bytes = dimx * sizeof(int);
    int *d_a = 0, *h_a = 0;//设备端和主机端的指针
    
    //在CPU端分配n个整型变量的空间
    h_a = (int*)malloc(num_bytes);
    //在GPU端分配n个整型变量的空间
    cudaMalloc((void**) &d_a, num_bytes);
    
    if(0 == h_a || 0 == d_a)
    {
        printf("Couldn't allocate memory\n");
        return 1;
    }
    
    //初始化GPU内存为0
    cudaMemset(d_a,0,num_bytes);
    
    //将数据从GPU拷贝到CPU
    cudaMemcpy(h_a, d_a, num_bytes, cudaMemcpyDeviceToHost);
    
    //打印变量
    for(int i=0; i < dimx; i++)
    {
        printf("%d\t", h_a[i]);
    }
    free(h_a);//主机端释放
    cudaFree(d_a);//设备端释放
    return 0;
}
  • 这里注意,cudaMemcpy是CPU阻塞的,也就是说,不执行完该函数CPU就不会一直向下执行。如果换成一个异步的版本,CPU会在拷贝还没有完成的时候就开始打印,打印的时候很多是错误的。

上面的示例代码缺少了GPU上面的并行计算,那么我们下一面将开始讲解如何写Kernel函数。