5.2 CUDA Histogram直方圖

什么是Histogramming

Histogramming是一種從大的數據集中提取典型特征和模式的方式.

在統計學中，直方圖（英語：Histogram）是一種對數據分布情況的圖形表示，是一種二維統計圖表，它的兩個坐標分別是統計樣本和該樣本對應的某個屬性的度量。

圖像直方圖（英語：Image Histogram）是用以表示數字圖像中亮度分布的直方圖，標繪了圖像中每個亮度值的像素數。可以借助觀察該直方圖了解需要如何調整亮度分布。這種直方圖中，橫坐標的左側為純黑、較暗的區域，而右側為較亮、純白的區域。因此，一張較暗圖片的圖像直方圖中的數據多集中於左側和中間部分；而整體明亮、只有少量陰影的圖像則相反。

很多數碼相機提供圖像直方圖功能，拍攝者可以通過觀察圖像直方圖了解到當前圖像是否過分曝光或者曝光不足。

計算機視覺領域常借助圖像直方圖來實現圖像的二值化。

顏色直方圖

在圖像處理和攝影領域中，顏色直方圖（英語：Color Histogram）指圖像中顏色分布的圖形表示。數字圖像的顏色直方圖覆蓋該圖像的整個色彩空間，標繪各個顏色區間中的像素數。

顏色直方圖本身可以針對任意色彩空間使用，但這一術語通常只用在諸如 RGB 和 HSV 的三維色彩空間，而針對灰度圖像時常使用亮度直方圖（英語：Intensity Histogram）這一術語。

 

直方圖例子

比如一個句子"Programming Massively Parallel Processors",我們構建一個字母出現頻率的直方圖.

A(4次), C(1), E(1), G(1), ...

Host端我們做統計的話,就會構建26大小的數組來表示a-z,然后遍歷這個句子,碰到不同的case進行累加操作.

Device端並行化: 把這個sentense 划分成多個block,每個block統計自己的情況,同意保存到一個數組里面.

但是當兩個thread同事在對A做+1的動作后,最后的結果是1,而正確的結果是2才對,這時候需要使用原子加操作:
unsigned int atomicAdd(unsigned int* address, unsigned int val);

其實在CUDA的device中執行計數的功能都需要使用原子操作才行.

 

優化histor kernel代碼:

不使用shared memory的情況下,執行原子操作,無論有多少個thread,在這個例子中最多同時執行的只有256,其他的都因為原子操作需要等待, 所以這里每個thread block使用一個 shared memory,  統計每個thread block,然后最后一起copy 給 histo

__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo)
{
    __shared__ unsigned int histo_private[256];
    if(threadId.x < 256)  //初始化shared histo
        histo_private[threadIdx.x] = 0;
    __syncthread();
    
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    // 步長是所有threads的數目
    int stride = blockDim.x * gridDim.x;
    while(i < size) {
        atomicAdd(&(private_histo[buffer[i]]), 1);
        i + = stride;
    }
    
    //等待所有線程執行完
    __syncthreads();
    
    if(threadIdx.x < 256){
        atomicAdd(&(histo[threadIdx.x]), private_histo[threadIdx.x]);
    }
}