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

Question

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

__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 by byte 到共享内存，它会比我慢 4 倍吗能够通过 unsigned int.

复制 unsigned int

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

编辑：

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

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)(short sizeofChild1, short sizeofChild2, short sizeofChild3...)

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

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

Answer 1

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

看你提供的代码，我大概会说，是的。你可能想使用 memcpy。编译器将通过这样做发出非常优化的字节复制循环。您可能还想研究更改加载的 ptxas 默认缓存行为以绕过 L1 缓存（因此 -Xptxas="--def-load-cache=cg" 选项）。它可能会提供更好的性能。
如果我从全局逐字节复制到共享内存，是否比我能够通过 unsigned int 复制 unsigned int 慢 4 倍。

您应该预料到内存吞吐量会降低。没有基准测试很难说多少。那是你的工作，如果你愿意的话
如果我一个字节一个字节地进行，是否有可能出现未对齐的内存访问？我想我读过字节访问总是对齐的。

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

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

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

Cuda efficiently copy from byte array to shared memory elements of different sizes

c++

cuda