DeepDriving | CUDA编程-02：初识CUDA编程

本文来源公众号“DeepDriving”，仅用于学术分享，侵权删，干货满满。

原文链接：CUDA编程-02：初识CUDA编程

上一篇文章DeepDriving | CUDA编程-01：搭建CUDA编程环境-CSDN博客介绍了如何搭建CUDA编程环境，从这篇文章开始正式开始介绍如何使用CUDA进行编程。

1 异构计算架构

如下图所示，一个典型的异构计算架构节点由一个多核CPU和一个或多个GPU组成，每个CPU和GPU都是独立的设备，它们之间通过PCIe总线连接。GPU作为CPU的协处理器用于执行一些并行计算任务。CPU适合用于做一些逻辑控制任务，而GPU则适合作为CPU的协处理器用于做一些大计算量的数据并行计算任务。

heterogeneous_architecture

一般我们称CPU为host，称GPU为device，相应地，一个异构计算的应用程序代码也被分为两部分：运行在CPU上的程序被称为host代码，运行在GPU上的程序被称为device代码。

heterogeneous_architecture

2 一个Hello World

在我们学习一种新的编程语言的时候，一般都是先从打印一句"Hello World"开始，今天开始学习CUDA编程，按照国际惯例，也从打印"Hello World"开始。

首先新建一个CUDA C源文件hello_world.cu，然后输入下面的内容：

#include <stdio.h>int main(void)
{printf("Hello World from CPU\n");return 0;
}

用nvcc进行编译生成可执行文件：

nvcc hello_world.cu -o hello_world

运行可执行文件hello_world，没有意外的话就能在终端看到打印出的"Hello World from CPU"了。这是在CPU上运行代码打印了这句话，那么怎么在GPU上打印呢？

要在GPU上运行程序，我们需要写一个能在GPU上执行的函数，然后在CPU上调用这个函数。来看一个例子：

#include <stdio.h>__global__ void HelloFromGPU(void)
{printf("Hello from GPU\n");
}int main(void)
{printf("Hello from CPU\n");HelloFromGPU<<<1, 5>>>();cudaDeviceReset();return 0;
}

在上面的这段代码里，我添加了一个用__global__函数类型限定符修饰的函数HelloFromGPU，然后在main函数中去调用它。在CUDA中，有以下3种函数类型限定符：

_global_

被__global__函数类型限定符修饰的函数被称为内核函数，该函数在host上被调用，在device上执行，只能返回void类型，不能作为类的成员函数。调用__global__修饰的函数是异步的，也就是说它未执行完就会返回。

_device_

被__device__函数类型限定符修饰的函数只能在device上被调用，在device上执行，用于在device代码中内部调用。

_host_

被__host__函数类型限定符修饰的函数只能在host上被调用，在host上执行，也就是host上的函数，__host__函数类型限定符可以省略。

与调用一般CPU函数不同的是，调用HelloFromGPU函数的时候会在后面写上<<< >>>，意思这是从host端到device端的内核函数调用，里面的参数是执行配置，用来说明使用多少线程来执行内核函数。一个内核函数是通过一组线程来执行的，所有线程执行的同样的代码，我这里设置的是用了5个线程来执行。程序编译成功后运行得到如下结果：

Hello from CPU
Hello from GPU
Hello from GPU
Hello from GPU
Hello from GPU
Hello from GPU

可以看到，内核函数里的打印语句被执行了5次。

3 用CUDA实现数组相加

一个典型的CUDA程序结构一般由以下5个步骤组成：

分配GPU内存；
从CPU内存中拷贝数据到GPU内存中；
调用CUDA内核函数执行程序指定的计算任务；
从GPU内存中把数据拷贝回CPU内存中；
释放GPU内存；

下面以一个数组相加的例子来展示以上过程，数组相加的计算过程比较简单：数组a和数组b中对应下标的元素相加然后存入数组c中。

vector_add

首先来看一下在CPU上实现数组相加的代码：

#include <iostream>void VectorAddCPU(const float *const a, const float *const b, float *const c,const int n) {for (int i = 0; i < n; ++i) {c[i] = a[i] + b[i];}
}int main(void) {// alloc memory for hostconst size_t size = 1024;float *ha = new float[size]();float *hb = new float[size]();float *hc = new float[size]();for (int i = 0; i < size; ++i) {ha[i] = i;hb[i] = size - i;}VectorAddCPU(ha, hb, hc, size);delete[] ha;delete[] hb;delete[] hc;return 0;
}

函数VectorAddCPU用了一个for循环在CPU上实现数组相加的过程。如果要用GPU来实现该过程，则调用CUDA的API按照前面说的5个步骤编写代码：

#include <cuda_runtime.h>
#include <iostream>__global__ void VectorAddGPU(const float *const a, const float *const b,float *const c, const int n) {int i = blockDim.x * blockIdx.x + threadIdx.x;if (i < n) {c[i] = a[i] + b[i];}
}int main(void) {// 分配CPU内存const size_t size = 1024;float *ha = new float[size]();float *hb = new float[size]();float *hc = new float[size]();for (int i = 0; i < size; ++i) {ha[i] = i;hb[i] = size - i;}// 分配GPU内存float *da = nullptr;float *db = nullptr;float *dc = nullptr;cudaMalloc((void **)&da, size);cudaMalloc((void **)&db, size);cudaMalloc((void **)&dc, size);// 把数据从CPU拷贝到GPUcudaMemcpy(da, ha, size, cudaMemcpyHostToDevice);cudaMemcpy(db, hb, size, cudaMemcpyHostToDevice);cudaMemcpy(dc, hc, size, cudaMemcpyHostToDevice);const int thread_per_block = 256;const int block_per_grid = (size + thread_per_block - 1) / thread_per_block;// 调用核函数VectorAddGPU<<<block_per_grid, thread_per_block>>>(da, db, dc, size);// 把数据从GPU拷贝回CPUcudaMemcpy(hc, dc, size, cudaMemcpyDeviceToHost);// 释放GPU内存cudaFree(da);cudaFree(db);cudaFree(dc);// 释放CPU内存delete[] ha;delete[] hb;delete[] hc;return 0;
}

在这段代码里，用到了几个CUDA的内存管理函数：

cudaMalloc

函数原型：

cudaError_t cudaMalloc(void** devPtr, size_t size)

该函数用于在GPU上分配指定大小的内存空间，类似于标准C语言中的malloc函数。

cudaMemcpy

函数原型：

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)

该函数用于内存拷贝，类似于标准C语言中的memcpy函数。数据拷贝的流向由参数kind指定，分为以下4种方式：

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

从它们的字面意思就能知道，如果是从host拷贝数据到device，那么kind参数应该设置为cudaMemcpyHostToDevice；如果是从device拷贝数据到host，那么kind参数应该设置为cudaMemcpyDeviceToHost。上面的代码体现了这一点。

cudaFree

udaError_t cudaFree(void* devPtr)

该函数用于释放已分配的GPU内存空间，类似于标准C语言中的free函数。

在上面的代码中，内核函数VectorAddGPU用于实现数组相加，与前面的代码中VectorAddCPU函数不同的是，VectorAddGPU函数里面并没有用for循环，因为在GPU中是将数据进行并行化划分，然后通过线程组去实现计算过程的，线程组中的每个线程都是执行c[i] = a[i] + b[i]这个计算过程。在这个程序里，数组的长度为1024，我设置了每个线程组中的线程数量为256，总共用了4个线程组去进行计算。

关于GPU中线程和线程组的相关知识，我将在下一篇文章中进行阐述。