多くの小さく不均等なサイズの配列の CUDA 削減

StackOverflow https://stackoverflow.com/questions/1773700

質問

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

SDK の並列リダクションの例は、単一の非常に大きな配列で動作し、そのサイズは都合よくブロックあたりのスレッド数の倍数であるように見えますが、私のケースはかなり異なります。

ただし概念的には、それぞれ 2 つのコンポーネントを含むオブジェクトが多数あります。 upper そして lower これらのコンポーネントにはそれぞれ、 x そして y 座標。つまり

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 つは、すべての配列をゼロ パディングして同じ長さにして、各オブジェクトに対してブロックのグループを動作させるというものでしたが、もしそれがうまくいくかどうかは、さまざまな問題があるようです。

前もって感謝します

役に立ちましたか？

解決

私が正しく理解している場合、

は、あなたが別の結果、1つの結果、Object1.lower.yにObject1.lower.xを削減し、その上にしたいです。任意の所与のオブジェクトに対して低減される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);

また、あなたは計算の能力1.2以上のために開発しているならば、これはかなり古いのGPUのような重要なようではないが、好ましくは、64バイト境界に合わせ、考慮の配列のアライメントを取る必要があります。

この例では、ブロックごとに4つのワープを起動する、すなわち128スレッド、および多くのブロックとして、あなたは、オブジェクトを持っているようになります。

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

他のヒント

トムの答えのフォローアップとして、次のことに言及したいと思います。 反りの軽減 によって簡単に実装できます カブ.

以下に実際の例を示します。

#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];
....
ライセンス： CC-BY-SA と 帰属
所属していません StackOverflow
scroll top