Jeremy Lin

上一篇最后有一个“Hello World”的例子，可是和C程序根本没差。现在我们来真正接触CUDA的代码到底要怎么写。

首先，Show the Code：

#include "cuda_runtime.h"
#include <stdio.h>

const int N = 10;

__global__ void add_Jeremy(int*a, int*b, int*c)
{
	int tid = blockIdx.x;
	if (tid < N)
	{
		c[tid] = a[tid] + b[tid];
	}
}

int main()
{
	int a[N], b[N], c[N];
	int *dev_a, *dev_b, *dev_c;

	cudaMalloc((void**)&dev_a, N*sizeof(int));
	cudaMalloc((void**)&dev_b, N*sizeof(int));
	cudaMalloc((void**)&dev_c, N*sizeof(int));

	for (int i = 0; i < N; i++)
	{
		a[i] = -i;
		b[i] = i * i;
	}

	cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);

	add_Jeremy<<<N,1>>>(dev_a, dev_b, dev_c);

	cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost);

	for (int i = 0; i < N; i++)
	{
		printf("%d + %d = %d\n", a[i], b[i], c[i]);
	}

	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	return 0;
}

上面的例子完成的功能是两个向量的相加，够简单~，所以在逻辑上，我们基本不用多说什么。

现在我们关注的是，CUDA的代码与一般的C/C++程序的差别到底在哪里？看上面的code，我想你应该可以看到如下几点差别：

（1）函数类型限定符 __global__

这个限定符告诉编译器，函数应该编译为在设备而不是主机运行，即函数add_Jeremy()将被交给编译设备代码的编译器，而main函数将被交给主机编译器。

————————————————————————————————————————

关于函数类型限定符的完整介绍如下：

__device__

使用__device__限定符声明的函数具有以下特征：

• 在设备执行；
• 仅可通过设备调用。

__global__

使用__global__限定符可将函数声明为内核。此类函数：

• 在设备上执行；
• 仅可通过主机调用。

__host__

使用__host__ 限定符声明的函数具有以下特征：

• 在主机上执行；
• 仅可通过主机调用。

仅使用 _host_ 限定符声明函数等同于不使用 _host_、_device_ 或 _global_ 限定符声明函数，这两种情况下，函数都将仅为主机进行编译。
但 _host_ 限定符也可与 _device_ 限定符一起使用，此时函数将为主机和设备进行编译。

一些限制：

1. _device_ 和 _global_ 函数不支持递归。
2. _device_ 和 _global_ 函数的函数体内无法声明静态变量。
3. _device_ 和 _global_ 函数不得有数量可变的参数。
4. _device_ 函数的地址无法获取，但支持 _global_ 函数的函数指针。
5. _global_ 和 _host_ 限定符无法一起使用。
6. _global_ 函数的返回类型必须为空。
7. _global_ 函数的调用是异步的，也就是说它会在设备执行完成之前返回。
8. _global_ 函数参数将同时通过共享存储器传递给设备，且限制为 256 字节。

————————————————————————————————————————

（2）内建变量  blockIdx

内建变量blockIdx是一个包含三个元素x,y,z的结构体，分别表示当前线程所在块在网格中x,y,z三个方向上的索引。

————————————————————————————————————————

关于内建变量的完整介绍如下：

gridDim

此变量的类型为dim3（dim3是一种整形向量类型，在定义类型为dim3的变量时，未指定的任何组件都被初始化为1），包含网格的维数。简单地讲，gridDim是一个包含三个元素x,y,z的结构体，分别表示网格在x,y,z三个方向上的尺寸，虽然其有三维，但是目前只能使用二维。

blockDim

此变量的类型也是dim3，是一个包含三个元素x,y,z的结构体，分别表示块在x,y,z三个方向上的尺寸，对应于执行配置中的第一个参数，对应于执行配置的第二个参数。

threadIdx

此变量的类型是uint3，是一个包含三个元素x,y,z的结构体，分别表示当前线程在其所在块中x,y,z三个方向上的索引。

warpSize

warpSzie表明warp的尺寸，在计算能力为1.0的设备中，这个值是24，在1.0以上的设备中，这个值是32.

————————————————————————————————————————

（3）内置函数

cudaMalloc((void**)&dev_Ptr, size_t size)

code中的cudaMalloc()是用来分配内存的，这个函数调用的行为非常类似于标准的C函数malloc()，但该函数的作用是告诉CUDA运行时在设备上分配内存。第一个参数是一个指针，指向用于保存新分配内存地址的变量，第二个参数是分配内存的大小。除了分配内存的指针不是作为函数的返回值外，这个函数的行为与malloc()是相同的，并且返回类型为void*。

一些要求：

• 可以将cudaMalloc()分配的指针传递给在设备上执行的函数；
• 可以在设备代码中使用cudaMalloc()分配的指针进行内存读/写操作；
• 可以将cudaMalloc()分配的指针传递给在主机上执行的函数；
• 不能在主机代码中使用cudaMalloc()分配的指针进行内存读/写操作。

cudaFree(void* dev_Ptr)

code中的cudaFree()用来释放cudaMalloc()分配的内存，这个函数的行为和free()的行为非常类似。

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

从上面我们可以知道，主机指针只能访问主机代码中的内存，而设备指针只能访问设备代码中的内存。因此，如果想要实现互相访问，则必须通过cudaMemcpy()函数来实现。这个函数的调用行为类似于标准C中的memcpy()，只不过多了一个参数来指定设备内存指针究竟是源指针还是目标指针。在上面的code中cudaMemcpy()的最后一个参数为cudaMemcpyDeviceToHost，这个参数将告诉运行时源指针是一个设备指针，而目标指针是一个主机指针。而cudaMemcpyHostToDevice将告诉运行时源指针位于主机上，而目标指针是位于设备上的。

本文地址：http://blog.csdn.net/linj_m/article/details/41345263

更多资源请 关注博客：LinJM-机器视觉  微博：林建民-机器视觉

CUDA学习日志：一个例子和编程接口