NVIDIA CUDA Fortran 编程指南

前言

本文档介绍了 CUDA Fortran，它是 Fortran 的一小组扩展，支持 CUDA 计算架构并基于其构建。

目标读者

本指南适用于精通 Fortran、C 和/或 C++ 语言编程的应用程序员、科学家和工程师。这些工具在各种操作系统上适用于 x86-64 和 Arm 服务器硬件平台。本指南假定读者熟悉基本操作系统用法。

组织结构

本文档的组织结构如下

简介: 包含一般介绍
编程指南: 作为 CUDA Fortran 的编程指南
参考: 描述 CUDA Fortran 语言参考
运行时 API: 描述 CUDA Fortran 和 CUDA 运行时 API 之间的接口
示例: 提供示例代码和简单示例的解释。

约定

本指南使用以下约定

斜体: 用于强调。
等宽字体: 用于文件名、目录、参数、选项、示例以及文本中的语言语句，包括汇编语言语句。
粗体: 用于命令。
[ item1 ]: 通常，方括号表示可选项目。在本例中，item1 是可选的。在 p/t-set 的上下文中，方括号是指定 p/t-set 所必需的。
{ item2 | item 3 }: 花括号表示需要选择。在本例中，您必须选择 item2 或 item3。
filename …: 省略号表示重复。可以出现零个或多个前面的项目。在本例中，允许多个文件名。
FORTRAN: Fortran 语言语句在本指南的文本中使用缩小的固定字号显示。
C/C++: C/C++ 语言语句在本指南的测试中使用缩小的固定字号显示。

NVIDIA HPC 编译器在 Linux 操作系统的 64 位变体上，以及各种 x86 兼容和 Arm 处理器上受支持。

1. 简介

欢迎使用 NVIDIA CUDA Fortran 2024 版本，它是 Fortran 的一小组扩展，支持 CUDA 计算架构并基于其构建。

图形处理单元或 GPU 已经发展成为可编程、高度并行的计算单元，具有非常高的内存带宽，以及在许多应用程序中的巨大潜力。 GPU 设计针对图形渲染中发现的计算进行了优化，但足够通用，可用于许多数据并行、计算密集型程序。

NVIDIA 引入了 CUDA^®，一种通用的并行编程架构，带有编译器和库来支持 NVIDIA GPU 的编程。 CUDA 配备了扩展的 C 编译器，这里称为 CUDA C，允许从高级语言直接编程 GPU。编程模型支持四个关键抽象：组织成线程组的协作线程、线程组内的共享内存和屏障同步，以及组织成网格的协调独立线程组。 CUDA 程序员必须将程序划分为可以在并行执行的粗粒度块。每个块被划分为细粒度线程，这些线程可以使用共享内存和屏障同步进行协作。正确设计的 CUDA 程序将在任何启用 CUDA 的 GPU 上运行，而与可用处理器内核的数量无关。

CUDA Fortran 包括 Fortran 2003 编译器和工具链，用于使用 Fortran 编程 NVIDIA GPU。 NVIDIA 2024 包括对 Linux 上 CUDA Fortran 的支持。 CUDA Fortran 类似于 NVIDIA 的 CUDA C 编译器。与基于 NVIDIA OpenACC 指令和编译器的模型相比，CUDA Fortran 是一种更低级别的显式编程模型，具有大量的运行时库组件，使专业程序员可以直接控制 GPGPU 编程的所有方面。

本文档中描述的 CUDA Fortran 扩展允许在 Fortran 程序中执行以下操作

声明在 GPU 设备内存中分配的变量
在 GPU 设备内存中分配动态内存
将数据从主机内存复制到 GPU 内存，然后再复制回来
编写在 GPU 上执行的子例程和函数
从主机调用 GPU 子例程
在主机上分配钉住内存
使用主机和 GPU 之间的异步传输
使用零复制和 CUDA 统一虚拟寻址功能。
通过纹理内存缓存访问只读数据。
使用内核循环指令自动生成 GPU 内核。
使用动态并行特性从设备上运行的其他 GPU 子例程启动 GPU 内核。
可重定位的设备代码：创建和链接设备库，并调用在其他模块和文件中定义的函数。
与 CUDA C 接口。
编程访问张量核心硬件。

2. 编程指南

本节通过 CUDA Fortran 编写的示例介绍 CUDA 编程模型。有关 CUDA Fortran 的参考，请参阅参考。

2.1. CUDA Fortran 主机和设备代码

所有 CUDA 程序，以及通常任何使用 GPU 进行计算的程序，都必须执行以下步骤

初始化并选择要在其上运行的 GPU。通常，这在程序中是隐式的，并默认为 NVIDIA 设备 0。
在 GPU 上为数据分配空间。
将数据从主机移动到 GPU，或者在某些情况下，在 GPU 上初始化数据。
从主机启动内核以在 GPU 上运行。
从 GPU 收集结果，以便从主机程序进行进一步分析或输出。
在步骤 2 中释放 GPU 上分配的数据。这可能会在主机程序退出时隐式执行。

这是一个简单的 CUDA Fortran 示例，它执行了所需的步骤

显式设备选择

主机代码

设备代码

program t1
use cudafor
use mytests
integer, parameter :: n = 100
integer, allocatable, device :: iarr(:)
integer h(n)
istat = cudaSetDevice(0)
allocate(iarr(n))
h = 0; iarr = h
call test1<<<1,n>>> (iarr)
h = iarr
print *,&
"Errors: ", count(h.ne.(/ (i,i=1,n) /))
deallocate(iarr)
end program t1

module mytests
contains
attributes(global)  &
subroutine test1( a )
integer, device :: a(*)
i = threadIdx%x
a(i) = i
return
end subroutine test1
end module mytests

在左侧的 CUDA Fortran 主机代码中，设备选择是显式的，通过第 7 行的 API 调用执行。提供的 cudafor 模块（在第 2 行中使用）包含到完整 CUDA 主机运行时库的接口，在本例中，公开了 cudaSetDevice() 的接口并确保正确调用它。在第 8 行，在设备上分配了一个数组。主机代码的第 9 行初始化主机和设备上的数据，并在第 10 行启动设备内核。内核的接口是显式的，在 Fortran 的意义上，因为包含内核的模块在第 3 行中使用。在主机代码的第 11 行，内核执行的结果被移回主机数组。GPU 数组的释放发生在第 14 行。

隐式设备选择

这是一个比前一个示例稍微复杂的 CUDA Fortran 示例。

主机代码

设备代码

program testramp
use cublas
use ramp
integer, parameter :: N = 20000
real, device :: x(N)
twopi = atan(1.0)*8
call buildramp<<<(N-1)/512+1,512>>>(x,N)
!$cuf kernel do
do i = 1, N
x(i) = 2.0 * x(i) * x(i)
end do
print *,"float(N) = ",sasum(N,x,1)
end program

module ramp
real, constant :: twopi
contains
attributes(global) &
subroutine buildramp(x, n)
real, device :: x(n)
integer, value :: n
real, shared :: term
if (threadidx%x == 1) term = &
twopi / float(n)
call syncthreads()
i = (blockidx%x-1)*blockdim%x &
+ threadidx%x
if (i <= n) then
x(i) = cos(float(i-1)*term)
end if
return
end subroutine
end module

在本例中，设备选择是隐式的，并默认为 NVIDIA 设备 0。主机代码中第 5 行的设备数组分配看起来是静态的，但实际上发生在程序初始化时。更大的数组大小在主机代码的第 7 行的内核启动中以及设备代码的第 10 行中都得到了处理。设备代码包含常量和共享数据的示例，这些数据在参考中描述。实际上，从主机代码启动了两个内核：一个显式提供并在第 10 行调用的内核，以及第二个内核，使用 CUDA Fortran 内核循环指令生成的，从第 11 行开始。最后，本示例演示了 cublas 模块的用法，该模块在主机代码的第 2 行中使用，并在第 12 行调用。

正如这两个示例所演示的，使用 GPU 的本节开头列出的所有步骤都包含在主机代码中。可以通过库调用和 CUDA Fortran 内核循环指令（如所示）或通过使用更高级别的基于指令的模型来编程 GPU，而无需编写任何内核和设备代码；但是，在像 CUDA 这样的更低级别的模型中编程，使程序员可以控制设备资源利用率和内核执行。

2.2. CUDA Fortran 内核

CUDA Fortran 允许定义 Fortran 子例程，这些子例程在从已调用并在主机上运行的 Fortran 程序或从 CUDA 5.0 开始在设备上调用时，在 GPU 上并行执行。这样的子例程称为设备内核或内核。

对内核的调用指定必须执行多少个内核的并行实例；每个实例将由不同的 CUDA 线程执行。CUDA 线程被组织成线程块，每个线程都有一个全局线程块索引，以及其线程块内的本地线程索引。

内核是使用子例程语句上的 attributes(global) 说明符定义的；内核是使用特殊的尖括号语法调用的，以指定线程块的数量和每个线程块内的线程数

! Kernel definition
attributes(global) subroutine ksaxpy( n, a, x, y )
    real, dimension(*) :: x,y
    real, value :: a
    integer, value :: n, i
    i = (blockidx%x-1) * blockdim%x + threadidx%x
    if( i <= n ) y(i) = a * x(i) + y(i)
end subroutine

! Host subroutine
subroutine solve( n, a, x, y )
    real, device, dimension(*) :: x, y
    real :: a
    integer :: n
    ! call the kernel
    call ksaxpy<<<n/64, 64>>>( n, a, x, y )
end subroutine

在这种情况下，对内核 ksaxpy 的调用指定 n/64 个线程块，每个线程块有 64 个线程。每个线程都被分配一个线程块索引，该索引通过内置的 blockidx 变量访问，线程索引通过 threadidx 访问。在本例中，每个线程执行常见的 SAXPY 循环操作的一次迭代。

2.3. 线程块

每个线程都被分配一个线程块索引，该索引通过内置的 blockidx 变量访问，线程索引通过 threadidx 访问。线程索引可以是一维、二维或三维索引。在 CUDA Fortran 中，每个维度的线程索引都从一开始。

同一线程块中的线程可以通过使用共享内存，并通过使用 SYNCTHREADS() 内联函数在屏障处同步来协作。块中的每个线程在调用 SYNCTHREADS() 时等待，直到所有线程都到达该调用。共享内存的作用类似于低延迟、高带宽的软件管理缓存内存。目前，线程块中线程的最大数量为 1024。

可以使用许多线程块调用内核，每个线程块都具有相同的线程块大小。线程块被组织成一维、二维或三维网格的块，因此每个线程在块内都有一个线程索引，在网格内有一个块索引。当调用内核时，尖括号 <<<>>> 语法中的第一个参数是网格大小，第二个参数是线程块大小。线程块必须能够独立执行；两个线程块可以并行执行，也可以依次执行，由同一核心或不同核心执行。

dim3 派生类型在 cudafor 模块中定义，可以用于声明主机代码中的变量，如果启动配置值不是标量，则可以方便地保存它们；例如

type(dim3) :: blocks, threads
...
blocks = dim3(n/256, n/16, 1)
threads = dim3(16, 16, 1)
call devkernel<<<blocks, threads>>>( ... )

2.4. 内存层次结构

CUDA Fortran 程序可以访问多个内存空间。在主机端，主机程序可以直接访问主机主内存中的数据。它还可以直接将数据复制到设备全局内存和从中复制数据；此类数据复制需要 DMA 访问设备，因此相对于主机内存而言速度较慢。主机还可以设置设备常量内存中的值，同样使用 DMA 访问实现。

在设备端，全局设备内存中的数据可以被所有线程读取或写入。常量内存空间中的数据由主机程序初始化；所有线程都可以读取常量内存中的数据。对常量内存的访问通常比对全局内存的访问更快，但它对线程是只读的，并且大小有限。同一线程块中的线程可以访问和共享共享内存中的数据；共享内存中的数据的生命周期为线程块。每个线程也可以有私有本地内存；线程本地内存中的数据可以实现为处理器寄存器，也可以在全局设备内存中分配；当线程本地数据限制为可以分配为处理器寄存器的少量标量时，通常可以获得最佳性能。

通过使用 cudafor 模块公开的 CUDA API，支持访问 CUDA 功能，例如映射内存、点对点内存访问和统一虚拟地址空间。用户应查看相关 CUDA 文档，了解这些功能的计算能力限制。有关设备数组映射的示例，请参阅映射内存示例。

从 CUDA 6.0 开始，托管或统一内存编程在某些平台上可用。有关统一内存编程的完整描述，请参阅CUDA C 编程指南的统一内存编程部分。托管内存提供了一个公共地址空间，并在主机和设备之间迁移数据，因为它被每组处理器使用。在主机端，数据驻留在主机主内存中。在设备端，它作为驻留在全局设备内存中访问。

2.5. 子例程 / 函数限定符

CUDA Fortran 中的子例程或函数具有一个附加属性，用于指定它是在主机上还是在设备上执行，以及如果是后者，它是从主机调用的内核，还是从另一个设备子程序调用的内核。

使用 attributes(host) 或默认使用 host 属性声明的子程序称为主机子程序。
使用 attributes(global) 或 attributes(device) 声明的子程序称为设备子程序。
使用 attributes(global) 声明的子例程也称为内核子例程。
从 cc70 硬件或更高版本开始支持使用 attributes(grid_global) 声明的子例程。这些内核中网格内的线程共同驻留在同一设备上，并且可以同步。

2.5.1. Attributes(host)

host 属性在子例程或函数语句上指定，声明子例程或函数将在主机上执行。这样的子程序只能从另一个主机子程序调用。默认值为 attributes(host)，如果没有指定 host、global 或 device 属性。

2.5.2. Attributes(global)

global 属性只能在子例程语句上指定；它声明子例程是一个内核子例程，将在设备上执行，并且只能使用包含尖括号语法和运行时映射参数的内核调用来调用。

2.5.3. Attributes(device)

device 属性在子例程或函数语句上指定，声明子程序将在设备上执行；这样的例程必须从具有 global 或 device 属性的子程序调用。

2.5.4. Attributes(host,device)

host,device 属性在子例程或函数语句上指定，声明子程序可以在主机和设备上执行；这样的例程可以从主机代码或从具有 global 或 device 属性的子程序调用。它通常用于小的、目标独立的功能。

2.5.5. Attributes(grid_global)

grid_global 属性只能在子例程语句上指定；它声明子例程是一个内核子例程，将在设备上执行，并且只能使用包含尖括号语法和运行时映射参数的内核调用来启动。启动内核是为了保证网格组中的所有线程都共同驻留在设备上。这允许在 cc70 硬件及更高版本上进行网格同步操作。

2.5.6. 限制

以下限制适用于子程序。

设备子程序不得包含具有 SAVE 属性或数据初始化的变量。
内核子例程也不能同时具有 device 或 host 属性。
对内核子例程的调用必须指定执行配置，如第 9 页“设备子程序中的预定义变量”中所述。这种调用是异步的，也就是说，进行调用的调用例程在设备完成内核子例程的执行之前继续执行。
设备子程序可能不包含在主机子例程或函数中，并且可能不包含任何子例程或函数。

2.6. 变量限定符

CUDA Fortran 中的变量具有一个新的属性，声明数据分配在哪个内存中。默认情况下，在模块或主机子程序中声明的变量分配在主机主内存中。最多可以为变量指定 device、managed、constant、shared 或 pinned 属性之一。

2.6.1. Attributes(device)

具有 device 属性的变量称为设备变量，并在设备全局内存中分配。

如果在模块中声明，则该模块中的任何子程序以及使用该模块的任何子程序都可以访问该变量。
如果在主机子程序中声明，则该子程序或包含在该子程序中的子程序可以访问该变量。

设备数组可以是显式形状数组、可分配数组，也可以是假定大小或假定形状的哑元数组。可分配设备变量具有动态生命周期，从分配时到释放时。其他设备变量具有整个应用程序的生命周期。

2.6.2. Attributes(managed)

从 CUDA 6.0 开始，在某些平台上，具有 managed 属性的变量称为托管变量。托管变量可以在主机和设备代码中使用。具有托管属性的变量在主机和设备之间迁移，具体取决于内存访问的来源。托管变量可以由主机读取和写入，但如果内核在设备上处于活动状态，则对托管变量的访问存在访问限制。在设备上，托管变量具有类似于设备变量的特性，但托管变量不能像设备变量那样从设备分配，因为设备变量可以从 CUDA 5.0 开始分配以支持动态并行性。

2.6.3. Attributes(constant)

具有 constant 属性的变量称为设备常量变量。设备常量变量在设备常量内存空间中分配。常量变量必须在模块的全局数据规范范围内声明。声明后，该模块中的任何子程序以及使用该模块的任何子程序都可以访问该变量。设备常量数据可能无法在任何设备子程序中分配或修改，但可以在主机子程序中修改。对常量内存的所有主机访问都必须通过使用或主机关联。设备常量变量可能不可分配，并且它们在设备常量内存中具有整个应用程序的生命周期。

2.6.4. Attributes(shared)

具有 shared 属性的变量称为设备共享变量或共享变量。共享变量只能在设备子程序中声明，并且只能在该子程序中访问，或者由作为参数传递给它的其他设备子程序访问。共享变量可能未进行数据初始化。共享变量在线程块的设备共享内存中分配，并具有线程块的生命周期。它可以被块中的所有线程读取或写入，尽管在一个线程中的写入只有在下次调用 SYNCTHREADS() 内联函数后才能保证对其他线程可见。

2.6.5. 属性(pinned)

具有 pinned 属性的变量称为pinned 变量。Pinned 变量必须是可分配数组。当分配 pinned 变量时，它将被分配在主机页锁定内存中。使用 pinned 变量的优点是，从页锁定内存复制到设备内存比从普通分页主机内存复制更快。某些操作系统或安装可能会限制页锁定内存的使用、可用性或大小；如果页锁定内存中的分配失败，则该变量将被分配在普通主机分页内存中，并且是异步移动所必需的。

2.6.6. 属性(texture)

在新 GPU 上，不再推荐或需要通过纹理内存接口读取值，并且 CUDA 12.0 中已删除对此功能的支持。

2.6.7. 属性(unified)

从 NVHPC 24.3 版本开始，在支持它的系统上，具有 unified 属性的变量称为unified 变量。与托管变量类似，unified 变量可以在主机代码和设备代码中使用。编译器将允许为期望设备变量的参数传递 unified 变量。具有 unified 属性的变量可能在主机和设备之间迁移，但根据驱动程序版本和设置，其迁移条件可能与托管变量不同。Unified 变量在主机系统内存中创建。与托管变量类似，当从主机和设备代码访问 unified 变量时，必须小心避免可能的竞争条件。

2.7. 设备子程序中的数据类型

以下固有数据类型允许在设备子程序和设备数据中使用

表 1. 固有数据类型
类型	类型种类
`integer`	1,2,4,8
`logical`	1,2,4,8
`real`	2,4,8
`double precision`	等效于 `real(kind=8)`
`complex`	4,8
`character(len=1)`	1

派生类型可以包含具有这些固有数据类型或其他允许的派生类型的成员。

2.7.1. 半精度浮点

在支持 CUDA 计算能力 6.0 及更高版本的 NVIDIA GPU 上，可以创建半精度浮点变量和数组。CUDA Fortran 提供对实数数据类型使用 kind 属性的支持；允许将数据声明为 real(2)。以下运算符支持此数据类型：+ , -, *, /, .lt., .le., .gt., .ge., .eq.,.ne.。当使用 real(2) 并以计算能力低于 6.0 的 GPU 为目标时，编译器将发出错误消息。

半精度表示为 IEEE 754 binary16。在用于表示浮点值的 16 位中，一位用于符号，五位用于指数，十位用于有效数。当遇到无法以该格式精确表示的值时，例如当添加两个 real(2) 数字时，IEEE 754 定义了舍入规则。在 real(2) 的情况下，默认规则是就近舍入，且当出现平局时舍入到偶数，这在 IEEE 754-2008 标准的 4.3.1 节中详细描述。此格式的动态范围很小，因此大于 65520 的值将四舍五入为无穷大。

2.8. 设备子程序中的预定义变量

设备子程序可以通过几个内置的只读变量访问块和网格索引以及维度。这些变量的类型为 dim3；模块 cudafor 将派生类型 dim3 定义如下

type(dim3)
    integer(kind=4) :: x,y,z
end type

这些预定义变量在主机子程序中不可访问。

变量 threadidx 包含其线程块内的线程索引；对于一维或二维线程块，threadidx%y 和/或 threadidx%z 组件的值为 1。
变量 blockdim 包含线程块的维度；blockdim 对于同一网格中的所有线程块都具有相同的值。
变量 blockidx 包含网格内的块索引；与 threadidx 一样，对于一维网格，blockidx%y 和/或 blockidx%z 的值为 1。
变量 griddim 包含网格的维度。
常量 warpsize 被声明为 integer 类型。线程以 32 个线程为一组执行，称为warps；warpsize 包含一个 warp 中的线程数，目前为 32。

2.9. 执行配置

对内核子例程的调用必须指定执行配置。执行配置定义了执行子例程的网格和线程块的维度和范围。它还可以指定动态共享内存范围（以字节为单位）和流标识符，以支持设备上的并发流执行。

内核子例程调用如下所示

call kernel<<<grid,block[,bytes][,streamid]>>>(arg1,arg2,...)

其中

grid 和 block 可以是整数表达式（对于一维网格和线程块），也可以是 type(dim3)（对于一维或二维网格和线程块）。
如果 grid 是 type(dim3)，则每个组件的值必须等于或大于 1，并且乘积通常受设备计算能力的限制。
如果 block 是 type(dim3)，则每个组件的值必须等于或大于 1，并且组件值的乘积必须小于或等于 1024。
bytes 的值必须是整数；它指定要为每个线程块分配的共享内存字节数，除了静态分配的共享内存之外。此内存用于线程块中的假定大小共享变量；有关更多信息，请参阅共享数据。如果未指定 bytes 的值，则将其值视为零。
streamid 的值必须是大于或等于零的整数；它指定此调用关联的流。可以使用对 cudaStreamCreate 的调用来创建非零流值。从 CUDA 7.0 开始，可以指定常量 cudaStreamPerThread，以便为每个 CPU 线程使用唯一的默认流。

2.10. 异步并发执行

使用 CUDA Fortran 进行异步并发执行有两个组成部分。

2.10.1. 并发主机和设备执行

当主机子程序调用内核子例程时，该调用实际上在内核子例程开始执行之前返回到主机程序。该调用可以被视为内核启动操作，其中启动实际上对应于将内核放置在设备执行的队列中。通过这种方式，主机可以继续执行，包括调用或排队更多内核以在设备上执行。通过调用运行时例程 cudaDeviceSynchronize，主机程序可以同步并等待所有先前启动或排队的内核。

程序员在使用并发主机和设备执行时必须小心；在主机程序读取或修改设备或常量数据的情况下，主机程序应与设备同步以避免错误的结果。

2.10.2. 并发流执行

涉及设备的操作，包括内核执行以及与设备内存之间的数据复制，都是使用流队列实现的。操作被放置在流队列的末尾，并且只有在该队列上的所有先前操作都已完成时才会启动。

应用程序可以通过使用多个流来管理更多并发性。每个用户创建的流管理其自己的队列；不同流队列上的操作可能会相对于它们放置在队列中的顺序无序执行，并且可能会彼此并发执行。

默认流（在未指定流标识符时使用）是流零；流零是特殊的，因为流零队列上的操作只有在所有队列上的所有先前操作都完成后才会开始，并且在流零操作完成之前，任何队列上的后续操作都不会开始。

2.11. 内核循环指令

CUDA Fortran 允许从包含一个或多个紧密嵌套循环的主机代码区域自动生成和调用内核。循环迭代到硬件的启动配置和映射通过熟悉的 CUDA chevron 语法在指令体中控制和指定。与任何内核一样，启动是异步的。程序可以使用 cudaDeviceSynchronize() 或 CUDA 事件来等待内核完成。

指令指定的循环中的工作在线程块和网格之间并行执行；程序员有责任确保并行执行是合法的并且产生正确的结果。此规则的一个例外是标量归约操作，例如对向量或矩阵中的值求和。对于这些操作，编译器处理最终归约内核的生成，并在内核中插入适当的同步。

2.11.1. 语法

内核指令的一般形式是

!$cuf kernel do[(n)] <<< grid, block  [optional stream] >>>

编译器将网格和块值指定的启动配置映射到最外层的 n 个循环，从循环 n 开始向外工作。网格和块值可以是整数标量或带括号的列表。或者，使用星号告诉编译器选择线程块形状和/或从线程块形状和循环限制计算网格形状。未映射到网格和块值的循环在每个线程上顺序运行。

有两种方法可以指定可选的 stream 参数

!$cuf kernel do[(n)] <<< grid, block, 0, streamid >>>

或

!$cuf kernel do[(n)] <<< grid, block, stream=streamid >>>

内核循环指令示例 1

!$cuf kernel do(2) <<< (*,*), (32,4) >>>
do j = 1, m
  do i = 1, n
    a(i,j) = b(i,j) + c(i,j)
  end do
end do

在此示例中，该指令定义了一个大小为 32x4 的二维线程块。

双重嵌套循环的主体被转换为内核主体

ThreadIdx%x 从 1 运行到 32，并映射到内部 i 循环。
ThreadIdx%y 从 1 运行到 4，并映射到外部 j 循环。

网格形状指定为 (*,*)，由编译器和运行时通过将循环行程计数 n 和 m 除以线程块大小来计算，因此计算所有迭代。

内核循环指令示例 2

!$cuf kernel do <<< *, 256 >>>
do j = 1, m
  do i = 1, n
    a(i,j) = b(i,j) + c(i,j)
  end do
end do

在 do 上没有显式的 n，计划仅适用于最外层循环，即默认值为 1。在这种情况下，只有外部 j 循环以线程块大小 256 并行运行。内部 i 维度在每个线程上顺序运行。

您可能会考虑如果交换两个循环，内核循环指令示例 2 中的代码是否会执行得更好。或者，您可以指定如下配置，其中线程以合并方式读取和写入矩阵。

!$cuf kernel do(2) <<< *, (256,1) >>>
do j = 1, m
  do i = 1, n
    a(i,j) = b(i,j) + c(i,j)
  end do
end do

内核循环指令示例 3

在 内核循环指令示例 2 中，每个块中的 256 个线程各自完成矩阵加法的一个元素。沿着 i 方向和跨 j 维度的所有工作的进一步扩展由到网格维度的映射处理。

要将更多工作“展开”到每个线程中，请为网格指定非星号值，如此处所示

!$cuf kernel do(2) <<< (1,*), (256,1) >>>
do j = 1, m
  do i = 1, n
    a(i,j) = b(i,j) + c(i,j)
  end do
end do

现在，线程块中的线程协同处理 i 方向上的所有值，以 256 递增。为每个 j 创建一个线程块。具体来说，j 循环映射到网格 x 维度，因为编译器跳过了 i 循环网格大小中的常量 1。在 CUDA 内置语言中，gridDim%x 等于 m。

2.11.2. CUF 内核指令的限制

以下限制适用于 CUF 内核指令

如果指令指定 n 个维度，则其后必须至少跟随相同数量的紧密嵌套 DO 循环。
紧密嵌套的 DO 循环必须具有不变的循环限制：下限、上限和增量必须相对于内核 do 中的任何其他循环保持不变。
不变的循环限制不能是来自数组表达式的值，除非这些数组具有托管属性。
在已映射到网格和块配置值的任何循环内部或之间，都不能有 GOTO 或 EXIT 语句。
循环体可以包含赋值语句、IF 语句、循环和 GOTO 语句。
循环内只允许使用 CUDA Fortran 数据类型。
如果设备代码中允许和支持 Fortran 固有函数，则可以使用它们。
允许使用设备特定的固有函数，例如 CUDA 原子函数，但需要显式使用来自 cudadevice 模块的接口才能正确编译。
应避免使用设备特定的固有函数，例如 syncthreads 和其他 warp 或块级协作、同步或屏障函数，除非在非常有限的情况下。
如果子例程和函数调用 attributes(device) 子程序与包含指令的代码在同一模块中，则允许调用。
在循环中使用或赋值的数组必须具有 device 或 managed 属性。
指令循环内不允许使用隐式循环和 F90 数组语法。
循环中使用或赋值的标量必须具有 device 属性，否则编译器将为每个线程制作该变量的设备副本，并在循环期间保持活动状态。归约的情况除外；当归约具有标量目标时，编译器会生成正确的同步操作序列，以在设备全局内存或主机中生成一个副本。

2.11.3. 求和示例

对两个设备数组执行点积的最简单指令形式利用了先前概述的标量使用属性。

