加载从全局内存

假设简单的内核是这样的：加载从全局内存

__global__ void fg(struct s_tp tp, struct s_param p) 

{ 

    const uint bid = blockIdx.y * gridDim.x + blockIdx.x; 
    const uint tid = threadIdx.x; 
    const uint idx = bid * blockDim.x + tid; 

    if(idx >= p.ntp) return; 

    double3 r = tp.rh[idx]; 

    double d = sqrt(r.x*r.x + r.y*r.y + r.z*r.z); 

    tp.d[idx] = d; 

}

这是真的：

double3 r = tp.rh[idx];

数据从全局内存加载到R参数。
r存储在寄存器中，或者在本地存储器中存在多个变量。
r未存储在共享内存中。
d是计算出来的，然后写回全局内存。
寄存器比其他存储器更快。
如果寄存器的空间已满（一些大的内核），局部存储器使用，并且访问速度较慢
当我需要的双打，是有什么办法可以加快步伐？例如，首先将数据加载到共享内存中，然后对其进行操作？

谢谢大家。

来源

2014-11-22 Henry

'double3 v'用于什么？它被赋予一个从未使用过的值。你的表述看起来很准确由于每个线程都从'tp.rh'中读取自己的值，读入共享内存没有任何好处。您可以通过在一个线程中处理多个数组元素来加快内核速度。 – 2014-11-22 09:22:20

我忘了删除它。现在没问题。 – Henry 2014-11-22 09:24:00

是的，这几乎都是真的。

•当我需要双打时，有什么方法可以加速吗？例如，首先将数据加载到共享内存中，然后对其进行操作？

当存在数据重用（通常由线程块中的多个线程加载相同的数据项）时，或者可能在专门使用共享内存时，使用共享内存很有用以帮助全球合并，例如在optimized matrix transpose期间。

数据重用意味着您不止一次地使用（加载）数据，并且为了共享内存的有用性，这意味着您正在通过多个线程多次加载它。如果您在单个线程中多次使用它，则只需将单个负载加上将其存储在寄存器中的编译器（自动）“优化”即可。

编辑 @Jez给出的答案对于优化加载有一些好的想法。我建议另一个想法是将您的AoS数据存储方案转换为SoA方案。数据存储转换是提高CUDA代码速度的常用方法。

您的s_tp结构，您没有显示，似乎有每个项目/结构的数量double数量的存储。如果您为这些数量中的每一个创建单独的数组，您将有机会获得最佳的加载/存储。事情是这样的：

__global__ void fg(struct s_tp tp, double* s_tp_rx, double* s_tp_ry, double* s_tp_rz, double* s_tp_d, struct s_param p) 

{ 

    const uint bid = blockIdx.y * gridDim.x + blockIdx.x; 
    const uint tid = threadIdx.x; 
    const uint idx = bid * blockDim.x + tid; 

    if(idx >= p.ntp) return; 

    double rx = s_tp_rx[idx]; 
    double ry = s_tp_ry[idx]; 
    double rz = s_tp_rz[idx]; 

    double d = sqrt(rx*rx + ry*ry + rz*rz); 

    s_tp_d[idx] = d; 

}

这种做法很可能会在你的设备代码的其他地方有优惠也为相似类型的使用模式。

来源

2014-11-22 16:24:11

这都是事实。

当我需要双打时，有什么方法可以加速吗？例如首先将数据加载到共享内存中，然后操作它们？

对于您给出的示例，您的实现可能不是最优的。你应该做的第一件事是比较一个参考内核的带宽，例如一个cudaMemcpy。如果差距很大，并且从缩小这一差距中获得的提速是非常重要的，那么优化也许是可能的。在你的内核

寻找有一些罢工我可能不理想的两件事情：

这里没有每个线程太多的工作。如果可能的话，每个线程处理多个元素可以提高性能。这部分是因为它避免了线程初始化/删除开销。
从double3加载不如从其他类型的加载效率。加载数据的最佳方式通常是每个线程使用128位加载。加载三个连续的64位值将会变慢，可能不会很多，但速度会更慢。

编辑：下面罗伯特Crovella的回答给出了一个很好的解决这需要在你的数据类型改变第二点。出于某种原因，我原本以为这不是一种选择，所以如果你只是改变你的数据类型，那么下面的解决方案可能是过顶的！

虽然为每个线程添加更多工作是一件相当简单的事情，但对于解决方案来说，优化内存访问模式（不更改数据类型）并不那么简单。幸运的是，有些库可以提供帮助。我认为使用CUB，尤其是BlockLoad集体，应该允许您更有效地加载。通过使用转置运算符为每个线程加载6 double项，您可以为每个线程处理两个元素，将它们打包成double2，并将它们正常存储。

来源

2014-11-22 18:23:04 Jez

加载从全局内存

回答

相关问题