3

CUDAで比較的小さいがサイズが異なる多数のアレイの平均/標準偏差を計算するための最良のアプローチを誰かが提案できるかどうか疑問に思っていますか？

SDKの並列削減の例は、単一の非常に大きなアレイで機能し、サイズはブロックあたりのスレッド数の倍数であると便利なようですが、私の場合はかなり異なります。

ただし、概念的には、それぞれに2つのコンポーネントが含まれるオブジェクトが多数あり、これらの各コンポーネントにはupperと座標があります。すなわちlowerxy

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構造体などに格納されていないことに注意してください。データは必要な方法で整理できます。重要なのは、各配列について、平均、標準偏差、そして最終的にはヒストグラムを計算する必要があり、1つの特定のオブジェクト内で、配列間の比率と差を計算する必要があるということです。

このデータをGPUデバイスに送信し、スレッドブロック階層を整理するにはどうすればよいですか？私が持っていたアイデアの1つは、すべての配列をゼロパッドして同じ長さにし、各オブジェクトでブロックのグループが機能するようにすることでしたが、このメソッドが機能する場合は、さまざまな問題があるようです。

前もって感謝します

4

2 に答える 2

1

私が正しく理解していれば、Object1.lower.xを1つの結果に、Object1.lower.yを別の結果に減らすというようになります。任意のオブジェクトに対して、縮小される4つの配列があり、すべて同じ長さです（オブジェクトの場合）。

これには多くの可能なアプローチがあります。影響を与える要因の1つは、システム内のオブジェクトの総数です。数が多いと思います。

最高のパフォーマンスを得るには、最適なメモリアクセスパターンが必要であり、発散を回避する必要があります。合同な配列の数は4であるため、以下のスレッドごとに1つの配列を実行するという単純なアプローチを採用すると、メモリアクセスが不十分になるだけでなく、h/wは各反復でスレッドをチェックする必要があります。ワープはループを実行する必要があります-ループを実行しないと無効になり、非効率になる可能性があります（特に、1つの配列が他の配列よりもはるかに長い場合）。

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

代わりに、各ワープで1つの配列を合計すると、効率が向上するだけでなく、隣接するスレッドが隣接する要素を読み取るため、メモリアクセスパターンも大幅に向上します[1]。

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

アレイのアラインメントも考慮に入れる必要があります。できれば64バイト境界でアラインメントしますが、コンピューティング機能1.2以上を開発している場合、これは古いGPUほど重要ではありません。

この例では、ブロックごとに4つのワープ、つまり128スレッド、およびオブジェクトと同じ数のブロックを起動します。

[1]好きなメモリ配置を選択できると言っていますが、array[0][0]がarray[1][0]の隣になるように配列をインターリーブすると便利な場合があります。これは、隣接するスレッドを意味するためです。隣接するアレイで動作し、合体したアクセスを取得できます。ただし、配列の長さは一定ではないため、これはおそらく複雑であり、短い配列のパディングが必要になります。

于 2009-11-21T13:07:13.913 に答える
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];
....
于 2015-07-30T05:10:20.063 に答える