0

私は最初のCUDAアプリケーションを作成しており、練習用にすべてのカーネルを自分で作成しています。

一部では、X_transpose*Xを単純に計算しています。

私はcudaMallocPitchとcudaMemcpy2Dを使用してきましたが、最初にXとX_transpose*X用にデバイスに十分なスペースを割り当てます。Xをデバイスにコピーします。カーネルは、X行列とX_transpose*Xの結果を書き込むためのスペースの2つの入力を受け取ります。

プロファイラーを使用すると、カーネルは元々、サイズ5000x6000のマトリックスで実行するのに104秒かかりました。カーネル内の行列の境界をチェックしないように、ブロックサイズの倍数になるように、ホスト上で行列にゼロを埋め込みます。32x32のブロックサイズを使用します。

グローバルメモリへの合体した読み取り/書き込みを最大化するためにいくつかの変更を加えましたが、これは非常に役立つようです。ビジュアルプロファイラーを使用してコードのリリースビルドをプロファイリングすると、カーネルの実行に4.27秒かかります。

matlabの実行の正確なタイミング（操作X'* X;のみ）を実行していませんが、約3秒のようです。CUDAを使用してmatlabよりもはるかに優れたスピードアップが得られることを望んでいました。

nvidiaビジュアルプロファイラーは私のカーネルの問題を見つけることができません。私はここのコミュニティがそれをより速くする方法についていくつかの提案があるかもしれないことを望んでいました。

カーネルコード：

__global__ void XTXKernel(Matrix X, Matrix XTX) {

//find location in output matrix
int blockRow = blockIdx.y;
int blockCol = blockIdx.x;

int row = threadIdx.y;
int col = threadIdx.x;

Matrix XTXsub = GetSubMatrix(XTX, blockRow, blockCol);
float Cvalue = 0;

for(int m = 0; m < (X.paddedHeight / BLOCK_SIZE); ++m) {

    //Get sub-matrix
    Matrix Xsub = GetSubMatrix(X, m, blockCol);
    Matrix XTsub = GetSubMatrix(X, m, blockRow);

    __shared__ float Xs[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float XTs[BLOCK_SIZE][BLOCK_SIZE];

    //Xs[row][col] = GetElement(Xsub, row, col);
    //XTs[row][col] = GetElement(XTsub, col, row);
    Xs[row][col] = *(float*)((char*)Xsub.data + row*Xsub.pitch) + col;
    XTs[col][row] = *(float*)((char*)XTsub.data + row*XTsub.pitch) + col;

    __syncthreads();

    for(int e = 0; e < BLOCK_SIZE; ++e)
        Cvalue += Xs[e][row] * XTs[col][e];

    __syncthreads();
}

//write the result to the XTX matrix
//SetElement(XTXsub, row, col, Cvalue);
((float *)((char*)XTXsub.data + row*XTX.pitch) + col)[0] = Cvalue;
}

私のマトリックス構造の定義：

struct Matrix {
matrixLocation location;
unsigned int width;             //width of matrix(# cols)
unsigned int height;            //height of matrix(# rows)
unsigned int paddedWidth;       //zero padded width
unsigned int paddedHeight;      //zero padded height
float* data;                    //pointer to linear array of data elements
size_t pitch;               //pitch in bytes, the paddedHeight*sizeof(float) for host, device determines own pitch
size_t size;                //total number of elements in the matrix
size_t paddedSize;          //total number of elements counting zero padding
};

よろしくお願いします。

編集：私は言及するのを忘れました、私はケプラーカード、GTX6704GBで実行しています。

4

1 に答える 1

2
  1. 16x16や8x8のような小さいブロックサイズの方が速い場合があります。このスライドは、ブロック/共有メモリの非正方形サイズが大きいほど、特定のマトリックスサイズの方が高速になる可能性があることも示しています。
  2. [BLOCK_SIZE][BLOCK_SIZE+1]共有メモリ割り当ての場合、バンクの競合を回避するためにを使用して、先頭のディメンションにダミー要素を追加します。
  3. を使用して、内側のforループを展開してみてください#pragma unroll

一方、十分な大きさのA'* Aの場合、MATLABGPUコードよりもはるかに高速になることはおそらくないでしょう。matlabのパフォーマンスのボトルネックは、カーネルのパフォーマンスではなく、呼び出しのオーバーヘッドであるためです。

cuBLASルーチンculas_gemm()は、行列の乗算に対して最高のパフォーマンスを発揮する可能性があります。あなたはそれとあなたのものを比較することができます。

MAGMAルーチンmagma_gemm()は、場合によってはcuBLASよりもパフォーマンスが高くなります。これはオープンソースプロジェクトです。また、彼らのコードからいくつかのアイデアを得るかもしれません。

于 2013-01-27T02:35:31.627 に答える