Windows 7, NVidia GeForce 425M.
行列の行の合計を計算する単純な CUDA コードを作成しました。行列は 1 次元表現 (float へのポインター) です。
コードのシリアル バージョンは次のとおりです (2予想どおり、ループがあります)。
void serial_rowSum (float* m, float* output, int nrow, int ncol) {
float sum;
for (int i = 0 ; i < nrow ; i++) {
sum = 0;
for (int j = 0 ; j < ncol ; j++)
sum += m[i*ncol+j];
output[i] = sum;
}
}
CUDA コード内で、行列を行単位でスイープするカーネル関数を呼び出します。以下は、カーネル呼び出しのスニペットです。
dim3 threadsPerBlock((unsigned int) nThreadsPerBlock); // has to be multiple of 32
dim3 blocksPerGrid((unsigned int) ceil(nrow/(float) nThreadsPerBlock));
kernel_rowSum<<<blocksPerGrid, threadsPerBlock>>>(d_m, d_output, nrow, ncol);
行の並列合計を実行するカーネル関数 (まだ1ループがあります):
__global__ void kernel_rowSum(float *m, float *s, int nrow, int ncol) {
int rowIdx = threadIdx.x + blockIdx.x * blockDim.x;
if (rowIdx < nrow) {
float sum=0;
for (int k = 0 ; k < ncol ; k++)
sum+=m[rowIdx*ncol+k];
s[rowIdx] = sum;
}
}
ここまでは順調ですね。シリアルとパラレル (CUDA) の結果は同じです。
nThreadsPerBlock要点は、パラメータを変更nThreadsPerBlockし32ても、CUDA バージョンの計算にはシリアル バージョンのほぼ 2 倍の時間がかかるということ1024です。
IMO、行列の次元は並列化を正当化するのに十分な大きさです: 90,000 x 1,000.
以下では、異なる を使用して、シリアル バージョンとパラレル バージョンの経過時間を報告しますnThreadsPerBlock。サンプルmsecの平均で報告された時間:100
マトリックス: nrow = 90000 x ncol = 1000
シリアル: ミリ秒単位のサンプルあたりの平均経過時間 (100サンプル): 289.18.
CUDA ( 32ThreadsPerBlock): サンプルあたりの平均経過時間 (ミリ秒) ( 100samples): 497.11.
CUDA ( 1024ThreadsPerBlock): サンプルあたりの平均経過時間 (ミリ秒) ( 100samples): 699.66.
念のため、 / が付いているバージョンが32最速1024 nThreadsPerBlock/最遅です。
ホストからデバイスにコピーするとき、およびその逆のときにある種のオーバーヘッドがあることは理解していますが、速度が遅いのは、最速のコードを実装していないためかもしれません。
私はCUDAの専門家にはほど遠いので:
このタスクの最速バージョンをコーディングしていますか? コードを改善するにはどうすればよいですか? カーネル関数のループを取り除くことはできますか?
どんな考えでも大歓迎です。
編集1
標準について説明しますが、 /のような値を持つ行の/操作にrowSum興味があります。とはいえ、一部のコメンテーターが示唆しているように、の列ベクトル トリックによる乗算を悪用することはできません。ANDOR(0;1}rowANDrowORcuBLAS1COL
編集2
ユーザーが提案したように、他のユーザーとここで承認されました：
独自の関数を記述しようとすることを忘れて、代わりに Thrust ライブラリを使用すると、魔法が起こります。