【问题标题】:Shared Memory in cuda fortran not working as expectedcuda fortran 中的共享内存无法按预期工作
【发布时间】:2013-07-02 00:37:07
【问题描述】:

我正在构建一个 cuda fortran 并且发生了一个奇怪的行为。我真的不明白为什么我的代码会这样运行,并感谢您的帮助。

似乎从未分配值 0,甚至循环 执行超出边界。

我尝试将 if 条件放在循环之后,但它也没有帮助。 谢谢你的帮助

    real, shared :: s_d_aaa_adk(0:15,0:15)
    real, shared :: s_d_bbb_adk(0:15,0:15)
    real, shared :: s_d_ccc_adk(0:15,0:15)

    d_k = (blockIdx%x-1)
    s_d_j = threadIdx%x-1
    s_d_l = threadIdx%y-1   

    if(d_k == kmax-1)then
        s_d_aaa_adk(s_d_j,s_d_l)  = 0 
        s_d_bbb_adk(s_d_j,s_d_l) = 0
        s_d_ccc_adk(s_d_j,s_d_l)  = 0       
    endif

    do d_k = 0, kmax-2              
        s_d_bbb_adk(s_d_j,s_d_l) = d_bbb(s_d_j,d_l,d_k+1)
        s_d_ccc_adk(s_d_j,s_d_l)  = d_ccc(d_j,s_d_l,d_k+1) 
        s_d_aaa_adk(s_d_j,s_d_l) = d_aaa(d_j,s_d_l,d_k+1)               
    end do  `

我将所有全局内存数组大小设置为 (16,16, kmax), 网格是 (128,1,1),块 (16,16,1),并且 内核启动为testkernell<<<grid,block>>>()

【问题讨论】:

  • kmax 的值是多少?
  • 我正在通过 z 切片一个 3d 到 2d 数组。 kmax 是 z 的最大值。我想把每个切片放在一个块中。
  • 在这种情况下可能不需要使用 do-loop。 x,y,z 中的 3d 数组的大小是多少?
  • 大小为 (64,64,32)。我发布的示例不是完整大小,因为我想先看看它在小数组中的表现。这就是我发布 (16,16,4) 的原因。所有 3 个数组的大小相同。

标签: cuda gpu shared-memory gpgpu


【解决方案1】:

由于您在 d_k 上调整 if 语句,它源自块索引:

d_k = (blockIdx%x-1)
if(d_k == kmax-1)then

这意味着网格中 128 个块中只有一个块会实际执行 if 语句,将那些特定的共享内存值设置为零。您的大多数块都不会执行 if 语句中的内容。

如果 kmax 恰好大于 128,那么您的任何块都不会执行 if 语句。

如果您希望在每个线程块中执行该 if 语句,则需要以块索引以外的其他内容为条件。

我会就如何重构代码提出一个建议,但我不清楚将数据加载到共享内存中你想要实现什么。例如,您的 do-loop 对我来说没有多大意义:

do d_k = 0, kmax-2              
    s_d_bbb_adk(s_d_j,s_d_l) = d_bbb(s_d_j,d_l,d_k+1)
    s_d_ccc_adk(s_d_j,s_d_l)  = d_ccc(d_j,s_d_l,d_k+1) 
    s_d_aaa_adk(s_d_j,s_d_l) = d_aaa(d_j,s_d_l,d_k+1)               
end do            ^     ^
                  |     |
         a given thread has specific values for these indices

您的 s_d_js_d_l 变量是线程索引。所以给定的线程会看到这个 do 循环,它会迭代地执行循环,将来自各种全局内存数组(d_bbbd_ccc 等)的连续值加载到 完全相同每个共享内存数组中的位置。

在我看来,您并不真正了解线程执行的工作原理。假设您是给定线程,将特定值分配给 s_d_js_d_l(和 d_k,尽管在将该变量重新用作循环索引时您正在覆盖块索引,这对我来说也很奇怪) ,然后看看你的代码执行是否有意义。

编辑:基于其他 cmets:

您已声明您的整体数据集大小 (x,y,z) 为 (64,64,32)。 您已经说过“我正在通过 z 切片......数组......我想将每个切片放在一个块中”

这将向我建议您应该在每个切片中启动一个块。或者,也许您有一个算法,将多个块分配给单个切片。无论如何,我将假设您希望所有切片数据(64、64)可用于分配给该切片的给定块。我现在假设您将启动 32 个区块。扩展到多个块在单个切片上工作的情况应该不难。我还将假设一个 32x32 线程块,而不是您指出的 16x16。如果您愿意,将其扩展为使用 16x16 应该不难。

你可能会这样做:

real, shared :: s_d_aaa_adk(0:63,0:63)
real, shared :: s_d_bbb_adk(0:63,0:63)
real, shared :: s_d_ccc_adk(0:63,0:63)

c above uses 48KB of shared mem, so assuming cc 2.0+ and cache config set accordingly

d_k = (blockIdx%x-1)
s_d_j = threadIdx%x-1
s_d_l = threadIdx%y-1   

c fill first quadrant
s_d_bbb_adk(s_d_j,s_d_l) = d_bbb(s_d_j,s_d_l,d_k+1)
s_d_ccc_adk(s_d_j,s_d_l) = d_ccc(s_d_j,s_d_l,d_k+1) 
s_d_aaa_adk(s_d_j,s_d_l) = d_aaa(s_d_j,s_d_l,d_k+1)
c fill second quadrant
s_d_bbb_adk(s_d_j+blockDim%x,s_d_l) = d_bbb(s_d_j+blockDim%x,s_d_l,d_k+1)
s_d_ccc_adk(s_d_j+blockDim%x,s_d_l) = d_ccc(s_d_j+blockDim%x,s_d_l,d_k+1) 
s_d_aaa_adk(s_d_j+blockDim%x,s_d_l) = d_aaa(s_d_j+blockDim%x,s_d_l,d_k+1)
c fill third quadrant
s_d_bbb_adk(s_d_j,s_d_l+blockDim%y) = d_bbb(s_d_j,s_d_l+blockDim%y,d_k+1)
s_d_ccc_adk(s_d_j,s_d_l+blockDim%y) = d_ccc(s_d_j,s_d_l+blockDim%y,d_k+1) 
s_d_aaa_adk(s_d_j,s_d_l+blockDim%y) = d_aaa(s_d_j,s_d_l+blockDim%y,d_k+1)
c fill fourth quadrant
s_d_bbb_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = d_bbb(s_d_j+blockDim%x,s_d_l+blockDim%y,d_k+1)
s_d_ccc_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = d_ccc(s_d_j+blockDim%x,s_d_l+blockDim%y,d_k+1) 
s_d_aaa_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = d_aaa(s_d_j+blockDim%x,s_d_l+blockDim%y,d_k+1)


c just guessing about what your intent was on filling with zeroes
c this just makes sure that one of the slices at the end gets zeroes
c instead of the values from the global arrays

if(d_k == kmax-1)then
c fill first quadrant
    s_d_bbb_adk(s_d_j,s_d_l) = 0
    s_d_ccc_adk(s_d_j,s_d_l) = 0
    s_d_aaa_adk(s_d_j,s_d_l) = 0
c fill second quadrant
    s_d_bbb_adk(s_d_j+blockDim%x,s_d_l) = 0
    s_d_ccc_adk(s_d_j+blockDim%x,s_d_l) = 0
    s_d_aaa_adk(s_d_j+blockDim%x,s_d_l) = 0
c fill third quadrant
    s_d_bbb_adk(s_d_j,s_d_l+blockDim%y) = 0
    s_d_ccc_adk(s_d_j,s_d_l+blockDim%y) = 0
    s_d_aaa_adk(s_d_j,s_d_l+blockDim%y) = 0
c fill fourth quadrant
    s_d_bbb_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = 0
    s_d_ccc_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = 0
    s_d_aaa_adk(s_d_j+blockDim%x,s_d_l+blockDim%y) = 0     
endif

【讨论】:

  • 谢谢罗伯特。我正在尝试在 CUDA 上像新人一样工作。我知道我现在不应该循环其他 blockId。我有一个 3D 数组 (x,y,z) 并想为每个 z 提取平面 (x,y) 并将这些值放在共享内存中。
  • 罗伯特,你真是太好了。感谢您抽出宝贵时间详细解释这一点。现在很清楚了。我将实现我的代码并让你知道。你是对的,我有一个算法,它会在每个切片中显示许多线程块。最后一个目标是将其移植到多 GPU 环境。
  • 实际上我想到您可能无法使用所有 48KB 的共享内存,因为系统将其中的一小部分(通常是几个字节)用于管理目的。所以我写的代码仍然不起作用。您没有足够的共享内存来使单个切片的所有数据可用于每个块,因此您必须考虑您的算法以及是否可以使用共享内存中的缩减集,或者操作其中的一部分全局内存不足。
  • 谢谢罗伯特。但是 48KB 是每个 SM 对吗?在这种情况下,多次使用内核应该有助于解决内存问题。我也在考虑流重叠发送和处理。
  • 是的,每个 SM 48KB。使用流将数据副本与计算重叠是个好主意。
猜你喜欢
  • 1970-01-01
  • 2017-11-30
  • 2016-07-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 2011-06-29
  • 1970-01-01
相关资源
最近更新 更多