使用 Fortran 和 CUDA 计算 PI

Posted 2023-02-16

【中文标题】使用 Fortran 和 CUDA 计算 PI

【英文标题】：Calculating PI with Fortran & CUDA

【发布时间】：2014-04-28 00:58:20

【问题描述】：

我正在尝试在 PGI 的 fortran 编译器中制作一个简单的程序。这个简单的程序将使用显卡使用“飞镖板”算法计算 pi。在与这个程序斗争了很长一段时间之后，我终于让它在大多数情况下都能正常工作。但是，我目前坚持正确传回结果。我必须说，这是一个相当棘手的调试程序，因为我不能再将任何打印语句推入子程序。该程序当前返回全零。我不确定发生了什么，但我有两个想法。两者我都不确定如何解决：

CUDA 内核没有运行？ 我没有正确转换值？ pi_parts = pi_parts_d

嗯，这是我当前程序的状态。所有带有 _d 的变量都代表 CUDA 准备的设备内存，其中所有其他变量（CUDA 内核除外）都是典型的 Fortran CPU 准备变量。现在有一些我已经注释掉的打印语句，我已经从 CPU Fortran 领域尝试过。这些命令是为了检查我是否真的正确地生成了随机数。至于 CUDA 方法，我目前已经注释掉了计算并将z 替换为静态等于1 只是为了看看会发生什么。

module calcPi
contains
    attributes(global) subroutine pi_darts(x, y, results, N)
        use cudafor
        implicit none
        integer :: id
        integer, value :: N
        real, dimension(N) :: x, y, results
        real :: z

        id = (blockIdx%x-1)*blockDim%x + threadIdx%x

        if (id .lt. N) then
            ! SQRT NOT NEEDED, SQRT(1) === 1
            ! Anything above and below 1 would stay the same even with the applied
            ! sqrt function. Therefore using the sqrt function wastes GPU time.
            z = 1.0
            !z = x(id)*x(id)+y(id)*y(id)
            !if (z .lt. 1.0) then
            !   z = 1.0
            !else
            !   z = 0.0
            !endif
            results(id) = z
        endif
    end subroutine pi_darts
end module calcPi

program final_project
    use calcPi
    use cudafor
    implicit none
    integer, parameter :: N = 400
    integer :: i
    real, dimension(N) :: x, y, pi_parts
    real, dimension(N), device :: x_d, y_d, pi_parts_d
    type(dim3) :: grid, tBlock

    ! Initialize the random number generaters seed
    call random_seed()

    ! Make sure we initialize the parts with 0
    pi_parts = 0

    ! Prepare the random numbers (These cannot be generated from inside the
    ! cuda kernel)
    call random_number(x)
    call random_number(y)

    !write(*,*) x, y

    ! Convert the random numbers into graphics card memory land!
    x_d = x
    y_d = y
    pi_parts_d = pi_parts

    ! For the cuda kernel
    tBlock = dim3(256,1,1)
    grid = dim3((N/tBlock%x)+1,1,1)

    ! Start the cuda kernel
    call pi_darts<<<grid, tblock>>>(x_d, y_d, pi_parts_d, N)

    ! Transform the results into CPU Memory
    pi_parts = pi_parts_d
    write(*,*) pi_parts

    write(*,*) 'PI: ', 4.0*sum(pi_parts)/N
end program final_project

编辑代码： 更改了多行以反映Robert Crovella 提到的修复。当前状态：cuda-memcheck 发现的错误在我的机器上显示：Program hit error 8 on CUDA API call to cudaLaunch。

如果有什么方法我可以用来测试这个程序，请告诉我。我正在投掷飞镖，看看它们在我目前使用 CUDA 的调试方式时落在哪里。不是最理想的，但在我找到另一种方法之前必须这样做。

愿 Fortran 诸神在这黑暗时刻怜悯我的灵魂。

【问题讨论】：

您有没有机会查看这个在 GPU 的帮助下对 π 进行蒙特卡罗计算的 Fortran 示例：pgroup.com/lit/articles/insider/v2n1a4.htm

【参考方案1】：

当我编译并运行你的程序时，我得到一个段错误。这是由于您传递给内核的最后一个参数 (N_d)：

call pi_darts<<<grid, tblock>>>(x_d, y_d, pi_parts_d, N_d)

由于N 是一个标量，内核希望直接使用它，而不是作为指针。因此，当您传递指向设备数据的指针 (N_d) 时，设置内核的过程会在尝试访问值 N 时生成 seg 错误（在主机代码中！），应直接将其传递为：

call pi_darts<<<grid, tblock>>>(x_d, y_d, pi_parts_d, N)

当我对您发布的代码进行更改时，我会得到实际的打印输出（而不是 seg 错误），它是一个由 1 和 0 组成的数组（256 个 1，后跟 144 个 0，总共N=400 个值），然后是计算出的 PI 值（在这种情况下恰好是 2.56 (4*256/400)，因为您已经使内核基本上是一个虚拟内核）。

这行代码也可能不是你想要的：

grid = dim3(N/tBlock%x,1,1)

如果N = 400 和tBlock%x = 256（来自前面的代码行），则计算结果为1（即grid 以(1,1,1) 结束，相当于一个线程块）。但是您真的想启动 2 个线程块，以便覆盖整个数据集范围（N = 400 个元素）。有很多方法可以解决这个问题，但为了简单起见，我们总是在计算中加 1：

grid = dim3((N/tBlock%x)+1,1,1)

在这种情况下，当我们启动的网格（就总线程而言）大于我们的数据集大小（512 个线程但在本例中只有 400 个数据元素）时，通常会进行 线程检查 em> 在我们内核的开头附近（在这种情况下，在 id 的初始化之后），以防止越界访问，如下所示：

if (id .lt. N) then

（以及对应的endif在内核代码的最后）这样，只有对应于实际有效数据的线程才被允许做任何工作。

通过上述更改，您的代码应该基本上可以正常运行，并且您应该能够将内核代码恢复为正确的语句并开始获得 PI 估计值。

请注意，您可以检查 CUDA API 的错误返回代码，也可以使用 cuda-memcheck 运行代码以了解内核是否进行越界访问。然而，这些都没有帮助解决这个特殊的段错误。

【讨论】：

非常感谢您对我的问题做出如此周到的回答，并解释了每个问题。我不再获得 sigabrt，但是我得到的都是 0。我在程序上运行了cuda-memcheck，它目前给了我：Program hit error 8 on CUDA API call to cudaLaunch。我一直无法追查到它来自什么地方或来自哪里。会不会是分配问题？再次感谢您的宝贵时间！

错误 8 是“无效的设备功能”。描述为：“请求的设备功能不存在或未针对正确的设备架构进行编译。”可能您没有为正确的设备架构进行编译。如果您仍然需要帮助，请说明您拥有的 GPU 以及您的编译命令行是什么。

我一直在使用pgf95 -o program.exe main.cuf 来编译我的应用程序。我的显卡是 nVidia GeForce 9800GTX+。我没有看到任何关于如何指定系统架构的示例。我一直认为编译器会选择当前运行的架构。我想不会吧？再次感谢您的大力帮助。

我建议你阅读pgi compiler user's guide的第7章。您的 9800GTX 是cc1.1 device。尝试将-ta=nvidia,cc11 标志添加到您的编译命令行。您还可以通过pgf95 -help 获得一些命令行帮助。我不确定，但我认为如果你不指定架构，默认可能是 cc2.0，它不会在你的设备上运行。

是的，我们开始了！设备架构需要设置为旧式设备！非常感谢！