优化 CUDA FDTD Fortran答案

【问题标题】：Optimizing CUDA FDTD Fortran优化 CUDA FDTD Fortran
【发布时间】：2012-03-09 00:33:23
【问题描述】：

我正在尝试使用 CUDA Fortran 优化此 FDTD 代码。我有三个带有输入、输出和常数的 3-D 立方体矩阵。

attributes (global) subroutine kernel_h(k,num_cells_x,num_cells_y,num_cells_z,Hx,Hy,Hz,Ex,Ey,Ez,Cbdx,Cbdy,Cbdz)
    implicit none
    integer :: idx,idy
    integer,value :: k,num_cells_x,num_cells_y,num_cells_z
    real(kind=8), intent(in), dimension(1:num_cells_x,1:num_cells_y,1:num_cells_z) :: Ex, Ey, Ez
    real(kind=8), intent(inout), dimension(1:num_cells_x,1:num_cells_y,1:num_cells_z) :: Hx, Hy, Hz
    real(kind=8), intent(in), constant, dimension(1:num_cells_x,1:num_cells_y,1:num_cells_z) :: Cbdx,Cbdy,Cbdz
    idx = threadIdx%x + ((blockIdx%x-1) * blockDim%x)
    idy = threadIdx%y + ((blockIdx%y-1) * blockDim%y)
    do while (idx < num_cells_x)
        Hz(idx,idy,k) = Hz(idx,idy,k) + ((Ex(idx,idy+1,k)-Ex(idx,idy,k))*Cbdy(idx,idy,k) + (Ey(idx,idy,k)-Ey(idx+1,idy,k))*Cbdx(idx,idy,k))
        Hx(idx,idy,k) = Hx(idx,idy,k) + ((Ey(idx,idy,k+1)-Ey(idx,idy,k))*Cbdz(idx,idy,k) + (Ez(idx,idy,k)-Ez(idx,idy+1,k))*Cbdy(idx,idy,k))
        Hy(idx,idy,k) = Hy(idx,idy,k) + ((Ez(idx+1,idy,k)-Ez(idx,idy,k))*Cbdx(idx,idy,k) + (Ex(idx,idy,k)-Ex(idx,idy,k+1))*Cbdz(idx,idy,k))
        idx = idx + (blockDim%x * gridDim%x)
        idy = idy + (blockDim%y * gridDim%y)
    end do
end subroutine kernel_h

我的内核启动是：

bdim=dim3(16,16,1)
gdim=dim3((num_cells_x+(bdim%x-1))/bdim%x,(num_cells_y+(bdim%y-1))/bdim%y,1)
do k=1,num_cells_z
 call kernel_h<<<gdim,bdim>>>(k,num_cells_x,num_cells_y,num_cells_z,Hx_d,Hy_d,Hz_d,Ex_d,Ey_d,Ez_d,Cbdx_d,Cbdy_d,Cbdz_d)
end do

我的问题是：为什么我不能加载超过 100x100x100 的矩阵？如果我尝试我会收到内核错误启动失败。我可以提高我的代码性能吗？我认为它可以以更好的方式编写。

【问题讨论】：

你有多少 VGA 内存？
设备名称：GeForce GTX 295 执行配置限制最大网格尺寸：65535 x 65535 x 1 最大块尺寸：512 x 512 x 64 每个块最大线程数：512 片外内存总全局内存 (B ): 938803200 总常量内存 (B): 65536 副本的最大内存间距 (B): 2147483647 集成: 每个多处理器没有片上内存共享内存 (B): 16384 每个多处理器的寄存器数量: 16384
该数组应该足够了。无论如何，CUDA Fortran 是 PGI 的专有格式，因此您应该寻求他们的支持。一些支持合同可能是您的许可的一部分。
这是否在您的主显示 GPU 上运行？内核是否运行了几秒钟？你在什么操作系统上运行？我猜你得到的是 TDR。
你好。我的运行执行时间是 16.62646 毫秒。 GPU是主GPU。我的系统是：Linux 2.6.32-39-server #86-Ubuntu SMP Mon Feb 13 23:15:11 UTC 2012 x86_64 GNU/Linux 什么是 TDR？你认为我可以在内核中展开循环吗？

标签： cuda fortran

【解决方案1】：

我猜你是越界访问了。

考虑一个 10x10x10 的体积 (x,y,z)。在这种情况下，您将启动一个 16x16 线程块。这些线程将访问一个 17x17 切片（因为模板半径为 1），这显然会超出边界。您需要禁用那些将访问越界的线程，并禁用那些将超出边界的线程以应用它们的模板。

考虑查看 CUDA SDK 中的 FDTD3D 示例。虽然它是用 C 语言编写的，但它说明了如何处理这个问题，还展示了如何使用共享内存来实现更加更高效的实现。

【讨论】：

k 的用法没问题 - 它是从主机端循环按值传递的。
感谢 talonmies，我错过了。相应地编辑了我的答案。