cuda从入门到精通（二）之函数关键字

最新推荐文章于 2025-03-19 09:30:41 发布

nudt_qxx

最新推荐文章于 2025-03-19 09:30:41 发布

阅读量3.1k

点赞数 8

分类专栏： CUDA C++ 并行计算 CUDA从入门到精通

C++ 同时被 3 个专栏收录

157 篇文章

订阅专栏

并行计算

8 篇文章

订阅专栏

CUDA从入门到精通

7 篇文章

订阅专栏

本文系转载，出处：浅析GPU计算——cuda编程

本系列博客配备github代码，本小节代码见：https://github.com/qixuxiang/cuda_zero_to_one/blob/master/run_gpu_or_cpu.cu

CUDA编程模型基础

在给出CUDA的编程实例之前，这里先对CUDA编程模型中的一些概念及基础知识做个简单介绍。CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行流程如下：

分配host内存，并进行数据初始化；
分配device内存，并从host将数据拷贝到device上；
调用CUDA的核函数在device上完成指定的运算；
将device上的运算结果拷贝到host上；
释放device和host上分配的内存。

关键字

之前我们讲解过，CPU是整个计算机的核心，它的主要工作是负责调度各种资源，包括其自身的计算资源以及GPU的计算计算资源。比如一个浮点数相乘逻辑，理论上我们可以让其在CPU上执行，也可以在GPU上执行。那这段逻辑到底是在哪个器件上执行的呢？cuda将决定权交给了程序员，我们可以在函数前增加修饰词来指定。

一般来说，我们只需要2个修饰词就够了，但是cuda却提供了3个——2个执行位置为GPU。这儿要引入一个“调用位置”的概念。父函数调用子函数时，父函数可能运行于CPU或者GPU，相应的子函数也可能运行于CPU或者GPU，但是这绝不是一个2*2的组合关系。因为GPU作为CPU的计算组件，不可以调度CPU去做事，所以不存在父函数运行于GPU，而子函数运行于CPU的情况。

由于GPU实际上是异构模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词开区别host和device上的函数，主要的三个函数类型限定词如下：

__global__：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
__device__：在device上执行，单仅可以从device中调用，不可以和__global__同时用。
__host__：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译。

关键字	调用位置	执行位置
`__host__`	CPU	CPU
`__global__`	CPU	GPU
`__device__`	GPU	GPU

注意

__global__描述的函数就是“被CPU调用，在GPU上运行的代码”，同时它也打通了__host__和__device__修饰的函数。
__global__既不能和__host__关键词一起用，也不能和__device__一起用。
__global__修饰的函数既是void返回类型，且是异步调用的。假设该函数有返回值，当函数返回时，接受返回值的变量可能已经被销毁了，所以设计其有返回值也没太多意义。

如果一段代码既需要运行于CPU，也要运行于GPU，怎么办？难道要写两次？当然不用，我们可以同时使用__host__和__device__修饰。这样编译器就会帮我们生成两份代码逻辑。

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

__host__ __device__ int run_on_cpu_or_gpu() {
	return 1;
}
 
__global__ void run_on_gpu() {
	printf("run_on_cpu_or_gpu GPU: %d\n", run_on_cpu_or_gpu());
}
 
int main() {

	printf("run_on_cpu_or_gpu CPU: %d\n", run_on_cpu_or_gpu());
	run_on_gpu<<<1, 1>>>();
	cudaDeviceReset();
	return 0;
}
/*
run_on_cpu_or_gpu CPU: 1
run_on_cpu_or_gpu GPU: 1
*/

对于上面的__global__修饰函数中<<<1, 1>>>,我们会在下一节cuda从入门到精通（三）之核函数详细讲到，这里你只需要把<<<1, 1>>>看成是一种并行计算的手段即可。

上面提过， __global__修饰的函数既是void返回类型，且是异步异步调用的。我们可以改写上面代码验证下：

#include "cuda.h" 
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
 
__host__ __device__ int run_on_cpu_or_gpu() {
	return 1;
}
 
__global__ void run_on_gpu() {
	printf("run_on_cpu_or_gpu GPU: %d\n", run_on_cpu_or_gpu());
}
 
int main() {

	printf("run_on_cpu_or_gpu CPU: %d\n", run_on_cpu_or_gpu());
	run_on_gpu<<<1, 1>>>();
	printf("will end\n");
	cudaDeviceReset();
	return 0;
}

如果上述代码都是同步执行的，那么最后一句输出应该是will end，然而实际输出却是：

run_on_cpu_or_gpu CPU: 1
will end
run_on_cpu_or_gpu GPU: 1

可见__global__修饰的函数的确是被异步执行的。

还有人会问，上面main函数怎么没有用修饰符修饰?cuda编程规定如果没有使用修饰符修饰的默认就是__host__类型。这种设计让大家熟悉的规则成为默认的规则，可以让更多第三方代码不用修改就直接被cuda编译器编译使用。

cuda是一个GPU编程环境，所以它对__device__修饰的函数进行了比较多的优化。比如它会根据它的规则，让某个__device__修饰函数成为内联函数（inline）。这些规则是程序员不可控，但是如果我们的确对是否内联有需求，cuda也提供了方式：使用__noinline__修饰函数不进行内联优化；使用 __forceinline__修饰函数强制进行内联优化，当然这两种修饰符不能同时使用。