Yinjun Zhu's Blog
首页 归档 标签 分类
Yinjun Zhu's Blog

目录

CUDA学习笔记——第一个CUDA程序

 收录于 编程语言
 Created on   约 1808 字   预计阅读 4 分钟 

安装

要利用 CUDA 写 GPU 并行程序，先要按照 CUDA 工具套件。Nvidia 官网上有详细介绍，这里就不展开讲了：[https:// developer.nvidia.com/cuda-toolkit](https:// developer.nvidia.com/cuda-toolkit)

需要注意的是，既然 CUDA 是基于 Nvidia GPU 的开发工具，所以 CUDA 程序必须运行在 Nvidia 的独立显卡上，在集显或其它品牌的独显上无法运行。

 

检查环境

安装完 CUDA Toolkit 后，可以检查是否安装成功：

1
2

> which nvcc
> /usr/local/cuda/bin/nvcc

尝试执行：

1
2

> nvcc
> nvcc fatal   : No input files specified; use option --help for more information

nvcc 已经安装成功，这里指没有 input file，即 CUDA 源代码文件。

如果 nvcc 找不到，需要手动将安装路径加到 PATH 环境变量中去。

nvcc 是 CUDA 编译器，它内部集成了 GCC 编译器，在编译 CUDA 程序时，会把 host code 和 device code 分开。host code，就是用标准 C 写的，nvcc 会用内嵌的 C 编译器去编译成符合 CPU 结构的机器码；device code 是用 CUDA C 的扩展的符号和语法写的函数，这类函数称为 kernel，nvcc 将它们编译为 GPU 指令。

 

GPU 上的 Hello World

创建源代码文件

创建一个文件，以 .cu 作为后缀名。在文件中输入以下代码：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

#include <stdio.h>

__global__ void hello_world_from_gpu (void)
{
	printf(“Hello World from GPU!\n”);
}

int main(void)
{
    printf(“Hello World from CPU!\n”);
    hello_world_from_gpu <<<1, 5>>>();
    cudaDeviceReset();
    return 0;
}

主函数中打印普通的语句，然后调用一个特殊函数hello_world_from_gpu，后面会讲。cudaDeviceReset会销毁和回收当前进程中的所有 GPU 相关的资源。

编译

保存好上述代码后，执行：

1

> nvcc -arch compute_50 hello.cu -o hello

可以看到，nvcc 的使用和 gcc 类似，gcc 的命令行参数也适用于 nvcc。-arch compute_50是用来指定当前运行在哪个架构的 GPU 上。可用的选项包括（不限于）：

  • compute_30, compute_32
  • compute_50, compute_52, compute_53
  • compute_60, compute_61, compute_62
  • compute_70, compute_72
  • …..

也可以用sm_50来表示。

执行

1
2
3
4
5
6
7

> ./hello
Hello World from CPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!

主程序中的语句打印了一次，而hello_world_from_gpu函数中的语句，打印了 5 次。为什么呢？

hello_world_from_gpu函数是一个特殊的函数，因为前面加了__global__修饰符，它表示这是一个 kernel，这里的 kernel 不是指操作系统的 kernel，而是 CUDA 编程中的术语——核函数。kernel 就是指将要在 GPU 上运行的程序。如同在多线程程序中，线程用函数来表达一样，GPU 上要运行的 kernel，也用函数来表达。启动一个 kernel 和普通函数调用类似，唯一区别就是要指定网格和线程块的维度：

1

hello_world_from_gpu <<<1, 5>>>();

<<< >>>中的数字或者变量，是 CUDA 的执行配置，表示网格和线程块的维度，用于指定将会启动多少线程去执行 kernel。CUDA 中网格和线程块的维度都是三维的，当然，也可以设置一维或二维，只需将其它维度设为 1。第一个参数 1， 表示网格的第一个维度是1，而其它维度默认是1，所以网格的维度是 1 *1 *1。5 表示线程块的第一个维度是 5，其它维度默认也是 1，所以线程块的维度是 5 * 1 * 1。

<<<1, 5>>>表示网格包含一个线程块，线程块包含 5 个线程。kernel 执行在这个网格中的线程上，所以一共有 5 个线程执行 kernel。

 

CUDA 程序结构

一般，典型的 CUDA 程序，执行以下几个步骤：

  1. 分配 GPU 内存；
  2. 将相关数据从 CPU 内存拷贝到 GPU 内存；
  3. 调用 CUDA kernel 来运行指定的计算；
  4. 将计算结果从 GPU 内存拷贝到 CPU 内存；
  5. 回收 GPU 内存。

在上述例子中，只执行了第三步：调用 kernel。后面的文章将会举例完整的 CUDA 程序。

 

和其它并行程序结构的不同

如果你用 C 语言写并行程序，大概率你用过 Pthreads 库，也可能用过 OpenMP，这是两个比较知名的并行开发技术，适用于大部分多核处理器架构和操作系统。Pthread 和 OpenMP 都属于基于共享内存系统的并行开发，在共享内存系统中，各个核能够共享访问计算机内存，理论上每个核都能读写内存的每个区域。Pthreads是 C 语言的扩展库，可以在 C 程序中使用扩展的类型定义、函数、宏；OpenMP包含了一个扩展库以及对 C 编译器的部分修改。

当你用 CUDA C 写程序时，实际上你写的是一串线性代码，在 CPU 上这串代码只在一个线程执行，而它的 kernel 函数，将被 GPU 上成千上万的线程执行，每个线程都执行相同的计算。

CUDA 把 GPU 底层细节隐藏起来，把硬件的执行细节抽象地表示。有三个主要的抽象表示：线程组的层次、内存组的层次、以及屏障同步机制。通过这些抽象表示，程序员只需基于简单的 C 语言的扩展，就可以操控相关的底层的硬件执行。

CPU + GPU 架构，如今已经是高性能计算的主流架构，“把数据并行工作交给 GPU，而任务并行工作交给 CPU”，算得上是并行计算的典范。

- 全文完 -

相关文章

「 您的赞赏是激励我创作和分享的最大动力! 」

 C/C++, 异构, 并发, 并行, 高性能计算, CUDA, GPU, 深度学习
返回 | 主页