Cuda 高效地从字节数组复制到不同大小的共享内存元素

Posted 2023-02-22

【中文标题】Cuda 高效地从字节数组复制到不同大小的共享内存元素

【英文标题】：Cuda efficiently copy from byte array to shared memory elements of different sizes

【发布时间】：2015-07-19 10:47:19

【问题描述】：

我有一个字节数组 (unsigned char *)，它表示内存中的树状数据结构。树的每个节点都包含不同大小的元素：1 bool 开头，n unsigned ints 和n 无符号短。我这样做是因为内存使用最少对我来说非常重要。不幸的是，当我尝试从全局内存复制到共享内存时，这会导致内存对齐问题：

__global__ void sampleerror(unsigned char * global_mem, unsigned int updated_idx...) 
    __shared__ unsigned int offsets[MAX_NUM_CHILDREN/2 +1]; 
    __shared__ unsigned int entries[ENTRIES_PER_NODE];
    __shared__ bool booleans[4];
    bool * is_last = &booleans[0];
    //First warp divergence here. We are reading in from global memory
    if (i == 0) 
        *is_last = (bool)global_mem[updated_idx];
    
    __syncthreads();

    if (*is_last) 
        //The number of entries in the bottom most nodes may be smaller than the size
        if (i < (size - 1)/entry_size) 
            entries[i] = *(unsigned int *)(&global_mem[updated_idx + 1 + i*sizeof(unsigned int)]);
        
     else 
        int num_entries = (size - 1 - sizeof(unsigned int) - sizeof(unsigned short))/(entry_size + sizeof(unsigned short));
        //Load the unsigned int start offset together with the accumulated offsets to avoid warp divergence
        if (i < ((num_entries + 1)/2) + 1) 
            offsets[i] = *(unsigned int *)(&global_mem[updated_idx + 1 * i*sizeof(unsigned int)]);
        
        __syncthreads();
        //Now load the entries
        if (i < num_entries) 
            entries[i] = *(unsigned int *)(&global_mem[updated_idx + 1 + (num_entries + 1)*sizeof(unsigned int) + i*sizeof(unsigned int)]);
        
    
    __syncthreads();

我得到未对齐的内存访问，因为我试图在此处（以及在 else 语句中）复制到共享内存：

entries[i] = *(unsigned int *)(&global_mem[updated_idx + 1 + i*sizeof(unsigned int)]);

因为 updated_idx + 1 不一定是对齐的。问题：

1) 如果我不想填充我的数据结构以很好地对齐整数，那么我唯一的选择是逐字节复制吗？

2) 如果我将 byte 通过 byte 从全局复制到共享内存，它是否会比我能够复制 unsigned int 慢 4 倍 by unsigned int。

3) 如果我逐字节进行，是否有可能获得未对齐的内存访问？我想我已经读过字节访问总是对齐的。

编辑：

我有一个 btree-ish 数据结构，其中每个节点都包含以下形式的有效负载：

struct Entry 
    unsigned int key;
    unsigned int next_level_offset;
    float prob1;
    float prob2;
 

为了搜索 btree，我只需要每个条目的关键信息，而不需要结构中的其余信息。因此，每个节点按以下方式折叠在一个字节数组中：

(bool is_last)(key1, key2, key3...)((offset, prob1 prob2 of key1), (offset, prob1 prob2 of key2), (offset, prob1 prob2 of key3))(unsigned int first_child_start_offset) (short sizeofChild1, short sizeofChild2, short sizeofChild3...)

很明显，如果 is_last 为 false，那么就不会存储任何 childrenOffsets。

我以这种方式排列数据的原因是每个节点的条目数可以是可变的，因此如果我将单独的东西存储在单独的数组中，我将不得不额外跟踪这些“元数据”的开始和结束索引" 数组，这将导致存储更多数据或在搜索期间必须使用状态机，我想避免这种情况。我相信对于每个节点的 bool 部分来说，它可以通过相对较少的工作来完成，但对于其他任何东西（比如偏移量）都不能。

【问题讨论】：

发布和实际可编译的最小版本代码会不会太难？当影响代码读取模式的许多变量未定义时，很难分析您的内核。

我可以发布我的代码的可编译版本，但您需要一个字节数组，其中包含实际的数据结构和构造它的方式，我不能轻易发布。我的代码有什么不清楚的地方？

条目的实际大小，对于初学者。 entry_size 来自哪里？是从全局内存读取还是从块和线程索引计算的 updated_idx。您还没有解释为什么无论如何都需要使用 AOS。为什么不使用 SOA 方法？

我已经用您需要的信息更新了我的帖子。如果不清楚并且您需要更多，请询问。

【参考方案1】：

如果我不想填充我的数据结构以很好地对齐整数，那么我唯一的选择是逐字节复制吗？

看看你提供的代码，我或多或少会说，是的。您可能想使用 memcpy。通过这样做，编译器将发出非常优化的字节复制循环。您可能还想调查更改了 ptxas 默认缓存行为以使负载绕过 L1 缓存（因此 -Xptxas="--def-load-cache=cg" 选项）。它可能会提供更好的性能。

如果我从全局逐字节复制到共享内存，会比我能够通过 unsigned int 复制 unsigned int 慢 4 倍吗？

您应该预计内存吞吐量会降低。没有基准测试很难说多少。如果你愿意，那就是你的工作

如果我一个字节一个字节地进行，是否有可能获得未对齐的内存访问？我想我已经读过字节访问总是对齐的。

对齐标准始终是字长。所以单字节字总是对齐的。但请记住，如果您将字节加载到共享内存缓冲区，然后尝试使用reinterpret_cast 读取未与共享字节数组对齐的较大字大小，您会遇到同样的问题。

您没有详细说明给定子树的大小。可能有一些模板技巧可用于将先验已知大小的字节加载扩展为一系列 32 位 char4 加载和 1 到 3 个尾随字节加载，以将给定的字节缓冲区大小放入内存。如果它适合您的数据结构设计，那应该会更高效。

【讨论】：

感谢您的详细解答！我的树通常与 gpu 内存允许的一样大（在我的情况下为 4 GB）。单个节点大约为 4 KB，我打算在处理它们之前将它们放在共享内存中。我想我会尝试限制可能的节点大小，以使内存 4 字节对齐。