CUDA合併內存加載行爲

Question

我正在處理一個結構數組，我希望每個塊都可以在共享內存中加載數組的一個單元。例如：塊0將在共享內存中加載數組[0]，塊1將加載數組[1]。

爲了做到這一點，我使用float *結構化數組來結合內存訪問。

我有兩個版本的代碼

1版

__global__ 
void load_structure(float * label){ 

    __shared__ float shared_label[48*16]; 
    __shared__ struct LABEL_2D* self_label; 


    shared_label[threadIdx.x*16+threadIdx.y] = 
      label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) +threadIdx.x*16+threadIdx.y]; 
    shared_label[(threadIdx.x+16)*16+threadIdx.y] = 
      label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) + (threadIdx.x+16)*16+threadIdx.y]; 
    if((threadIdx.x+32)*16+threadIdx.y < sizeof(struct LABEL_2D)/sizeof(float)) { 
    shared_label[(threadIdx.x+32)*16+threadIdx.y] = 
      label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) +(threadIdx.x+32)*16+threadIdx.y]; 
    } 

    if(threadIdx.x == 0){ 
    self_label = (struct LABEL_2D *) shared_label; 
    } 
    __syncthreads(); 
    return; 
} 

... 

dim3 dimBlock(16,16); 
load_structure<<<2000,dimBlock>>>((float*)d_Label; 

計算時間：0.740032毫秒

2版

__global__ 
void load_structure(float * label){ 

    __shared__ float shared_label[32*32]; 
    __shared__ struct LABEL_2D* self_label; 

    if(threadIdx.x*32+threadIdx.y < *sizeof(struct LABEL_2D)/sizeof(float)) 
    shared_label[threadIdx.x*32+threadIdx.y] = 
       label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float)+threadIdx.x*32+threadIdx.y+]; 


    if(threadIdx.x == 0){ 
     self_label = (struct LABEL_2D *) shared_label; 
    } 
    __syncthreads(); 
    return; 
} 

dim3 dimBlock(32,32); 
load_structure<<<2000,dimBlock>>>((float*)d_Label); 

計算時間：2.559264毫秒

在這兩個版本我使用了NVIDIA探查和全局負載效率爲8％。

我有兩個問題： 1 - 我不明白爲什麼有時間差異。 2 - 我的電話是否合併？

我使用2.1計算能力（32線/套）

Answer 1

我解決了我的問題，訪問內存模式在以前的版本中是不正確的。 閱讀cuda最佳實踐指南的第6.2.1段後，我發現如果它們對齊，訪問速度會更快。

爲了調整我的訪問模式，我在結構中添加了一個「假」變量，以便可以將結構大小除以128（現金大小線）。使用此策略我獲得了良好的性能：爲了將2000結構加載到2000塊中，它只需要0.16ms。

這裏是代碼的版本：

struct TEST_ALIGNED{ 
    float data[745]; 
    float aligned[23]; 
}; 


__global__ 
void load_structure_v4(float * structure){ 

    // Shared structure within a block 
    __shared__ float s_structure[768]; 
    __shared__ struct TEST_ALIGNED * shared_structure; 

    s_structure[threadIdx.x] = 
    structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x]; 
    s_structure[threadIdx.x + 256] = 
    structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x + 256]; 
    if(threadIdx.x < 745) 
     s_structure[threadIdx.x + 512] = 
      structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x + 512]; 
    if(threadIdx.x == 0) 
     shared_structure = (struct TEST_ALIGNED*) s_structure; 

    __syncthreads(); 

    return; 
} 

dim3 dimBlock(256); 
load_structure_v4<<<2000,dimBlock>>>((float*)d_test_aligned); 

我現在還在找工作的優化，我將它張貼在這裏，如果我找到了一些。

Answer 2

您的全球負載不合並的視頻卡。 8％相當低，你可能做的最差的是3％。

我相信這樣做的主要原因是您基於threadIdx.x和threadIdx.y進行索引的方式。讓我們看看這行代碼從第二個內核（第一內核具有類似的問題）：

shared_label[threadIdx.x*32+threadIdx.y] = label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float)+threadIdx.x*32+threadIdx.y]; 

特別是考慮這個索引：

threadIdx.x*32+threadIdx.y 

CUDA經紗在X的順序分組， Y，Z。這意味着在一個warp中快速變化的索引將首先在X索引上，然後在Y上，然後在Z上。因此，例如，如果我有一個16x16線程塊，第一個warp將會有threadIdx.x跨越的線程從0到15，而threadIdx.y只跨越0到1.在這種情況下，相鄰的線程大多會有相鄰的threadIdx.x索引。

您的代碼的結果是，由於您的索引，您已經破壞了合併。如果你能調整你的裝載和存儲使用這種類型的索引：

threadIdx.y*32+threadIdx.x 

你會突然看到在您的全局負載效率顯著改善。 （你的共享內存使用情況可能會更好。）

我意識到你有兩個問題，當我想到第一個問題時，我感到困惑。你已經告訴我們「計算時間」是約。第二次執行時間延長4倍，但大概你指的是compute_interpolation內核，除此之外，你還沒有顯示任何細節，除了在第二種情況下你啓動了4倍的線程。也許這裏沒有神祕。你沒有顯示任何代碼。使用內核將大量內容加載到共享內存中，然後退出也無濟於事。共享內存內容不會從一個內核調用持續到下一個內核調用。