rsum = 0.0
!$cuf kernel do <<< *, * >>>
do i = 1, n
    rsum = rsum + x(i)* y(i)
end do

对于归约，编译器会识别标量的使用并仅生成一个最终结果。

此 CUF 内核之后可以在同一子程序中跟另一个 CUF 内核

!$cuf kernel do <<< *, * >>>
do i = 1, n
    rsum= x(i) * y(i)
    z(i) = rsum
end do

在此 CUF 内核中，编译器将 rsum 识别为应在每个线程上本地分配的标量临时变量。但是，在此循环之后在主机上使用 rsum 是未定义的。

2.11.4. 显式归约

CUDA Fortran 编译器通常在简单循环中很好地识别归约。当由于复杂的控制流或其他问题而未被编译器检测到归约时，从 21.7 版本开始，可以使用类似于 OpenACC 和 OpenMP 编程模型中使用的语法来指定显式归约。

value = 0.0
!$cuf kernel do <<< *, * >>> reduce(+:value)
do i = 1, n
    a(i) = real(int(a(i) * 100.0 - 50.0),kind=4)
    if (a(i) .ge. 0.0) then
        value = value + a(i)
    else
        value = value + a(i) + 50.0
    end if
end do

reduce 和 reduction 关键字都被接受。通常，OpenACC Fortran 中接受的所有数据类型和归约类型都在 CUF 内核中接受。这包括 Fortran integer 类型的 +、*、max、min、iand、ior 和 ieor；Fortran real 类型的 +、*、max、min；Fortran complex 类型的 +，以及 Fortran logical 类型的 .and.、.or.。

2.12. 使用 Fortran 模块

现代 Fortran 使用模块来打包全局数据、定义、派生类型和接口块。在 CUDA Fortran 中，这些模块可以用于轻松地在主机代码和设备代码之间传递数据和定义。本节包含一些使用 Fortran 模块的示例。

2.12.1. 从其他模块访问数据

在以下示例中，一组模块在一个文件中定义，并由另一个模块访问。

从其他模块访问数据。

在一个文件 moda.cuf 中，您可以定义一组模块

module moda
  real, device, allocatable :: a(:)
end module

module modb
  real, device, allocatable :: b(:)
end module

在另一个模块或文件 modc.cuf 中，您可以定义另一个模块，该模块使用两个模块 moda 和 modb

module modc
  use moda
  use modb
  integer, parameter :: n = 100
  real, device, allocatable :: c(:)
  contains
    subroutine vadd()
    !$cuf kernel do <<<*,*>>>
    do i = 1, n
      c(i) = a(i) + b(i)
    end do
    end subroutine
end module

在主机程序中，您使用顶层模块，并获取 n 的定义和 vadd 的接口。您还可以重命名设备数组，使其与主机命名约定不冲突

program t
use modc, a_d => a, b_d => b, c_d => c
real a,b,c(n)
allocate(a_d(n),b_d(n),c_d(n))
a_d = 1.0
b_d = 2.0
call vadd()
c = c_d
print *,all(c.eq.3.0)
end

2.12.2. 从其他模块调用例程

从 CUDA 5.0 开始，除了能够访问在另一个模块中声明的数据之外，您还可以调用包含在另一个模块中的设备函数。在以下示例中，文件 ffill.cuf 包含一个设备函数来填充数组

使用可重定位设备代码从其他模块调用例程。

module ffill
  contains
    attributes(device) subroutine fill(a)
    integer, device :: a(*)
    i = (blockidx%x-1)*blockdim%x + threadidx%x
    a(i) = i
    end subroutine
end module

要生成可重定位设备代码，请使用 -⁠gpu=rdc 标志编译此文件

% nvfortran -cuda -gpu=rdc -c ffill.cuf

现在编写另一个模块和测试程序，调用此模块中的子例程。由于您正在调用 attributes(device) 子例程，因此您不使用 chevron 语法。为方便起见，文件 tfill.cuf 中包含一个重载的 Fortran sum 函数，在本例中，该函数采用一维整数设备数组。

module testfill
    use ffill
    contains
    attributes(global) subroutine Kernel(arr)
        integer, device :: arr(*)
        call fill(arr)
    end subroutine Kernel

    integer function sum(arr)
        integer, device :: arr(:)
        sum = 0
        !$cuf kernel do <<<*,*>>>
        do i = 1, size(arr)
          sum = sum + arr(i)
        end do
    end function sum
end module testfill

program tfill
use testfill
integer, device :: iarr(100)
iarr = 0
call Kernel<<<1,100>>>(iarr)
print *,sum(iarr)==100*101/2
end program tfill

此文件也需要使用 -⁠gpu=rdc 标志编译，然后可以与先前的目标文件链接

% nvfortran -cuda -gpu=rdc tfill.cuf ffill.o

-⁠gpu=rdc 选项已成为许多版本的默认设置。-⁠gpu=nordc 标志将覆盖当前默认设置。

2.12.3. 声明设备指针和目标数组

最近，NVIDIA 增加了对指向设备数据的 F90 指针的支持。目前，这仅限于在模块作用域中声明的指针。可以通过模块关联访问指针，也可以将其传递到全局子例程中。associated() 函数在设备代码中也受支持。以下代码显示了使用 F90 指针的许多示例。这些指针也可以在 CUF 内核中使用。

在 CUDA Fortran 模块中声明设备指针和目标数组

module devptr
! currently, pointer declarations must be in a module
  real, device, pointer, dimension(:) :: mod_dev_ptr
  real, device, pointer, dimension(:) :: arg_dev_ptr
  real, device, target,  dimension(4) :: mod_dev_arr
  real, device, dimension(4) :: mod_res_arr
contains
  attributes(global) subroutine test(arg_ptr)
    real, device, pointer, dimension(:) :: arg_ptr
    ! copy 4 elements from one of two spots
    if (associated(arg_ptr)) then
      mod_res_arr = arg_ptr
    else
      mod_res_arr = mod_dev_ptr
    end if
  end subroutine test
end module devptr

program test
use devptr
real, device, target, dimension(4) :: a_dev
real  result(20)

a_dev = (/ 1.0, 2.0, 3.0, 4.0 /)

! Pointer assignment to device array declared on host,
! passed as argument.  First four result elements.
arg_dev_ptr => a_dev
call test<<<1,1>>>(arg_dev_ptr)
result(1:4) = mod_res_arr

!$cuf kernel do <<<*,*>>>
do i = 1, 4
  mod_dev_arr(i) = arg_dev_ptr(i) + 4.0
  a_dev(i)       = arg_dev_ptr(i) + 8.0
end do

! Pointer assignment to module array, argument nullified
! Second four result elements
mod_dev_ptr => mod_dev_arr
arg_dev_ptr => null()
call test<<<1,1>>>(arg_dev_ptr)
result(5:8) = mod_res_arr

! Pointer assignment to updated device array, now associated
! Third four result elements
arg_dev_ptr => a_dev
call test<<<1,1>>>(arg_dev_ptr)
result(9:12) = mod_res_arr

!$cuf kernel do <<<*,*>>>
do i = 1, 4
  mod_dev_arr(i) = 25.0 - mod_dev_ptr(i)
  a_dev(i)       = 25.0 - arg_dev_ptr(i)
end do

! Non-contiguous pointer assignment to updated device array
! Fourth four element elements
arg_dev_ptr => a_dev(4:1:-1)
call test<<<1,1>>>(arg_dev_ptr)
result(13:16) = mod_res_arr

! Non-contiguous pointer assignment to updated module array
! Last four elements of the result
nullify(arg_dev_ptr)
mod_dev_ptr => mod_dev_arr(4:1:-1)
call test<<<1,1>>>(arg_dev_ptr)
result(17:20) = mod_res_arr

print *,all(result==(/(real(i),i=1,20)/))
end

2.12.4. 声明纹理

在新 GPU 上，不再推荐或需要通过纹理内存接口读取值，并且 CUDA 12.0 中已删除对此功能的支持。

CUDA 纹理内存可以通过特殊的纹理属性来获取，该属性归因于指向具有 target 属性的设备数据的 F90 指针。在 CUDA Fortran 中，纹理目前仅用于只读数据，这些数据通过纹理缓存传递。由于有单独的硬件来支持此缓存，因此在许多情况下，使用纹理属性可以提高性能，尤其是在线程之间的访问不规则且不连续的情况下。以下简单示例演示了此功能

在 CUDA Fortran 模块中声明纹理

module memtests
  real(8), texture, pointer :: t(:)  ! declare the texture
  contains
    attributes(device) integer function bitrev8(i)
    integer ix1, ix2, ix
    ix = i
    ix1 = ishft(iand(ix,z'0aa'),-1)
    ix2 = ishft(iand(ix,z'055'), 1)
    ix = ior(ix1,ix2)
    ix1 = ishft(iand(ix,z'0cc'),-2)
    ix2 = ishft(iand(ix,z'033'), 2)
    ix = ior(ix1,ix2)
    ix1 = ishft(ix,-4)
    ix2 = ishft(ix, 4)
    bitrev8 = iand(ior(ix1,ix2),z'0ff')
    end function bitrev8

    attributes(global) subroutine without( a, b )
    real(8), device :: a(*), b(*)
    i = blockDim%x*(blockIdx%x-1) + threadIdx%x
    j = bitrev8(threadIdx%x-1) + 1
    b(i) = a(j)
    return
    end subroutine

    attributes(global) subroutine withtex( a, b )
    real(8), device :: a(*), b(*)
    i = blockDim%x*(blockIdx%x-1) + threadIdx%x
    j = bitrev8(threadIdx%x-1) + 1
    b(i) = t(j)  ! This subroutine accesses a through the texture
    return
    end subroutine
end module memtests

program t
use cudafor
use memtests
real(8), device, target, allocatable :: da(:)
real(8), device, allocatable :: db(:)
integer, parameter :: n = 1024*1024
integer, parameter :: nthreads = 256
integer, parameter :: ntimes = 1000
type(cudaEvent) :: start, stop
real(8) b(n)

allocate(da(nthreads))
allocate(db(n))

istat = cudaEventCreate(start)
istat = cudaEventCreate(stop)

db = 100.0d0
da = (/ (dble(i),i=1,nthreads) /)

call without<<<n/nthreads, nthreads>>> (da, db)
istat = cudaEventRecord(start,0)
do j = 1, ntimes
  call without<<<n/nthreads, nthreads>>> (da, db)
end do
istat = cudaEventRecord(stop,0)
istat = cudaDeviceSynchronize()
istat = cudaEventElapsedTime(time1, start, stop)
time1 = time1 / (ntimes*1.0e3)
b = db
print *,sum(b)==(n*(nthreads+1)/2)

db = 100.0d0
t => da  ! assign the texture to da using f90 pointer assignment

call withtex<<<n/nthreads, nthreads>>> (da, db)
istat = cudaEventRecord(start,0)
do j = 1, ntimes
  call withtex<<<n/nthreads, nthreads>>> (da, db)
end do
istat = cudaEventRecord(stop,0)
istat = cudaDeviceSynchronize()
istat = cudaEventElapsedTime(time2, start, stop)
time2 = time2 / (ntimes*1.0e3)
b = db
print *,sum(b)==(n*(nthreads+1)/2)

print *,"Time with    textures",time2
print *,"Time without textures",time1
print *,"Speedup with textures",time1 / time2

deallocate(da)
deallocate(db)
end

2.13. CUDA Fortran 条件编译

本节介绍在应用程序中启用 CUDA Fortran 的几种方法，同时最大限度地减少为维护单个 CPU/GPU 代码库而进行的更改。

如果在编译中启用了 CUDA Fortran，可以通过在命令行上指定 -⁠cuda 或使用 .cuf 或 .CUF 扩展名重命名文件来启用，那么对于以 !@cuf 哨兵开头的源代码行，该行其余部分显示为语句，否则整行都是注释。

如果在编译中启用了 CUDA Fortran，可以通过在命令行上指定 -⁠cuda 来启用，并且预处理通过 -⁠Mpreprocess 编译器选项或通过在文件名扩展名（.CUF、.F90 等）中使用大写字母来启用，则定义 _CUDA 宏。

如果在编译中启用了 CUDA Fortran，则 CUF 内核指令（由 !$cuf kernel 表示）将为该循环生成设备代码。否则，代码将在 CPU 上运行。

变量重命名可以通过上述技术的组合以及 use ..., only: Fortran 语句来重命名模块变量来完成。例如，您可以结合 CUF 哨兵或 _CUDA 宏，使用 use device_declaration_mod, only : a => a_dev, b => b_dev 来重命名模块中包含的设备数组。Fortran associate 块可以类似地使用，并提供对变量重命名更精细的控制。

此示例显示了许多这些技术，并且可以在启用或不启用 CUDA Fortran 的情况下编译和运行。

program p
!@cuf use cudafor
real a(1000)
!@cuf attributes(managed) :: a
real b(1000)
!@cuf real, device :: b_dev(1000)
b = 2.0
!@cuf b_dev = b
!@cuf associate(b=>b_dev)
!$cuf kernel do(1) <<<*,*>>>
do i = 1, 1000
    a(i) = real(i) * b(i)
end do
!@cuf end associate
#ifdef _CUDA
print *,"GPU sum passed? ",sum(a).eq.1000*1001
#else
print *,"CPU sum passed? ",sum(a).eq.1000*1001
#endif
end program

2.14. 构建 CUDA Fortran 程序

当文件名使用 CUDA Fortran 扩展名时，NVIDIA Fortran 编译器支持 CUDA Fortran。.cuf 扩展名指定该文件是自由格式的 CUDA Fortran 程序；也可以使用 .CUF 扩展名，在这种情况下，程序在编译之前由预处理器处理。要编译固定格式的程序，请添加命令行选项 -⁠Mfixed。可以通过添加 -⁠cuda 命令行选项在任何 Fortran 源代码文件中启用 CUDA Fortran 扩展。重要的是要记住，如果您使用 -⁠cuda 命令行选项编译文件，则还必须使用 -⁠cuda 命令行选项链接该文件。如果您使用 -⁠cuda 编译，但不使用 -⁠cuda 链接，您将收到对符号 cuda_compiled 的未定义引用。

要更改从默认值使用的 CUDA Toolkit 版本，请指定 -⁠cuda -⁠gpu=cudaX.Y；必须安装 CUDA Toolkit 版本 X.Y。

默认情况下生成可重定位设备代码。您可以通过指定 -⁠cuda -⁠gpu=nordc 来覆盖此选项。

如果您正在使用 CUDA 内核循环指令的许多实例，即 CUF 内核，您可能需要添加 -⁠Minfo 开关来验证是否在您期望的位置生成了 CUDA 内核，以及您是否遵循了前面章节中概述的限制。

2.15. 托管和统一内存选项和互操作性

CUDA Fortran 是 Fortran 开发人员可用的几种 GPU 编程模型之一。其他模型使用了 CUDA Fortran 开发人员可能会觉得有用的编译器选项。有关现在支持的内存模型的通用讨论，请参见 NVIDIA HPC Compiler User’s Guide，可在线访问 docs.nvidia.com/hpc-sdk。

多年来一直支持的一个 nvfortran 编译器选项是 -⁠gpu=mem:managed。这在 stdpar 编程模型中尤其有用。对于 CUDA Fortran，此选项的作用是对所有 Fortran 可分配数据使用 cudaMallocManaged()，实际上是将可分配数组视为它们具有托管属性（请参阅 2.6.2）。这允许这些数组在全局子例程和 CUF 内核中使用，并传递到通常采用设备数组的库函数中。

这样做的一个缺点是，当数组通过子例程级别传递时，Fortran 编译器可能会丢失信息，即该数组最初是可分配的。因此，托管属性行为可能会丢失；它在顶层函数中以一种方式工作，但在您真正需要的叶函数中则不工作。有一些方法可以解决这个问题，但它们通常是不希望对代码进行的更改。与 OpenACC 或 stdpar 不同，CUDA Fortran 没有隐式数据移动。所有操作都是显式的，在开发人员的控制之下，通过数据属性、赋值语句和 API 调用（如 cudaMemcpy）。

从 23.11 版本开始，在支持 HMM/ATS 和统一内存的系统上，NVHPC 编译器现在支持一个名为 -⁠gpu=mem:unified 的选项。这类似于 managed 选项，但它不仅适用于可分配数据，而且适用于所有主机数据：可分配数据、本地堆栈数据和全局静态数据。所有程序数据都可以在 GPU 上访问。在某些方面，对于 CUDA Fortran 而言，此选项所做的只是消除编译器错误和警告，这些错误和警告指出在需要设备数据的地方使用了主机数据。CPU 和 GPU 访问之间来回的底层数据移动由操作系统和 CUDA 驱动程序处理，并且不需要单独的主机和设备数据副本。

当然，已经针对两个离散内存进行调优，并利用异步操作、多个流以及 CPU 和 GPU 并发操作的 CUDA 程序，在使用这些选项时可能会遇到竞争条件，并且现在只有一个数据副本，而不是两个。为了帮助调试这些问题，NVCOMPILER_ACC_SYNCHRONOUS 环境变量现在接受一个位字段。将此环境变量的值设置为 2 将在每个 CUF 内核的末尾插入一个同步点，将其设置为 4 将在每次全局内核启动的末尾插入一个同步点。

以下是一个简单的 CUDA Fortran 程序的示例，该程序现在可以在使用 -⁠gpu=mem:unified 选项编译时作用于统一内存

module m1
integer, parameter :: N = 5
integer :: m(N)

contains
  attributes(global) subroutine g1( a )
    integer :: a(*)
    i = threadIdx%x
    if (i .le. N) a(i) = m(i)
    return
  end subroutine g1
end module m1

program t1
  use m1
  use cudafor
  integer :: istat, a(N)
  m = [ ((i),i=1,N) ]  ! Init global data
  call g1 <<<1,N>>> (a)
  istat = cudaDeviceSynchronize()
  print *,a(1:N)
end program t1

请注意，我们添加了对 cudaDeviceSynchronize() 的调用，因为统一数据在设备上读取和写入，从主机打印，并且全局内核启动仍然相对于主机是异步的。

此外，从 24.3 版本开始，同样在支持 HMM/ATS 和统一内存的系统上，NVHPC CUDA Fortran 编译器支持 unified 属性。一般来说，具有 unified 属性的数据的行为类似于 managed 数据，但是它是使用系统内存分配的，而不是使用 cudaMallocManaged()。该属性允许程序员在逐个变量的基础上启用统一内存，并且不需要使用 -⁠gpu=mem:unified 进行编译。

当使用 Managed Memory Model、Unified Memory Model 或变量的 managed 或 unified 属性时，请参考 cudaMemAdvise 或 cudaMemPrefetchAsync 以获取内存提示，这些提示已证明在许多情况下可以提高应用程序性能。

3. 参考

本节是 CUDA Fortran 语言参考。

3.1. 新的子例程和函数属性

CUDA Fortran 为子例程和函数添加了新的属性。本节介绍如何指定新属性、它们的含义和限制。

子例程可以具有 host、global 或 device 属性，或者可以同时具有 host 和 device 属性。函数可以具有 host 或 device 属性，或者两者都具有。这些属性使用子例程或函数语句上的 attributes(attr) 前缀指定；如果子程序语句上没有 attributes 前缀，则使用默认规则，如下节所述。

3.1.1. Host 子例程和函数

host 属性可以在子例程或函数语句上显式指定，如下所示

attributes(host) subroutine sub(...)
attributes(host) integer function func(...)
integer attributes(host) function func(...)

host 属性前缀可以位于任何其他允许的子例程或函数前缀说明符（recursive、pure、elemental、function return datatype）之前或之后。具有 host 属性的子例程或函数称为 host 子例程或函数，或主机子程序。主机子程序被编译为在主机处理器上执行。没有 attributes 前缀的子程序默认具有 host 属性。

3.1.2. Global 和 Grid_Global 子例程

global 和 grid_global 属性可以在子例程语句上显式指定，如下所示

attributes(global) subroutine sub(...)

attributes(grid_global) subroutine subg(...)

函数不能具有 global 属性。具有 global 属性或 grid_global 属性的子例程称为内核子例程。内核子例程不能是递归的、纯的或元素级的，因此不允许使用其他子例程前缀。内核子例程被编译为在设备上执行的内核，以便从主机例程中使用执行配置调用。内核子例程不能包含在另一个子例程或函数中，也不能包含任何其他子程序。grid_global 子例程在 cc70 硬件或更高版本上受支持，并指定内核应以允许网格中的所有线程同步的方式启动。

启动边界可以可选地指定为 global 子例程定义的一部分，以向编译器提供优化提示。这主要有助于寄存器使用、溢出和低级代码生成中使用的占用率启发式方法。有关更多信息，请参阅 CUDA C 编程指南。CUDA Fortran 中使用的形式是

attributes(global) launch_bounds(maxTPB, minBPM) subroutine sub(...)

其中 maxTPB 是 maxThreadsPerBlock，应用程序将永远启动的最大每块线程数，minBPM 是期望的每多处理器最小驻留块数。这两个值都必须是数字常量。

从 23.3 版本开始，Hopper (cc90) 和更高版本的目标启用了线程块集群支持。要指定集群的维度，请使用 cluster_dims 语法并指定每个 x、y 和 z 维度。值必须是数字常量。有关更多信息，请参阅 CUDA C 编程指南。例如，CUDA Fortran 中的这个例子

attributes(global) cluster_dims(2,2,1) subroutine sub(...)

将设置集群中 2x2（x 和 y）的线程块集。使用尖括号语法的对这些内核的启动将在调用站点进行适当调整。

此外，作为 Hopper 支持的一部分，launch_bounds 语法已扩展为接受第三个参数，即集群大小的上限。

3.1.3. Device 子例程和函数

device 属性可以在子例程或函数语句上显式指定，如下所示

attributes(device) subroutine sub(...)
attributes(device) datatype function func(...)
datatype attributes(device) function func(...)

具有 device 属性的子例程或函数称为设备子程序。设备子程序被编译为在设备上执行，并且可以从内核子例程或其他设备子程序调用。设备子程序也可以是递归的、纯的或元素级的。具有 device 属性的子例程或函数可以与调用者位于不同的文件或作用域中，但您必须使用可重定位的设备代码链接，并提供显式接口。否则，设备例程应与调用者位于同一模块中。

3.1.4. 内核子例程和设备子程序的限制

具有 device 或 global 属性的子例程或函数必须满足以下限制

它不能包含另一个子程序。
它不能包含在另一个子例程或函数中。
内核子例程不能是递归的，也不能在子例程语句上具有 recursive 前缀。
内核子例程不能是纯的或元素级的，也不能在子例程语句上具有 pure 或 elemental 前缀。

有关更多信息，请参阅设备代码。

3.2. 变量属性

CUDA Fortran 为变量和数组添加了新的属性。本节介绍如何指定新属性以及它们的含义和限制。

在主机子程序中声明的变量可以具有三个新属性之一：它们可以声明为位于设备全局内存、managed 内存或 pinned 内存中。

模块中的变量可以声明为位于设备全局内存、managed 内存空间或 constant 内存空间中。

在设备程序单元中声明的变量可以具有三个新属性之一：它们可以声明为位于设备全局内存、constant 内存空间、线程块共享内存中，或者没有任何附加属性，它们将被分配在线程本地内存中。出于性能和可用性原因，value 属性也可以用于标量哑元，以便它们按值传递，而不是 Fortran 默认的按引用传递参数。

3.2.1. Device 数据

具有 device 属性的变量或数组被定义为驻留在设备全局内存中。device 属性可以使用 attributes 语句指定，也可以作为类型声明语句的属性。以下示例声明了两个数组，a 和 b，作为大小为 100 的设备数组。

real :: a(100)
attributes(device) :: a
real, device :: b(100)

以下规则适用于 device 数据

可分配的 device 数组动态分配设备全局内存。
Device 变量和数组可以出现在模块中，但不能出现在 Common 块或 Equivalence 语句中。
除非是可分配的，否则派生类型的成员不能具有 device 属性。
Device 变量和数组可以作为实参传递给主机和设备子程序；在这种情况下，子程序接口必须是显式的（在 Fortran 意义上），并且匹配的哑元也必须具有 device 属性。
除非是可分配的，否则在主机子程序中声明的 Device 变量和数组不能具有 Save 属性。

在主机子程序中，device 数据只能以下列方式使用

在声明语句中
在 Allocate 和 Deallocate 语句中
作为 Allocated intrinsic 函数的参数
作为数据传输赋值语句中的源或目标
作为内核子例程的实参
作为另一个主机子程序或运行时 API 调用的实参
作为主机子程序中的哑元

Device 数组可以具有 allocatable 属性，或者可以具有可调整的范围。

3.2.2. Managed 数据

具有 managed 属性的变量或数组由统一内存系统管理，并在主机主内存和设备全局内存之间迁移。managed 属性可以使用 attributes 语句指定，也可以作为类型声明语句的属性。Managed 局部数组可以具有显式形状、是自动的或可分配的。Managed 哑元数组可以是假定大小或假定形状的。以下示例声明了两个数组，a 和 b，作为大小为 100 的 managed 数组，并分配了第三个数组 c，大小为 200。

real :: a(100)
attributes(managed) :: a
real, managed :: b(100)
real, allocatable, managed :: c(:)
. . .
allocate(c(200))

以下规则适用于主机上的 managed 数据

Managed 变量和数组可以出现在主机子程序和模块中，但不能出现在 Common 块或 Equivalence 语句中。
除非是可分配的，否则在主机子程序中声明的 Managed 变量和数组不能具有 Save 属性。
派生类型可以具有 managed 属性。
派生类型的成员可以具有 managed 属性。
Managed 派生类型也可以包含可分配的 device 数组。
Managed 变量和数组可以作为实参传递给其他主机子程序；如果子程序接口被重载，则通用匹配优先级是首先匹配另一个 managed 哑元，其次匹配具有 device 属性的哑元，最后匹配没有（或 host）属性的哑元。
将非 managed 实参传递给 managed 哑元将导致编译错误（如果接口是显式的），否则会导致意外行为。
Managed 变量和数组可以像 device 变量和数组一样作为实参传递给 global 子例程。
默认情况下，managed 数据以全局作用域分配，即传递给 cudaMallocManaged 的标志是 cudaMemAttachGlobal。
可以通过调用 ``cudaStreamAttachMemAsync`` 来更改 managed 变量的作用域。
可以通过调用 cudaforSetDefaultStream 将单个 managed 变量与给定的流关联起来。
所有后续分配的 managed 变量也可以通过调用 cudaforSetDefaultStream 与给定的流关联起来。
当正在运行的内核在设备上的同一作用域内访问 managed 数据时，在主机上访问 managed 数据将导致分段错误或竞争条件。

以下规则适用于设备上的 managed 数据

managed 属性可以用于哑元。
Managed 数据被视为好像它是 device 数据一样。
不支持在设备上分配或释放 managed 数据。

注意

即使您的应用程序仅使用单个 GPU，如果您在具有多个非对等互连启用的 GPU 的系统上运行，managed 内存也将被分配为零复制内存，并且性能将因此受到影响。一种解决方法是设置环境变量 CUDA_VISIBLE_DEVICES 以便仅看到一个 GPU，或者通过设置 CUDA_MANAGED_FORCE_DEVICE_ALLOC 强制在 GPU 上分配。CUDA C 编程指南在统一内存编程部分中有关于此的更多详细信息。

3.2.3. Unified 数据

具有 unified 属性的变量或数组可以从主机和设备代码访问。数据是否迁移或跨内存总线读取和写入受 CUDA 驱动程序和设置的控制。unified 属性可以使用 attributes 语句指定，也可以作为类型声明语句的属性。Unified 数组必须是局部的、固定大小的、自动的或可分配的。unified 属性不允许用于全局模块数组。局部 F90 指针可以具有 unified 属性，然后指向具有 target 属性的全局数组。此指针可以作为内核参数传递到设备代码内核中，类似于 CUDA C++ 中的统一内存支持。

以下规则适用于主机上的 unified 数据

Unified 变量和数组只能出现在主机子程序中，而不能出现在模块中。
派生类型可以具有 unified 属性。
派生类型的成员可以具有 unified 属性。
Unified 派生类型也可以包含可分配的 device 数组，这对于深层数据结构很有用。
当正在运行的内核在设备上访问 unified 数据时，在主机上访问 unified 数据可能会导致竞争条件。

以下规则适用于设备上的 unified 数据

Unified 数据被视为好像它是 device 数据一样。
不支持在设备上分配或释放 unified 数据。

当存在通用的、重载的接口公开时，下表可能有助于解释如何计算实参与哑元的匹配。对于每对参数，返回一个距离。最小距离（小于无穷大）获胜。

