CUDA编程 1

这段时间想重新来学习一下cuda编程

1. 第一个程序

hello,world!

#include <iostream>

__global__ void kernel(void) {
}

int main(void)
{
	kernel<<<1,1>>>();
	printf("Hello,world!\n");
	return 0;
}

我怀疑你在骗我，cuda编程这么简单的吗？这不是C吗？哈哈，相比较于C，确实多了两个值得注意的地方：

1. 一个空的kernel()，并且带有修饰符__global__
2. 对这个空函数的调用，并且带有修饰字符 <<<1,1>>>

global这个修饰符，告诉编译器，函数应该编译为在device而不是在host上运行，函数kernel() 交给编译设备代码的编译器，而main函数将被交给主机编译器

因此这个看上去有些奇怪的函数调用实际上表示调用设备代码，但是为什么要用尖括号，或者数字呢？尖括号表示要将一些参数传递给运行时系统。这些参数并不是传递给设备代码的参数，而是告诉运行时如何启动设备代码。

1.1 传递参数

我们现在对hello world进行修改

#include <iostream>
#include "book.h"

__global__ void add(int a, int b, int *c)
{
	*c = a + b;
}

int main(void) {
	int c;
	int *dev_c;
	HANDLE_ERROR (cudaMalloc( (void**) &dev_c, sizeof(int)));
	
	add<<<1,1>>> (2,7,dev_c);
	
	HANDLE_ERROR( cudaMemcpy (&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost ));
	printf("2+7=%d\n", c);
	cudaFree (dev_c);

	return 0;
}

这里增加了多行代码，在这些代码中包含了两个概念：

1. 可以像调用c函数那样将参数传递给核函数
2. 当设备执行任何有用的操作时，都需要分配内存，例如将计算值返回给host 主机
3. 简单解释一下cudaMemcpyHostToDevice将告诉运行时相反的含义，即源地址位于host，而目标指针位于device。此外还可以通过传递参数 cudaMemcpyDeviceToDevice 来告诉运行时这两个指针都是位于device上。如果源指针和目标指针都位于host上，那么可以直接调用标准C的memcpy()函数.

1.2 查询设备

这些api主要用来查询GPU的信息

1. cudaGetDeviceCount() ,可以对每个设备进行迭代，并查询各个设备的相关信息。CUDA运行时将返回一个cudaDeviceProp的结构，其中包含了设备的相关属性。我们可以获得哪些属性：
   设备查询的代码：

#include "../common/book.h"

int main(void) {
	cudaDeviceProp prop;
	
	int count;
	HANDLE_ERROR(cudaGetDeviceCount(&count));

	for (int i=0; i<count; i++){
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, i));
		// 对设备的属性执行某些操作
	printf(i, prop.minor, prop.clockRate);
	}
}

1.3 设备属性的使用

根据在cudaGetDeviceCount()和cudaGetDeviceProperties 中返回的结果，我们可以对每个设备进行迭代，并且查找主版本号大于1，或者主版本号为1且次版本号大于等于3的设备。但是这种迭代操作执行起来有些繁琐，所以cuda运行时提供了一种自动方式来执行这个迭代操作。首先，找出我们希望设备拥有的属性并将这些属性填充到一个cudaDeviceProp结构。

cudaDeviceProp prop;
memset(&prop, 0, sizeof(cudaDeviceProp));
prop.major = 1;
prop.minor = 3; 

再填充完这个cudaDeviceProp结构后，将其传递给cudaChooseDevice()，这样cuda运行时将查找是否存在某个设备满足这些条件。cudaChooseDevice() 函数将返回一个设备ID，然后我们可以将这个ID传递给cudaSetDevice()。随后，所有的设备操作都将在这个设备上执行。

#include ".../common/book.h"

int main(void){

	cudaDeviceProp prop;
	int dev;
	
	HANDLE_ERROR(cudaGetDevice(&dev));

	memset( &prop, 0, sizeof(cudaDeviceProp));
	prop.major = 1;
	prop.minor = 3; 
	HANDLE_ERROR(cudaChooseDevice(&dev, &prop));
	printf(dev);
	HANDLE_ERROR(cudaSetDevice(dev));
}

小结：
本次主要时学习了添加关键字 global 将告诉编译器把该函数放在GPU上运行。为了使用GPU的专门内存，我们还学习了与C的malloc()、memcopy()和free等api对应的cuda api

参考：GPU高性能编程CUDA实战