本系列博客配备github代码 ,本小节代码见:https://github.com/qixuxiang/cuda_zero_to_one/blob/master/async_copy.cu
和
https://github.com/qixuxiang/cuda_zero_to_one/blob/master/array_sum.cu
知道了CUDA编程基础,我们就来个简单的实战:利用CUDA编程实现两个向量的加法。在实现之前,先简单介绍一下CUDA编程中内存管理API。首先是在device上分配内存的cudaMalloc
函数。
cudaError_t cudaMalloc(void** devPtr, size_t size);
这个函数和C语言中的malloc类似,但是在device上申请一定字节大小的显存,其中devPtr是指向所分配内存的指针。同时要释放分配的内存使用cudaFree函数,这和C语言中的free函数对应。另外一个重要的函数是负责host和device之间数据通信的cudaMemcpy
函数:
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)
其中src指向数据源,而dst是目标区域,count是复制的字节数,其中kind控制复制的方向:cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost
及cudaMemcpyDeviceToDevice
,如cudaMemcpyHostToDevice
将host上数据拷贝到device上。
需要指出的是cudaMemcpy
是阻塞式的API,也就是CPU端代码在调用该API时,只有当该API完成拷贝之后,CPU才能继续处理后面的任务。这有一个好处就是保证了计算结果已经完全从GPU端拷贝到了CPU。同时CUDA也提供了非阻塞拷贝的APIcudaMemcpyAsync
, 非阻塞拷贝也称为异步拷贝,指的是该API在拷贝完成之前就返回,使得CPU可以继续处理后续的代码。异步拷贝API使得CPU与GPU之间的数据拷贝与CPU计算的并发称为可能。如果该API与CUDA中流(Stream)相结合使用,也可以实现数据的拷贝与GPU计算进行并发执行,这一点会在流与并发这一部分进行介绍。
在host和device之间异步拷贝数据的一个简单例子如下:
#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__device__ int devData;
__host__ __device__ int run_on_cpu_or_gpu() {
return 1;
}
__global__ void run_on_gpu() {
printf("run_on_cpu_or_gpu GPU: %d\n", run_on_cpu_or_gpu());
}
int main() {
int val = run_on_cpu_or_gpu();
cudaMemcpyToSymbol(devData, &val, sizeof(int));
printf("run_on_cpu_or_gpu CPU: %d\n", run_on_cpu_or_gpu());
cudaMemcpyFromSymbol(&val, devData, sizeof(int));
run_on_gpu<<<1, 1>>>();
cudaDeviceReset();
return 0;
}
现在我们来实现一个向量加法的实例,这里grid和block都设计为1-dim,首先定义kernel如下:
// 两个向量加法kernel,grid和block均为一维
__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];
}
}
然后按照CUDA程序的执行流程继续编写代码:
1. 分配host内存,并进行数据初始化;
2. 分配device内存,并从host将数据拷贝到device上;
3. 调用CUDA的核函数在device上完成指定的运算;
4. 将device上的运算结果拷贝到host上;
5. 释放device和host上分配的内存。
代码如下:
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;
}
在这里可以附一个完整的利用CUDA 并行化思想来对数组进行求和和CPU求和的对比程序:
// 相关 CUDA 库
#include "cuda_runtime.h"
#include "cuda.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <cstdlib>
using namespace std;
const int N = 100;
// 块数
const int BLOCK_data = 3;
// 各块中的线程数
const int THREAD_data = 10;
// CUDA初始化函数
bool InitCUDA()
{
int deviceCount;
// 获取显示设备数
cudaGetDeviceCount (&deviceCount);
if (deviceCount == 0)
{
cout << "找不到设备" << endl;
return EXIT_FAILURE;
}
int i;
for (i=0; i<deviceCount; i++)
{
cudaDeviceProp prop;
if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 获取设备属性
{
if (prop.major>=1) //cuda计算能力
{
break;
}
}
}
if (i==deviceCount)
{
cout << "找不到支持 CUDA 计算的设备" << endl;
return EXIT_FAILURE;
}
cudaSetDevice(i); // 选定使用的显示设备
return EXIT_SUCCESS;
}
// 此函数在主机端调用,设备端执行。
__global__
static void Sum (int *data,int *result)
{
// 取得线程号
const int tid = threadIdx.x;
// 获得块号
const int bid = blockIdx.x;
int sum = 0;
// 有点像网格计算的思路
for (int i=bid*THREAD_data+tid; i<N; i+=BLOCK_data*THREAD_data)
{
sum += data[i];
}
// result 数组存放各个线程的计算结果
result[bid*THREAD_data+tid] = sum;
}
int main ()
{
// 初始化 CUDA 编译环境
if (InitCUDA()) {
return EXIT_FAILURE;
}
cout << "成功建立 CUDA 计算环境" << endl << endl;
// 建立,初始化,打印测试数组
int *data = new int [N];
cout << "测试矩阵: " << endl;
for (int i=0; i<N; i++)
{
data[i] = rand()%10;
cout << data[i] << " ";
if ((i+1)%10 == 0) cout << endl;
}
cout << endl;
int *gpudata, *result;
// 在显存中为计算对象开辟空间
cudaMalloc ((void**)&gpudata, sizeof(int)*N);
// 在显存中为结果对象开辟空间
cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data*THREAD_data);
// 将数组数据传输进显存
cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice);
// 调用 kernel 函数 - 此函数可以根据显存地址以及自身的块号,线程号处理数据。
Sum<<<BLOCK_data,THREAD_data,0>>> (gpudata,result);
// 在内存中为计算对象开辟空间
int *sumArray = new int[THREAD_data*BLOCK_data];
// 从显存获取处理的结果
cudaMemcpy (sumArray, result, sizeof(int)*THREAD_data*BLOCK_data, cudaMemcpyDeviceToHost);
// 释放显存
cudaFree (gpudata);
cudaFree (result);
// 计算 GPU 每个线程计算出来和的总和
int final_sum=0;
for (int i=0; i<THREAD_data*BLOCK_data; i++)
{
final_sum += sumArray[i];
}
cout << "GPU 求和结果为: " << final_sum << endl;
// 使用 CPU 对矩阵进行求和并将结果对照
final_sum = 0;
for (int i=0; i<N; i++)
{
final_sum += data[i];
}
cout << "CPU 求和结果为: " << final_sum << endl;
getchar();
return 0;
}