表 2. 属性化参数匹配距离值
哑元	实参 None (host)	实参 Device	实参 Managed	实参 Unified	实参 OpenACC use_ device	实参 None (gpu= mem: unified)	实参 None (gpu= mem: managed)
`None(host)`	0	INF	3	3	1	3	3
`Device`	INF	0	2	2	0	2	2
`Managed`	INF	INF	0	1	INF	1	0
`Unified`	INF	INF	1	0	INF	0	1

应该注意的是，CUDA Fortran 中提供的 Fortran 主机模块（例如本章后面讨论的 CUDA 库提供的模块）包含接口，在这些接口中，哑元是主机或设备。因此，对于 device、managed 或 unified 的实参，仍然首选对设备数据进行操作的实现。

还要记住，匹配是基于数据的使用方式和位置。在分配和释放数据时必须小心，通常属性必须在这两个操作期间完全匹配。下一节中讨论的 pinned 属性具有相同的要求。

3.2.4. Pinned 数组

具有 pinned 属性的可分配数组将在特殊的页锁定主机内存中分配，当此类内存可用时。使用 pinned 内存的优点是设备和 pinned 内存之间的传输更快，并且可以是异步的。具有 pinned 属性的数组可以在模块或主机子程序中声明。pinned 属性可以使用 attributes 语句指定，也可以作为类型声明语句的属性。以下示例声明了两个数组，p 和 q，作为 pinned 可分配数组。

real :: p(:)
allocatable :: p
attributes(pinned) :: p
real, allocatable, pinned :: q(:)

Pinned 数组可以作为参数传递给主机子程序，无论接口是否显式，或者哑元是否具有 pinned 和 allocatable 属性。在释放数组的位置，数组的声明仍然必须具有 pinned 属性，否则释放可能会失败。

3.2.5. Constant 数据

具有 constant 属性的变量或数组被定义为驻留在设备 constant 内存空间中。constant 属性可以使用 attributes 语句指定，也可以作为类型声明语句的属性。以下示例声明了两个数组，c 和 d，作为大小为 100 的 constant 数组。

real :: c(100)
attributes(constant) :: c
real, constant :: d(100)

以下规则适用于 constant 数据

Constant 变量和数组可以出现在模块中，但不能出现在 Common 块或 Equivalence 语句中。出现在模块中的 Constant 变量可以通过主机和设备子程序中的 use 语句访问。
Constant 数据不能具有 Pointer、Target 或 Allocatable 属性。
派生类型的成员不能具有 constant 属性。
具有 constant 属性的数组必须具有固定大小。
Constant 变量和数组可以作为实参传递给主机和设备子程序，只要子程序接口是显式的，并且匹配的哑元也具有 constant 属性。Constant 变量不能在主机子程序和设备 global 子程序之间作为实参传递。
在设备子程序中，具有 constant 属性的变量和数组不能被赋值或修改。
在主机子程序中，具有 constant 属性的数据可以被读取和写入。

在主机子程序中，具有 constant 属性的数据只能以下列方式使用

作为 USE 语句中的命名实体。
作为数据传输赋值语句中的源或目标
作为另一个主机子程序的实参
作为主机子程序中的哑元

3.2.6. Shared 数据

具有 shared 属性的变量或数组被定义为驻留在线程块的共享内存空间中。shared 变量或数组只能在设备子程序内部声明和使用。shared 属性可以使用 attributes 语句指定，也可以作为类型声明语句的属性。以下示例声明了两个数组，s 和 t，作为大小为 100 的 shared 数组。

real :: c(100)
attributes(shared) :: c
real, shared :: d(100)

以下规则适用于 shared 数据

Shared 数据不能具有 Pointer、Target 或 Allocatable 属性。
Shared 变量不能出现在 Common 块或 Equivalence 语句中。
派生类型的成员不能具有 shared 属性。
Shared 变量和数组可以从设备子程序传递到另一个设备子程序作为实参，只要接口是显式的并且匹配的哑元具有 shared 属性。

不是哑元的 Shared 数组可以声明为假定大小的数组；也就是说，shared 数组的最后一个维度可以有一个星号作为其上限

real, shared :: x(*)

这样的数组具有特殊的意义。它的大小在运行时由对内核的调用确定。调用内核时，执行配置中的 bytes 参数的值用于指定为每个线程块动态分配的共享内存字节数。此内存用于该线程块中的假定大小的共享内存数组；如果存在多个假定大小的共享内存数组，则它们都隐式地等价，从相同的共享内存地址开始。程序员在编码时必须考虑到这一点。

Shared 数组可以声明为 Fortran 自动数组。对于自动数组，边界被声明为包含常量、参数、blockdim 变量和按值传递的整数参数的表达式。自动数组的分配也来自通过尖括号启动配置指定的动态区域。如果声明了多个自动数组，则编译器和运行时会管理到动态区域的偏移量。程序员必须在尖括号启动配置共享内存值中提供足够的字节数，以覆盖全局子程序中声明的所有自动数组。

