我在 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>>>()
因为你在
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_j
和 s_d_l
变量是线程索引。所以一个给定的线程会看到这个 do 循环,它会迭代地执行循环,将来自各种全局内存数组(d_bbb
、d_ccc
等)的连续值加载到每个共享内存中的完全相同位置数组。
在我看来你并不真正理解线程执行是如何工作的。假设你是一个给定的线程,为
s_d_j
和s_d_l
分配特定的值(和d_k
,虽然你在重新使用该变量作为循环索引时覆盖了块索引,这对我来说也很奇怪) ,然后看看你的代码执行是否有意义。
编辑:基于其他评论:
您已声明您的整体数据集大小 (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