目录
- 一、上机环境
- 二、核函数的概念
- 三、入门 hello gpu(单线程 && 多线程)
- 四、线程索引
这学期在上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