By Ying Nie

更新日期:

文章目录
  1. 1. cuda简介
  2. 2. 核函数(kernel)
    1. 2.1. 声明与调用
    2. 2.2. 参数传递
    3. 2.3. 线程并行
  3. 3. 总结

前言
两个月前，我接下了移植实验室三维重建系统到gpu上执行的任务，以期实现实时重建的效果。现在，我将结合近两个月的学习更新一系列博文，以总结cuda的核心知识点。关于cuda学习资料，我推荐nvidia官方出版的《gpu高性能编程——cuda实战》与《cuda编程指南》两书，本系列博文也将其作为参考。

cuda简介

cuda架构：是nvidia推出的通用并行计算架构，它包含了一个统一的着色器流水线，使得执行通用计算的程序能够对芯片上的每个alu进行排列,也就是每一个alu单元都可以并行起来进行大规模计算。
cuda c：则是专门为这种架构设计的编程语言。它基于c语言的，因此只要有c语言基础的同学都能很快上手。
编程模型：cuda允许程序员定义成为内核（kernel）的c语言函数，在调用此类函数时，他将由n个不同的cuda线程并行执行n次。

核函数(kernel)

声明与调用

首先我们来看两段代码：

代码1：标准c

1
2
3
4
5
6
void add(){ }
int main (void){
	add();
	return 0;	
}

代码2：cuda c

1
2
3
4
5
6
__global__ void add_kernel(){ }
int main (void){
	add_kernel<<<1,n>>>(); //kernel
	return 0;	
}

这两段代码，有两个有区别的地方：

  • kernel函数add_kernel()定义时,带有修饰符__global__;
  • 该函数调用时，带有修饰符<<<1,n>>>。

区别解释：

  • __global__声明说明符，表示该函数在gpu上运行，属于设备代码；
  • <<<1,n>>>修饰符，则告诉gpu如何调用设备代码，指定每次并行的cuda线程数。

其实cuda c就是这么简单，它将设备代码传给nvcc编译，其它代码则交给主机编译器。对主机来说，cuda就这样变成透明的了，是不是很机智？

参数传递

根据以上说明，初步了解了kernel的声明与调用, 那么应该如何传递参数给核函数呢？

问1：如何传递参数给核函数呢？
答1： 可以像c函数一样将参数传递给核函数，但kernel中只能调用显存中的数据。因此我们必须为任何有用数据分配显存空间，并将返回值传给主机。

代码3：带参数的cuda c

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
__global__ void add_kernel(int a, int b, int *c){
	int i = threadidx.x；
	c[i] = a[i] + b[i] ;
 }
int main (void){
	int h_c;
	int *dev_c;
	cudamalloc((void**)&dev_c, sizeof(int)); //分配显存空间
	
	add_kernel<<<1,n>>>（dev_a, dev_b, dev_c); //假设dev_a, dev_b已赋值
	cudamemcpy(&h_c, dev_c, sizeof(int), cudamemcpydevicetohost); //将显存数据传回主机	
	cudafree(dev_c);
	return 0;	
}

代码3的作用是，将大小为n的向量a和向量b相加，并将结果存入c。main函数里多了三行代码cudamalloc()/cudafree()和cudamemcpy();它们均是cuda提供的标准函数，分别负责分配/设备设备空间和主机与设备的数据交换。

函数原型：
cudaerror_t cudamalloc(void** devptr, size_t count);
cudaerror_t cudafree(void* devptr);
cudaerror_t cudamemcpy(void* dst, const void* src, size_t count, enum cudamemcpykind kind);
说明：cudamemcpy()从src指向的存储器区域将count个字符复制到dst指向的区域；
其中kind是cudamemcpyhosttohost、udamemcpyhosttodevice、cudamemcpydevicetohost或cudamemcpydevicetodevice之一。

线程并行

学到这里，我们已经知道如何声明和调用核函数了，并且已经能正确传递参数给它，但依然有个问题待解决。

问2：应该如何定义核函数，以使线程并行执行并准确获取所需参数呢？
答2：执行内核的每个线程都会被分配一个独特的线程id，可通过内置的threadidx变量在内核中访问此id。我们可以利用线程id去该线程索引所需的参数值；

依然以代码3为例，add_kernel()实际所做的工作也就是将a,b两组数据的对应元素两两相加。该函数中的代码:
int i = threadidx.x作用在于获取线程id(以0开始计数)；
c[i] = a[i] + b[i]则表示，每个线程将线程id作为索引号，获取该线程对应处理的元素；
这样就能使n个线程并行起来，每个线程实现其中一组元素的加法操作，从而避免单线程循环n次计算，大大提高运算速度。

总结

总结起来，
编程模型： 定义称为kernel的c语言函数，kernel将有n个不同的cuda线程并行执行n次；
声明与调用： __global__ void kernel ( void ){};
kernel<<<...>>> (); //<<<...>>> 指定每次调用的cuda线程数。
参数传递：核函数与c函数参数传递形式一样，但核函数参数必须为设备存储区域的数据。
线程并行：每个线程利用线程id去索引合适的元素进行计算。

如果看到这里，对并行编程还不太明白，没关系。我将在下一篇博文详细介绍cuda并行编程以及线程层次结构。

文章目录
  1. 1. cuda简介
  2. 2. 核函数(kernel)
    1. 2.1. 声明与调用
    2. 2.2. 参数传递
    3. 2.3. 线程并行
  3. 3. 总结