Cuda的有效地从字节数组复制到不同的尺寸

Question

我有一个字节阵列（无符号字符*）表示像存储器数据结构树共享存储器元件。树的每个节点包含不同大小的元素： 布尔开头，Ñ无符号整型 S和Ñ无符号短秒。我这样做是因为使用内存最少对我来说非常重要。不幸的是，导致内存对齐问题，当我试图访问从全局内存复制到共享内存中：

__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不necessarely对准。问题：

1）如果我不希望我的垫数据结构可以很好地对齐的整数，以字节我唯一的选择复制字节？

2）如果我通过字节从全局共享存储器复制字节，这将是比慢4倍，如果我能复制unsigned int类型通过unsigned int类型。

3）是否有可能得到未对齐的内存访问，如果我按字节做字节？我想我已经读过字节访问总是对齐的。

编辑：

我有一个二叉树的十岁上下的数据结构，其中每个节点包含的形式，有效载荷：

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

用于搜索B树我只要求从每个条目的关键信息，而不是结构中的其余信息。因此，每个节点按以下方式在字节数组中折叠：

（bool is_last）（key1，key2，key3 ...）（（offset1，prob1 prob1），offset2，key2的prob1 prob2） ，（偏移量，key1的prob1 prob2））（unsigned int first_child_start_offset）（short sizeofChild1，short sizeofChild2，short sizeofChild3 ...）

显然，如果is_last为false，那么只会存储没有childrenOffset。

我以这种方式布置数据的原因是，每个节点的条目数量可以变化，因此如果我将单独的数据存储在单独的数组中，我将不得不保留对这些“开始和结束索引”元数据“数组，这会导致更多数据被存储或者在搜索期间不得不使用状态机，这是我想避免的。我相信对于每个节点的布尔部分都可以做相对较少的工作，但是对于其他任何东西（比如偏移量）都可以。

Answer 1

1. 如果我不希望我的垫数据结构可以很好地对齐的整数，以字节我唯一的选择复制字节？

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

2. 如果我从全局复制逐个字节的共享内存，是否比我能够通过unsigned int复制unsigned int慢4倍。

   您应该期望减少内存吞吐量。如果没有基准测试，很难说多少。这是你的工作，如果你是如此倾向

3. 如果我正在逐字节地进行内存访问是否可能得到错位？我想我已经读过字节访问总是对齐的。

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

你还没有给出关于给定子树大小的更多细节。可能有一些模板技巧可用于将先前已知大小的字节负载扩展为一系列32位char4负载，并带有1到3个尾随字节负载，以获得给予字节缓冲区大小的内存。如果它适合您的数据结构设计，那应该更具性能。