2009-11-20 91 views
3

我想知道是否有人可以建議最好的方法來計算大量相對較小的平均值/標準偏差,但是在CUDA中的不同大小的陣列CUDA減少許多小的,不相等大小的陣列

SDK中平行下降示例工作一個非常大的陣列上,似乎大小爲方便每塊的線程數的倍數,但我的情況有點不同:

概念,然而,我具有大量的對象,每個對象包含兩個組件,分別爲upperlower,並且這些組件中的每一個都具有xy座標。即

upper.x, lower.x, upper.y, lower.y 

這些陣列中的每一個的長度爲大約800但它的對象之間變化(未在對象內),例如

Object1.lower.x = 1.1, 2.2, 3.3 
Object1.lower.y = 4.4, 5.5, 6.6 
Object1.upper.x = 7.7, 8.8, 9.9 
Object1.upper.y = 1.1, 2.2, 3.3 

Object2.lower.x = 1.0, 2.0, 3.0, 4.0, 5.0 
Object2.lower.y = 6.0, 7.0, 8.0, 9.0, 10.0 
Object2.upper.x = 11.0, 12.0, 13.0, 14.0, 15.0 
Object2.upper.y = 16.0, 17.0, 18.0, 19.0, 20.0 

請注意以上僅僅是我代表陣列的方式和我的數據沒有存儲在C結構或類似的東西:數據可以以任何方式,我需要進行組織。重點是,對於每個陣列,需要計算均值,標準偏差和最終直方圖,並且在一個特定對象內,需要計算數組之間的比率和差異。

我應該如何將這些數據發送到GPU設備並組織我的線程塊層次結構?我想到的一個想法是將所有數組填充爲零,以使它們具有相同的長度,並且在每個對象上都有一組塊,但是如果它能夠工作,似乎存在各種各樣的問題。

在此先感謝

回答

1

如果我理解正確的,你想Object1.lower.x減少到一個結果,Object1.lower.y另一個結果等。對於任何給定的對象,都有四個數組要縮小,所有長度都相等(對於該對象)。

有很多可能的方法來解決這個問題,其中一個影響因素是系統中的對象總數。我會假設這個數字很大。

爲了獲得最佳性能,您希望獲得最佳內存訪問模式,並且希望避免發散。因爲一致數組的數量是4,所以如果你採用了每個線程在下面做一個數組的簡單方法,那麼你不僅會遭受較差的內存訪問,而且h/w還需要檢查每次迭代的線程數warp需要執行循環 - 那些不會被禁用的可能是低效的(特別是如果一個數組比其他數據長得多)。

for (int i = 0 ; i < myarraylength ; i++) 
    sum += myarray[i]; 

相反,如果你每經總結一個數組,那麼不僅會更加高效,而且你的內存訪問模式會更好,因爲相鄰的線程將讀取的相鄰元素[1]。

for (int i = tidwithinwarp ; i < warparraylength ; i += warpsize) 
{ 
    mysum += warparray[i]; 
} 
mysum = warpreduce(mysum); 

你也應該採取陣列的排列考慮,最好是64字節邊界上對齊但如果你是在計算能力1.2或更高版本開發,那麼這是不是很上了年紀的GPU一樣重要。

在這個例子中,你將啓動每個塊四個warp,即128個線程和儘可能多的塊。

[1]你的確可以說你可以選擇你喜歡的任何內存佈局,通常交錯數組使得array [0] [0]緊挨array [1] [0]是有用的,因爲這會意味着相鄰的線程可以在相鄰的陣列上運行並獲得合併訪問。但是由於數組的長度不是常數,這可能很複雜,需要填充較短的數組。

+0

謝謝,我決定去每個陣列1塊 – zenna 2009-11-23 01:10:51

1

作爲湯姆回答的後續,我想提一下,減量可以通過CUB輕鬆實現。

這裏是一個工作實例:

#include <cub/cub.cuh> 
#include <cuda.h> 

#include "Utilities.cuh" 

#include <iostream> 

#define WARPSIZE 32 
#define BLOCKSIZE 256 

const int N = 1024; 

/*************************/ 
/* WARP REDUCTION KERNEL */ 
/*************************/ 
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) { 

    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    unsigned int warp_id = threadIdx.x/WARPSIZE; 

    // --- Specialize WarpReduce for type float. 
    typedef cub::WarpReduce<float, WARPSIZE> WarpReduce; 

    // --- Allocate WarpReduce shared memory for (N/WARPSIZE) warps 
    __shared__ typename WarpReduce::TempStorage temp_storage[BLOCKSIZE/WARPSIZE]; 

    float result; 
    if(tid < N) result = WarpReduce(temp_storage[warp_id]).Sum(indata[tid]); 

    if(tid % WARPSIZE == 0) outdata[tid/WARPSIZE] = result; 
} 

/********/ 
/* MAIN */ 
/********/ 
int main() { 

    // --- Allocate host side space for 
    float *h_data  = (float *)malloc(N * sizeof(float)); 
    float *h_result  = (float *)malloc((N/WARPSIZE) * sizeof(float)); 

    float *d_data;  gpuErrchk(cudaMalloc(&d_data, N * sizeof(float))); 
    float *d_result; gpuErrchk(cudaMalloc(&d_result, (N/WARPSIZE) * sizeof(float))); 

    for (int i = 0; i < N; i++) h_data[i] = (float)i; 

    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice)); 

    sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    gpuErrchk(cudaMemcpy(h_result, d_result, (N/WARPSIZE) * sizeof(float), cudaMemcpyDeviceToHost)); 

    std::cout << "output: "; 
    for(int i = 0; i < (N/WARPSIZE); i++) std::cout << h_result[i] << " "; 
    std::cout << std::endl; 

    gpuErrchk(cudaFree(d_data)); 
    gpuErrchk(cudaFree(d_result)); 

    return 0; 
} 

在該示例中,創建長度N的陣列,並且其結果是連續32元素的總和。所以

result[0] = data[0] + ... + data[31]; 
result[1] = data[32] + ... + data[63]; 
....