优化寄存器/L1中数据的每线程复制和0填充

Optimizing per-thread copying and 0-padding of data in registers/L1

本文关键字：复制 线程 填充 寄存器 L1 数据 优化      更新时间：2023-10-16

我正在编写一个内核，它让每个线程用构成低位字节的数据填充一个变量，并填充其余部分(假设字节序很小)。这是在多个线程之间重复而非一致地完成的，因为有些线程可能有更多的字节要复制到其变量中，而填充量较小，而有些线程则需要复制的字节较少，填充量较大。结果和未填充的数据要么在寄存器中(对于较小的大小)，要么在共享内存中，要么在本地内存中(应该由L1覆盖)。

换句话说，假设每个线程都执行：

T padded;
pad_high_bytes(padded, my_low_bytes, my_num_low_bytes);
do_stuff(padded);

我们有：

template <typename T>
__device__ __forceinline__
void pad_high_bytes(
T&                                result,
const unsigned char* __restrict__ low_bytes, 
unsigned                          num_low_bytes);

如果T很大(比如说struct { int data[50]; })，那么我想我可能应该只使用CUDA的设备代码memcpy()。然而，通常情况并非如此——T的大小通常为4或8，低字节数通常在1到3之间，甚至0也不罕见。

很明显，我可以在字节上循环，并抱着最好的希望。我也可以"在int上循环"，只要还有超过4个字节要复制，然后在剩余的字节上循环。但是：

  1. 我能做得更好吗
  2. 我应该以某种方式解释对齐吗？还是麻烦太多了
  3. 复制和填充在代码中应该完全分开，还是应该以某种方式将它们组合在一起
  4. 切换到memcpy()的sizeof(T)的合理值是多少

记住，函数是在结果类型上模板化的，所以如果你只对较小/较大的T有意见，那会很有帮助。

我真的很难理解Q1-3是如何真正负责的，它可以被解释为"我的话是这个模糊描述的任务的最佳设计吗"。所以我甚至都不打算尝试。

Q4应负责：

设备侧memcpy(或设备侧cudaMemcpy，它只是memcpy的一个薄薄的包装器)总是发出一个执行逐字节复制的循环。当您在编译时知道要复制的类型的大小时，总是可以通过自己编写复制循环来做得更好，该循环利用关于类型大小的先验知识(受对齐约束等)。如果你知道类型的大小和要复制的单词的数量，那么除了字节大小的事务之外，你还可以利用循环展开来做得更好

如果您不知道这两件事中的任何一件，那么memcpy仍然是最佳选择，因为它简化了代码，并为工具链中的幕后惯用优化开辟了可能性。我唯一反对的一次是，如果你有机会将其他操作与副本融合在一起，在这种情况下，自己做一些事情可能仍然有意义。

相关文章：