attributes(global) subroutine sub(A, n,
integer, value :: n, nb
real, shared :: s(nb*blockdim%x,nb)

如果 shared 数组不是哑元，也不是假定大小或自动的，则它必须是固定大小的。在这种情况下，shared 数组的分配不是来自启动配置中指定的动态分配的共享内存区域，而是静态地在函数中声明的。如果全局例程仅使用固定大小的 shared 数组，或者根本不使用，则在启动时不需要指定共享内存量。

3.2.7. 纹理数据

在新 GPU 上，不再推荐或需要通过纹理内存接口读取值，并且 CUDA 12.0 中已删除对此功能的支持。

只读实数和整数 device 数据可以通过纹理内存在设备子程序中访问，方法是将 F90 指针变量分配给底层设备数组。要以这种方式使用纹理内存，请按照以下步骤操作

向包含设备代码的模块声明部分添加声明，以便该声明可以通过主机关联对设备子程序可用，并通过主机或 use 关联对主机代码可用
```
real, texture, pointer :: t(:)
```
在您的主机代码中，将 target 属性添加到您希望通过纹理内存访问的 device 数据
- 更改：real, device :: a(n)
- 改为：real, target, device :: a(n)
target 属性是标准 F90/F2003 语法，用于表示可能被另一个实体“指向”的数组或其他数据结构。
通过在您的主机代码中使用 F90 指针赋值运算符将纹理声明绑定到 device 数组。像下面这样的简单表达式执行所有底层的 CUDA 纹理绑定操作。
```
t => a
```
可以通过主机关联引用 t 的 CUDA Fortran 设备代码现在可以访问 t 的元素，而无需更改语法。

在以下示例中，针对 a 的 t 的访问通过纹理缓存进行。
```
! Vector add, s through device memory, t is through texture memory
i = threadIdx%x + (blockIdx%x-1)*blockDim%x
s(i) = s(i) + t(i)
```

3.2.8. Value 哑元

在设备子程序中，按照 Fortran 的规则，哑元默认情况下按引用传递。这意味着实参必须存储在设备全局内存中，并且参数的地址传递给子程序。标量参数可以按值传递，就像在 C 中一样，方法是将 value 属性添加到变量声明中。

attributes(global) subroutine madd( a, b, n )
      real, dimension(n,n) :: a, b
      integer, value :: n

在这种情况下，n 的值可以从主机传递，而无需驻留在设备内存中。对应于哑元 a 和 b 的变量数组必须在调用之前设置为驻留在设备上。

3.3. 分配设备内存、Pinned 内存和 Managed 内存

本节介绍 Allocate 语句的扩展，专门用于动态分配 device 数组、主机 pinned 数组、managed 数组以及其他支持的用于分配 CUDA Fortran 特有的内存的方法。

3.3.1. 分配设备内存

Device 数组可以具有 allocatable 属性。这些数组在主机子程序中使用 Allocate 语句动态分配，并使用 Deallocate 语句动态释放。如果在主机子程序中声明的 device 数组没有 Save 属性，它将在子程序返回时自动释放。

real, allocatable, device :: b(:)
allocate(b(5024),stat=istat)
...
if(allocated(b)) deallocate(b)

可以使用 Fortran 2003 可分配标量功能在设备上分配标量变量。要使用这些功能，请在主机上将标量声明和初始化为

integer, allocatable, device :: ndev
allocate(ndev)
ndev = 100

该语言还支持创建等效于自动和局部 device 数组的能力，而无需使用 allocate 语句。这些数组的生命周期也将与 Fortran 语言通常的子程序相同

subroutine vfunc(a,c,n)
    real, device :: adev(n)
    real, device :: atmp(4)
    ...
end subroutine vfunc   ! adev and atmp are deallocated

以这种方式声明的自动和局部数组（不包含 allocatable 属性）不能具有 Save 属性。

3.3.2. 使用运行时例程分配设备内存

对于熟悉 CUDA C 编程环境的程序员，我们提供了 CUDA 内存管理运行时例程的 Fortran 接口。这些函数返回的内存将绕过某些 Fortran 可分配属性（例如自动释放），因此这些数组更像是 C 语言中 malloc 分配的区域。不支持对给定数组混合使用标准 Fortran allocate/deallocate 和运行时 Malloc/Free。

cudaMalloc 函数可用于分配受支持的固有数据类型的单维数组，而 cudaFree 可用于释放它。

real, allocatable, device :: v(:)
istat = cudaMalloc(v, 100)
...
istat = cudaFree(v)

有关内存管理运行时例程的完整列表，请参阅内存管理。

3.3.3. 分配固定内存

具有 pinned 属性的可分配数组使用 Allocate 语句动态分配。如果可用，编译器将生成代码以在主机页锁定内存中分配数组。如果此类内存空间不可用或已耗尽，则编译器会在正常的分页主机内存中分配数组。否则，pinned 可分配数组的工作方式和行为与主机上的任何其他可分配数组相同。

real, allocatable, pinned :: p(:)
allocate(p(5000),stat=istat)
...
if(allocated(p)) deallocate(p)

为了确定页锁定内存的分配是否成功，allocate 语句添加了一个额外的 PINNED 关键字。它返回一个逻辑成功值。

logical plog
allocate(p(5000), stat=istat, pinned=plog)
if (.not. plog) then
. . .

3.3.4. 分配托管内存

托管数组可能具有也可能不具有 allocatable 属性。这些数组都像设备数组一样动态分配。

real, allocatable, managed :: b(:)
allocate(b(5024),stat=istat)
...
if(allocated(b)) deallocate(b)

CUDA Fortran 支持创建等效于自动和局部托管数组的功能，而无需使用 allocate 语句。这些数组的生命周期也将与 Fortran 语言通常的子程序生命周期相同。

subroutine vfunc(a,c,n)
    real, managed :: aman(n)
    real, managed :: atmp(4)
    ...
end subroutine vfunc   ! aman and atmp are deallocated

3.3.5. 使用运行时例程分配托管内存

cudaMallocManaged 函数可用于分配受支持的固有数据类型的单维托管数组，而 cudaFree 可用于释放它。

use cudafor
real, allocatable, managed :: v(:)
istat = cudaMallocManaged(v, 100, cudaMemAttachHost)
...
istat = cudaFree(v)

有关内存管理运行时例程的完整列表，请参阅内存管理。

3.3.6. 异步分配设备内存

从 CUDA 11.2 开始，可以使用 Allocate 语句在主机子程序中异步地在指定的流上动态分配可分配的设备数组。

real, allocatable, device :: b(:)
integer(kind=cuda_stream_kind) :: istream
...
allocate(b(5024),stream=istream)

这些数组也可以使用 Deallocate 语句动态释放。没有必要或不允许在释放期间指定流。如果在主机子程序中声明的设备数组没有 Save 属性，它将在子程序返回时自动释放。给定上面的分配，此语句将在 istream 指定的流上释放数组 b。

if(allocated(b)) deallocate(b)

使用带有流的 Allocate 语句声明的数组与该流相关联，就像为设备数据和流的组合调用了 cudaforSetDefaultStream 函数一样。为了在此流之外的操作中使用此数据，用户应首先调用 cudaStreamSynchronize 以阻止主机执行，直到所有流操作完成。

3.3.7. 使用运行时例程异步分配设备内存

cudaMallocAsync 函数可用于在给定流上异步地分配受支持的固有数据类型的单维数组，而 cudaFreeAsync 可用于异步地释放它。

real, allocatable, device :: v(:)
integer(kind=cuda_stream_kind) :: istream
istat = cudaMallocAsync(v, 100, istream)
...
istat = cudaFreeAsync(v, istream)

有关内存管理运行时例程的完整列表，请参阅内存管理。

3.3.8. 控制设备数据是否为托管内存

从 HPC SDK 编译器版本 21.9 开始，可以更改 CUDA Fortran 设备数据分配行为，以实际分配托管内存而不是设备内存，并且可能无需更改任何代码。

这对于超额订阅可用的 GPU 内存非常有用，并允许操作系统和驱动程序根据需要将内存分页到 GPU 和从 GPU 分页，无论是作为实验还是为了运行比通常可用内存更大的问题规模。

所有 CUDA Fortran 设备分配都通过一个小的包装层，然后再进行实际的 CUDA API 调用。通过设置环境变量

NVCOMPILER_CUDAFOR_DEVICE_IS_MANAGED=1

本节前两个小节形式的分配，分配设备内存 和 使用运行时例程分配设备内存 最终将调用 cudaMallocManaged 而不是 cudaMalloc。此外，还添加了一些预取提示，以使从 GPU（当前设备）访问新分配的数据最为有效。

3.4. 主机和设备内存之间的数据传输

本节介绍在主机和设备内存之间传输数据的方法。

3.4.1. 使用赋值语句进行数据传输

您可以使用主机子程序中的简单赋值语句将变量和数组从主机内存复制到设备内存。默认情况下，使用赋值语句读取或写入设备、托管或常量数据隐式使用 CUDA 流零。这意味着此类数据复制是同步的，并且数据复制会等待直到所有先前的内核和数据复制完成。或者，您可以使用 cudaforSetDefaultStream 调用将一个或多个设备和托管变量与特定流关联。在此调用发生后，对这些变量的赋值语句将在指定的流上异步运行。

有关赋值语句的具体信息

左侧是设备变量或设备数组或数组切片，右侧是主机变量或主机数组或数组切片的赋值语句，将数据从主机内存复制到设备全局内存。
左侧是主机变量或主机数组或数组切片，右侧是设备变量或设备数组或数组切片的赋值语句，将数据从设备全局内存复制到主机内存。
赋值语句的两侧都是设备变量或设备数组或数组切片，则在两个设备变量或数组之间复制数据。

同样，您可以使用简单的赋值语句来复制或分配具有 constant 属性的变量或数组。

有关赋值语句和托管数据的具体信息

左侧是托管变量或托管数组，右侧是符合的标量常量、主机变量、主机数组或数组切片的赋值语句，使用 cudaMemcpy、memset 或类似操作将数据从主机内存复制到设备全局内存。
左侧是托管数组切片，右侧是任何主机变量的赋值语句，使用生成的主机代码复制数据。
左侧是托管变量、托管数组或数组切片，右侧是设备变量或设备数组或数组切片的赋值语句，使用 cudaMemcpy 或类似操作将数据从设备全局内存复制到主机内存。
右侧是托管变量或托管数组，左侧是主机变量、主机数组或数组切片的赋值语句，使用 cudaMemcpy 或类似操作将数据从设备全局内存复制到主机内存。
右侧是托管数组切片，左侧是任何主机或托管变量的赋值语句，使用生成的主机代码复制数据。
右侧是托管变量、托管数组或数组切片，左侧是设备变量或设备数组或数组切片的赋值语句，使用 cudaMemcpy 复制数据并从设备访问数据。

有关托管内存的 Memcpy 和 Memset 行为的更多信息，请参见 CUDA C 编程指南 的统一内存编程部分。

3.4.2. 表达式中的隐式数据传输

一些有限的数据传输可以包含在表达式中。一般来说，经验法则是所有算术或运算都必须在主机上进行，这通常只允许一个设备数组出现在表达式的右侧。会生成临时数组以容纳设备数据的主机副本（如果需要）。例如，如果 a、b 和 c 是符合的主机数组，而 adev、bdev 和 cdev 是符合的设备数组，则以下表达式是合法的

a = adev

adev = a

b = a + adev

c = x * adev + b

以下表达式是不合法的，因为它们要么错误地暗示了实际计算发生的位置，要么以另一种方式编写会更有效，或者两者兼而有之

c = adev + bdev

adev = adev + a

b = sqrt(adev)

语言支持元素传输，但性能较差。也支持数组切片，其性能取决于切片的大小、切片中连续数据的量以及实现方式。

3.4.3. 使用运行时例程进行数据传输

对于熟悉 CUDA C 编程环境的程序员，我们提供了 CUDA 内存管理运行时例程的 Fortran 接口。这些函数可以将数据从主机传输到设备，从设备传输到主机，或者从一个设备数组传输到另一个设备数组。

cudaMemcpy 函数可用于在主机和 GPU 之间复制数据

real, device :: wrk(1024)
real cur(512)
istat = cudaMemcpy(wrk, cur, 512)

对于熟悉 CUDA C 例程的人来说，Memcpy 例程的 kind 参数在 Fortran 中是可选的，因为数组的属性已显式声明。Fortran 运行时例程的参数中表示的计数以数据类型元素而不是字节为单位表示。

有关内存管理运行时例程的完整列表，请参阅内存管理。

3.5. 调用内核子例程

调用内核子例程必须给出调用的执行配置。执行配置给出网格和线程块的大小和形状，这些网格和线程块执行该函数，以及用于假定大小的共享内存数组和关联流的共享内存量。

执行配置在调用语句中的子例程名称之后指定；它具有以下形式

<<< grid, block, bytes, stream >>>

grid 是一个整数，type(dim3) 类型的值，或 *。如果 grid 是一个整数，则将其转换为 dim3(grid,1,1)。如果是 type(dim3)，则乘积 grid%x*grid%y*grid%z 给出要启动的线程块数。此乘积必须小于或等于设备支持的最大块数。启动 grid_global 子例程内核会对块数施加进一步的限制。将网格设置为 * 指示运行时通过调用 cudaOccupancyMaxActiveBlocksPerMultiprocessor() 来计算块数，这会考虑 grid_global（或不考虑）。将单个网格 dim3 x、y 或 z 值设置为 -1 也采用相同的运行时查询路径。
block 是一个整数，或 type(dim3) 类型。如果是 type(dim3)，则每个线程块的线程数为 block%x*block%y*block%z，这必须小于或等于设备支持的最大值。如果 block 是一个整数，则将其转换为 dim3(block,1,1)。
bytes 是可选的；如果存在，则它必须是标量整数，并指定要为每个线程块分配的共享内存字节数，以用于假定大小的共享内存数组。有关更多信息，请参阅共享数据。如果未指定，则使用零值。
stream 是可选的；如果存在，则它必须是一个整数，并且值为零，或调用 cudaStreamCreate 返回的值。请参阅第 41 页的第 4.5 节。它指定此调用排队的流。可以指定流常量值 cudaStreamPerThread。这将为每个 CPU 线程使用唯一的流。

例如，内核子例程

attributes(global) subroutine sub( a )

可以像这样调用

call sub <<< DG, DB, bytes >>> ( A )

如果 grid 或 block 参数大于允许的最大尺寸，或者如果 bytes 大于可用的共享内存，则函数调用将失败。共享内存也可能被内核中的固定大小共享内存声明以及其他专用用途（例如函数参数和执行配置参数）消耗。

3.6. 设备代码

3.6.1. 允许的数据类型

具有 device、constant 或 shared 属性的变量和数组，或在设备子程序中声明的变量和数组，仅限于本节中描述的类型。它们可以具有下表中的任何固有数据类型。

表 3. 设备代码固有数据类型
类型	类型种类
`integer`	1,2,4(默认),8
`logical`	1,2,4(默认),8
`real`	2,4(默认),8
`double precision`	等效于 `real(kind=8)`
`complex`	4(默认),8
`character(len=1)`	1 (默认)

此外，它们可以是派生类型，其中派生类型的成员具有允许的固有数据类型之一，或另一种允许的派生类型。

系统模块 cudafor 包含派生类型 dim3 的定义，定义为

type(dim3)
    integer(kind=4) :: x,y,z
end type

3.6.2. 内置变量

几个 CUDA Fortran 只读预定义变量在设备代码中可用。它们的声明如下

type(dim3) :: threadidx, blockdim, blockidx, griddim
integer(4), parameter :: warpsize = 32

变量 threadidx 包含其线程块内的线程索引；对于一维或二维线程块，threadidx%y 和/或 threadidx%z 组件的值为 1。
变量 blockdim 包含线程块的维度；blockdim 对于同一网格中的所有线程都具有相同的值；对于一维或二维线程块，blockdim%y 和/或 blockdim%z 组件的值为 1。
变量 blockidx 包含网格内的块索引；与 threadidx 一样，对于一维网格，blockidx%y 的值为 1。blockidx%z 的值始终为 1。blockidx 的值对于同一线程块中的所有线程都相同。
变量 griddim 包含网格的维度。griddim 的值对于同一网格中的所有线程都相同；对于一维网格，griddim%y 和 griddim%z 的值为 1。
变量 threadidx、blockdim、blockidx 和 griddim 仅在设备子程序中可用。
常量 warpsize 包含一个 warp 中的线程数。当前定义为 32。

3.6.3. Fortran 内在函数

本节列出了设备子程序中允许的 Fortran 内在函数。

需要使用系统模块 wmma 来调用使用 real(2) 数据类型的数学和一些数值内在函数。有关哪些内在函数仅通过 wmma 模块可用的信息，可以在 WMMA 模块描述部分找到。

表 4. Fortran 数值和逻辑内在函数
名称	参数数据类型	名称	参数数据类型
`abs`	integer, real(2,4,8), complex	`int`	integer, real(2,4,8), complex
`aimag`	complex	`logical`	logical
`aint`	real(4,8)	`max`	integer, real(2,4,8)
`anint`	real(4,8)	`min`	integer, real(2,4,8)
`ceiling`	real(4,8)	`mod`	integer, real(4,8)
`cmplx`	real(2,4,8) or (real,real)	`modulo`	integer, real(4,8)
`conjg`	complex	`nint`	real(4,8)
`dim`	integer, real(4,8)	`real`	integer, real(2,4,8), complex
`floor`	real(4,8)	`sign`	integer, real(4,8)

表 5. Fortran 数学内在函数
名称	参数数据类型	名称	参数数据类型
`acos`	real(2,4,8)	`cosh`	real(2,4,8)
`acosh`	real(4,8)	`erf`	real(4,8)
`asin`	real(2,4,8)	`erfc`	real(4,8)
`asinh`	real(4,8)	`exp`	real(2,4,8), complex
`atan`	real(2,4,8)	`gamma`	real(4,8)
`atanh`	real(4,8)	`hypot`	(real(4,8),real(4,8))
`atan2`	(real,real)	`log`	real(2,4,8), complex
`bessel_j0`	real(4,8)	`log10`	real(2,4,8)
`bessel_j1`	real(4,8)	`log_gamma`	real(4,8)
`bessel_jn`	(int,real(4,8))	`sin`	real(2,4,8), complex
`bessel_y0`	real(4,8)	`sinh`	real(2,4,8)
`bessel_y1`	real(4,8)	`sqrt`	real(2,4,8), complex
`bessel_yn`	(int,real(4,8))	`tan`	real(2,4,8)
`cos`	real(2,4,8), complex	`tanh`	real(2,4,8)

表 6. Fortran 数值查询内在函数
名称	参数数据类型	名称	参数数据类型
`bit_size`	integer	`precision`	real(2,4,8), complex
`digits`	integer, real(2,4,8)	`radix`	integer, real(2,4,8)
`epsilon`	real(2,4,8)	`range`	integer, real(2,4,8), complex
`huge`	integer, real(2,4,8)	`selected_int_kind`	integer
`maxexponent`	real(2,4,8)	`selected_real_kind`	(integer,integer)
`minexponent`	real(2,4,8)	`tiny`	real(2,4,8)

表 7. Fortran 位操作内在函数
名称	参数数据类型	名称	参数数据类型
`btest`	integer	`ishft`	integer
`iand`	integer	`ishftc`	integer
`ibclr`	integer	`leadz`	integer
`ibits`	integer	`mvbits`	integer
`ibset`	integer	`not`	integer
`ieor`	integer	`popcnt`	integer
`ior`	integer	`poppar`	integer

表 8. Fortran 归约和数组内在函数
名称	参数数据类型	名称	参数数据类型
`all`	logical	`maxval`	integer, real(2,4,8)
`any`	logical	`minloc`	integer, real(4,8)
`count`	logical	`minval`	integer, real(2,4,8)
`dot_product`	real(4,8)	`norm2`	real(4,8)
`matmul`	real(4,8), complex	`product`	integer, real(4,8), complex
`maxloc`	integer, real(4,8)	`sum`	integer, real(4,8), complex

3.6.4. 同步函数

本节介绍设备子程序中支持的同步函数和子例程。

同步函数

同步函数控制线程块执行期间各种线程的同步。

syncthreads
syncthreads_count
syncthreads_and
syncthread_or

syncwarp
threadfence
threadfence_block
threadfence_system

有关这些函数的详细信息，请参阅线程管理。

SYNCTHREADS

syncthreads 内在子例程充当单个线程块中所有线程的屏障同步；它没有参数

subroutine syncthreads()

有时，块内的线程访问共享内存或全局内存中的相同地址，从而为某些内存访问创建潜在的先写后读、先读后写或先写后写危害。为了避免这些潜在问题，请使用 syncthreads() 在内核中指定同步点。此内在函数充当屏障，块中的所有线程都必须等待，然后任何线程才能继续进行。块内的线程通过同步其执行来协调内存访问，从而进行协作和共享数据。

线程块中的每个线程在 syncthreads 调用处暂停，直到所有线程都到达该调用。如果线程块中的任何线程发出对 syncthreads 的调用，则所有线程也必须到达并执行相同的调用语句，否则内核将无法正确完成。

SYNCTHREADS_AND

integer syncthreads_and(int_value)

syncthreads_and，像 syncthreads 一样，充当一个屏障，块中的所有线程都必须等待，然后任何线程才能继续进行。此外，syncthreads_and 评估块中所有线程的整数参数 int_value，并且当且仅当 int_value 对所有线程的求值结果为非零时，才返回非零值。

SYNCTHREADS_COUNT

integer syncthreads_count(int_value)

syncthreads_count，像 syncthreads 一样，充当一个屏障，块中的所有线程都必须等待，然后任何线程才能继续进行。此外，syncthreads_count 评估块中所有线程的整数参数 int_value，并返回 int_value 的求值结果为非零的线程数。

SYNCTHREADS_OR

integer syncthreads_or(int_value)

syncthreads_or，像 syncthreads 一样，充当一个屏障，块中的所有线程都必须等待，然后任何线程才能继续进行。此外，syncthreads_or 评估块中所有线程的整数参数 int_value，并且当且仅当 int_value 对任何线程的求值结果为非零时，才返回非零值。

SYNCWARP

subroutine syncwarp(int_mask)

syncwarp 将导致 warp 内的所有执行线程（在掩码参数中指定）到达一个屏障，此时掩码中的所有线程都必须执行 syncwarp，然后任何线程才能继续进行。

内存栅栏

一般来说，当一个线程按特定顺序向内存发出写入序列时，其他线程可能会以不同的顺序看到这些内存写入的效果。您可以使用 threadfence()、threadfence_block() 和 threadfence_system() 创建内存栅栏以强制执行排序。

例如，假设您使用内核在一个调用中计算 N 个数字的数组的总和。每个块首先对数组的子集求和，并将结果存储在全局内存中。当所有块都完成时，最后一个完成的块从全局内存中读取每个部分和，并将它们相加以获得最终结果。为了确定哪个块最后完成，每个块原子地递增一个计数器，以指示它已完成计算并存储其部分和。如果在存储部分和和递增计数器之间没有放置栅栏，则计数器可能会在部分和存储之前递增。

THREADFENCE

subroutine threadfence()

threadfence 充当内存栅栏，创建一个等待。通常，当一个线程按特定顺序向内存发出写入序列时，其他线程可能会以不同的顺序看到这些内存写入的效果。threadfence() 是强制执行特定顺序的一种方法。在 threadfence() 之前，调用线程进行的所有全局和共享内存访问对于以下项都是可见的

线程块中的所有线程（对于共享内存访问）
设备中的所有线程（对于全局内存访问）

THREADFENCE_BLOCK

subroutine threadfence_block()

threadfence_block 充当内存栅栏，创建一个等待，直到在 threadfence_block() 之前，调用线程进行的所有全局和共享内存访问对于线程块中的所有线程对于所有访问都是可见的。

THREADFENCE_SYSTEM

subroutine threadfence_system()

threadfence_system 充当内存栅栏，创建一个等待，直到在 threadfence_system() 之前，调用线程进行的所有全局和共享内存访问对于以下项都是可见的

线程块中的所有线程（对于共享内存访问）
设备中的所有线程（对于全局内存访问）
主机线程（对于页锁定主机内存访问）

threadfence_system() 仅受计算能力为 2.0 或更高的设备支持。

3.6.5. Warp-Vote 操作

新的 warp-vote 和 warp 匹配操作已添加到 NVIDIA CUDA Fortran 中。旧版本出于遗留原因仍然保留；它们将调用更新的功能，并使用掩码指定 warp 中的所有线程。

ALLTHREADS

allthreads 函数是一个 warp-vote 操作，带有一个标量逻辑参数

if( allthreads(a(i)<0.0) ) allneg = .true.

函数 allthreads 为当前 warp 中的所有线程评估其参数。当且仅当参数的值对于 warp 中的所有线程都为 .true. 时，函数的值才为 .true.。

ANYTHREAD

anythread 函数是一个 warp-vote 操作，带有一个标量逻辑参数

if( anythread(a(i)<0.0) ) allneg = .true.

函数 anythread 为当前 warp 中的所有线程评估其参数。当且仅当参数的值对于 warp 中的所有线程都为 .false. 时，函数的值才为 .false.。

BALLOT

ballot 函数是一个 warp-vote 操作，带有一个整数参数

unsigned integer ballot(int_value)

函数 ballot 为 warp 的所有线程评估参数 int_value，并返回一个整数，当且仅当 int_value 对 warp 的第 N 个线程的求值结果为非零时，才设置该整数的第 N 位。

此函数仅受计算能力为 2.0 的设备支持。

示例

if( ballot(int_value) ) allneg = .true.

ACTIVEMASK

unsigned integer activemask()

activemask 函数返回调用 warp 中所有当前活动线程的 32 位整数掩码。当调用 activemask 时，如果 warp 中的第 N 个通道处于活动状态，则设置第 N 位。

ALL_SYNC

integer all_sync(int_mask, int_predicate)

all_sync 函数评估掩码中所有非退出线程的谓词参数，如果谓词对所有线程都为非零，则返回非零值。

ANY_SYNC

integer any_sync(int_mask, int_predicate)

any_sync 函数评估掩码中所有非退出线程的谓词参数，如果谓词对任何线程都为非零，则返回非零值。

BALLOT_SYNC

unsigned integer ballot_sync(int_mask, int_predicate)

ballot_sync 函数为调用 Warp 中由掩码设置的所有非退出线程评估谓词参数。如果第 N 个线程的谓词非零，则在第 N 个通道中设置第 N 位。

MATCH_ALL_SYNC

unsigned integer match_all_sync(int_mask, value, int_predicate)

match_all_sync 函数对掩码参数指定的 Warp 内所有线程的值执行广播和比较。如果所有线程都具有相同的值，则返回 int_mask，否则返回 0。在前一种情况下，int_predicate 设置为 true，在后一种情况下设置为 false。此函数目前接受 value 类型为 integer(4)、integer(8)、real(4) 或 real(8)。

MATCH_ANY_SYNC

unsigned integer match_any_sync(int_mask, value)

match_any_sync 函数对掩码参数指定的 Warp 内所有线程的值执行广播和比较。它返回一个掩码，其中包含与 value 具有相同值的线程。此函数目前接受 value 类型为 integer(4)、integer(8)、real(4) 或 real(8)。

3.6.6. 使用缓存提示的加载和存储函数

这些加载和存储函数可以提供对缓存行为的更精细控制，并充当优化提示。它们不会更改程序的内存一致性行为。这些函数和子例程可以对大多数支持的数据类型进行操作，包括 integer(4)、integer(8)、real(2)、real(4)、real(8)、complex(4) 和 complex(8)。还支持 dimension(4) 的 integer(4) 和 real(4)，以及 dimension(2) 的 integer(8) 和 real(8)，即 128 位加载和存储。

缓存加载函数包括

表 9. 使用缓存提示的加载函数
函数	缓存行为
`value = \__ldca(mem)`	`在所有级别缓存`
`value = \__ldcg(mem)`	`在全局级别缓存`
`value = \__ldcs(mem)`	`缓存流式传输，访问一次`
`value = \__ldlu(mem)`	`最后使用`
`value = \__ldcv(mem)`	`不缓存，视为易失性`

缓存存储子例程包括

表 10. 使用缓存提示的存储子例程
子例程	缓存行为
`call \__stwb(mem, value)`	`缓存写回所有一致性级别`
`call \__stcg(mem, value)`	`在全局级别缓存`
`call \__stcs(mem, value)`	`缓存流式传输，访问一次`
`call \__stwt(mem, value)`	`缓存写通`

3.6.7. 原子函数

原子函数读取和写入其第一个操作数的值，该操作数必须是共享内存（具有 shared 属性）或设备全局内存（具有 device 属性）中的变量或数组元素。只有计算能力为 1.1 及更高版本的设备才支持原子函数。如果第一个参数具有 shared 属性，则需要计算能力 1.2 或更高版本。某些 real(4) 和 real(8) 原子函数可能需要计算能力 2.0 及更高版本。

即使同一或不同线程块中的多个线程尝试在没有任何同步的情况下读取和更新同一位置，原子函数也会返回正确的值。

算术和按位原子函数

这些原子函数读取并返回第一个参数的值。它们还会根据函数将该值与第二个参数的值组合，并将组合后的值存储回第一个参数的位置。对于 atomicadd、atomicsub、atomicmax、atomicmin 和 atomicexch，数据类型可以是 integer(4)、integer(8)、real(4) 或 real(8)。对于 atomicand、atomicor 和 atomicxor，仅支持 integer(4) 参数。

注意

这些函数中的每一个的返回值都是第一个参数 mem。

这些函数包括

表 11. 算术和按位原子函数
函数	附加原子更新
`atomicadd( mem, value )`	`mem = mem + value`
`atomicsub( mem, value )`	`mem = mem − value`
`atomicmax( mem, value )`	`mem = max(mem,value)`
`atomicmin( mem, value )`	`mem = min(mem,value)`
`atomicand( mem, value )`	`mem = iand(mem,value)`
`atomicor( mem, value )`	`mem = ior(mem,value)`
`atomicxor( mem, value )`	`mem = ieor(mem,value)`
`atomicexch( mem, value )`	`mem = value`

计数原子函数

这些原子函数读取并返回第一个参数的值。它们还会将第一个参数与第二个参数进行比较，并根据比较结果将新值存储回第一个参数的位置。这些函数旨在实现循环计数器，向上或向下计数到第二个参数中指定的最大值。两个参数的类型都必须为 integer(kind=4)。

注意

这些函数中的每一个的返回值都是第一个参数 mem。

这些函数包括

表 12. 计数原子函数
函数	附加原子更新
`atomicinc( mem, imax )`	if (mem<imax) then mem = mem+1 else mem = 0 endif
`atomicdec( mem, imax )`	if (mem<imax .and. mem>0) then mem = mem-1 else mem = imax endif

比较和交换原子函数

此原子函数读取并返回第一个参数的值。它还会将第一个参数与第二个参数进行比较，如果第一个和第二个参数相等，则原子地将新值存储回第一个参数的位置。所有三个参数的类型必须相同，可以是 integer(kind=4)、integer(kind=8)、real(kind=4) 或 real(kind=8)。

注意

此函数的返回值是第一个参数 mem。

函数为

表 13. 比较和交换原子函数
函数	附加原子更新
`atomiccas(mem,comp,val)`	if (mem == comp) then mem = val endif

3.6.8. Fortran I/O

NVIDIA Fortran 编译器包括对 GPU 设备代码中 PRINT 语句的有限支持。Fortran GPU 运行时库在 CUDA Fortran 和 OpenACC 之间共享，用于 NVIDIA GPU 目标，它会缓冲输出并在一个操作中打印整行。支持 Integer、character、logical、real 和 complex 数据类型。

底层的 CUDA printf 实现将内核启动中的 print 语句数量限制为 4096。用户在使用此功能时应考虑此限制。

3.6.9. PRINT 示例

通过添加编译器选项 -cuda=charstring，CUDA Fortran 设备代码中现在也提供对字符字符串、字符子字符串、字符变量和字符串赋值的一些有限支持。这是一个简短的示例

attributes(global) subroutine printtest()
character*12 c
i = threadIdx%x
if (i/2*2.eq.i) then
  c = "Even Thread:"
else
  c = " Odd Thread:"
endif
print *,c,c(6:11),i
end subroutine

3.6.10. Shuffle 函数

CUDA Fortran 设备代码可以访问计算能力 3.x 的 shuffle 函数。这些函数允许访问 Warp（称为通道）内线程之间的变量。在 CUDA Fortran 中，通道使用 Fortran 的从 1 开始的编号方案。

__shfl()

__shfl() 返回由 ID 为 srcLane 的线程持有的 var 的值。如果 srcLane 超出 1:width 的范围，则返回线程自身的 var 值。width 参数在所有 shuffle 函数中都是可选的，默认值为 32，即当前的 Warp 大小。


integer(4) function __shfl(var, srcLane, width)
   integer(4) var, srcLane
   integer(4), optional :: width


integer(8) function __shfl(var, srcLane, width)
   integer(8) :: var
   integer(4) :: srcLane
   integer(4), optional :: width


real(4) function __shfl(var, srcLane, width)
   real(4) :: var
   integer(4) :: srcLane
   integer(4), optional :: width


real(8) function __shfl(var, srcLane, width)
   real(8) :: var
   integer(4) :: srcLane
   integer(4), optional :: width


__shfl_up()
__shfl_up() 通过从调用线程的线程 ID 中减去 delta 来计算源通道 ID。返回结果线程 ID 持有的 var 值；实际上，var 在 Warp 中向上移动 delta 个通道。
源通道索引不会环绕 width 的值，因此较低的 delta 通道保持不变。
integer(4) function __shfl_up(var, delta, width)
   integer(4) var, delta
   integer(4), optional :: width


integer(8) function __shfl_up(var, delta, width)
   integer(8) :: var
   integer(4) :: delta
   integer(4), optional :: width


real(4) function __shfl_up(var, delta, width)
   real(4) :: var
   integer(4) :: delta
   integer(4), optional :: width


real(8) function __shfl_up(var, delta, width)
   real(8) :: var
   integer(4) :: delta
   integer(4), optional :: width


__shfl_down()
__shfl_down() 通过将 delta 添加到调用线程的线程 ID 来计算源通道 ID。返回结果线程 ID 持有的 var 值：这具有将 var 在 Warp 中向下移动 delta 个通道的效果。源通道的 ID 号不会环绕 width 的值，因此较高的 delta 通道保持不变。
integer(4) function __shfl_down(var, delta, width)
   integer(4) var, delta
   integer(4), optional :: width


integer(8) function __shfl_down(var, delta, width)
   integer(8) :: var
   integer(4) :: delta
   integer(4), optional :: width


real(4) function __shfl_down(var, delta, width)
   real(4) :: var
   integer(4) :: delta
   integer(4), optional :: width


real(8) function __shfl_down(var, delta, width)
   real(8) :: var
   integer(4) :: delta
   integer(4), optional :: width


__shfl_xor()
__shfl_xor() 使用 ID-1 通过对调用线程的通道 ID 与 laneMask 执行按位异或运算来计算源通道 ID。返回结果通道 ID 持有的 var 值。如果结果通道 ID 超出 width 允许的范围，则返回线程自身的 var 值。此模式实现了一种蝶形寻址模式，例如树归约和广播中使用的模式。
integer(4) function __shfl_xor(var, laneMask, width)
   integer(4) var, laneMask
   integer(4), optional :: width


integer(8) function __shfl_xor(var, laneMask, width)
   integer(8) :: var
   integer(4) :: laneMask
   integer(4), optional :: width


real(4) function __shfl_xor(var, laneMask, width)
   real(4) :: var
   integer(4) :: laneMask
   integer(4), optional :: width


real(8) function __shfl_xor(var, laneMask, width)
   real(8) :: var
   integer(4) :: laneMask
   integer(4), optional :: width


以下是使用 __shfl_xor() 计算 Warp 内每个线程的变量贡献总和的示例
j = . . .
k = __shfl_xor(j,1);  j = j + k
k = __shfl_xor(j,2);  j = j + k
k = __shfl_xor(j,4);  j = j + k
k = __shfl_xor(j,8);  j = j + k
k = __shfl_xor(j,16); j = j + k



3.6.11. 限制
本节列出了设备子程序中可能出现的语句和功能的限制。

不允许使用递归子例程和函数。
不允许使用 PAUSE 语句。
大多数输入/输出语句完全不允许使用：READ、FORMAT、NAMELIST、OPEN、CLOSE、BACKSPACE、REWIND、ENDFILE、INQUIRE。
为计算能力 2.0 及更高版本编译时，可以使用指向默认单元的列表定向 PRINT 和 WRITE 语句；禁止所有其他 PRINT 和 WRITE 用法。
不允许使用备用返回规范。
不允许使用 ENTRY 语句。
不支持浮点异常处理。
不支持第 3.6.3 节中未列出的 Fortran 内部函数。
不支持 Cray 指针。



3.7. 主机代码
主机子程序可以使用内部函数，例如 sizeof 内部函数，来查找 Fortran 数据结构的大小（以字节为单位）。

3.7.1. SIZEOF 内部函数
调用 sizeof(A)，其中 A 是变量或表达式，将返回保存 A 值所需的字节数。
integer(kind=4) :: i, j
j = sizeof(i)    ! this assigns the value 4 to j





3.8. Fortran 设备模块
NVIDIA 默认提供一个设备模块，该模块允许访问和接口到许多 CUDA 设备内置例程。
要显式访问此模块，请执行以下操作之一

将此行添加到您的 Fortran 程序中
use cudadevice



将此行添加到您的 C 程序中
#include <cudadevice.h>




您可以在 CUDA Fortran 全局和设备子程序、CUF 内核以及 Fortran 和 C 中的 NVIDIA Accelerator 计算区域中使用这些例程。此外，NVIDIA HPC 编译器附带这些例程的主机代码实现，尽管这些实现并非专门针对主机进行优化。在 CUDA Fortran 全局和设备子程序以外的用途中，您必须在主机子程序单元中显式使用该模块。
CUDA 内置例程 列出了可用的 CUDA 内置例程

表 14. CUDA 内置例程







__brev
__brevll
clock
clock64

__clz
__clzll
__cosf
cospi

cospif
__dadd_rd
__dadd_rn
__dadd_ru

__dadd_rz
__ddiv_rd
__ddiv_rn
__ddiv_ru

__ddiv_rz
__dmul_rd
__dmul_rn
__dmul_ru

__dmul_rz
__double2float_rd
__double2float_rn
__double2float_ru

__double2float_rz
__double2hiint
__double2int_rd
__double2int_rn

__double2int_ru
__double2int_rz
__double2loint
__double2ll_rd

__double2ll_rn
__double2ll_ru
__double2ll_rz
__double2uint_rd

__double2uint_rn
__double2uint_ru
__double2uint_rz
__double2ull_rd

__double2ull_rn
__double2ull_ru
__double2ull_rz
__double_as_longlong

__drcp_rd
__drcp_rn
__drcp_ru
__drcp_rz

__dsqrt_rd
__dsqrt_rn
__dsqrt_ru
__dsqrt_rz

__exp10f
__expf
__fadd_rd
__fadd_rn

__fadd_ru
__fadd_rz
__fdiv_rd
__fdiv_rn

__fdiv_ru
__fdiv_rz
fdivide
fdividef

__fdividef
__ffs
__ffsll
__float2half_rn

__float2int_rd
__float2int_rn
__float2int_ru
__float2int_rz

__float2ll_rd
__float2ll_rn
__float2ll_ru
__float2ll_rz

__float_as_int
__fma_rd
__fma_rn
__fma_ru

__fma_rz
__fmaf_rd
__fmaf_rn
__fmaf_ru

__fmaf_rz
__fmul_rd
__fmul_rn
__fmul_ru

__fmul_rz
__frcp_rd
__frcp_rn
__frcp_ru

__frcp_rz
__fsqrt_rd
__fsqrt_rn
__fsqrt_ru

__fsqrt_rz
__half2float
__hiloint2double
__int2double_rn

__int2float_rd
__int2float_rn
__int2float_ru
__int2float_rz

__int_as_float
__ll2double_rd
__ll2double_rn
__ll2double_ru

__ll2double_rz
__ll2float_rd
__ll2float_rn
__ll2float_ru

__ll2float_rz
__log10f
__log2f
__logf

__longlong_as_double
__mul24
__mulhi
__popc

__popcll
__powf
__sad
__saturatef

sincos
sincosf
sincospi
sincospif

__sinf
sinpi
sinpif
__tanf

__uint2double_rn
__uint2float_rd
__uint2float_rn
__uint2float_ru

__uint2float_rz
__ull2double_rd
__ull2double_rn
__ull2double_ru

__ull2double_rz
__ull2float_rd
__ull2float_rn
__ull2float_ru

__ull2float_rz
__umul24
__umulhi
__usad




3.8.1. LIBM 设备模块
NVIDIA 还提供了一个设备模块，该模块提供到标准 libm 函数的接口，这些函数不在 Fortran 内部函数库中。
要访问此模块，请将此行添加到您的 Fortran 子程序中
use libm


这些接口在 libm 设备模块中定义

表 15. CUDA 设备 libm 例程








名称
参数数据类型

名称
参数数据类型



cbrt,cbrtf
real(8),real(4) 返回 real

llround,llroundf
real(8),real(4) 返回 integer

ceil,ceilf
real(8),real(4) 返回 real

lrint,lrintf
real(8),real(4) 返回 integer

copysign,copysignf
2*real(8),real(4) 返回 real

lround,lroundf
real(8),real(4) 返回 integer

expm1,expm1f
real(8),real(4) 返回 real

logb,logbf
real(8),real(4) 返回 real

exp10,exp10f
real(8),real(4) 返回 real

log1p,log1pf
real(8),real(4) 返回 real

exp2,exp2f
real(8),real(4) 返回 real

log2,log2f
real(8),real(4) 返回 real

fabs,fabsf
real(8),real(4) 返回 real

modf,modff
2*real(8),real(4) 返回 real

floor,floorf
real(8),real(4) 返回 real

nearbyint,nearbyintf
real(8),real(4) 返回 real

fma,fmaf
3*real(8),real(4) 返回 real

nextafter,nextafterf
2*real(8),real(4) 返回 real

fmax,fmaxf
2*real(8),real(4) 返回 real

remainder,remainderf
2*real(8),real(4) 返回 real

fmin,fminf
2*real(8),real(4) 返回 real

remquo,remquof
2*real(8),real(4) integer 返回 real

frexp,frexpf
real(8),real(4) integer 返回 real

rint,rintf
real(8),real(4) 返回 real

ilogb,ilogbf
real(8),real(4) 返回 real

scalbn,scalbnf
real(8),real(4) integer 返回 real

ldexp,ldexpf
real(8),real(4) integer 返回 real

scalbln,scalblnf
real(8),real(4) integer 返回 real

llrint,llrintf
real(8),real(4) 返回 integer

trunc,truncf
real(8),real(4) 返回 real



这是一个使用 LIBM 设备模块的简单示例
attributes(global) subroutine testlibm( a, b )
  use libm
  real, device :: a(*), b(*)
  i = threadIdx%x
  b(i) = cbrt(a(i))
  end subroutine




3.8.2. 协作组设备模块
在支持 CUDA 计算能力 7.0 及更高版本的 NVIDIA GPU 上，NVIDIA 提供了一个设备模块，该模块提供到协作组功能的接口，NVIDIA 在 CUDA 9.0 中开始提供此功能。在我们的 23.3 版本中，协作组模块还支持 Hopper (cc90) 和更新架构的线程块集群编程。
要访问此模块，请将此行添加到您的 Fortran 子程序中
use cooperative_groups


这是一个使用 cooperative_groups 设备模块的简单示例，该模块启用了协作网格内核
attributes(grid_global) subroutine g1(a,b,n,some_offset)
  use cooperative_groups
  real, device :: a(n), b(n)
  integer, value :: n, some_offset
  type(grid_group) :: gg
  gg = this_grid()
  do i = gg%rank, n, gg%size
    a(i) = min(max(a(i),0.0),100.0) + 0.5
  end do
  call syncthreads(gg)
  do i = gg%rank, n, gg%size
    j = i + some_offset
    if (j.gt.n) j = j - n
    b(i) = a(i) + a(j)
  end do
return
end subroutine


目前，小于或等于线程块大小的协作组的功能有限。在即将发布的版本中将添加更多功能。目前，模块中定义了以下类型：grid_group、thread_group、coalesced_group 和 cluster_group。每种类型都有两个公共成员：size 和 rank。syncthreads 子例程在 cooperative_groups 模块中被重载，以将类型作为参数，从而适当地同步该组中的线程。支持的最小代码序列是
协作组等于线程块
. . .
 use cooperative_groups
 type(thread_group) :: tg
 tg = this_thread_block()
 call syncthreads(tg)


协作组等于 Warp
. . .
 use cooperative_groups
 type(coalesced_group) :: wg
 wg = this_warp()
 call syncthreads(wg)


协作组等于线程块集群
. . .
 use cooperative_groups
 type(cluster_group) :: clg
 clg = this_cluster()
 call syncthreads(clg)


线程块集群的主要优势是利用分布式共享内存，这使得可以将更大部分的数据保存在靠近处理元件的位置。我们建议在访问相邻共享内存时使用 Cray 指针语法，以尽可能降低寄存器压力。这是一个简短的示例
attributes(global) cluster_dims(2,1,1) subroutine t1(rnks)
  use cooperative_groups
  integer, device :: rnks(32,*)
  type(cluster_group) :: clg  ! Defined in cooperative_groups
  integer, shared :: smem(*)
  integer, shared :: dmem(*); pointer(pmem,dmem)
  i = threadIdx%x; j = blockIdx%x
  clg = this_cluster()        ! Defined in cooperative_groups
  nrank = clg%rank
  rnks(i,j) = clg%rank        ! Initialize rnks to 1 or 2
  call syncthreads(clg)       ! Sync both blocks
  if (nrank.eq.1) then ! Get a pointer to the other
    pmem = cluster_map_shared_rank(smem, 2)
  else
    pmem = cluster_map_shared_rank(smem, 1)
  end if
  dmem(i) = 100*nrank + i ! Write to the other blocks shared memory
  call syncthreads(clg)   ! Sync both blocks
  rnks(i,j) = rnks(i,j) + smem(i) ! Read what the other block wrote
end subroutine


协作组模块还定义了新的 shfl_sync() 函数。这些函数与本文档前面部分讨论的 shfl() 函数类似，但采用额外的 mask 作为第一个参数。32 位掩码参数指定 Warp 中哪些线程参与 shuffle 操作，并且对于大多数用例，可以作为值为 z'ffffffff' 的 integer(4) 传递。请注意，如果您将旧版 shfl() 函数与 CUDA 9.0 或更高版本一起使用，我们将隐式使用掩码为 z'ffffffff' 的 shlf_sync()。如果您的 Warp 内存在线程分歧，则这可能不正确。在这种情况下，请使用新的 shfl_sync() 函数并提供正确的掩码，该掩码可以使用 ballot() 设备函数生成。


3.8.3. WMMA（Warp 矩阵乘法累加）模块
在支持 CUDA 计算能力 7.0 及更高版本的 NVIDIA GPU 上，NVIDIA 包含一个设备模块，该模块提供到矩阵运算的接口，这些矩阵运算利用 Tensor Core 来加速矩阵问题。这使使用 Fortran 的科学程序员能够利用 real(2) 矩阵运算。
要访问该模块，请将此行添加到您的 Fortran 子程序中
use wmma


wmma 模块中提供的 API 例程包括 C = Matmul(A, B) 形式的矩阵乘法运算，其中

A 是维度为 A(m,k) 的二维 real(2) 数组
B 是维度为 B(k,n) 的二维 real(2) 数组
C 是维度为 C(m,n) 的二维 real(2) 或 real(4) 数组

使用 Fortran kind 属性，可以声明和使用半精度格式的数据。有关表示形式和使用要求的详细信息，请参见 半精度浮点数 部分。
这是一个使用 wmma 设备模块使用单个 Warp 线程执行矩阵乘法的简单示例。有两个 16×16 real(2) 矩阵相乘并累加到 16×16 real(4) 矩阵中
#include "cuf_macros.CUF"
    module m
    integer, parameter :: wmma_m = 16
    integer, parameter :: wmma_n = 16
    integer, parameter :: wmma_k = 16

    contains
    ! kernel for 16 x16 matrices (a, b, and c) using wmma
    ! Should be launched with one block of 32 threads
    attributes(global) subroutine wmma_single(a, b, c)
    use wmma
    implicit none
    real(2), intent(in) :: a(wmma_m,*) , b(wmma_k,*)
    real(4) :: c(wmma_m,*)
    WMMASubMatrix(WMMAMatrixA, 16, 16, 16, Real, WMMAColMajor) :: sa
    WMMASubMatrix(WMMAMatrixB, 16, 16, 16, Real, WMMAColMajor) :: sb
    WMMASubMatrix(WMMAMatrixC, 16, 16, 16, Real, WMMAKind4) :: sc
    integer :: lda, ldb, ldc

    lda = wmma_m
    ldb = wmma_k
    ldc = wmma_m

    sc = 0.0_4
    call wmmaLoadMatrix(sa, a(1, 1), lda)
    call wmmaLoadMatrix(sb, b(1, 1), ldb)
    call wmmaMatMul(sc, sa, sb, sc)
    call wmmaStoreMatrix(c(1, 1), sc, ldc)

    end subroutine wmma_single
    end module m


调用站点如下所示，以使用单个 Warp 线程调用
call wmma_single<<<1,32>>>(ah_d, bh_d, c_d)


对于这个简单的示例，作为参数传递到内核的矩阵与 WMMA 子矩阵的大小相同。因此，为了执行矩阵乘法，我们只需将 C WMMA 子矩阵初始化为 0.0，将 A 和 B 矩阵从全局内存加载到 WMMA 子矩阵，对子矩阵执行矩阵乘法，并将结果从 WMMA 子矩阵存储到全局内存。
您可能已经注意到线程索引 threadIdx 在此代码中根本没有出现。这突出了从此示例中得出的重要概念：Warp 中的线程协同工作以完成这些任务。因此，在处理 WMMA 子矩阵时，我们执行的是 Warp 级编程，而不是线程级编程。此内核使用单个包含 32 个线程的 Warp 启动，但我们的每个 WMMA 子矩阵都有 16×16 或 256 个元素。当初始化语句
sc = 0.0_4


执行时，每个线程将 16×16 子矩阵中的 8 个元素设置为零。线程到子矩阵元素的映射对于此操作和涉及 WMMA 子矩阵的其他操作是不透明的 - 从编程的角度来看，我们只关注 Warp 线程组对 WMMA 子矩阵执行的操作。
将 A 和 B 从全局内存加载到 WMMA 子矩阵的语句
call wmmaLoadMatrix(sa, a(1, 1), lda)
    call wmmaLoadMatrix(sb, b(1, 1), ldb)


也协同工作。在这些调用中，WMMA 子矩阵被指定为第一个参数，第二个参数包含要加载到 WMMA 子矩阵的全局（或共享）内存中图块的左上角元素的地址。全局（或共享）内存中矩阵的前导维度是第三个参数。请注意，传递给 wmmaLoadMatrix() 的参数对于 Warp 中的所有线程都是相同的。由于 Warp 中元素到线程的映射是不透明的，因此每个线程只需传递 16×16 矩阵中第一个元素的地址以及前导维度作为第三个参数，并且加载操作在 Warp 中的线程之间分配。
WMMA 子矩阵上的矩阵乘法由语句
call wmmaMatMul(sc, sa, sb, sc)


执行，该语句再次由 Warp 线程组协同执行。这里在 wmmaMatMul() 调用中对第一个和最后一个参数使用了相同的累加器子矩阵，这就是为什么需要将其初始化为零的原因。
wmmaStoreMatrix() 调用
call wmmaStoreMatrix(c(1, 1), sc, ldc)


类似于之前的 wmmaLoadMatrix 调用，但此处第一个参数是全局（或共享）内存中图块的左上角元素的地址，第二个参数是要存储其值的 WMMA 子矩阵。当使用累加器 (WMMAMatrixC) 参数调用 wmmaLoadMatrix() 和 wmmaStoreMatrix() 时，有一个可选的第四个参数，用于指定存储顺序。在 CUDA Fortran 中，默认值为 WMMAColMajor 或列优先存储顺序。
关于 wmmaLoadMatrix() 和 wmmaStoreMatrix() 例程的参数的最后一点说明。要求矩阵的前导维度（由这些例程的第三个参数指定）必须是 16 字节的倍数（例如，8 个 real(2) 字或 4 个 real(4) 字）。
有关数据声明和 wmma 操作的更多详细信息，请访问 使用 CUDA Fortran 进行 Tensor Core 编程 https://devblogs.nvidia.com/tensor-core-programming-using-cuda-fortran/
wmma 模块还提供对以下半精度数学内部函数的访问，并且需要 use wmma 才能访问它们：abs、sin、cos、tan、acos、asin、atan、atan2、sinh、cosh、tanh、log、log10、exp 和 sqrt。预计在未来的版本中，这些内部函数将无需提及 use wmma 即可使用。



3.9. Fortran 主机模块
NVIDIA 为 CUDA Fortran 提供的主要 Fortran 模块名为 cudafor。此模块包含对下一章中列出的 CUDA 运行时 API 的所有受支持接口。此外，它还包含对下面章节中描述的一些 Fortran 数组内部函数的接口。
还有许多其他 Fortran 模块接口到 CUDA 库。这些库在 NVIDIA Fortran CUDA 接口 文档中进行了详尽的描述。这些库包括用于计算的库，如 CUBLAS、 CUFFT 和 CUSPARSE，用于通信的库，如 NCCL、 NVSHMEM，以及用于分析的库，如 NVTX。
我们将在本章中描述的另一个主机模块是 CUTENSOR。它已在名为 cutensorex 的模块中进行了扩展，并包含许多 Fortran 数组内部函数的重载接口，其中一些接口调用 NVIDIA CUTENSOR 库，而另一些接口则不调用，但它们使用相同的延迟评估技术。这些实现对设备（或托管）数据进行操作，并从主机调用。

3.9.1. CUDAFOR 中重载的 Fortran 归约内部函数
当使用 cudafor 模块时，SUM、MAXVAL、MINVAL、MAXLOC 和 MINLOC Fortran 内部函数被重载以接受来自主机代码的设备或托管数组。如果使用掩码可选参数，则掩码参数必须是设备逻辑数组，或者包含托管操作数和常量的表达式，即掩码必须可在主机上计算，但可在设备上读取。与标准 Fortran 中一样，掩码形状和大小（如果存在）必须与数据数组一致。
这是一个完整的示例，它在 GPU 上执行 sum 和 maxval 归约
program multidimred
use cudafor
real(8), managed :: a(5,5,5,5,5)
real(8), managed :: b(5,5,5,5)
real(8) :: c
call random_number(a)
do idim = 1, 5
  b = sum(a, dim=idim)
  c = max(maxval(b), c)
end do
print *,"Max along any dimension",c
end program


数组切片也受支持。虽然在 GPU 上运行效率可能较低，但功能仍然非常强大，并且对调试很有用。
real(4), managed :: a(n,m)
reslt(ix) = sum(a(2:n-1,:))
reslt(ix) = sum(a(:,3:m-2))
reslt(ix) = sum(a(n2:n,m2:m))
reslt(ix) = sum(a(1:n3,1:m3))
reslt(ix) = sum(a(n2:n3,m2:m3))


默认情况下，设备上支持的内部归约操作将在设备上对（足够大的）托管数组执行。在某些情况下，人们可能希望在主机上对托管数据执行归约操作。这可以使用“use”语句的重命名功能来完成，例如
program reductionRename
use cudafor, gpusum => sum
implicit none
integer, managed :: m(3000)
integer :: istat
m = 1
istat = cudaDeviceSynchronize()
write(*,*) sum(m)    ! executes on host
write(*,*) gpusum(m) ! executes on device
end program


从 NVHPC 23.1 版本开始，所有五个函数，SUM、MAXVAL、MINVAL、MAXLOC 和 MINLOC 现在都可以接受可选的流参数。如果通过调用 cudaforSetDefaultStream 设置了唯一的每线程默认流，则归约操作将拾取该流并在该流上运行。鉴于新的、更简单的功能，从 23.1 版本开始，也已删除对 cudaforReductionSetStream() 和 cudaforReductionGetStream() 的支持。例如
integer(kind=cuda_stream_kind) :: istrm
        x = sum(a, stream=istrm)


现在是在特定流上运行求和归约的最简单方法。
以下章节将更详细地描述每个函数，包括当前的支持和限制。

3.9.1.1. Fortran SUM 内部函数
SUM 的重载接口位于 cudafor 模块中。它可以返回标量（最常见的情况），或者如果使用了可选的 dim 参数，则返回数组。支持 real(4)、real(8)、integer(4) 和 integer(8) 数据类型。复数类型可能会在未来的版本中添加。输入数组可以在一到七维之间。两种形式是
function sum ( array, mask, stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    logical(4), device, optional, intent(in) :: mask(...)
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    type, intent(out) :: res ! same type as array


function sum ( array, dim, mask, stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    integer(4), intent(in) :: dim
    logical(4), device, optional, intent(in) :: mask(...)
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    type, allocatable, managed, intent(out) :: res(...) ! same type as array
                                               ! rank is one less than array




3.9.1.2. Fortran MAXVAL 内部函数
MAXVAL 的重载接口（返回数组中元素的最大值）位于 cudafor 模块中。它可以返回标量（最常见的情况），或者如果使用了可选的 dim 参数，则返回数组。支持 real(4)、real(8)、integer(4) 和 integer(8) 数据类型。输入数组可以在一到七维之间。两种形式是
function maxval ( array, mask, stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    logical(4), device, optional, intent(in) :: mask(...)
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    type, intent(out) :: res ! same type as array


function maxval ( array, dim, mask, stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    integer(4), intent(in) :: dim
    logical(4), device, optional, intent(in) :: mask(...)
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    type, allocatable, managed, intent(out) :: res(...) ! same type as array
                                               ! rank is one less than array




3.9.1.3. Fortran MINVAL 内部函数
MINVAL 的重载接口（返回数组中元素的最小值）位于 cudafor 模块中。它可以返回标量（最常见的情况），或者如果使用了可选的 dim 参数，则返回数组。支持 real(4)、real(8)、integer(4) 和 integer(8) 数据类型。输入数组可以在一到七维之间。两种形式是
function minval ( array, mask, stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    logical(4), device, optional, intent(in) :: mask(...)
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    type, intent(out) :: res ! same type as array


function minval ( array, dim, mask, stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    integer(4), intent(in) :: dim
    logical(4), device, optional, intent(in) :: mask(...)
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    type, allocatable, managed, intent(out) :: res(...) ! same type as array
                                               ! rank is one less than array




3.9.1.4. Fortran MAXLOC 内部函数
MAXLOC 的重载接口（返回一个索引数组，从 1 开始，标识数组中首先出现的最大值元素）位于 cudafor 模块中。函数结果的大小等于输入数组的秩，并且是一个整数主机数组。支持 real(4)、real(8)、integer(4) 和 integer(8) 数据类型。输入数组可以在一到七维之间。dim 参数仅支持一维数组，在这种情况下，结果是一个标量，而不是大小为 1 的数组。还有可选的 kind、back 和 stream 参数，前两个是标准 Fortran，后者是 CUDA Fortran 扩展。
function maxloc ( array, mask, kind, back stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    logical(4), device, optional, intent(in) :: mask(...)
    integer, optional, intent(in) :: kind
    logical, optional, intent(in) :: back
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    integer, intent(out) :: res(*) ! Size of res is equal to rank of array




3.9.1.5. Fortran MINLOC 内部函数
MINLOC 的重载接口（返回一个索引数组，从 1 开始，标识数组中首先出现的最小值元素）位于 cudafor 模块中。函数结果的大小等于输入数组的秩，并且是一个整数主机数组。支持 real(4)、real(8)、integer(4) 和 integer(8) 数据类型。输入数组可以在一到七维之间。dim 参数仅支持一维数组，在这种情况下，结果是一个标量，而不是大小为 1 的数组。还有可选的 kind、back 和 stream 参数，前两个是标准 Fortran，后者是 CUDA Fortran 扩展。
function minloc ( array, mask, kind, back stream ) result(res)
    type, device :: array(...)  ! type is real or integer, kind = 4 or 8
    logical(4), device, optional, intent(in) :: mask(...)
    integer, optional, intent(in) :: kind
    logical, optional, intent(in) :: back
    integer(kind=cuda_stream_kind), optional, intent(in) :: stream
    integer, intent(out) :: res(*) ! Size of res is equal to rank of array





3.9.2. Fortran 排序子例程模块
通常，为了获得最佳性能，我们建议使用 CUDA Thrust、nvcc 编译器生成排序例程，并从 Fortran 调用这些函数。从 23.5 版本开始，我们还在 CUDA Fortran 库中包含基本的排序子例程，这些子例程易于使用，并且可能提供“足够好”的性能。
可以通过在代码中添加 use sort 来访问库函数的接口，重载的排序子例程名为 fsort()。该库为 integer(4)、integer(8)、real(4) 和 real(8) 数组提供了基数排序实现。子例程可以接受主机、托管或设备数组。子例程还可以接受索引数组，以返回排序置换。下面列出了其他可选参数。
这是一个在 GPU 上对实数数组进行排序的简单示例
program sortit
use sort
real(4), managed :: a(1000)
call random_number(a)
call fsort(a, 1000)
print *,all(a(1:999) .le. a(2:1000))
end program


主机和设备功能分为四种类型的调用，每种调用的参数如下
! Host arrays, no indices
subroutine fsort(array, n, stream)
type(kind) :: array(*)  ! Type is integer or real, kind is 4 or 8
integer(kind) :: n      ! kind is 4 or 8
integer(kind=cuda_stream_kind), optional :: stream
end subroutine


! Host arrays, with indices
subroutine fsort(array, indices, n, init_index, stream)
type(kind) :: array(*)  ! Type is integer or real, kind is 4 or 8
integer(4) :: indices(*)
integer(kind) :: n      ! kind is 4 or 8
logical(4), optional :: init_index  ! Flag to initialize the indices to 1..n
integer(kind=cuda_stream_kind), optional :: stream
end subroutine


! Managed or device arrays, no indices
subroutine fsort(array, n, workspace, worksize, stream)
type(kind), device :: array(*)  ! Type is integer or real, kind is 4 or 8
integer(kind) :: n      ! kind is 4 or 8
type(kind), device, optional :: workspace(*)  ! Same type as array
integer(8), optional :: worksize              ! Size of workspace in elements
integer(kind=cuda_stream_kind), optional :: stream
end subroutine


! Managed or device arrays, with indices
subroutine fsort(array, indices, n, init_index, workspace, worksize, stream)
type(kind), device :: array(*)  ! Type is integer or real, kind is 4 or 8
integer(4), device :: indices(*)
integer(kind) :: n      ! kind is 4 or 8
logical(4), optional :: init_index  ! Flag to initialize the indices to 1..n
type(kind), device, optional :: workspace(*)  ! Same type as array
integer(8), optional :: worksize              ! Size of workspace in elements
integer(kind=cuda_stream_kind), optional :: stream
end subroutine


在没有提供 workspace 参数的情况下，子例程将使用 cudaMalloc() 或 cudaMallocAsync() 分配临时工作空间，具体取决于 CUDA 版本支持以及是否指定了流。对于采用 worksize 参数的子例程，避免临时分配所需的工作空间量大致为 N 个元素，kind 等于 4 或 8，再加上最多 2 MBytes 以上的空间。例如，对大小为 1000 万的 integer(4) 数组进行排序将使用大约 42 MBytes 或 1050 万个元素的工作空间。


3.9.3. CUTENSOREX 中重载的 Fortran 归约内部函数
当从主机代码使用 cutensorex 模块时，ALL、ANY 和 COUNT Fortran 内部函数被重载以接受设备或托管数组。由于这三个函数仅对掩码进行操作，因此选择了一种不同的策略，使这些函数更灵活，并识别和高效地评估常用的掩码表达式。
使用与 cutensorex 中用于 matmul()、spread()、transpose() 和 reshape() 的相同的延迟求值和赋值技术，从 23.1 版本开始，我们现在支持更多的 F90 数组内部操作。
这三个函数不调用 cuTensor 库，而是在先前为这些包装器开发的软件基础设施之上构建和扩展。
首先，以下是为延迟求值识别的掩码表达式
对于 A、B、x、dx、alpha、beta
A is a device array of real(4), real(8), integer(4), or integer(8)
B is a device array with the same type as A.
A and B are 1-3 dimensional, (conforming arrays)
x is a scalar with the same type as A
dx is a device scalar with the same type as A
alpha and beta are host scalars with the same type as A


在一个内核启动中，我们支持以下掩码表达式
A .relop. B
A .relop. x
A .relop. dx
abs(A) .relop. B
abs(A) .relop. x
abs(A) .relop. dx
(A +/- B) .relop. x
(A +/- B) .relop. dx
abs(A +/- B) .relop. x
abs(A +/- B) .relop. dx
(alpha*A + beta*B) .relop. x
(alpha*A + beta*B) .relop. dx
abs(alpha*A + beta*B) .relop. x
abs(alpha*A + beta*B) .relop. dx


对于 relop in EQ、NE、LE、LT、GE、GT
为了方便起见，一个例外是，如果操作是“A .relop. x”，如果 A 的 kind=8，则 x 的 kind 可以为 4
在大多数情况下，B 数组也可以是 spread() 或 transpose() 内部函数的结果，以使 B 符合 A 的形状。目前不支持掩码操作数的一般 reshape() 支持。

3.9.3.1. CUTENSOREX 中重载的逻辑数组赋值
上面章节中逻辑表达式的结果可以赋值给 logical(4) 类型的数组。例如
A 和 B 是符合规范的 real(4) 类型设备数组，x 是 real(4) 标量，L 是 logical(4) 类型设备数组
block; use cutensorex
L = A .LT. B
L = ABS(A) .GE. 1.0
L = ABS(A - B) .LE. x
end block


当然，可以使用任何方法生成逻辑数组：CUDA 内核、CUF 内核，或者从主机计算/复制。这些功能作为便利提供，但请注意，如果掩码在多次使用中是恒定的，则可能更快地计算一次并将其传递给这些函数，而不是多次重新评估它。


3.9.3.2. Fortran ALL 内部函数
ALL 的接口位于 cutensorex 模块中。
Fortran 数组归约 ALL 在掩码的每个元素都为真时返回真，否则返回假。掩码可以是逻辑数组、数组切片或上述任何逻辑表达式。目前不支持 ALL() 的可选 dim 参数。
logical(4) function all ( mask )
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions


例如，如果 A 和 B 是具有设备或托管属性的符合规范的数组，而 X 是相同类型的标量
IF (ALL(A .EQ. B)) PRINT *,"PASSED"
IF (ALL(ABS(A - B) .GT. X)) CALL REDO()




3.9.3.3. Fortran ANY 内部函数
ANY 的接口位于 cutensorex 模块中。
Fortran 数组归约 ANY 在掩码的任何元素为真时返回真，如果所有元素都为假，则返回假。掩码可以是逻辑数组、数组切片或上述任何逻辑表达式。目前不支持 ANY() 的可选 dim 参数。
logical(4) function any ( mask )
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions


例如，如果 A 和 B 是具有设备或托管属性的符合规范的数组，而 X 是相同类型的标量
IF (ANY(A .EQ. B)) PRINT *,"FAILED"
IF (ANY(ABS(A) .GT. X)) CALL REDO()




3.9.3.4. Fortran COUNT 内部函数
COUNT 的接口位于 cutensorex 模块中。
Fortran 数组归约 COUNT 返回掩码的真元素数量。掩码可以是逻辑数组、数组切片或上述任何逻辑表达式。目前不支持 COUNT() 的可选 dim 参数。
integer function count ( mask )
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions


例如，如果 A 和 B 是符合规范的 real(4) 类型数组，具有设备或托管属性，EPS 是 real(4) 标量，ICNT1 和 ICNT2 是整数标量
ICNT1 = COUNT(A .EQ. B)
ICNT2 = COUNT(ABS(A - B) .LE. EPS)





3.9.4. CUTENSOREX 中重载的 Fortran 数组内部函数
本节列出了 cutensorex 模块中可用的其他重载函数。与上一节类似，当从主机代码使用 cutensorex 模块时，这些 Fortran 内部函数接受设备或托管数组。
本节中的前五个函数也接受掩码参数，并接受上一节中描述的相同掩码数组或表达式。这组中更复杂的函数使用本文中描述的扫描算法：Single-pass Parallel Prefix Scan with Decoupled Look-Back，作者 Duane Merrill 和 Michael Garland。
第二组函数调用 cuTensor 或 cuRand 库，并在此处包含以求完整性。它们之前已在 NVIDIA Fortran CUDA 接口 文档中记录。

3.9.4.1. Fortran MERGE 内部函数
MERGE 的接口位于 cutensorex 模块中。
Fortran merge() 内部函数是基于掩码评估的元素选择。它接受三个参数，一个“真”值数组、一个或多个“假”值和一个掩码。merge() 内部函数可以将上述形式的掩码表达式作为参数，或 logical(4) 设备数组。在当前实现中，只有第二个参数（假选择）可以是标量。仅支持 real(4)、real(8)、integer(4) 和 integer(8) 数组，且仅适用于 1 - 3 维数组。tsource 参数和 mask 参数必须是符合规范的数组，如果 fsource 是数组，则也必须符合规范。
function merge ( tsource, fsource, mask ) result(res)
    type, intent(in) :: tsource(...)  ! type is real or integer, kind = 4 or 8
    type, intent(in) :: fsource(...)  ! type same as tsource, array or scalar
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions
    type, intent(out) :: res(...)     ! type, kind, rank same as tsource


例如：对于 integer(4) 类型的数组 A、B、C 和 integer(4) 类型的标量 K
C = MERGE(A, B, A .GT. B)
C = MERGE(A, 0, ABS(A) .LT. K)




3.9.4.2. Fortran PACK 内部函数
PACK 的接口位于 cutensorex 模块中。
Fortran pack() 内部函数对于将多维数组中的选定数据收集到秩为 1 的数组中非常有用。pack 内部函数的独特之处在于，输出数组的大小在函数完全评估之前是未知的。此 Fortran pack() 内部函数是一个高效的并行实现，可以将上面指定的掩码表达式作为掩码参数。
目前，作为我们强调性能的一部分，我们不重新分配 LHS 目标以适应结果；用户有责任确保 LHS 目标数组足够大。
仅支持 real(4)、real(8)、integer(4) 和 integer(8) 数组，且仅适用于 1 - 3 维数组。此实现不支持 pack() 的 vector 可选参数。此实现确实添加了一个新的可选参数 “count”，它可以返回通过掩码结果的计数，基本上是写入 LHS 结果的元素数量
function pack ( array, mask, count ) result(res)
    type, intent(in) :: array(...)    ! type is real or integer, kind = 4 or 8
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions
    integer, optional, intent(in) :: count
    type, intent(out) :: res(*)       ! type is same as array, rank is 1-D


例如：A 和 B 是 real(4) 类型的设备数组，x 是 real(4) 类型的标量。C 和 D 是相同类型的设备数组，其中 C 符合 A 和 B，D 是 1 维数组
D = PACK(C, ABS(A - B) .GT. x)
D = PACK(C, MASK=(A .EQ. B), COUNT=ICNT)




3.9.4.3. Fortran PACKLOC 函数
PACKLOC 的接口位于 cutensorex 模块中。
Fortran packloc() 函数是 PACK 内部函数的扩展，但不接受源数组。相反，它生成一个打包的索引数组或位置数组，其中掩码评估为真。此 Fortran packloc() 函数使用与 PACK 相同的有效并行实现，并且可以将上面指定的掩码表达式作为掩码参数。
目前，作为我们强调性能的一部分，我们不重新分配 LHS 目标以适应结果；用户有责任确保 LHS 目标数组足够大。
目前仅支持涉及一维数组的掩码表达式。此实现确实支持可选参数 “count”，它可以返回通过掩码结果的计数，基本上是写入 LHS 结果的元素数量
与 maxloc 和其他位置函数类似，索引从 1 开始，并且不受传递给函数的数组的非单位步幅或下界的影响。
function packloc ( mask, count ) result(res)
    logical, intent(in) :: mask(:)   ! mask is a 1D logical array or supported expression
    integer(4), optional, intent(out) :: count
    integer(4), intent(out) :: res(*)


例如：A 和 B 是 real(4) 类型的设备数组，x 是 real(4) 类型的标量。D 是 integer(4) 设备数组
D = PACKLOC(ABS(A - B) .GT. x, COUNT=ICNT)




3.9.4.4. Fortran UNPACK 内部函数
UNPACK 的接口位于 cutensorex 模块中。
Fortran unpack() 内部函数是 pack() 的补充，可以将上面指定的掩码表达式作为掩码参数。当前 unpack() 的实现中，与 field 参数相关的一些限制。在此实现中，field 参数是可选的，如果省略，则 LHS 目标被视为 field。如果 field 参数是标量，则 unpack 按照标准工作。如果 field 参数是数组，则掩码操作必须是逻辑数组，而不是掩码表达式，当然，掩码和 field 的大小和形状必须符合规范。
仅支持 real(4)、real(8)、integer(4) 和 integer(8) 数组，且仅适用于 1 - 3 维数组。输出数组和掩码参数必须是符合规范的数组。
function unpack ( array, mask, field ) result(res)
    type, intent(in) :: array(*)      ! type is real or integer, kind = 4 or 8
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions
    type, optional, intent(in) :: field(...)  ! array or scalar
    type, intent(out) :: res(...)     ! type same as array, rank same as mask


例如：对于 real(4) 类型的设备数组 A 和 B，x 和 y 是 real(4) 类型的标量，C 和 D 是相同类型的数组，其中 C 符合 A 和 B，D 是 1 维数组
C = UNPACK(D, ABS(A - B) .GT. x)
C = UNPACK(D, MASK=(ABS(A - B) .GT. x), FIELD=y)




3.9.4.5. Fortran COUNT_PREFIX 内部函数
COUNT_PREFIX 的接口位于 cutensorex 模块中。
count_prefix 函数在高性能 Fortran (HPF) 中定义。它计算数组存储顺序中真掩码值的运行计数。可选的逻辑参数 EXCLUSIVE 指定当前掩码结果不影响当前输出，仅影响后续计数。另一个可选参数整数 DIM 指定仅在特定维度上计算多维数组的计数。
完整的函数声明和参数列表是
function count_prefix ( mask, dim, exclusive ) result(res)
    logical, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions
    integer, optional, intent(in) :: dim  ! 1 - 3 depending on rank of mask
    logical, optional, intent(in) :: exclusive ! Default is .false. (inclusive)
    integer, intent(out) :: res(...)  ! same size and rank as mask


例如：对于 real(4) 类型的 A、B 和 x，integer(4) 类型的 C
C = COUNT_PREFIX(A .GT. 0)
C = COUNT_PREFIX(A .EQ. B, DIM=1)
C = COUNT_PREFIX(MASK=ABS(A - B) .LE. x, DIM=2, EXCLUSIVE=.true.)


HPF 还指定了 SEGMENT 可选参数，但当前版本中没有该功能。


3.9.4.6. Fortran SUM_PREFIX 内部函数
SUM_PREFIX 的接口位于 cutensorex 模块中。
sum_prefix 函数也在 HPF 中定义。它计算数组元素值的运行总和，对于这些值，相应的掩码为真，按数组存储顺序排列。可选的逻辑参数 EXCLUSIVE 指定数组值不影响当前输出，仅影响后续总和。另一个可选参数整数 DIM 指定仅在特定维度上计算多维数组的计数。对于此函数，MASK 也是可选的；如果没有它，则每个数组元素都参与总和。
完整的函数声明和参数列表是
function sum_prefix ( array, mask, dim, exclusive ) result(res)
    type, intent(in) :: array(...)    ! type is real or integer, kind = 4 or 8
    logical, optional, intent(in) :: mask(...)  ! mask is 1 - 3 dimensions
    integer, optional, intent(in) :: dim  ! 1 - 3 depending on rank of array
    logical, optional, intent(in) :: exclusive ! Default is .false. (inclusive)
    type, intent(out) :: res(...)  ! same size and rank as array


例如：对于 real(4) 类型的 A、B、C、D 和 x
D = SUM_PREFIX(C, A .GT. 0)
D = SUM_PREFIX(C, MASK=(A .NE. B), DIM=2)
D = SUM_PREFIX(C, MASK=ABS(A - B) .LE. x, EXCLUSIVE=.true.)




3.9.4.7. Fortran RESHAPE 内部函数
RESHAPE 的接口位于 cutensorex 模块中。此函数也在 NVIDIA Fortran CUDA 接口 文档的 cuTensor 章节中进行了详细记录。
Fortran reshape() 内部函数更改数组的形状，并可能置换维度和布局。它的调用方式为
D = alpha * func(reshape(A, shape=[...], order=[...]))
数组 A 和 D 的类型可以是 real(2)、real(4)、real(8)、complex(4) 或 complex(8)。A 和 D 的秩（维度数）可以是 1 到 7。alpha 值应与 A 的类型相同，或者与 func(reshape(A)) 的类型相同（如果不同）。上面引用的 Fortran CUDA 接口文档中列出了可以应用于 reshape 结果的接受函数。目前不支持 F90 reshape 函数的 pad 参数。除了初始化和设置 cuTENSOR 描述符外，此 Fortran 调用还映射到 cutensorPermutation()。
function reshape ( source, shape, order ) result(res)
    type, intent(in) :: source(...)  ! type is real or complex
    integer, intent(in) :: shape(:)
    integer, optional, intent(in) :: order(*)
    type, intent(out) :: res(...)     ! type, kind same as source


! Example to switch the 2nd and 3rd dimension layout
D = reshape(a,shape=[ni,nk,nj], order=[1,3,2])
! Same example, take the absolute value and scale by 2.5
D = 2.5 * abs(reshape(a,shape=[ni,nk,nj], order=[1,3,2]))




3.9.4.8. Fortran TRANSPOSE 内部函数
TRANSPOSE 的接口位于 cutensorex 模块中。此函数也在 NVIDIA Fortran CUDA 接口 文档的 cuTensor 章节中进行了详细记录。
Fortran transpose() 内部函数转置矩阵（二维数组）。它的调用方式为
D = alpha * func(transpose(A))
数组 A 和 D 的类型可以是 real(2)、real(4)、real(8)、complex(4) 或 complex(8)。A 和 D 的秩（维度数）为 2。应用缩放（alpha 参数）或将函数应用于转置结果是可选的。alpha 值应与 A 的类型相同，或者与 func(transpose(A)) 的类型相同（如果不同）。上面引用的 Fortran CUDA 接口文档中列出了可以应用于转置结果的接受函数。除了初始化和设置 cuTENSOR 描述符外，此 Fortran 调用还映射到 cutensorPermutation()。
! Example of transpose
D = transpose(A)
! Same example, take the absolute value and scale by 2.5
D = 2.5 * abs(tranpose(A))


transpose() 函数也作为上述掩码表达式中 “B” 参数的一部分受支持。例如，如果 A 是一个 2 维 mxn 数组，B 是 nxm
ICNT = COUNT(A .GT. TRANSPOSE(B))




3.9.4.9. Fortran SPREAD 内部函数
SPREAD 的接口位于 cutensorex 模块中。此函数也在 NVIDIA Fortran CUDA 接口 文档的 cuTensor 章节中进行了记录。
Fortran spread() 内部函数通过指定的维度将数组的秩增加 1，并在新维度上广播值。它的调用方式为
D = alpha * func(spread(A, dim=i, ncopies=n))
数组 A 和 D 的类型可以是 real(2)、real(4)、real(8)、complex(4) 或 complex(8)。A 和 D 的秩（维度数）可以是 1 到 7。alpha 值应与 A 的类型相同。上面引用的 Fortran CUDA 接口文档中列出了可以应用于 spread 结果的接受函数。除了初始化和设置 cuTENSOR 描述符外，此 Fortran 调用还映射到 cutensorPermutation()。
! Example to add and broadcast values over the new first dimension
D = spread(A, dim=1, ncopies=n1)
! Same example, take the absolute value and scale by 2.5
D = 2.5 * abs(spread(A, dim=1, ncopies=n1))


spread() 函数也作为上述掩码表达式中 “B” 参数的一部分受支持。例如，如果 A 是一个 2 维 mxn 数组，B 是长度为 m 的 1 维数组
ICNT = COUNT(A .GT. SPREAD(B, dim=2, ncopies=n))




3.9.4.10. Fortran MATMUL 内部函数
MATMUL 的接口位于 cutensorex 模块中。此函数也在 NVIDIA Fortran CUDA 接口 文档的 cuTensor 章节中进行了详细记录。
Fortran matmul() 内部函数执行矩阵乘法，即张量收缩的一个实例。matmul 的任一操作数都可以是置换数组，reshape()、transpose() 或 spread() 调用的结果。cuTENSOR 库目前不支持将元素函数应用于数组操作数，但可以缩放结果和累加器。以下是一些受支持的形式
D = matmul(A, B)
D = matmul(permute(A), B)
D = matmul(A, permute(B))
D = matmul(permute(A), permute(B))
D = C + matmul(A, B)
D = C - matmul(A, B)
D = alpha * matmul(A, B) + beta * C
数组 A、B、C 和 D 的类型可以是 real(2)、real(4)、real(8)、complex(4) 或 complex(8)。在任何置换之后，A、B、C 和 D 的秩（维度数）必须为 2。数组 C 和 D 当前必须具有相同的形状、步幅和类型。alpha 值应与 A 和 B 的类型相同。beta 值应与 C 的类型相同。Fortran 包装器不进行类型转换，但 cuTENSOR 可能会进行类型转换。对数组一致性的编译时检查是有限的。对不支持的组合的其他运行时检查可能来自 Fortran 包装器或 cuTENSOR。除了初始化和设置 cuTENSOR 描述符外，Fortran 对 Matmul 的支持还映射到 cutensorContraction()。
! Example to multiply two matrices together
D = matmul(A, B)
! Same example, accumulate into C
C = C +  matmul(A, B)
! Same example, transpose the first argument
C = C + matmul(transpose(A), B)




3.9.4.11. Fortran DOT_PRODUCT 内部函数
DOT_PRODUCT 的接口位于 cutensorex 模块中。此函数也在 NVIDIA Fortran CUDA 接口 文档的 cuTensor 章节中进行了详细记录。
Fortran dot_product() 内部函数执行两个向量的点积，即张量收缩的一个特定实例。在标准形式中，它返回与输入参数类型相同的标量，并且赋值 LHS 上的目标必须具有设备或托管属性。dot_product 的任一操作数都可以是置换数组，reshape() 调用的结果（创建一维数组）。请注意，复数变量的 Fortran dot_product 定义对第一个参数执行共轭。以下是一些受支持的形式
S = dot_product(A, B)
S = dot_product(reshape(T, shape=[m*n]), B)
S = dot_product(ABS(A), B)
ZC = dot_product(ZX,ZY) ! BLAS ZDOTC 等效
ZU = dot_product(CONJG(ZX),ZY) ! BLAS ZDOTU 等效
输入数组的类型可以是 real(2)、real(4)、real(8)、complex(4) 或 complex(8)。除了初始化和设置 cuTENSOR 描述符外，Fortran 对 DOT_PRODUCT 的支持还映射到 cutensorContraction()。
此实现已扩展为在高层次上公开更多的 cutensorContraction() 功能。DOT_PRODUCT 的扩展接口接受多维数组和一个 dim 参数。点积将仅沿指定维度计算，从而生成秩比输入数组少 1 的数组。
与 MATMUL 一起提供的许多相同的置换、函数和累加操作也与 DOT_PRODUCT 一起提供。以下是一些示例：对于 NxN 矩阵 A 和 B，长度为 N 的向量 X 和 V，以及标量 alpha
X = dot_product(A, B, dim=1)
X = X + dot_product(A, transpose(B), dim=1)
X = X - alpha * dot_product(spread(V, dim=1, ncopies=N), B, dim=2)


3.9.4.12. Fortran RANDOM_NUMBER 固有函数
RANDOM_NUMBER 的接口实际上是在 curandex 模块中，但该模块包含在/使用于 cutensorex 模块内。
Fortran 子例程 RANDOM_NUMBER 返回介于 0.0 和 1.0 之间的随机数。此接口作为一种便利提供，尚未经过广泛的测试。当您传递具有 device 或 managed 属性的数组时，该子例程将调用 cuRAND 库函数来生成值。还完成了一些额外的工作来支持类型 real(2)、real(4)、real(8)、complex(4) 和 complex(8)，其中一些是非标准的。仅支持 1 - 3 维数组。
例如，如果 A 是一个实数数组并且具有 device 或 managed 属性
block; use cutensorex
CALL RANDOM_NUMBER(A)
end block


curandex 模块使用的默认生成器是 CURAND_RNG_PSEUDO_XORWOW。
curandex 模块中提供了一些辅助函数来微调 cuRAND 库随机数生成器，这些函数应该是自明的
integer(4) function curandExSetCurandGenerator(g)
type(curandGenerator) :: g
end function

function curandExGetCurandGenerator() result(s)
type(curandGenerator) :: s
end function

integer(4) function curandExSetStream(stream)
integer(kind=cuda_stream_kind), value :: stream
end function

function curandExGetStream() result(s)
integer(kind=cuda_stream_kind) :: s
end function





3.9.5. 其他 CUDA 库主机模块
有关 CUDA 库的 Fortran 接口的更多详细信息，请参阅 NVIDIA Fortran CUDA 接口 文档。本节讨论了一些更常用的接口和库。
NVIDIA 提供了一个模块，该模块定义了从 NVIDIA CUDA Fortran 到 CUBLAS 库的接口。通过在 CUDA Fortran 主机代码程序单元中放置以下语句，可以访问这些接口。
use cublas


这些接口目前有三种形式

重载的传统 BLAS 接口，它将设备数组作为参数而不是主机数组，即：
call saxpy(n, a, x, incx, y, incy)


其中参数 x 和 y 具有 device 属性。

可移植的旧版 CUBLAS 接口，它直接与 CUBLAS 版本 < 4.0 接口，即：
call cublasSaxpy(n, a, x, incx, y, incy)


其中参数 x 和 y 必须具有 device 属性。

新的 CUBLAS 4.0+ 接口，可以访问新库的所有功能。
这些接口都是函数调用的形式，将句柄作为第一个参数，并通过引用传递许多标量参数和结果，即：
istat = cublasSaxpy_v2(h, n, a, x, incx, y, incy)


在 saxpy 的情况下，用户现在可以选择让 a 驻留在主机或设备上。传统上返回标量的函数，例如 sdot() 和 isamax()，现在采用额外的参数来返回结果。传统上采用 character*1 参数的函数，例如 t 或 n 来控制转置，现在采用在 cublas 模块中定义的整数值。
为了支持第三种形式，在 cublas 模块中定义了一个名为 cublasHandle 的派生类型。您可以使用以下命令定义此类型的变量
type(cublasHandle) :: h


通过将其传递给 cublasCreate 函数来初始化它。


当使用 CUBLAS 4.0 及更高版本时，cublas 模块会从串行和 OpenMP 并行区域正确生成前两种形式的句柄。
允许混合使用这三种形式。要访问 cublas 模块内部使用的句柄，请使用
h = cublasGetHandle()


还支持以下形式“istat = cublasGetHandle(h)”。
istat = cublasGetHandle(h)


支持 cublasHandle 类型的赋值以及相等性和不等性测试。
cublas 模块中定义的 CUDA 4.0+ 辅助函数
integer function cublasCreate(handle)
integer function cublasDestroy(handle)
integer function cublasGetVersion(handle, version)
integer function cublasSetStream(handle, stream)
integer function cublasGetStream(handle, stream)
integer function cublasGetPointerMode(handle, mode)
integer function cublasSetPointerMode(handle, mode)


有关演示 cublas 模块、cublasHandle 类型以及三种调用形式的用法的示例，请参阅 Cublas 模块示例。
NVIDIA 提供了另一个模块，该模块定义了从 NVIDIA CUDA Fortran 到 CUFFT 库的接口。通过在 CUDA Fortran 主机代码程序单元中放置以下语句，可以访问这些接口。
use cufft


这是一个使用 cufft 接口的一些代码示例
program cufft2dTest
  use cufft
  integer, parameter :: n=450
  complex :: a(n,n),b(n,n)
  complex, device :: a_d(n,n), b_d(n,n)
  integer :: plan, ierr
  real, dimension(3) :: res, exp

  a  = 1; a_d  = a

  ierr = cufftPlan2D(plan,n,n,CUFFT_C2C)
  ierr = ierr + cufftExecC2C(plan,a_d,b_d,CUFFT_FORWARD)
  ierr = ierr + cufftExecC2C(plan,b_d,b_d,CUFFT_INVERSE)

  b = b_d
  res(1) = maxval(abs(a-b/(n*n)))
  print *,'Max error C2C: ', res(1)


该发行版还包含一个模块，该模块定义了从 NVIDIA CUDA Fortran 到 CUSPARSE 库的接口。通过在 CUDA Fortran 主机代码程序单元中放置以下语句，可以显式访问这些接口。
use cusparse


除了函数接口之外，cusparse 模块中还定义了几个重要的派生类型和常量。这是一个它们用法的示例
! Compile with "nvfortran testLevel3.cuf -cudalib=cusparse"
program testLevel3
  use cudafor
  use cusparse

  implicit none

  integer, parameter :: nd = 20 ! # rows/cols in dense matrix

  type(cusparseHandle) :: h
  type(cusparseMatDescr) :: descrA
  type(cusparseSolveAnalysisInfo) :: saInfo
  integer :: status, version, mode, i

  ! D-data
  ! dense
  real(8) :: DAde(nd,nd), DBde(nd,nd), DCde(nd,nd), DmaxErr
  real(8), device :: DAde_d(nd,nd), DBde_d(nd,nd), DCde_d(nd,nd)
  ! csr
  real(8) :: csrValDA(nd)
  real(8), device :: csrValDA_d(nd)
  real(8) :: Dalpha, Dbeta
  real(8), device :: Dalpha_d, Dbeta_d

  ! integer data common to all data types
  integer :: nnz
  integer :: nnzPerRowA(nd), csrRowPtrA(nd+1), csrColIndA(nd)
  integer, device :: nnzPerRowA_d(nd), csrRowPtrA_d(nd+1), csrColIndA_d(nd)

  ! initalization

  status = cusparseCreate(h)
  status = cusparseGetVersion(h, version)
  write(*,*) '... version:', version

  status = cusparseCreateMatDescr(descrA)
  status = cusparseSetMatType(descrA, CUSPARSE_MATRIX_TYPE_GENERAL)
  status = cusparseSetMatIndexBase(descrA, CUSPARSE_INDEX_BASE_ONE)
  status = cusparseCreateSolveAnalysisInfo(saInfo)

  ! Initialize matrix (Identity)

  DAde = 0.0
  i = 1, nd
      DAde(i,i) = 1.0
  end do
  DAde_d = DAde
  call random_number(DBde)
  DBde_d = DBde

  ! convert from dense to csr
  status = cusparseDnnz_v2(h, CUSPARSE_DIRECTION_ROW, nd, nd, descrA, &
      DAde_d, nd, nnzPerRowA_d, nnz)
  status = cusparseDdense2csr(h, nd, nd, descrA, DAde_d, nd, nnzPerRowA_d, &
      csrValDA_d, csrRowPtrA_d, csrColIndA_d)

  ! csrmm HPM
  Dalpha = 1.0
  Dbeta = 0.0
  status = cusparseDcsrmm(h, CUSPARSE_OPERATION_NON_TRANSPOSE, nd, nd, nd, &
      nnz, Dalpha, descrA, csrValDA_d, csrRowPtrA_d, csrColIndA_d, DBde_d, &
      nd, Dbeta, DCde_d, nd)
  if (status /= CUSPARSE_STATUS_SUCCESS) write (*,*) 'CSRMM Error:',status

  DCde = DCde_d
  DmaxErr = maxval(abs(DCde-DBde))

  status = cusparseDestroy(h)
  write(*,*) 'cusparseDestroy', status, DmaxErr

end program testLevel3



4. 运行时 API
系统模块 cudafor 定义了运行时 API 例程的接口。
有关本章中每个例程的用途和功能的完整说明，请参阅 docs.nvidia.com/cuda/cuda-runtime-api。
大多数运行时 API 例程都是返回错误代码的整数函数；如果调用成功，则返回零值，如果出现错误，则返回非零值。要解释错误代码，请参阅 错误处理。
除非提供了特定的 kind，否则 plain integer 类型表示 integer(4)，plain real 类型表示 real(4)。

4.1. 初始化
不需要显式初始化；运行时在首次调用运行时例程或分配设备数组时初始化并连接到设备。

提示
在进行计时运行时，请注意初始化可能会增加一些开销。



4.2. 设备管理
使用本节中的函数进行设备管理。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的设备管理部分。

4.2.1. cudaChooseDevice
integer function cudaChooseDevice ( devnum, prop )
     integer, intent(out) :: devnum
     type(cudadeviceprop), intent(in) :: prop


cudaChooseDevice 将最符合 prop 中给出的属性的设备编号分配给其第一个参数。


4.2.2. cudaDeviceGetAttribute
integer function cudaDeviceGetAttribute ( val, attribute, devicenum )
     integer, intent(out) :: val
     integer, intent(in)  :: attribute
     integer, intent(in)  :: devicenum


cudaDeviceGetAttribute 返回有关设备的信息。返回的特定信息由指定设备编号的 attribute 参数确定。


4.2.3. cudaDeviceGetCacheConfig
integer function cudaDeviceGetCacheConfig ( cacheconfig )
     integer, intent(out) :: cacheconfig


cudaDeviceGetCacheConfig 返回当前设备的首选缓存配置。当前可能的缓存配置定义为 cudaFuncCachePreferNone、cudaFuncCachePreferShared 和 cudaFuncCachePreferL1。
cudaDeviceGetCacheConfig 在 CUDA 5.0 开始的设备代码中可用。


4.2.4. cudaDeviceGetLimit
integer function cudaDeviceGetLimit( val, limit )
   integer(kind=cuda_count_kind) :: val
   integer :: limit


cudaGetDeviceGetLimit 在 val 中返回限制的当前大小。当前可能的限制参数是 cudaLimitStackSize、cudaLimitPrintfSize 和 cudaLimitMallocHeapSize。
cudaGetDeviceGetLimit 在 CUDA 5.0 开始的设备代码中可用。


4.2.5. cudaDeviceGetSharedMemConfig
integer function cudaDeviceGetSharedMemConfig ( config )
integer, intent(out) :: config


cudaDeviceGetSharedMemConfig 返回当前设备上共享内存库的当前大小。此例程适用于具有可配置共享内存库的设备，并且从 CUDA 4.2 开始受支持。当前可能的共享内存配置定义为 cudaSharedMemBankSizeDefault、cudaSharedMemBankSizeFourByte 和 cudaSharedMemBankSizeEightByte。


4.2.6. cudaDeviceGetStreamPriorityRange
integer function cudaDeviceGetStreamPriorityRange ( leastpriority, greatestpriority )
integer, intent(out) :: leastpriority, greatestpriority


cudaDeviceGetStreamPriorityRange 返回与当前上下文和设备的最小和最大流优先级相对应的数值。


4.2.7. cudaDeviceReset
integer function cudaDeviceReset()


cudaDeviceReset 重置附加到当前进程的当前设备。


4.2.8. cudaDeviceSetCacheConfig
integer function cudaDeviceSetCacheConfig ( cacheconfig )
     integer, intent(in) :: cacheconfig


cudaDeviceSetCacheConfig 设置当前设备的首选缓存配置。当前可能的缓存配置定义为 cudaFuncCachePreferNone、cudaFuncCachePreferShared 和 cudaFuncCachePreferL1。


4.2.9. cudaDeviceSetLimit
integer function cudaDeviceSetLimit( limit, val )
  integer :: limit

  integer(kind=cuda_count_kind) :: val


cudaGetDeviceSetLimit 将当前设备的限制设置为 val。当前可能的限制参数是 cudaLimitStackSize、cudaLimitPrintfSize 和 cudaLimitMallocHeapSize。


4.2.10. cudaDeviceSetSharedMemConfig
integer function cudaDeviceSetSharedMemConfig ( config )
    integer, intent(in) :: config


cudaDeviceSetSharedMemConfig 设置当前设备上共享内存库的大小。此例程适用于具有可配置共享内存库的设备，并且从 CUDA 4.2 开始受支持。当前可能的共享内存配置定义为 cudaSharedMemBankSizeDefault、cudaSharedMemBankSizeFourByte 和 cudaSharedMemBankSizeEightByte。


4.2.11. cudaDeviceSynchronize
integer function cudaDeviceSynchronize()


cudaDeviceSynchronize 阻止当前设备，直到所有先前的请求任务都已完成。
cudaDeviceSynchronize 在 CUDA 5.0 开始的设备代码中可用。
cudaDeviceSynchronize 已从 NVHPC 22.11 版本开始的 CUDA Fortran 设备代码中删除，因为它在 CUDA 编程模型中不再受支持。


4.2.12. cudaGetDevice
integer function cudaGetDevice( devnum )
    integer, intent(out) :: devnum


cudaGetDevice 将与此主机线程关联的设备编号分配给其第一个参数。
cudaGetDevice 在 CUDA 5.0 开始的设备代码中可用。


4.2.13. cudaGetDeviceCount
integer function cudaGetDeviceCount( numdev )
    integer, intent(out) :: numdev


cudaGetDeviceCount 将可用设备的数量分配给其第一个参数。
cudaGetDeviceCount 在 CUDA 5.0 开始的设备代码中可用。


4.2.14. cudaGetDeviceProperties
integer function cudaGetDeviceProperties( prop, devnum )
     type(cudadeviceprop), intent(out) :: prop
     integer, intent(in) :: devnum


cudaGetDeviceProperties 返回给定设备的属性。
cudaGetDeviceProperties 在 CUDA 5.0 开始的设备代码中可用。


4.2.15. cudaSetDevice
integer function cudaSetDevice( devnum )
    integer, intent(in) :: devnum


cudaSetDevice 选择要与此主机线程关联的设备。


4.2.16. cudaSetDeviceFlags
integer function cudaSetDevice( flags )
    integer, intent(in) :: flags


cudaSetDeviceFlags 记录 CUDA 运行时如何与此主机线程交互。


4.2.17. cudaSetValidDevices
integer function cudaSetValidDevices( devices, numdev )
    integer :: numdev, devices(numdev)


cudaSetValidDevices 设置 CUDA 执行的有效设备列表，优先级顺序如 devices 数组中指定的那样。



4.3. 线程管理
有时，块内的线程访问共享内存或全局内存中的相同地址，从而为某些内存访问创建潜在的先写后读、先读后写或先写后写危害。为了避免这些潜在问题，请使用本节中的函数进行线程管理。这些函数已从 CUDA 4.0 开始弃用。

4.3.1. cudaThreadExit
integer function cudaThreadExit()


cudaThreadExit 显式清理与主机线程关联的所有运行时相关的 CUDA 资源。任何后续的 CUDA 调用或操作都将重新初始化运行时。
调用 cudaThreadExit 是可选的；当主机线程退出时，它会被隐式调用。


4.3.2. cudaThreadSynchronize
integer function cudaThreadSynchronize()


cudaThreadSynchronize 阻止主机子程序的执行，直到所有先前的内核和操作完成。如果先前的操作之一失败，它可能会返回错误条件。

注意
此函数已被弃用，因为其名称不能反映其行为。其功能与非弃用函数 cudaDeviceSynchronize() 相同，您应该改用该函数。




4.4. 错误处理
使用本节中的函数进行错误处理。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的错误处理部分。

4.4.1. cudaGetErrorString
function cudaGetErrorString( errcode )
    integer(4), intent(in) :: errcode
    character*(*) :: cudaGetErrorString


cudaGetErrorString 返回与给定错误代码关联的消息字符串。


4.4.2. cudaGetLastError
integer function cudaGetLastError()


cudaGetLastError 返回最近从此主机线程中的任何运行时调用返回的错误代码。


4.4.3. cudaPeekAtLastError
integer function cudaPeekAtLastError()


cudaPeekAtLastError 返回 CUDA 运行时产生的最后一个错误代码，而不会像 cudaGetLastError 一样将错误代码重置为 cudaSuccess。



4.5. 流管理
使用本节中的函数进行流管理。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的流管理部分。

4.5.1. cudaforGetDefaultStream
integer(kind=cuda_stream_kind) function cudaforGetDefaultStream( devptr )


cudaforGetDefaultStream 返回已通过调用 cudaforSetDefaultStream 与线程、托管变量或设备变量关联的默认流。devptr 可以是 设备代码固有数据类型 中指定的支持类型的任何托管或设备标量或数组。devptr 参数是可选的；如果未指定，则该函数返回绑定到线程的流，或零（默认流）。
从 cudaforGetDefaultStream 返回的流值可以用作其他 CUDA 库的参数，例如例程 cublasSetStream()、cufftSetStream() 和 cusparseSetStream()。


4.5.2. cudaforSetDefaultStream
integer function cudaforSetDefaultStream( devptr, stream )
    integer(kind=cuda_stream_kind), intent(in) :: stream


cudaforSetDefaultStream 为由该 CPU 线程启动的对托管数据或设备数据执行的所有后续高级 CUDA Fortran 操作设置默认流。受托管数据影响的特定操作是通过 Fortran allocate 语句进行的分配、赋值（memset 和 memcpy 类型）、CUF Kernel 和全局内核启动以及 sum()、maxval() 和 minval() 缩减操作。devptr 可以是 设备代码固有数据类型 中指定的支持类型的任何托管或设备标量或数组。devptr 参数是可选的；如果未指定，则该函数将指定的流绑定到在该线程上执行的所有后续的、允许的高级操作。


4.5.3. cudaStreamAttachMemAsync
integer function cudaStreamAttachMemAsync( stream, devptr, length, flags )
    integer(kind=cuda_stream_kind), intent(in) :: stream
    integer(kind=cuda_count_kind), optional, intent(in) :: length
    integer, optional, intent(in) :: flags


cudaStreamAttachMemAsync 启动流操作，以将地址 devptr 处的托管分配附加到指定的流。devptr 可以是 设备代码固有数据类型 中指定的支持类型的任何托管标量或数组。参数 len 是可选的，但目前必须为零。flags 参数必须是 cudaMemAttachGlobal、cudMemAttachHost 或 cudMemAttachSingle。
cudaStreamAttachMemAsync 从 CUDA 6.0 开始可用。


4.5.4. cudaStreamCreate
integer function cudaStreamCreate( stream )
    integer(kind=cuda_stream_kind), intent(out) :: stream


cudaStreamCreate 创建一个异步 stream 并将其标识符分配给其第一个参数。


4.5.5. cudaStreamCreateWithFlags
integer function cudaStreamCreateWithFlags( stream, flags )
    integer(kind=cuda_stream_kind), intent(out) :: stream
    integer, intent(in) :: flags


cudaStreamCreateWithFlags 创建一个异步 stream 并将其标识符分配给其第一个参数。flags 的有效值是 cudaStreamDefault 或 cudaStreamNonBlocking。
cudaStreamCreateWithFlags 在 CUDA 5.0 开始的设备代码中可用。


4.5.6. cudaStreamCreateWithPriority
integer function cudaStreamCreateWithPriority( stream, flags, priority )
    integer(kind=cuda_stream_kind), intent(out) :: stream
    integer, intent(in) :: flags, priority


cudaStreamCreateWithPriority 创建一个异步 stream 并将其标识符分配给其第一个参数。flags 的有效值是 cudaStreamDefault 或 cudaStreamNonBlocking。较低的优先级值表示较高的优先级。较高优先级流中的工作可能会抢占已在较低优先级流中执行的工作。


4.5.7. cudaStreamDestroy
integer function cudaStreamDestroy( stream )
    integer(kind=cuda_stream_kind), intent(in) :: stream


cudaStreamDestroy 释放与给定 stream 关联的任何资源。
cudaStreamDestroy 在 CUDA 5.0 开始的设备代码中可用。


4.5.8. cudaStreamGetPriority
integer function cudaStreamGetPriority( stream, priority )
    integer(kind=cuda_stream_kind), intent(in) :: stream
    integer, intent(out) :: priority


cudaStreamGetPriority 查询并在 priority 中返回给定流的优先级。


4.5.9. cudaStreamQuery
integer function cudaStreamQuery( stream )
    integer(kind=cuda_stream_kind), intent(in) :: stream


cudaStreamQuery 测试是否已完成排队到所选 stream 的所有操作；如果所有操作都已完成，则返回零（成功），如果未完成，则返回 cudaErrorNotReady 值。如果某些异步操作失败，它也可能返回另一个错误条件。


4.5.10. cudaStreamSynchronize
integer function cudaStreamSynchronize( stream )
    integer(kind=cuda_stream_kind), intent(in) :: stream


cudaStreamSynchronize 阻止主机子程序的执行，直到与给定 stream 关联的所有先前内核和操作完成。它可能会返回来自先前异步操作的错误代码。


4.5.11. cudaStreamWaitEvent
integer function cudaStreamWaitEvent( stream, event, flags )
    integer(kind=cuda_stream_kind) :: stream
    type(cudaEvent), intent(in) :: event
    integer :: flags


cudaStreamWaitEvent 阻止在 stream 上提交的所有工作执行，直到 event 报告完成。
cudaStreamWaitEvent 在 CUDA 5.0 开始的设备代码中可用。



4.6. 事件管理
使用本节中的函数来管理事件。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的事件管理部分。

4.6.1. cudaEventCreate
integer function cudaEventCreate( event )
    type(cudaEvent), intent(out) :: event


cudaEventCreate 创建一个事件对象，并将事件标识符分配给其第一个参数


4.6.2. cudaEventCreateWithFlags
integer function cudaEventCreateWithFlags( event, flags )
    type(cudaEvent), intent(out) :: event
    integer :: flags


cudaEventCreateWithFlags 使用指定的标志创建一个事件对象。当前支持的标志是 cudaEventDefault、cudaEventBlockingSync 和 cudaEventDisableTiming。
cudaEventCreateWithFlags 在 CUDA 5.0 开始的设备代码中可用。


4.6.3. cudaEventDestroy
integer function cudaEventDestroy( event )
    type(cudaEvent), intent(in) :: event


cudaEventDestroy 销毁与事件对象关联的资源。
cudaEventDestroy 在 CUDA 5.0 开始的设备代码中可用。


4.6.4. cudaEventElapsedTime
integer function cudaEventElapsedTime( time, start, end)
    real :: time
    type(cudaEvent), intent() :: start, end


cudaEventElapsedTime 计算两个事件之间经过的时间（以毫秒为单位）。如果任一事件尚未记录，则返回 cudaErrorInvalidValue。此函数仅对在流零上记录的事件有效。


4.6.5. cudaEventQuery
integer function cudaEventQuery( event )
    type(cudaEvent), intent(in) :: event


cudaEventQuery 测试是否已记录 event。如果事件已记录，则返回成功（零），如果未记录，则返回 cudaErrorNotReady。如果尚未为此事件调用 cudaEventRecord，则返回 cudaErrorInvalidValue。


4.6.6. cudaEventRecord
integer function cudaEventRecord( event, stream )
    type(cudaEvent), intent(in) :: event
    integer, intent(in) :: stream


cudaEventRecord 向给定 stream 发出操作以记录 event。在流中的所有先前操作完成后记录事件。如果 stream 为零，则在所有流中的所有先前操作完成后记录事件。
cudaEventRecord 在 CUDA 5.0 开始的设备代码中可用。


4.6.7. cudaEventSynchronize
integer function cudaEventSynchronize( event )
    type(cudaEvent), intent(in) :: event


cudaEventSynchronize 阻塞直到 event 已被记录。如果尚未为此事件调用 cudaEventRecord，则返回值 cudaErrorInvalidValue。



4.7. 执行控制
CUDA Fortran 不支持所有复制 chevron 语法功能的 API 例程。提供了更高版本的 CUDA 提供的附加功能。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nv idia.com/cuda/cuda-runtime-api 上的执行控制部分。

4.7.1. cudaFuncGetAttributes
integer function cudaFuncGetAttributes( attr, func )
    type(cudaFuncAttributes), intent(out) :: attr
    external :: func


cudaFuncGetAttributes 获取由 func 参数命名的函数的属性，该参数必须是全局函数。
cudaFuncGetAttributes 在 CUDA 5.0 开始的设备代码中可用。


4.7.2. cudaFuncSetAttribute
integer function cudaFuncSetAttribute( func, attribute, value )
    external :: func
    integer  :: attribute
    integer  :: value


cudaFuncSetAttribute 设置由 func 参数命名的函数的属性，该参数必须是全局函数。


4.7.3. cudaFuncSetCacheConfig
integer function cudaFuncSetCacheConfig( func, cacheconfig )
    character*(*) :: func
    integer :: cacheconfig


cudaFuncSetCacheConfig 设置由 func 参数命名的函数的首选缓存配置，该参数必须是全局函数。当前可能的缓存配置定义为 cudaFuncCachePreferNone、cudaFuncCachePreferShared 和 cudaFuncCachePreferL1。


4.7.4. cudaFuncSetSharedMemConfig
integer function cudaFuncSetSharedMemConfig( func, cacheconfig )
character*(*) :: func
integer :: cacheconfig


cudaFuncSetSharedMemConfig 设置由 func 参数命名的函数的共享内存库大小，该参数必须是全局函数。此例程用于具有可配置共享内存库的设备，并从 CUDA 4.2 开始受支持。当前可能的共享内存配置定义为 cudaSharedMemBankSizeDefault、cudaSharedMemBankSizeFourByte 和 cudaSharedMemBankSizeEightByte。


4.7.5. cudaSetDoubleForDevice
integer function cudaSetDoubleForDevice( d )
    real(8) :: d


cudaSetDoubleForDevice 将参数 d 设置为适用于不支持双精度算术的设备的内部表示形式。


4.7.6. cudaSetDoubleForHost
integer function cudaSetDoubleForHost( d )
    real(8) :: d


cudaSetDoubleForHost 将来自不支持双精度算术的设备上的内部表示形式的参数 d 设置为正常的宿主机表示形式。



4.8. Occupancy
占用率例程将全局子例程作为参数，并返回值，这些值与占用率相关，可用于内核启动配置。
CUDA Fortran 扩展了 chevron 语法以接受 * 参数，该参数将在运行时（即在后台）调用这些函数。如果内核被多次启动，则此便利性可能不是期望的，因为它确实为每次调用都引入了一些开销。
显式或通过 * 语法使用占用率调用在启动 grid_global 内核时特别有用，因为启动参数必须调整大小以适合当前设备。
在本节中使用函数进行显式占用率计算。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的占用率部分。

4.8.1. cudaOccupancyMaxActiveBlocksPerMultiprocessor
integer function cudaOccupancyMaxActiveBlocksPerMultiprocessor( numBlocks, func, blockSize, dynamicSMemSize )
    integer :: numBlocks
    external :: func
    integer :: blockSize
    integer :: dynamicSMemSize


cudaOccupancyMaxActiveBlocksPerMultiprocessor 返回占用率，以每个多处理器块的数量表示，给定由 func 参数命名的全局子例程、内核预期启动的块大小（线程数）以及内核预期启动的动态共享内存量（以字节为单位）。


4.8.2. cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
integer function cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( numBlocks, func, blockSize, dynamicSMemSize, flags )
    integer :: numBlocks
    external :: func
    integer :: blockSize
    integer :: dynamicSMemSize
    integer :: flags


cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags 返回占用率，以每个多处理器块的数量表示，给定由 func 参数命名的全局子例程、内核预期启动的块大小（线程数）以及内核预期启动的动态共享内存量（以字节为单位），以及指定的标志。


4.8.3. cudaOccupancyMaxPotentialClusterSize
integer function cudaOccupancyMaxPotentialClusterSize( csize, func, config )
    integer :: csize
    external :: func
    type(cudaLaunchConfig) :: config


cudaOccupancyMaxPotentialClusterSize 返回可以启动的最大集群大小，给定输入内核 func 和 config 参数中指定的启动配置。


4.8.4. cudaOccupancyMaxActiveClusters
integer function cudaOccupancyMaxActiveClusters( maxc, func, config )
    integer :: maxc
    external :: func
    type(cudaLaunchConfig) :: config


cudaOccupancyMaxActiveClusters 返回目标设备上可以共存的最大集群数。集群大小可以在 config 参数中设置。



4.9. Memory Management
许多内存管理例程可以将设备数组作为参数。有些还可以将 C 类型作为参数，这些类型通过 Fortran 2003 iso_c_binding 模块提供，以简化与现有 CUDA C 代码的接口。
CUDA Fortran 通过提供在 cudafor 模块中定义的 C 设备指针 TYPE(C_DEVPTR) 扩展了 F2003 派生类型 TYPE(C_PTR)。一致地使用 TYPE(C_PTR) 和 TYPE(C_DEVPTR)，以及 Fortran 设备数组和宿主机数组之间的一致性检查应该是有益的。
目前，可以通过使用 iso_c_binding 子例程 c_f_pointer 的扩展，从 TYPE(C_DEVPTR) 构造 Fortran 设备数组。在 CUDA Fortran 下，c_f_pointer 将采用 TYPE(C_DEVPTR) 作为第一个参数，可分配设备数组作为第二个参数，形状作为第三个参数，并有效地将分配转移到 Fortran 数组。类似地，还定义了一个函数 C_DEVLOC()，它将创建一个 TYPE(C_DEVPTR)，该类型保存 Fortran 设备数组参数的 C 地址。当将来支持设备数据的正确 Fortran 指针时，这两个功能都可能会发生变化。
在本节中使用函数进行内存管理。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的内存管理部分。

4.9.1. cudaFree
integer function cudaFree(devptr)


cudaFree 释放设备上的数据。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何可分配设备数组。或者，devptr 可以是 TYPE(C_DEVPTR) 类型。
cudaFree 在 CUDA 5.0 及更高版本的设备代码中可用。


4.9.2. cudaFreeArray
integer function cudaFreeArray(carray)
    type(cudaArrayPtr) :: carray


cudaFreeArray 释放设备上分配的数组。


4.9.3. cudaFreeAsync
integer function cudaFreeAsync(devptr, stream)


cudaFreeAsync 在指定的流上异步释放设备上的数据。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何可分配设备数组。或者，devptr 可以是 TYPE(C_DEVPTR) 类型。stream 参数是 kind=cuda_stream_kind 的整数。
cudaFreeAsync 从 CUDA 11.2 开始可用。


4.9.4. cudaFreeHost
integer function cudaFreeHost(hostptr)
    type(C_PTR) :: hostptr


cudaFreeHost 释放使用 cudaMalloHost 在宿主机上分配的已锁定内存。


4.9.5. cudaGetSymbolAddress
integer function cudaGetSymbolAddress(devptr, symbol)
    type(C_DEVPTR) :: devptr
    type(c_ptr) :: symbol


cudaGetSymbolAddress 在 devptr 参数中返回设备上 symbol 的地址。symbol 可以通过字符串设置为外部设备名称。
以下代码序列从 CUDA C 内核初始化全局设备数组“vx”
type(c_ptr) :: csvx
type(c_devptr) :: cdvx
real, allocatable, device :: vx(:)
csvx = 'vx'
Istat = cudaGetSymbolAddress(cdvx, csvx)
Call c_f_pointer(cdvx, vx, 100)
Vx = 0.0




4.9.6. cudaGetSymbolSize
integer function cudaGetSymbolSize(size, symbol)
    integer :: size
    type(c_ptr) :: symbol


cudaGetSymbolSize 将变量 size 设置为 symbol 引用的全局或常量内存空间中设备区域的大小。


4.9.7. cudaHostAlloc
integer function cudaHostAlloc(hostptr, size, flags)
    type(C_PTR) :: hostptr
    integer :: size, flags


cudaHostAlloc 在宿主机上分配已锁定内存。它在 hostptr 中返回页面锁定分配的地址，如果内存不可用，则返回错误。Size 以字节为单位。flags 参数允许指定影响分配的不同选项。正常的 iso_c_binding 子例程 c_f_pointer 可用于将 type(c_ptr) 移动到 Fortran 指针。


4.9.8. cudaHostGetDevicePointer
integer function cudaHostGetDevicePointer(devptr, hostptr, flags)
    type(C_DEVPTR) :: devptr
    type(C_PTR) :: hostptr
    integer :: flags


cudaHostGetDevicePointer 返回指向与宿主机上的已锁定内存对应的设备内存地址的指针。hostptr 是通过 cudaHostAlloc() 分配的已锁定内存缓冲区。它在 devptr 中返回一个地址，该地址可以传递给设备上运行的内核并由其读取和写入。flags 参数是为将来的版本提供的。正常的 iso_c_binding 子例程 c_f_pointer 可用于将 type(c_devptr) 移动到设备数组。


4.9.9. cudaHostGetFlags
integer function cudaHostGetFlags(flags, hostptr)
    integer :: flags
    type(C_PTR) :: hostptr


cudaHostGetFlags 返回与宿主机指针关联的标志。


4.9.10. cudaHostRegister
integer function cudaHostRegister(hostptr, count, flags)
    integer :: flags
    type(C_PTR) :: hostptr


cudaHostRegister 根据 flags 参数，对与宿主机指针关联的内存进行页面锁定，其大小由 count 参数提供。


4.9.11. cudaHostUnregister
integer function cudaHostRegister(hostptr)
    type(C_PTR) :: hostptr


cudaHostUnregister 取消映射与宿主机指针关联的内存，并使其再次可分页。参数 hostptr 必须与 cudaHostRegister 一起使用的参数相同。


4.9.12. cudaMalloc
integer function cudaMalloc(devptr, count)


cudaMalloc 在设备上分配数据。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何可分配的一维设备数组。count 以元素为单位。或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。
cudaMalloc 在 CUDA 5.0 及更高版本的设备代码中可用。


4.9.13. cudaMallocArray
integer function cudaMallocArray(carray, cdesc, width, height)
    type(cudaArrayPtr) :: carray
    type(cudaChannelFormatDesc) :: cdesc
    integer :: width, height


cudaMallocArray 在设备上分配数据数组。


4.9.14. cudaMallocAsync
integer function cudaMallocAsync(devptr, count, stream)


cudaMallocAsync 在指定的流上异步分配设备上的数据。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何可分配的一维设备数组。count 以元素为单位。或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。stream 参数是 kind=cuda_stream_kind 的整数。
cudaFreeAsync 从 CUDA 11.2 开始可用。


4.9.15. cudaMallocManaged
integer function cudaMallocManaged(devptr, count, flags)


cudaMallocManaged 分配将由统一内存系统管理的数据。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何可分配的一维托管数组。count 以元素为单位。或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。flags 参数必须是 cudaMemAttachGlobal 或 cudaMemAttachHost。
cudaMallocManaged 从 CUDA 6.0 开始可用。


4.9.16. cudaMallocPitch
integer function cudaMallocPitch(devptr, pitch, width, height)


cudaMallocPitch 在设备上分配数据。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何可分配的二维设备数组。width 以元素数为单位。height 是一个整数。
cudaMallocPitch 可能会填充数据，填充后的宽度在变量 pitch 中返回。Pitch 是 kind=cuda_count_kind 的整数。devptr 也可以是 TYPE(C_DEVPTR) 类型，在这种情况下，整数值以字节表示。


4.9.17. cudaMalloc3D
integer function cudaMalloc3D(pitchptr, cext)
    type(cudaPitchedPtr), intent(out) :: pitchptr
    type(cudaExtent), intent(in) :: cext


cudaMalloc3D 在设备上分配数据。pitchptr 是在 cudafor 模块中定义的派生类型。cext 也是一个派生类型，它保存已分配数组的范围。或者，pitchptr 可以是 允许的数据类型 中指定的受支持类型的任何可分配的三维设备数组。


4.9.18. cudaMalloc3DArray
integer function cudaMalloc3DArray(carray, cdesc, cext)
    type(cudaArrayPtr) :: carray
    type(cudaChannelFormatDesc) :: cdesc
    type(cudaExtent) :: cext


cudaMalloc3DArray 在设备上分配数组数据。


4.9.19. cudaMemAdvise
integer function cudaMemAdvise(devptr, count, advice, device)


cudaMemAdvise 向统一内存子系统提供有关指定内存范围的预期使用模式的建议。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何托管内存标量或数组。count 以元素为单位。或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。
cudafor 模块中定义的 advice 当前可能的值为 cudaMemAdviseSetReadMostly、cudaMemAdviseUnsetReadMostly、cudaMemAdviseSetPreferredLocation、cudaMemAdviseUnsetPreferredLocation、cudaMemAdviseSetAccessedBy 和 cudaMemAdviseUnsetAccessedBy。
device 参数指定目标设备。为设备传入 cudaCpuDeviceId（在 cudafor 模块中定义为参数）将为 CPU 设置建议。


4.9.20. cudaMemcpy
integer function cudaMemcpy(dst, src, count, kdir)


cudaMemcpy 将数据从一个位置复制到另一个位置。dst 和 src 可以是 设备代码固有数据类型 中指定的受支持类型的任何设备或宿主机、标量或数组。count 以元素为单位。kdir 可能是可选的；有关更多信息，请参阅 使用运行时例程进行数据传输。如果指定了 kdir，则它必须是定义的枚举 cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost 或 cudaMemcpyDeviceToDevice 之一。或者，dst 和 src 可以是 TYPE(C_DEVPTR) 或 TYPE(C_PTR) 类型，在这种情况下，count 以字节为单位。
cudaMemcpy 在 CUDA 5.0 及更高版本的设备代码中可用。


4.9.21. cudaMemcpyArrayToArray
integer function cudaMemcpyArrayToArray(dsta, dstx, dsty,
                        srca, srcx, srcy, count, kdir)
    type(cudaArrayPtr) :: dsta, srca
    integer :: dstx, dsty, srcx, srcy, count, kdir


cudaMemcpyArrayToArray 将数组数据复制到设备和从设备复制数组数据。


4.9.22. cudaMemcpyAsync
integer function cudaMemcpyAsync(dst, src, count, kdir, stream)


cudaMemcpyAsync 将数据从一个位置复制到另一个位置。dst 和 src 可以是 设备代码固有数据类型 中指定的受支持类型的任何设备或宿主机、标量或数组。count 以元素为单位。kdir 可能是可选的；有关更多信息，请参阅 使用运行时例程进行数据传输。如果指定了 kdir，则它必须是定义的枚举 cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost 或 cudaMemcpyDeviceToDevice 之一。或者，dst 和 src 可以是 TYPE(C_DEVPTR) 或 TYPE(C_PTR) 类型，在这种情况下，count 以字节为单位。
此函数仅对页面锁定的宿主机内存进行操作。复制可以与流关联，方法是传递非零流参数；否则，stream 参数是可选的，默认为零。
cudaMemcpyAsync 在 CUDA 5.0 及更高版本的设备代码中可用。


4.9.23. cudaMemcpyFromArray
integer function cudaMemcpyFromArray(dst, srca, srcx, srcy, count, kdir)
    type(cudaArrayPtr) :: srca
    integer :: dstx, dsty, count, kdir


cudaMemcpyFromArray 将数组数据复制到设备和从设备复制数组数据。


4.9.24. cudaMemcpyFromSymbol
integer function cudaMemcpyFromSymbol(dst, symbol, count, offset, kdir, stream)
    type(c_ptr) :: symbol
    integer :: count, offset, kdir
    integer, optional :: stream


cudaMemcpyFromSymbol 将数据从 symbol 引用的全局或常量内存空间中的设备区域复制到宿主机上的目标位置。dst 可以是 允许的数据类型 中指定的受支持类型的任何宿主机标量或数组。count 以元素为单位。


4.9.25. cudaMemcpyFromSymbolAsync
integer function cudaMemcpyFromSymbolAsync(dst, symbol, count, offset, kdir, stream)
    type(c_ptr) :: symbol
    integer :: count, offset, kdir
    integer, optional :: stream


cudaMemcpyFromSymbolASYNC 将数据从 symbol 引用的全局或常量内存空间中的设备区域复制到宿主机上的目标位置。dst 可以是 允许的数据类型 中指定的受支持类型的任何宿主机标量或数组。count 以元素为单位。
cudaMemcpyFromSymbolASYNC 相对于宿主机是异步的。此函数仅对页面锁定的宿主机内存进行操作。复制可以与流关联，方法是传递非零流参数。


4.9.26. cudaMemcpyPeer
integer function cudaMemcpyPeer(dst, dstdev, src, srcdev, count)


cudaMemcpyPeer 将数据从一个设备复制到另一个设备。dst 和 src 可以是 设备代码固有数据类型 中指定的受支持类型的任何设备标量或数组。count 以元素为单位。或者，dst 和 src 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。


4.9.27. cudaMemcpyPeerAsync
integer function cudaMemcpyPeerAsync(dst, dstdev, src, srcdev, count, stream)


cudaMemcpyPeerAsync 将数据从一个设备复制到另一个设备。dst 和 src 可以是 设备代码固有数据类型 中指定的受支持类型的任何设备标量或数组。count 以元素为单位。或者，dst 和 src 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。复制可以与流关联，方法是传递非零流参数。


4.9.28. cudaMemcpyToArray
integer function cudaMemcpyToArray(dsta, dstx, dsty, src, count, kdir)
    type(cudaArrayPtr) :: dsta
    integer :: dstx, dsty, count, kdir


cudaMemcpyToArray 将数组数据复制到设备和从设备复制数组数据。


4.9.29. cudaMemcpyToSymbol
integer function cudaMemcpyToSymbol(symbol, src, count, offset, kdir)
   type(c_ptr) :: symbol
   integer :: count, offset, kdir


cudaMemcpyToSymbol 将数据从源复制到 symbol 引用的全局或常量内存空间中的设备区域。src 可以是 设备代码固有数据类型 中指定的受支持类型的任何宿主机标量或数组。count 以元素为单位。


4.9.30. cudaMemcpyToSymbolAsync
integer function cudaMemcpyToSymbolAsync(symbol, src, count, offset, kdir, stream)
    type(c_ptr) :: symbol
    integer :: count, offset, kdir
    integer, optional :: stream


cudaMemcpyToSymbolAsync 将数据从源复制到 symbol 引用的全局或常量内存空间中的设备区域。src 可以是 允许的数据类型 中指定的受支持类型的任何宿主机标量或数组。count 以元素为单位。
此函数仅对页面锁定的宿主机内存进行操作。复制可以与流关联，方法是传递非零流参数。


4.9.31. cudaMemcpy2D
integer function cudaMemcpy2D(dst, dpitch, src, spitch, width, height, kdir)


cudaMemcpy2D 将数据从一个位置复制到另一个位置。dst 和 src 可以是 设备代码固有数据类型 中指定的受支持类型的任何设备或宿主机数组。width 和 height 以元素为单位。与 Fortran 程序员可能认为的内存布局相反，并且为了保持与 CUDA C 的兼容性，宽度指定前导维度中连续元素的数量，高度是此类连续部分的数量。kdir 可能是可选的；有关更多信息，请参阅 使用运行时例程进行数据传输。如果指定了 kdir，则它必须是定义的枚举 cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost 或 cudaMemcpyDeviceToDevice 之一。或者，dst 和 src 可以是 TYPE(C_DEVPTR) 或 TYPE(C_PTR) 类型，在这种情况下，width 和 height 以字节为单位。
cudaMemcpy2D 在 CUDA 5.0 及更高版本的设备代码中可用。


4.9.32. cudaMemcpy2DArrayToArray
integer function cudaMemcpy2DArrayToArray(dsta, dstx, dsty,
                     srca, srcx, srcy, width, height, kdir)
    type(cudaArrayPtr) :: dsta, srca
    integer :: dstx, dsty, srcx, srcy, width, height, kdir


cudaMemcpy2DArrayToArray 将数组数据复制到设备和从设备复制数组数据。


4.9.33. cudaMemcpy2DAsync
integer function cudaMemcpy2DAsync(dst, dpitch, src, spitch, width,
              height, kdir, stream)


cudaMemcpy2D 将数据从一个位置复制到另一个位置。dst 和 src 可以是 设备代码固有数据类型 中指定的受支持类型的任何设备或宿主机数组。width 和 height 以元素为单位。与 Fortran 程序员可能认为的内存布局相反，并且为了保持与 CUDA C 的兼容性，宽度指定前导维度中连续元素的数量，高度是此类连续部分的数量。kdir 可能是可选的；有关更多信息，请参阅 使用运行时例程进行数据传输。如果指定了 kdir，则它必须是定义的枚举 cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost 或 cudaMemcpyDeviceToDevice 之一。或者，dst 和 src 可以是 TYPE(C_DEVPTR) 或 TYPE(C_PTR) 类型，在这种情况下，width 和 height 以字节为单位。
此函数仅对页面锁定的宿主机内存进行操作。复制可以与流关联，方法是传递非零 stream 参数，否则，stream 参数是可选的，默认为零。
cudaMemcpy2DAsync 在 CUDA 5.0 及更高版本的设备代码中可用。


4.9.34. cudaMemcpy2DFromArray
integer function cudaMemcpy2DFromArray(dst, dpitch, srca, srcx, srcy,
                               width, height, kdir)
    type(cudaArrayPtr) :: srca
    integer :: dpitch, srcx, srcy, width, height, kdir


cudaMemcpy2DFromArray 将数组数据复制到设备和从设备复制数组数据。


4.9.35. cudaMemcpy2DToArray
integer function cudaMemcpy2DToArray(dsta, dstx, dsty, src,
                     spitch, width, height, kdir)
    type(cudaArrayPtr) :: dsta
    integer :: dstx, dsty, spitch, width, height, kdir


cudaMemcpy2DToArray 将数组数据复制到设备和从设备复制数组数据。


4.9.36. cudaMemcpy3D
integer function cudaMemcpy3D(p)
    type(cudaMemcpy3DParms) :: p


cudaMemcpy3D 按照派生类型 p 中保存的数据的指定，将元素从一个 3D 数组复制到另一个 3D 数组。


4.9.37. cudaMemcpy3DAsync
integer function cudaMemcpy3D(p, stream)
    type(cudaMemcpy3DParms) :: p
    integer :: stream


cudaMemcpy3DAsync 按照派生类型 p 中保存的数据的指定，将元素从一个 3D 数组复制到另一个 3D 数组。
此函数仅对页面锁定的宿主机内存进行操作。复制可以与流关联，方法是传递非零 stream 参数。


4.9.38. cudaMemGetInfo
integer function cudaMemGetInfo( free, total )
    integer(kind=cuda_count_kind) :: free, total


cudaMemGetInfo 返回设备上可用于分配的可用和总内存量。返回的值单位为字节。


4.9.39. cudaMemPrefetchAsync
integer function cudaMemPrefetchAsync(devptr, count, device, stream)


cudaMemPrefetchAsync 将内存预取到指定的目的设备。devptr 可以是 设备代码固有数据类型 中指定的受支持类型的任何托管内存标量或数组。count 以元素为单位。或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位。
device 参数指定目标设备。 stream 参数指定将预取操作排队到哪个流。
为设备传入在 cudafor 模块中定义为参数的 cudaCpuDeviceId，会将数据预取到 CPU 内存。


4.9.40. cudaMemset
integer function cudaMemset(devptr, value, count)


cudaMemset 将某个位置或数组设置为指定值。 devptr 可以是设备标量或设备代码固有数据类型中指定的支持类型的数组，请参阅 设备代码固有数据类型。 value 的类型和种类必须匹配。 count 以元素为单位。 或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位，并且使用 value 的最低字节。


4.9.41. cudaMemsetAsync
integer function cudaMemsetAsync(devptr, value, count, stream)


cudaMemsetAsync 将某个位置或数组设置为指定值。 devptr 可以是设备标量或设备代码固有数据类型中指定的支持类型的数组，请参阅 设备代码固有数据类型。 value 的类型和种类必须匹配。 count 以元素为单位。 或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，count 以字节为单位，并且使用 value 的最低字节。 内存设置操作与指定的流相关联。


4.9.42. cudaMemset2D
integer function cudaMemset2D(devptr, pitch, value, width, height)


cudaMemset2D 将数组设置为指定值。 devptr 可以是设备代码固有数据类型中指定的支持类型的设备数组，请参阅 设备代码固有数据类型。 value 的类型和种类必须匹配。 pitch、width 和 height 以元素为单位。 或者，devptr 可以是 TYPE(C_DEVPTR) 类型，在这种情况下，pitch、width 和 height 以字节为单位，并且使用 value 的最低字节。 与 Fortran 程序员可能认为的内存布局相反，为了保持与 CUDA C 的兼容性，width 指定前导维度中连续元素的数量，height 是此类连续部分的数量。


4.9.43. cudaMemset3D
integer function cudaMemset3D(pitchptr, value, cext)
    type(cudaPitchedPtr) :: pitchptr
    integer :: value
    type(cudaExtent) :: cext


cudaMemset3D 将数组的元素设置为指定值，每个维度的范围由 cext 指定，该数组是使用 cudaMalloc3D 分配的。



4.10. 统一寻址和对等设备内存访问
使用本节中的函数从同一进程和线程管理多个设备。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的统一寻址和对等设备内存访问部分。

4.10.1. cudaDeviceCanAccessPeer
integer function cudaDeviceCanAccessPeer( canAccessPeer, device, peerDevice )
    integer :: canAccessPeer, device, peerDevice


如果 device 参数可以访问 peerDevice 参数指定的设备中的内存，则 cudaDeviceCanAccessPeer 在 canAccessPeer 中返回值 1。


4.10.2. cudaDeviceDisablePeerAccess
integer function cudaDeviceDisablePeerAccess ( peerDevice )
    integer :: peerDevice


cudaDeviceDisablePeerAccess 禁用当前设备访问 peerDevice 参数指定的设备上的内存的能力。


4.10.3. cudaDeviceEnablePeerAccess
integer function cudaDeviceEnablePeerAccess ( peerDevice, flags )
    integer :: peerDevice, flags


cudaDeviceEnablePeerAccess 启用当前设备访问 peerDevice 参数指定的设备上的内存的能力。 目前，标志必须为零。


4.10.4. cudaPointerGetAttributes
integer function cudaPointerGetAttributes( attr, ptr )
    type(cudaPointerAttributes), intent(out) :: attr


cudaPointerGetAttributes 在属性类型中返回设备或主机指针的属性。 ptr 可以是设备代码中指定的支持类型的任何主机或设备标量或数组，请参阅 允许的数据类型。 它也可以是 C_PTR 或 C_DEVPTR 类型。 它可能具有主机、设备、托管或固定属性。



4.11. 版本管理
使用本节中的函数进行版本管理。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的版本管理部分。

4.11.1. cudaDriverGetVersion
integer function cudaDriverGetVersion(iversion)
    integer :: iversion


cudaDriverGetVersion 以 iversion 返回已安装的 CUDA 驱动程序的版本号。 如果未安装驱动程序，则它以 iversion 返回 0。
如果 iversion 参数为 NULL，则此函数自动返回 cudaErrorInvalidValue。


4.11.2. cudaRuntimeGetVersion
integer function cudaRuntimeGetVersion(iversion)
    integer :: iversion


cudaRuntimeGetVersion 以 iversion 返回已安装的 CUDA 运行时的版本号。
如果 iversion 参数为 NULL，则此函数自动返回 cudaErrorInvalidValue。



4.12. 性能分析管理
使用本节中的函数进行性能分析管理。
有关此处列出的每个例程的用途和功能的完整说明，请参阅 https://docs.nvda.net.cn/cuda/cuda-runtime-api 上的 Profiler Control 部分。

4.12.1. cudaProfilerStart
integer function cudaProfilerStart()


cudaProfilerStart 启用活动性能分析工具的性能分析收集。


4.12.2. cudaProfilerStop
integer function cudaProfilerStop()


cudaProfilerStop 禁用活动性能分析工具的性能分析收集。



4.13. CUDA 图形管理
使用本节中的函数进行 CUDA 图形管理，从 CUDA Fortran 捕获和重放 CUDA 图形。
有关此处列出的每个例程的用途和功能的完整说明，请参阅此处的 CUDA 图形部分 https://docs.nvda.net.cn/cuda/cuda-c-programming-guide。
CUDA Fortran 提供了三种用于管理 cuda 图形的类型
TYPE cudaGraph
  TYPE(C_PTR)  :: graph
END TYPE cudaGraph

TYPE cudaGraphExec
  TYPE(C_PTR)  :: exec
END TYPE cudaGraphExec

TYPE cudaGraphNode
  TYPE(C_PTR)  :: node
END TYPE cudaGraphNode



4.13.1. cudaGraphCreate
integer function cudaGraphCreate( graph, flags )
    type(cudagraph), intent(out) :: graph
    integer :: flags


cudaGraphCreate 创建一个空的 graph。


4.13.2. cudaGraphDestroy
integer function cudaGraphDestroy( graph )
    type(cudagraph) :: graph


cudaGraphDestroy 释放与给定 graph 关联的任何资源。


4.13.3. cudaGraphExecDestroy
integer function cudaGraphExecDestroy( graphExec )
    type(cudagraphexec) :: graphExec


cudaGraphExecDestroy 释放与给定 graphExec 关联的任何资源。


4.13.4. cudaGraphInstantiate
integer function cudaGraphInstantiate( graphExec, graph, flags )
    type(cudagraphexec), intent(out) :: graphExec
    type(cudagraph), intent(in) :: graph
    integer, intent(in) :: flags


cudaGraphInstantiate 从指定的 graph 实例化 graphExec 对象。


4.13.5. cudaGraphLaunch
integer function cudaGraphLaunch( graphExec, stream )
    type(cudagraphexec) :: graphExec
    integer(kind=cuda_stream_kind), intent(in) :: stream


cudaGraphLaunch 在指定的 stream 上开始异步图形启动或重放。


4.13.6. cudaStreamBeginCapture
integer function cudaStreamBeginCapture( stream, mode )
    integer(kind=cuda_stream_kind), intent(in) :: stream
    integer, intent(in) :: mode


cudaStreamBeginCapture 在指定的 stream 上开始图形捕获。


4.13.7. cudaStreamEndCapture
integer function cudaStreamEndCapture( stream, graph )
    integer(kind=cuda_stream_kind), intent(in) :: stream
    type(cudagraph), intent(out) :: graph


cudaStreamEndCapture 结束在指定的 stream 上的图形捕获，并提供一个 cudagraph 以供进一步使用。


4.13.8. cudaStreamIsCapturing
integer function cudaStreamIsCapturing( stream, status )
    integer(kind=cuda_stream_kind), intent(in) :: stream
    integer, intent(out) :: status


cudaStreamIsCapturing 查询指定 stream 中的捕获 status。




5. 示例
本节包含带源代码的示例。

5.1. 矩阵乘法示例
此示例演示了一个计算两个矩阵 A 和 B 的乘积 C 的程序，如下所示

每个线程块计算 C 的一个 16x16 子矩阵；
块内的每个线程计算子矩阵的一个元素。

选择子矩阵大小，以便块中线程的数量是 warp 大小 (32) 的倍数，并且小于每个线程块的最大线程数 (512)。
结果的每个元素都是 A 的一行与 B 的一列的乘积。该程序通过累积子矩阵乘积来计算乘积；它读取 A 的块子矩阵和 B 的块子矩阵，累积子矩阵乘积，然后按行移动到 A 的下一个子矩阵，并按列移动到 B 的下一个子矩阵。该程序将 A 和 B 的子矩阵缓存在快速共享内存中。
为简单起见，该程序假定矩阵大小是 16 的倍数，并且尚未针对执行时间进行高度优化。

5.1.1. 源代码列表
矩阵乘法
! start the module containing the matmul kernel
module mmul_mod
 use cudafor
contains
 ! mmul_kernel computes A*B into C where
 ! A is NxM, B is MxL, C is then NxL
 attributes(global) subroutine mmul_kernel( A, B, C, N, M, L )
  real :: A(N,M), B(M,L), C(N,L)
  integer, value :: N, M, L
  integer :: i, j, kb, k, tx, ty
  ! submatrices stored in shared memory
  real, shared :: Asub(16,16), Bsub(16,16)
  ! the value of C(i,j) being computed
  real :: Cij
  ! Get the thread indices
  tx = threadidx%x
  ty = threadidx%y


! This thread computes C(i,j) = sum(A(i,:) * B(:,j))
i = (blockidx%x-1) * 16 + tx
j = (blockidx%y-1) * 16 + ty
Cij = 0.0
! Do the k loop in chunks of 16, the block size
do kb = 1, M, 16
  ! Fill the submatrices
  ! Each of the 16x16 threads in the thread block
  ! loads one element of Asub and Bsub
   Asub(tx,ty) = A(i,kb+ty-1)
  Bsub(tx,ty) = B(kb+tx-1,j)
  ! Wait until all elements are filled
  call syncthreads()
  ! Multiply the two submatrices
  ! Each of the 16x16 threads accumulates the
  ! dot product for its element of C(i,j)
  do k = 1,16
    Cij = Cij + Asub(tx,k) * Bsub(k,ty)
  enddo
  ! Synchronize to make sure all threads are done
  ! reading the submatrices before overwriting them
  ! in the next iteration of the kb loop
  call syncthreads()
enddo
! Each of the 16x16 threads stores its element
! to the global C array
C(i,j) = Cij
end subroutine mmul_kernel


! The host routine to drive the matrix multiplication
 subroutine mmul( A, B, C )
  real, dimension(:,:) :: A, B, C
  ! allocatable device arrays
  real, device, allocatable, dimension(:,:) :: Adev,Bdev,Cdev
  ! dim3 variables to define the grid and block shapes
  type(dim3) :: dimGrid, dimBlock


! Get the array sizes
N = size( A, 1 )
M = size( A, 2 )
L = size( B, 2 )
! Allocate the device arrays
allocate( Adev(N,M), Bdev(M,L), Cdev(N,L) )



! Copy A and B to the device
Adev = A(1:N,1:M)
Bdev(:,:) = B(1:M,1:L)


! Create the grid and block dimensions
dimGrid = dim3( N/16, L/16, 1 )
dimBlock = dim3( 16, 16, 1 )
call mmul_kernel<<<dimGrid,dimBlock>>>( Adev, Bdev, Cdev, N, M, L)


 ! Copy the results back and free up memory
 C(1:N,1:L) = Cdev
 deallocate( Adev, Bdev, Cdev )
end subroutine mmul
end module mmul_mod




5.1.2. 源代码描述
此源代码模块 mmul_mod 有两个子例程。 主机子例程 mmul 是内核例程 mmul_kernel 的包装器。
MMUL
此主机子例程有两个输入数组 A 和 B，以及一个输出数组 C，作为假定形状数组传递。 该例程执行以下操作

它确定 N、M 和 L 中矩阵的大小。
它分配设备内存数组 Adev、Bdev 和 Cdev。
它使用数组赋值将数组 A 和 B 复制到 Adev 和 Bdev。
它填充 dimGrid 和 dimBlock 以保存网格和线程块大小。
它调用 mmul_kernel 以在设备上计算 Cdev。
它将 Cdev 从设备内存复制回 C。
它释放设备内存数组。

由于数据复制操作是同步的，因此在复制操作和内核启动之间不需要额外的同步。
MMUL_KERNEL
此内核子例程具有两个设备内存输入数组 A 和 B，一个设备内存输出数组 C，以及三个标量，用于给出数组大小。 执行此例程的线程是线程块中协作的 16x16 个线程之一。 此例程计算特定 i 和 j 值的 A(i,:)*B(:,j) 的点积，具体取决于块和线程索引。
它执行以下操作

它确定此线程的线程索引。
它确定 i 和 j 索引，它正在计算 ``C(i,j)`` 的哪个元素。
它初始化一个标量，它将在其中累积点积。
它以大小为 16 的块逐步遍历数组 A 和 B。
对于每个块，它执行以下步骤

它将 A 和 B 的子矩阵的一个元素加载到共享内存中。
它同步以确保块中的所有线程都加载了两个子矩阵。
它累积其子矩阵的行和列的点积。
它再次同步以确保所有线程在开始下一个块之前都已完成读取子矩阵。


最后，它将计算出的值存储到 C 的正确元素中。




5.2. 映射内存示例
此示例演示了 cudafor 模块中支持的 CUDA API 的使用，用于将页面锁定主机内存映射到设备的地址空间中。 它利用 iso_c_binding c_ptr 类型和 cudafor c_devptr 类型来与 C 例程接口，然后 Fortran c_f_pointer 调用将类型映射到 Fortran 数组。
映射内存
module atest
  contains
    attributes(global) subroutine matrixinc(a,n)
      real, device :: a(n,n)
      integer, value :: n
      i = (blockidx%x-1)*10 + threadidx%x
      j= (blockidx%y-1)*10 + threadidx%y
      if ((i .le. n) .and. (j .le. n)) then
        a(i,j) = a(i,j) + 1.0
      endif
      return
    end subroutine
end module

program test
use  cudafor
use  atest
use, intrinsic :: iso_c_binding

type(c_ptr) :: a
type(c_devptr) :: a_d
real, dimension(:,:), pointer :: fa
real, dimension(:,:), allocatable, device :: fa_d
type(dim3) :: blcks, thrds

istat= cudaSetDeviceFlags(cudadevicemaphost)

istat = cudaHostAlloc(a,100*100*sizeof(1.0),cudaHostAllocMapped)

! can move the c_ptr to an f90 pointer
call c_f_pointer(a, fa, (/ 100, 100 /) )

! update the data on the host
do j = 1, 100
  do i = 1, 100
    fa(i,j)= real(i) + j*;100.0
  end do
end do

! get a device pointer to the same array
istat= cudaHostGetDevicePointer(a_d, a, 0)

! can move the c_devptr to an device allocatable array
call c_f_pointer(a_d, fa_d, (/ 100, 100 /) )
!
blcks = dim3(10,10,1)
thrds = dim3(10,10,1)
!
call matrixinc <<<blcks, thrds>>>(fa_d, 100)

! need to synchronize
istat = cudaDeviceSynchronize()
!
do j = 1, 100
  do i = 1, 100
    if (fa(i,j) .ne. (real(i) + j*100.0 + 1.0)) print *,"failure",i,j
  end do
end do
!
istat = cudaFreeHost(a)
end




5.3. Cublas 模块示例
此示例演示了 cublas 模块、cublasHandle 类型、cublas 调用的三种形式以及映射的固定内存的使用，所有这些都在多线程 OpenMP 程序的框架内。
Cublas 模块
program tdot
! Compile with "nvfortran -mp tdot.cuf -cudalib=cublas -lblas
! Set OMP_NUM_THREADS environment variable to run with
! up to 2 threads, currently.
!
use cublas
use cudafor
use omp_lib
!
integer, parameter :: N = 10000
real*8 x(N), y(N), z
real*8, device, allocatable :: xd0(:), yd0(:)
real*8, device, allocatable :: xd1(:), yd1(:)
real*8, allocatable :: zh(:)
real*8, allocatable, device :: zd(:)
integer, allocatable :: istats(:), offs(:)
real*8 reslt(3)
type(C_DEVPTR) :: zdptr
type(cublasHandle) :: h

! Max at 2 threads for now
nthr = omp_get_max_threads()
if (nthr .gt. 2) nthr = 2
call omp_set_num_threads(nthr)
! Run on host
call random_number(x)
call random_number(y)
z = ddot(N,x,1,y,1)
print *,"HostSerial",z


! Create a pinned memory spot
!$omp PARALLEL private(i,istat)
    i = omp_get_thread_num()
    istat = cudaSetDeviceFlags(cudaDeviceMapHost)
    istat = cudaSetDevice(i)
!$omp end parallel
allocate(zh(512),align=4096)
zh = 0.0d0
istat = cudaHostRegister(C_LOC(zh(1)), 4096, cudaHostRegisterMapped)
istat = cudaHostGetDevicePointer(zdptr, C_LOC(zh(1)), 0)
call c_f_pointer(zdptr, zd, 512 )


! CUDA data allocation, run on one card, blas interface
allocate(xd0(N),yd0(N))
xd0 = x
yd0 = y
z = ddot(N,xd0,1,yd0,1)
ii = 1
reslt(ii) = z
ii = ii + 1
deallocate(xd0)
deallocate(yd0)


! Break up the array into sections
nsec = N / nthr
allocate(istats(nthr),offs(nthr))
offs = (/ (i*nsec,i=0,nthr-1) /)

! Allocate and initialize the arrays
!$omp PARALLEL private(i,istat)
    i = omp_get_thread_num() + 1
    if (i .eq. 1) then
        allocate(xd0(nsec), yd0(nsec))
        xd0 = x(offs(i)+1:offs(i)+nsec)
        yd0 = y(offs(i)+1:offs(i)+nsec)
    else
        allocate(xd1(nsec), yd1(nsec))
        xd1 = x(offs(i)+1:offs(i)+nsec)
        yd1 = y(offs(i)+1:offs(i)+nsec)
    endif
!$omp end parallel


! Run the blas kernel using cublas name
!$omp PARALLEL private(i,istat,z)
    i = omp_get_thread_num() + 1
    if (i .eq. 1) then
        z = cublasDdot(nsec,xd0,1,yd0,1)
    else
        z = cublasDdot(nsec,xd1,1,yd1,1)
    endif
    zh(i) = z
!$omp end parallel


z = zh(1) + zh(2)
reslt(ii) = z
ii = ii + 1

zh = 0.0d0


! Now write to our pinned area with the v2 blas
!$omp PARALLEL private(h,i,istat)
    i = omp_get_thread_num() + 1
    h = cublasGetHandle()
    istat = cublasSetPointerMode(h, CUBLAS_POINTER_MODE_DEVICE)
    if (i .eq. 1) then
        istats(i) = cublasDdot_v2(h, nsec, xd0, 1, yd0, 1, zd(1))
    else
        istats(i) = cublasDdot_v2(h, nsec, xd1, 1, yd1, 1, zd(2))
    endif
    istat = cublasSetPointerMode(h, CUBLAS_POINTER_MODE_HOST)
    istat = cudaDeviceSynchronize()
!$omp end parallel


z = zh(1) + zh(2)
reslt(ii) = z

print *,"Device, 3 ways:",reslt

! Deallocate the arrays
!$omp PARALLEL private(i)
    i = omp_get_thread_num() + 1
    if (i .eq. 1) then
        deallocate(xd0,yd0)
    else
        deallocate(xd1,yd1)
    endif
!$omp end parallel
deallocate(istats,offs)

end




5.4. CUDA 设备属性示例
此示例演示了如何从 CUDA Fortran 访问设备属性。
CUDA 设备属性
! An example of getting device properties in CUDA Fortran
! Build with
!   nvfortran cufinfo.cuf
!
program cufinfo
use cudafor
integer istat, num, numdevices
type(cudadeviceprop) :: prop
istat = cudaGetDeviceCount(numdevices)
do num = 0, numdevices-1
  istat = cudaGetDeviceProperties(prop, num)
  call printDeviceProperties(prop, num)
end do
end
!
subroutine printDeviceProperties(prop, num)
use cudafor
type(cudadeviceprop) :: prop
integer num
ilen = verify(prop%name, ' ', .true.)
write (*,900) "Device Number: "      ,num
write (*,901) "Device Name: "        ,prop%name(1:ilen)
write (*,903) "Total Global Memory: ",real(prop%totalGlobalMem)/1e9," Gbytes"
write (*,902) "sharedMemPerBlock: "  ,prop%sharedMemPerBlock," bytes"
write (*,900) "regsPerBlock: "       ,prop%regsPerBlock
write (*,900) "warpSize: "           ,prop%warpSize
write (*,900) "maxThreadsPerBlock: " ,prop%maxThreadsPerBlock
write (*,904) "maxThreadsDim: "      ,prop%maxThreadsDim
write (*,904) "maxGridSize: "        ,prop%maxGridSize
write (*,903) "ClockRate: "          ,real(prop%clockRate)/1e6," GHz"
write (*,902) "Total Const Memory: " ,prop%totalConstMem," bytes"
write (*,905) "Compute Capability Revision: ",prop%major,prop%minor
write (*,902) "TextureAlignment: "   ,prop%textureAlignment," bytes"
write (*,906) "deviceOverlap: "      ,prop%deviceOverlap
write (*,900) "multiProcessorCount: ",prop%multiProcessorCount
write (*,906) "integrated: "         ,prop%integrated
write (*,906) "canMapHostMemory: "   ,prop%canMapHostMemory
write (*,906) "ECCEnabled: "         ,prop%ECCEnabled
write (*,906) "UnifiedAddressing: "  ,prop%unifiedAddressing
write (*,900) "L2 Cache Size: "      ,prop%l2CacheSize
write (*,900) "maxThreadsPerSMP: "   ,prop%maxThreadsPerMultiProcessor
900 format (a,i0)
901 format (a,a)
902 format (a,i0,a)
903 format (a,f5.3,a)
904 format (a,2(i0,1x,'x',1x),i0)
905 format (a,i0,'.',i0)
906 format (a,l0)
return
end




5.5. CUDA 异步内存传输示例
此示例演示了如何使用 CUDA Fortran 中的 CUDA API 执行与设备的异步复制。
CUDA 异步内存传输
! This code demonstrates strategies hiding data transfers via
! asynchronous data copies in multiple streams

module kernels_m
contains
  attributes(global) subroutine kernel(a, offset)
    implicit none
    real :: a(*)
    integer, value :: offset
    integer :: i
    real :: c, s, x
    i = offset + threadIdx%x + (blockIdx%x-1)*blockDim%x
    x = threadIdx%x + (blockIdx%x-1)*blockDim%x
    s = sin(x); c = cos(x)
    a(i) = a(i) + sqrt(s**2+c**2)
  end subroutine kernel
end module kernels_m

program testAsync
  use cudafor
  use kernels_m
  implicit none
  integer, parameter :: blockSize = 256, nStreams = 8
  integer, parameter :: n = 16*1024*blockSize*nStreams
  real, pinned, allocatable :: a(:)
  real, device :: a_d(n)
  integer(kind=cuda_Stream_Kind) :: stream(nStreams)
  type (cudaEvent) :: startEvent, stopEvent, dummyEvent
  real :: time
  integer :: i, istat, offset, streamSize = n/nStreams
  logical :: pinnedFlag
  type (cudaDeviceProp) :: prop

  istat = cudaGetDeviceProperties(prop, 0)
  write(*,"(' Device: ', a,/)") trim(prop%name)

  ! allocate pinned  host memory
  allocate(a(n), STAT=istat, PINNED=pinnedFlag)
  if (istat /= 0) then
     write(*,*) 'Allocation of a failed'
     stop
  else
     if (.not. pinnedFlag) write(*,*) 'Pinned allocation failed'
     end if

  ! create events and streams
  istat = cudaEventCreate(startEvent)
  istat = cudaEventCreate(stopEvent)
  istat = cudaEventCreate(dummyEvent)
  do i = 1, nStreams
     istat = cudaStreamCreate(stream(i))
  enddo


! baseline case - sequential transfer and execute
a = 0
istat = cudaEventRecord(startEvent,0)

a_d = a
call kernel<<<n/blockSize, blockSize>>>(a_d, 0)
a = a_d
istat = cudaEventRecord(stopEvent, 0)
istat = cudaEventSynchronize(stopEvent)
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
write(*,*) 'Time for sequential transfer and execute (ms): ', time
write(*,*) '  max error: ', maxval(abs(a-1.0))

! asynchronous version 1: loop over {copy, kernel, copy}
a = 0
istat = cudaEventRecord(startEvent,0)

do i = 1, nStreams
   offset = (i-1)*streamSize
   istat = cudaMemcpyAsync(a_d(offset+1),a(offset+1),streamSize,stream(i))
   call kernel<<<streamSize/blockSize, blockSize, &
                 0, stream(i)>>>(a_d,offset)
   istat = cudaMemcpyAsync(a(offset+1),a_d(offset+1),streamSize,stream(i))
enddo
istat = cudaEventRecord(stopEvent, 0)
istat = cudaEventSynchronize(stopEvent)
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
write(*,*) 'Time for asynchronous V1 transfer and execute (ms): ', time
write(*,*) '  max error: ', maxval(abs(a-1.0))

! asynchronous version 2:
! loop over copy, loop over kernel, loop over copy
a = 0
istat = cudaEventRecord(startEvent,0)
do i = 1, nStreams
   offset = (i-1)*streamSize
   istat = cudaMemcpyAsync(a_d(offset+1),a(offset+1),streamSize,stream(i))
enddo
do i = 1, nStreams
   offset = (i-1)*streamSize
   call kernel<<<streamSize/blockSize, blockSize, &
                 0, stream(i)>>>(a_d,offset)
enddo
do i = 1, nStreams
   offset = (i-1)*streamSize
   istat = cudaMemcpyAsync(a(offset+1),a_d(offset+1),streamSize,stream(i))
enddo
istat = cudaEventRecord(stopEvent, 0)
istat = cudaEventSynchronize(stopEvent)
istat = cudaEventElapsedTime(time, startEvent, stopEvent)
write(*,*) 'Time for asynchronous V2 transfer and execute (ms): ', time
write(*,*) '  max error: ', maxval(abs(a-1.0))

! cleanup
istat = cudaEventDestroy(startEvent)
istat = cudaEventDestroy(stopEvent)
istat = cudaEventDestroy(dummyEvent)

do i = 1, nStreams
   istat = cudaStreamDestroy(stream(i))
enddo
deallocate(a)

end program testAsync




5.6. 托管内存示例
此示例演示了在 OpenMP 程序中使用 CUDA 托管内存。 在主程序中，为每个 OpenMP 线程创建一个流。 调用 cudaforSetDefaultStream 以将其设置为所有后续高级语言构造的默认流。 默认流在 CUF 内核的启动配置中显式使用，也用作线程的输入参数以进行同步。 一旦发生 cudaStreamSynchronize，即使其他线程可能正在内核启动的中间，此线程也可以安全地访问主机上的托管数据，在本例中是在 any() 函数中。
CUDA Fortran 中的托管内存和 OpenMP
program ompcuf
use cudafor
use omp_lib
integer(kind=cuda_stream_kind) :: mystream

!$omp parallel private(istat,mystream)
istat = cudaStreamCreate(mystream)
istat = cudaforSetDefaultstream(mystream)
call ompworker()
!$omp end parallel
end

subroutine ompworker()
use cudafor
use omp_lib
real, managed :: a(10000)
j = omp_get_thread_num()
a = real(j)

!$cuf kernel do <<< *, *, stream=cudaforGetDefaultStream() >>>
do i = 1, 10000
  a(i) = a(i) + 1.0
end do
istat = cudaStreamSynchronize(cudaforGetDefaultStream())

if (any(a.ne.real(j+1))) then
  print *,"Found error on ",j
else
  print *,"Looks good on ",j
endif
end




5.7. WMMA Tensor Core 示例
此示例演示了使用 NVIDIA Volta tensor core 执行 real(2) 矩阵乘法。 结果是一个 real(4) 矩阵。 此示例利用了 cuf_macros.CUF 中的定义，该文件在 NVIDIA 包的 examples 目录中提供。 Fortran tensor core 编程中当前使用的实际派生类型可能会在以后发生更改，但始终支持这些宏。 该程序展示了 Fortran real(2) 数据类型在主机和设备代码中的使用。 更多示例，重点介绍了采用 WMMASubMatrix 类型的重载设备函数，以及使用 real(2) 数据向量以提高性能的函数，可以在 NVIDIA 包的 examples 目录中找到。
CUDA Fortran 中的 Tensor Core 编程
#include "cuf_macros.CUF"

module params
  integer, parameter :: m = 16
  integer, parameter :: n = 16
  integer, parameter :: k = 16
end module

module mod1
use params  ! Define matrix m, n, k
contains
  attributes(global) subroutine test1(a,b,c)
  use wmma
  real(2), device :: a(m,k)
  real(2), device :: b(k,n)
  real(4), device :: c(m,n)
  WMMASubMatrix(WMMAMatrixA, 16, 16, 16, Real, WMMAColMajor) :: sa
  WMMASubMatrix(WMMAMatrixB, 16, 16, 16, Real, WMMAColMajor) :: sb
  WMMASubMatrix(WMMAMatrixC, 16, 16, 16, Real, WMMAKind4)    :: sc
  sc = 0.0
  call wmmaLoadMatrix(sa, a(1,1), m)
  call wmmaLoadMatrix(sb, b(1,1), k)
  call wmmaMatmul(sc, sa, sb)
  call wmmaStoreMatrix(c(1,1), sc, m)
  end subroutine
end module

program main
use cudafor
use mod1
real(2), managed :: a(m,k)
real(2), managed :: b(k,n)
real(4), managed :: c(m,n)
a = real(1.0,kind=2)
b = 2.0_2
c = 0.0
call test1 <<<1,32>>> (a,b,c)
istat = cudaDeviceSynchronize()
print *,all(c.eq.2*k)
end program




5.8. OpenACC 互操作性示例
此示例演示了在同一程序中一起使用 CUDA Fortran 和 OpenACC 的两种方式，包括数据共享和控制流。 在最低级别，我们稍微修改了 BLAS daxpy 子例程，方法是将其插入模块中，使其成为 OpenACC 向量例程，并添加 OpenACC 向量循环指令。 第二个文件包含纯 CUDA Fortran，一个全局子例程，它为每个线程块使用相同的参数调用 daxpy。 在最高级别，我们有一个 Fortran 主程序，它使用 OpenACC 进行数据管理，但通过 host_data 指令调用 CUDA Fortran 全局函数和重载的 CUDA Fortran 归约。 此指令指示编译器传递相应的设备指针，这些指针由 OpenACC 运行时隐式管理，用于 x 和 y 参数。
混合使用 CUDA Fortran 和 OpenACC
      module daxpy_mod
      contains
      subroutine daxpy(n,da,dx,incx,dy,incy)
!$acc routine vector nohost
!
!     constant times a vector plus a vector.
!     uses unrolled loops for increments equal to one.
!     jack dongarra, linpack, 3/11/78.
!     modified 12/3/93, array(1) declarations changed to array(*)
!
      integer, value :: n, incx, incy
      double precision, value ::  da
      double precision dx(*),dy(*)
      integer i,ix,iy
!
      if(n.le.0)return
      if (da .eq. 0.0d0) return
      if(incx.eq.1.and.incy.eq.1) then
!
!        code for both increments equal to 1
!
!$acc loop vector
        do i = 1, n
          dy(i) = dy(i) + da*dx(i)
        end do
      else
!
!        code for unequal increments or equal increments
!          not equal to 1
!
!$acc loop vector
        do i = 1, n
          if(incx.lt.0) then
            ix = 1 + (-n+i) * incx
          else
            ix = 1 + (i-1) * incx
          end if
          if(incy.lt.0) then
            iy = 1 + (-n+i) * incy
          else
            iy = 1 + (i-1) * incy
          end if
          dy(iy) = dy(iy) + da*dx(ix)
        end do
      end if
      return
      end
      end module daxpy_mod


module mdaxpy
use daxpy_mod
contains
  attributes(global) subroutine mdaxpy(x,y,n)
  integer, value :: n
  real(8), device :: x(n), y(n,n)
  real(8) :: a
  a = 0.5d0
  j = blockIdx%x
  call daxpy(n, a, x, 1, y(1,j), 1)
  return
  end subroutine
end module


在 CUDA 代码中必须小心，其中编程模型允许线程应用方式的很大灵活性，以遵循 OpenACC 在调用 OpenACC 例程时的期望。 从线程块中的每个线程调用 OpenACC 向量例程，传递相同的参数，通常是安全的。 从 CUDA 线程调用 OpenACC 顺序例程也是安全的。 这通常是一个新功能，尚未经过全面测试。
program tdaxpy
! Compile with "nvfortran -cuda daxpy.F mdaxpy.CUF tdaxpy.F90"
use cudafor
use mdaxpy
integer, parameter :: n = 100
real(8) :: x(n), y(n,n)
x = 2.0d0
y = 3.0d0
!$acc data copyin(x), copy(y)
!$acc host_data use_device(x,y)
call mdaxpy <<<n, n>>> (x, y, n)
print *,sum(y),maxval(y).eq.minval(y)
!$acc end host_data
!$acc end data
end program


有很多从 OpenACC 计算区域内调用 CUDA 代码的示例。 NVIDIA 包中的 examples 目录中有几个，来自 Fortran、C 和 C++。 还有许多使用 OpenACC host_data 指令的示例。 有关该指令和其他指令的更多信息，请参见 OpenACC 规范。
声明
声明
所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“资料”）均“按原样”提供。NVIDIA 不对这些资料作任何明示、暗示、法定或其他方面的保证，并且明确声明不承担任何关于不侵权、适销性和适用于特定用途的暗示保证。
所提供的信息据信是准确可靠的。但是，NVIDIA Corporation 对使用此类信息或因使用此类信息而可能导致的侵犯第三方专利或其他权利的后果不承担任何责任。未通过暗示或其他方式授予 NVIDIA Corporation 任何专利权下的许可。本出版物中提及的规范如有更改，恕不另行通知。本出版物取代并替换以前提供的所有其他信息。未经 NVIDIA Corporation 明确书面批准，NVIDIA Corporation 产品不得用作生命支持设备或系统中的关键组件。
商标
NVIDIA、NVIDIA 徽标、CUDA、CUDA-X、GPUDirect、HPC SDK、NGC、NVIDIA Volta、NVIDIA DGX、NVIDIA Nsight、NVLink、NVSwitch 和 Tesla 是 NVIDIA Corporation 在美国和其他国家/地区的商标和/或注册商标。其他公司和产品名称可能是与其关联的各自公司的商标。