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個尾隨字節負載，以獲得給予字節緩衝區大小的內存。如果它適合您的數據結構設計，那應該更具性能。