《计算机视觉与CUDA加速》 01 入门CUDA运算

Author : Kevin
Copyright : Kevin Ren from HIT
Description : CUDA compute reciprocals

https://blog.csdn.net/h_l_dou/article/details/84199350

1.1 主机和设备

/*
 ============================================================================
 Name        : myFirst.cu
 Author      : Kevin
 Version     :
 Copyright   : Kevin Ren from HIT
 Description : CUDA compute reciprocals
 ============================================================================
 */

#include <cstdio>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>


//define of kernel function to add
__global__ void gpuAdd(int d_a,int d_b,int *d_c){
	*d_c = d_a + d_b;
}

int main(void){
	// host
	int h_c;
	// device
	int *d_c;
	cudaMalloc((void**)&d_c,sizeof(int));
	gpuAdd << <1,1>> >(1,4,d_c);
	cudaMemcpy(&h_c,d_c,sizeof(int),cudaMemcpyDeviceToHost);
	printf("1+4=%d\n",h_c);
	cudaFree(d_c);

	return 0;
}

在上面代码中一共有两个函数，其中$gpuAdd$ 被$\_\_global\_\_$修饰，可以看出这是在GPU外设上执行的，而main函数则是在cpu上执行的。24行我们在创建了变量h_c作为主变量，我们需要把这个变量放到GPU上去运算，就需要想ANSIC一样在GPU上申请空间。也就是$cudaMalloc（）$函数。随后，如果计算结果需要被运用到cpu上，我们就需要把结果从GPU上拷贝回来，也就是上面29行的$cudaMemcpy()$函数。随后我们要释放内存，否则可能显存会溢出，导致无法运算。

1.2 什么是kernel？

使用C语言撰写的用于GPU设备上的代码我们叫做kernel（核心），他们是从host代码中通过kernel call产生的。 通常来讲，kernel就是产生很多并行的数据块和程序线，在GPU上运行。Kernel code和标准C语言非常像，只是这些代码可以再GPU上生成很多并行的线。 kernel的写法看起来有些奇怪。

|kernel << <number of blocks, number of thread per block, size of shared memory> >> (parameters)

简单地说，核的结构从kernel名称开始，随后是 << < 标志，然后就是一些核内参数，在> >>之后，我们输入传入核的参数。

在我们第一节的例子里，我们使用了这样一个核函数：

gpuAdd << <1,1>> >(1,4,d_c);

表示我们需要一块单线程的区域，大小程序自己看着办，传入参数是1,4，和一个指针。这个指针指向的是gpu上的内存区域（显存区域），这也是为什么我们最后要用cudaMemcpy从显存指针拷贝到host上。

【注意，显存指针不能直接指向host，否则会导致程序崩溃】

可以理解，他们都不在一个硬件上，本身联系是不大的。

上面这个例子我们使用了一区块单线程，这显然无法体现cuda加速的优势。那么我们如何进行并行运算呢？

1.3 设置kernel call参数【核心调用】

进行多线程运算，我们首先要进行kernel call的参数设置。

第一，我们要设置区块数和线程数，通常来件，我们设置每个区块的线程数不超过512或者1024。每个区块中的参数都可以通过他们所在的连续存储空间进行共享。我们只能设置数量，不能控制具体的区块和线程位置。

假如我们在一个块里设置了500个线，那么我们需要这样做：

gpuAdd<< <1,500>> >(1,4,d_c)

当然我们也可以选择两个区块，每个区块250线。

对于我们利用cuda进行计算机视觉的图像处理加速来说，我们通常需要避免使用一个维度的区块或者线。

GPU支持三维栅格[简单的说就是三维矩阵]的区块和线，例如：

mykernel<< <dim3(Nbx, Nby, Nbz), dim3(Ntx, Nty, Ntz)> >>()

上面的$N_{b_x},N_{b_y},N_{b_z}$表示在三个方向上需要的block数量。对于线来说也是一样的。如果高维度的参数空缺，那么gpu会把他们默认为1 。

所以举个例子，我们要处理一个图片，你可以创造一个$16\times 16$的区块，每一个都包含$16\times 16$的线，那么你的核心调用应该这样写：

