内容概览

1. CUDA 并行计算简介
2. CUDA-Fortran 开发环境配置
3. CUDA 基本概念
4. 第一个 CUDA-Fortran 程序

1 CUDA并行计算简介

1.1 并行计算简史

• 随着计算机硬件的发展，CPU不能再以提高单核时钟速度的方式来提高性能，导致其横向发展，也就出现了越来越多的多核处理器。
• 相比于图形处理单元（Graphics Processing Unit），其并行规模显得特别苍白无力。
• 使用GPU做通用计算被称为（GPGPU，General Purpose GPU），这需要面对应用程序接口进行编程。
• CUDA架构不仅包含应伟达GPU硬件，还包含一个软件编程环境。

1.2 CUDA 并行计算

    GPU内部可以产生大量线程，上下文切换开销非常小。
    进行数据级并行，是一种细粒度并行。

2 CUDA-Fortran 开发环境配置

作为开发工具，首先要做的就是配置环境，笔者自己的笔记本电脑为联想Y430P，安装有 XUbuntu 18.04.1 64位操作系统，显卡为GForce GTX 850M，显存2G。

• 安装系统这里不再赘述。

2.1 配置 CUDA 环境

1 更换国内源，因为更新快。见下图。

方法：系统设置 -> 软件和更新 -> Ubuntu 软件 -> 下载自（选择一个，推荐清华）

换源

2 换完之后，一般会刷新缓存，但是还是建议手动刷新一次，并更新。

$ sudo apt uodate && sudo apt upgrade -y 

3 更新安装完之后需要勾选Nvidia显卡驱动，见下图。

方法：系统设置 -> 软件和更新 -> 附加驱动程序 -> 选择驱动（推荐闭源驱动）

勾选N卡驱动

4 安装nvidia-cuda-toolkit
安装命令如下；图中是仓库中搜索的结果。

$ sudo apt install nvidia-cuda-toolkit

软件仓库中的nvidia-cuda-toolkit

5 安装完之后重启电脑，检测环境。

$ nvidia-smi # NVIDIA System Management Interface program，可以监测CUDA设备
$ nvcc --version # 输出NVIDIA CUDA C 编译器版本

正常输出，就表明环境配置正确，可以干活了。如上面两条命令，在笔者电脑上的输出分别如下。

命令 nvidia-smi 输出
命令 nvcc --version 输出

2.2 安装配置 CUDA-Fortran 编译器

PGI（前身为The Portland Group，Inc。）是一家为高性能计算系统生产一套商用Fortran，C和C ++ 编译器的公司。2013年7月29日，NVIDIA公司收购了The Portland Group，Inc。波特兰集团（或PGI）的名称现在被称为NVIDIA公司生产的软件开发工具品牌。

之后，PGI发布免费使用的社区版，学习使用还是蛮不错的，笔者自己也使用该编译器，点击跳转到下载页。

PGI Community Edition

我们要使用CUDA-Fortran，就需要安装相应编译器，安装过程此处不再叙述，和一般 linux 软件安装套路一样，可能以后补充上，如果读者需要，可以留言，我会及时更新文章。
另外，安利一本书《OpenACC并行编程实战》(何沧平)，这本书里第九章有详细的 PGI 编译器安装教程，此外这本书讲 OpenACC，非常基础，笔者自己也有一本。

3 CUDA 基本概念

如果你看到了这里，那么恭喜你，说明你已经配置好了开发环境，并且可以开始在CUDA-Fortran的世界里翱翔了。言归正传，在进行 CUDA-Fortran 开发之前，我们需要了解一些基本概念，不过不用担心，这些概念并不是很复杂。

• 主机：指CPU及其内存。
• 设备：指GPU及其内存。
• CPU代码：指一个仅用到CPU的实现。
• 内核：指一个在设备上执行但被主机调用的子例行程序。
• 执行配置：子例程名字与参数列表之间三尖号内的一组参数，如call sub<<<m，n>>>( arguments )。

另外，在 CUDA 这样的混合编程模型中，必须解决的一个问题是主机与设备间的同步。我们将要在示例代码中使用的赋值语句是阻塞传递的，有利于隐式同步CPU和GPU。

4 第一个 CUDA-Fortran 程序

在这一章中，我们要编写第一个 CUDA-Fortran 程序，编译并且执行他。

好了，我们先来看代码。

module simpleOps_m
    implicit none 
  contains 
    attributes(global) subroutine increment( a, b )
!       |         |  
!       |         + --- 指明这个代码运行在设备上,但需要从主机上调用。
!       |               global 描述作用域;从主机和设备均可看到本例程
!       + ------------- 指明子例行程序的属性 
        integer,intent(inout) :: a(:)
        integer,value         :: b 
        integer               :: i
        i = threadIdx%x
!               |
!               + --- GPU 线程同时执行子例程。
!                     每一个线程通过在所有设备代码中可用的内置变量 threadIdx 来识别自已
!                     并将该变量用作数组下标。
        a(i) = a(i) + b
    end subroutine increment
end module simpleOps_m

program incrementTestCPU
    use simpleOps_m
    use cudafor 
    implicit none 
    integer,parameter :: n = 256
    integer           :: a(n), b
    integer,device    :: a_d(n)
    
    a = 1; b = 3
    
    a_d = a ! 阻塞传递数据，隐式同步
    call increment<<<1,n>>>( a_d, b )
    !             |       |   |   |
    !             |       |   |   + --- 第二个内核标量参数 b 驻留主机内存，需要 ## 传值 ##
    !             |       |   + ------- 第一个内核数组参数 a_d 驻留设备内存
    !             + ----- + ----------- 执行配置，此中分别是<<<线程块数量，块内线程数量>>> 
    a = a_d ! 阻塞传递数据，隐式同步
    
    if( any( a .ne. 4 ) ) then 
        write(*,*) "**** Program Failed ****"
    else 
        write(*,*) "**** Program Passed ****"
    end if 
end program incrementTestCPU

读者可能在第一次看到这段代码时，会对其中一些地方感到迷惑，不过不要紧，除去少部分 CUDA-Fortran 的内容外，其余部分代码和代码含义与普通的 Fortran 程序没有什么大的区别。

读者可以仔细阅读一下代码中的注释，以及回顾一下 CUDA 基本概念。

现在，我们编译这段代码。

$ pgfortran -Mcuda file_name.f90 

然后执行编译好的程序

$ ./a.out

最后你会看到这样的结果。

第一个 CUDA-Fortran 程序编译及执行结果

今天就到这里了，下次再更其他内容，谢谢阅读。