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

Question

我有一个字节数组 (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 部分来说，它可以通过相对较少的工作来完成，但对于其他任何东西（比如偏移量）都不能。

Answer 1

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

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

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

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

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

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

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