CUDA將GPU的內(nèi)存模型暴露給開發(fā)人員，包括全局內(nèi)存、常量/紋理內(nèi)存、共享內(nèi)存、本地內(nèi)存、寄存器，不同類型內(nèi)存的讀取和訪問的模式有所差別。在不合理的訪問模式下，全局內(nèi)存訪問可能發(fā)生“分區(qū)沖突”（partion camping），其類似于共享內(nèi)存中的bank conflict，只不過粒度較大（資料[1]中介紹的架構(gòu)下分區(qū)寬度為256字節(jié)，而bank寬度通常為4或8字節(jié)）。全局內(nèi)存按照256字節(jié)劃分為多個分區(qū)，所有針對全局內(nèi)存的訪問操作由不同的分區(qū)完成，如果多個內(nèi)存訪問操作地址落在同一個分區(qū)中，這些訪問操作將被串行處理，對性能有較大的影響（全局內(nèi)存訪問本身就是高延遲的操作）。下圖是一個分區(qū)總數(shù)為8的全局內(nèi)存的分區(qū)情況。

以全局內(nèi)存分區(qū)數(shù)量為8為例，下面圖片給出了發(fā)生和不發(fā)生”分區(qū)沖突“的全局內(nèi)存訪問情況。

在發(fā)生”分區(qū)沖突“時，SM-1到SM-30的全局內(nèi)存訪問操作完全變成串行訪問（全部由分區(qū)1處理）。我們下面分別給出發(fā)生和不發(fā)生”分區(qū)沖突“的核函數(shù)示例，通過執(zhí)行該核函數(shù)可以對”分區(qū)沖突”對性能的影響有大致了解。

不發(fā)生“分區(qū)沖突”

//TYPEcanbea2-,4-oran8-byteword
__global__ void readBenchmark(TYPE *d_arr){
    //assignuniquepartitionstoblocks,
    int numOfPartitions=8;
    int curPartition=blockIdx.x%numOfPartitions;
    int partitionSize=256;//256bytes
    int elemsInPartition=partitionSize/sizeof(int);
    //jumptouniquepartition
    int startIndex=elemsInPartition*curPartition;
    TYPE readVal=0;
    //Loopcounter’x’ensurescoalescing
    for(int x=0;x<ITERATIONS;x+=16){
    /*offsetguaranteestorestrictthe
    indextothesamepartition*/
        int offset=((threadIdx.x+x)%elemsInPartition);
        int index=startIndex+offset;
        //Readfromglobalmemorylocation
        readVal=d_arr[index];
    }
    /*Writeoncetomemorytopreventtheabove
    codefrombeingoptimizedout*/
    d_arr[0]=readVal;
}

發(fā)生“分區(qū)沖突”

//TYPEcanbea2-,4-oran8-byteword
__global__ void readBenchmarkPC(TYPE *d_arr){
    int partitionSize=256;//256bytes
    int elemsInPartition=partitionSize/sizeof(TYPE);
    TYPE readVal=0;
    //Loopcounter’x’ensurescoalescing.
    for(int x=0;x<ITERATIONS;x+=16){
        /*allblocksreadfromasinglepartition
        tosimulatePartitionCamping*/
        int index=((threadIdx.x+x)%elemsInPartition);
        //Readfromglobalmemorylocation
        readVal=d_arr[index];
    }
    /*Writeoncetomemorytopreventtheabove
    codefrombeingoptimizedout*/
    d_arr[0]=readVal;
}

具體執(zhí)行配置：網(wǎng)格配置為256x1，線程塊配置為32x32，數(shù)據(jù)類型（TYPE）為整型，數(shù)據(jù)個數(shù)為256x8，迭代次數(shù)為4096x4096，設(shè)備為RTX2080ti。下面是執(zhí)行結(jié)果。

從圖中我們可以看到，即使在第一個核函數(shù)執(zhí)行更多指令的情況下，“分區(qū)沖突”還是使核函數(shù)的性能下降了4倍左右。

注意：

在編譯時需要禁用一級緩存 ，否則讀操作可能由緩存完成而不訪問全局內(nèi)存，從而無法觀察到“分區(qū)沖突”現(xiàn)象。

參考資料

1. 書籍《cuda C權(quán)威編程指南》
2. 論文《Bounding the Effect of Partition Camping in GPU Kernels》