2

4800x9600 マトリックスの行を一緒に追加しようとすると、マトリックス 1x9600 になります。

私が行ったことは、4800x9600 をそれぞれ長さ 4800 の 9,600 の行列に分割することです。次に、4800 要素に対して削減を実行します。

問題は、これが本当に遅いことです...

誰にも何か提案がありますか？

基本的に、MATLAB の sum(...) 関数を実装しようとしています。

これは私が正常に動作することを確認したコードです。それは本当に遅いだけです:

void reduceRows(Matrix Dresult,Matrix DA)
{
        //split DA into chunks
        Matrix Dchunk;
        Dchunk.h=1;Dchunk.w=DA.h;
        cudaMalloc((void**)&Dchunk.data,Dchunk.h*Dchunk.w*sizeof(float));

        Matrix DcolSum;
        DcolSum.h=1;DcolSum.w=1;
        //cudaMalloc((void**)&DcolSum.data,DcolSum.h*DcolSum.w*sizeof(float));

        int i;
        for(i=0;i<DA.w;i++)   //loop over each column
        {
                //printf("%d ",i);
                cudaMemcpy(Dchunk.data,&DA.data[i*DA.h],DA.h*sizeof(float),cudaMemcpyDeviceToDevice);
                DcolSum.data=&Dresult.data[i];
                reduceTotal(DcolSum,Dchunk);
        }
        cudaFree(Dchunk.data);
}

行列は次のように定義されます。

typedef struct{
        long w;
        long h;
        float* data;
}Matrix;

ReduceTotal() は、標準の NVIDIA リダクションを呼び出すだけで、Dchunk のすべての要素を合計し、その答えを DcolSum に入れます。

答えが見つからない場合は、CPUでこれをすべて実行しようとしています... ;(

よろしくお願いします。

4

3 に答える 3

3

各列をループする代わりに、列を並列化します。4600 個のスレッドのそれぞれが、その列の 9600 個のエントリを合計し、その合計を結果ベクトルの適切な場所に配置します。

Cuda の操作を簡単にするライブラリを探している場合は、Thrust を強くお勧めします: http://code.google.com/p/thrust/

Thrust を使用して、行列のポインターをデバイス メモリに保持するファンクターを作成し、それを一連の列インデックスにマップします。ファンクタの operator() はインデックスを取り、行列のその列のすべてを合計して、合計を返します。次に、メモリ コピー (または直接の CUDA 呼び出しさえも) なしで、thrust::device_vector に合計を配置します。

ファンクターは次のようになります。

struct ColumnSumFunctor {
    const Matrix matrix;

    // Make a functor to sum the matrix
    ColumnSumFunctor(const Matrix& matrix);

    // Compute and return the sum of the specified column
    __device__
    int operator()(const int& column) const;
};
于 2010-07-27T03:28:05.413 に答える
1

削減はGPGPUの非常に基本的な操作であり、高速であると想定されており、9600倍の削減も遅くないはずです。

どのグラフィックカードを使用していますか？

4800要素の配列を1つの結果に減らすたびに、9600配列に分割することをお勧めします。reduceTotalの代わりに、CUDPPを使用して削減操作を実行することをお勧めします。CUDPPはCUDAのSTLのようなものです。パフォーマンスを考慮して実装されています。

http://code.google.com/p/cudpp/

于 2010-07-23T13:30:25.720 に答える
0

問題は、9600X2 カーネルを起動していることだと思います。これは、単一のカーネルとして表現するのが簡単なアルゴリズムです。

それを実装する最も単純な方法は、メモリを結合しませんが、現在行っている方法よりも高速になる可能性があります。

単純な方法で動作するようになったら、メモリ読み取りを結合します。たとえば、ブロック内のすべてのスレッドが 16 個の連続する float を共有メモリに読み取り、syncthreads を実行し、関連する 16 個の float をレジスタに蓄積し、synthreads を繰り返します。

Computing SDK には、リダクション手法の例がたくさんあります。

于 2010-07-23T15:46:05.820 に答える