mykernel << < dim3(16,16), dim3(16, 16)> >>()

总而言之，核心调用(kernel call)的参数设置在创建cuda运算的时候非常重要。我们应该选择合适的大小，合理的利用GPU资源。

接下来我们会讲一些重要的CUDA函数。

1.4 CUDA API函数

在我们的cuda加法例子中，我们接触到了一下有趣的函数和关键字，他们不属于C和C++项目。下面我们就详细的讲解一下这些关键字和函数。

1.4.1 __global__

CUDA中有三个标准关键字，他们分别是

__global__
__device__
__host__

从他们的名字我们基本可以看穿这些关键字的意义。

不过首先我要强调一点，这些关键字都只能被host调用，不能写在kernel code里。

如果你需要这个函数在GPU上运行，那么你就需要给他限定为device，二host关键字是用来定义那些只能被host调用的函数。如果不写的话，编译器默认所有函数都是host函数。这些关键字都可以在定义函数的时候同时被使用。如果同时使用，这些函数在编译的时候会生成两个函数，一个给host用，一个给device用。

1.4.2 cudaMalloc

我们来看一下他的函数原型

cudaMalloc（void ** d_pointer, size_t size）
// e.g. cudaMalloc((void**)&d_c,sizeof(int));

例子中给出，我们申请了一块区域，大小为一个int型，地址就是d_c指向的地方。

1.4.3 cudaMemcpy

这个函数也类似C语言的标准函数Memcpy。它被用来从一个block中复制内存中的数据回host。函数原型如下

cudaMemcpy(void ** dst_ptr, const void * src_ptr, size_t size, enum cudaMemcpyK)
// e.g. cudaMemcpy(&h_c,d_c,sizeof(int),cudaMemcpyDeviceToHost);

这四个参数分别是

• 数据位置指针
• 主机位置指针
• 区域大小
• 复制方式（对象）

1.4.4 cudaFree

释放内存，函数原型如下

cudaFree（void * d_ptr）

小结一下，cuda还有很多其他的API函数，以上三个是我们最常用的。其他的细节请大家移步NVIDIA官网的CUDA programming guide

1.5 向CUDA传递参数

1.5.1 传值

	gpuAdd << <1,1>> >(1,4,d_c);

传值固然简单名了，但是对于函数来说，他们的权限太低，我们很容易想到c++中传递引用的特效。

1.5.2 传引

我们对先前的加法函数进行一些修改：

__global__ void gpuAdd(int *d_a,int *d_b,int *d_c){
	*d_c = *d_a + *d_b;
}

这里要注意的是，如果是传地址过去，我们在申请空间的时候也要为指针们留位置。

int main(void){
	// host
	int h_a,h_b,h_c;
	h_a = 2;
	h_b = 3;
	// device
	int *d_a,*d_b,*d_c;
	cudaMalloc((void**)&d_a,sizeof(int));
	cudaMalloc((void**)&d_b,sizeof(int));
	cudaMalloc((void**)&d_c,sizeof(int));
    // copy host variable to devices
	cudaMemcpy(d_a,&h_a,sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(d_b,&h_b,sizeof(int),cudaMemcpyHostToDevice);

	gpuAdd << <1,1>> >(d_a,d_b,d_c);
    // copy answer in devices to host
	cudaMemcpy(&h_c,d_c,sizeof(int),cudaMemcpyDeviceToHost);
    
	printf("%d+%d=%d\n",h_a,h_b,h_c);
    
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);

	return 0;
}

我们可以看到，h_a, h_b, h_c 是host内存里的变量，他们就像普通的C语言一样被定义。另一方面，d_a, d_b, d_c 也是属于host内存里的。但是他们指向的是GPU显存中的地址。我们通过下图来了解。

我们知道，d系列的变量就是host和device之间的桥梁，他们存在与host内存，指向device显存。

【注意，当你通过地址传递参数的时候，一定要注意这些指针都是指向device的，否则可能导致程序崩溃。】

小结:通过这一章的完整介绍，我们大概知道了cuda加速中host和device之间的关系。我们将在下一章为大家介绍device中的线。

尽请期待。