【CUDA编程笔记】（1）CUDA C 编程入门

这学期在上GPU并行计算的课程，大作业是CUDA C编程，所以写一些笔记记录一下学习过程。

参考资料：《CUDA编程 基础与实践》樊哲勇 清华大学出版社

一、上机环境

学校机房，windows，VS2019，CUDA10.2，1080Ti

二、核函数的概念

GPU只是一个设备，想要工作的话就需要CPU主机给它下达命令。所以一个真正利用了GPU的CUDA程序包含主机代码和设备代码（需要设备执行的代码）两部分。其中主机对设备的调用是通过核函数实现的，下面的程序中用 __ global __ 修饰的就是核函数，除了前缀之后，其他部分与c++相同。

一般有以下三个形式的前缀：

__ device __ 在GPU上调用，在GPU上执行
__ global __ 在CPU上调用，在GPU上执行
__ host __ 在CPU上调用，在CPU上执行

一个典型的CUDA程序结构为：

int main()
{
	主机代码
	核函数的调用
	主机代码
	return 0;
}

三、入门 hello gpu（单线程 && 多线程）

线程数=网格大小*线程块大小

一个核函数的全部线程块构成一个网格（grid），所以网格大小就是线程块的个数（block size）。线程块大小表示一个线程块中的线程数（thread）。

单线程：

#include <stdio.h>

//kernal function，用于实现主机对设备的调用
//核函数的返回类型必须是空类型，即void
__global__ void hello_gpu() {
	printf("hello gpu\n");
}

int main()
{
	//<<<网格大小，线程块大小>>>，线程数=网格大小*线程块大小
	hello_gpu << <1, 1 >> > ();
	cudaDeviceSynchronize;   //同步主机与设备，促使缓存区刷新。
	return 0;
}


//output:
hello gpu

多线程：

#include <stdio.h>

__global__ void hello_gpu() {
	printf("hello gpu\n");
}

int main()
{
	//线程数 = grid_size * block_size = 6，所以执行6次核函数
	hello_gpu << <2, 3 >> > ();
	cudaDeviceSynchronize;   
	return 0;
}


//output:
hello gpu
hello gpu
hello gpu
hello gpu
hello gpu
hello gpu

四、线程索引

#include <stdio.h>

__global__ void hello_gpu() {
	const unsigned bid = blockIdx.x;    //一个网格中的线程块指标
	const unsigned tid = threadIdx.x;   //一个线程块中的线程指标
	printf("hello gpu from block %d、thread %d\n",bid,tid);
}

int main()
{
	hello_gpu << <2, 3 >> > ();
	cudaDeviceSynchronize;   
	return 0;
}


//output(first run):
hello gpu from block 1、thread 0
hello gpu from block 1、thread 1
hello gpu from block 1、thread 2
hello gpu from block 0、thread 0
hello gpu from block 0、thread 1
hello gpu from block 0、thread 2


//output(second run):
hello gpu from block 0、thread 0
hello gpu from block 0、thread 1
hello gpu from block 0、thread 2
hello gpu from block 1、thread 0
hello gpu from block 1、thread 1
hello gpu from block 1、thread 2

可以看到每次执行的结果都是不同的，即有时是第0个线程块先完成计算，有时是第1个线程块先完成计算。这反映CUDA程序执行时的一个重要特征：每个线程块的计算都是相互独立的。

程序中的 blockIdx是一个 unit3类型的 struct，有x,y,z三个成员，表示一个三维网络（类似于多维数组）的线程。我们这里默认的网格和线程块都是一维的，所以调用 blockIdx.x