执行函数_PyCUDA(1)-执行核函数

本文将介绍如何在GPU上并行地执行一段简单的Hello World代码，并介绍什么是核函数。

https://github.com/YouQixiaowu/CUDA-Programming-with-Python​github.com

Hello World

首先让我们从GPU中打印Hello World，现在无需知道每一行代码都在做什么，请先复制到你的编辑器中并运行它。

import pycuda.autoinit
from pycuda.compiler import SourceModule
kernel_code = r"""
__global__ void hello_from_gpu(void)
{
    printf("Hello World from the GPU!n");
}
"""
mod = SourceModule(kernel_code)
hello_from_gpu = mod.get_function("hello_from_gpu")
hello_from_gpu(block=(1,1,1))

如果你的CUDA环境没有问题，那么下面内容将会输出在你的屏幕。

Hello World from the GPU!

逐行讲解Python代码

第一部分

import pycuda.autoinit                           # 以自动的方式对pycuda进行初始化

不同于普通的Python模块，PyCUDA在执行其他方法之前，需要对其进行初始化。我们目前无需知道他究竟做了什么，使用他提供的全自动初始化就好，这足以满足大多数用户的需求。如果你有其他需求，也可以通过手动方式进行初始化。

第二部分

from pycuda.compiler import SourceModule         # 导入一个类，这个类将对CUDA代码进行编译。
kernel_code = r"""
__global__ void hello_from_gpu(void)
{
    printf("Hello World from the GPU!n");
}
"""
mod = SourceModule(kernel_code)                  # 编译上面定义的源代码

这一段对于Python来说其实只有三句代码，中间那个字符串比较长，它只是Python里面简单的字符串类型（这没有什么特别的）。这里面的内容是即将被编译成GPU程序的CUDA代码。我们将在介绍完Python的内容后再来介绍CUDA语法的内容。

这一部分的第一句代码是从pycuda.compiler里面导入一个类（class），这个类以储存着CUDA代码的字符串作为构造参数，去构造一个对象。在这个对象的构造函数中，代码将被编译成CUDA程序。

如果示例代码是第一次被运行，它将会慢一些，而且在程序正常输出之前还会打印下面内容。

UserWarning: The CUDA compiler succeeded, but said the following:
kernel.cu

这也说明了我们的CUDA代码的确是在第一次运行时被编译的。

第三部分

hello_from_gpu = mod.get_function("hello_from_gpu")
hello_from_gpu(block=(1,1,1))

这两句是示例代码的最后两句。第一行可以理解为从GPU程序中获得一个函数执行入口，参数就是这个函数的名字。第二行执行了这个函数，就像是在执行普通的Python函数一样，它的参数我们暂时不去讨论。

小结

目前我们完成了Python部分代码的讲解，这个过程很简单，让我们整理一下这个流程。

1. 初始化PyCUDA模块。
2. 导入一个用于编译的类。
3. 定义一个存储着CUDA代码的字符串。
4. 构造一个类对象（编译CUDA代码，使之成为可以被执行的GPU程序）。
5. 从类对象中获得一个函数的执行入口。
6. 通过这个入口，执行GPU程序。

逐行讲解CUDA代码

第一部分

我们将Python声明字符串内容取出，单独讲解。

__global__ void hello_from_gpu(void)
{
    printf("Hello World from the GPU!n");
}

__global__是一个CUDA的关键字，它出现在一个函数的前面，限定这是一个CUDA核函数（kernel function)，而不是C++中的函数。除此外定义函数的方式与C语言无异，这里不再赘述。

第二部分

下面我们来关注核函数的执行。

hello_from_gpu(block=(1,1,1))

请注意，上面核函数的参数列表是void，而我们在Python中调用该函数时却给予了一个参数block，该参数的格式必须是一个长度为3元组，而且元组元素的类型为int。那这三个参数是什么意义呢，让我们来实践一下。

我们先来尝试修改如下参数。

hello_from_gpu(block=(3,1,1))

你将看到的结果。

Hello World from the GPU!
Hello World from the GPU!
Hello World from the GPU!

该函数被执行了3遍！！！

再来尝试一下这个。

hello_from_gpu(block=(3,4,1))

还有这个。

hello_from_gpu(block=(3,4,5))

你可以分别可以看到12行输出与60行输出！

我想你一定发现了，核函数的执行次数就是里面的数字的乘积。那么你可能要有一个疑问并行(3,4,5)为什么不直接写60呢？这是由于并行经常被用于处理2D、3D问题，这样写参数就很方便（下篇文章中详细说明原因）。

这些线程就像一个格点长方体，参数block指定了每一个边的格点数量

小结

参数block决定着有多少个线程在执行核函数。那这些函数是串行（依次执行）还是并行（同时运行）执行呢？那么我们应该如何验证呢？我们将在下一篇文章中通过实践讨论这个问题。