【发布时间】：2017-11-12 17:23:50
【问题描述】：
我正在处理存储在 3D 数组中的大数据。这是我所做的一个内核示例（在 CPU 的 for 循环中调用）：
attributes(global) subroutine mykernel (A,B,C,p,nx,ny,nz)
real,dimension(:,:,:),device :: A,B
real,dimension(:),device :: C
real,device :: p
integer,device :: nx,ny,nz
xInd = blockDim.x * (blockIdx.x-1) + threadIdx.x;
yInd = blockDim.y * (blockIdx.y-1) + threadIdx.y;
zInd = blockDim.z * (blockIdx.z-1) + threadIdx.z;
if (xInd<=nx) then
if (yInd<=ny) then
if (zInd<=nz) then
A(xInd,yInd,zInd)=(A(xInd,yInd+1,zInd)-A(xInd,yInd,zInd))*p-(B(xInd,yInd,zInd+1)-C(yInd)+B(xInd+1,yInd,zInd))*p+C(yInd+1)
end if
end if
end if
end subroutine mykernel
当我启动内核时，一切似乎都很好，GPU 结果与 CPU 结果相同……但就时间而言，性能并不是很好。
我认为这是由于这里的内存访问，但我不确定。我会将我的 3D 数组放在共享内存中，但我正在处理 nxnynz > 1M 数据，因此共享内存中没有足够的空间。
所以我的以下问题是关于性能问题的，有大量数据：
- 是否应该将 3D 阵列展平为 1D 阵列？我会得到提升吗？
- 是否可以在不使用全局或共享内存的情况下读取（内存访问）大型数据数组？
- 在这种情况下，性能问题的其他可能性是什么？
【问题讨论】：
-
假设在 DO 循环等中调用此函数，您将需要内联以控制函数调用开销。即使那样，您的条件语句似乎也会阻止重要的优化，例如矢量化。在没有有效线性化数组访问的优化的情况下，您可能会看到一维展平的增益。我认为在其他问题得到解决之前，您不会看到内存模型的性能。
-
所以我需要在编译器中添加一个选项来内联我的内核？我对内联的理解是将函数调用替换为他自己的代码。 if 语句在这里很昂贵？如果是这样，我会尝试在 1D 中展平，谢谢！在处理大数组时，如果有比全局或共享内存更优化的东西，我会很感兴趣。