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

Question

我有一个字节数组(unsigned char *)代表内存中的树状数据结构。树的每个节点包含不同大小的元素:1 bool 在开头，n unsigned int 和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 倍吗 通过 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, key1 的 prob1 prob2), (offset, key2 的 prob1 prob2), (offset, key3 的 prob1 prob2))(unsigned int first_child_start_offset) (短 sizeofChild1，短 sizeofChild2，短 sizeofChild3 ...)

显然，如果 is_last 为 false，则不会存储任何 childrenOffsets。

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

Answer 1

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

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

2. 如果我从全局逐字节复制到共享内存，是否会比我能够逐个无符号整数复制无符号整数慢 4 倍。

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

3. 如果我逐字节进行，是否可能出现未对齐的内存访问？我想我已经读到字节访问总是对齐的。

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

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