performance - OpenCL と CUDA の間の異常なパフォーマンス ギャップ

Question

CUDA で単純なタイル行列乗算をコーディングしました。こんな感じです：

__global__ void matrixMultiplyShared(float * A, float * B, float * C,
                         int numARows, int numAColumns,
                         int numBRows, int numBColumns,
                         int numCRows, int numCColumns) {

    __shared__ float ds_A[TILE_WIDTH][TILE_WIDTH];
    __shared__ float ds_B[TILE_WIDTH][TILE_WIDTH];

    int bx = blockIdx.x; int by = blockIdx.y;
    int tx = threadIdx.x; int ty = threadIdx.y;

    int row = by * TILE_WIDTH + ty;
    int col = bx * TILE_WIDTH + tx;

    float Cvalue = 0.0;

// Loop over the M and N tiles required to compute the Pd element
    for (int m = 0; m < (numAColumns-1)/TILE_WIDTH+1; ++m) {
        if(row<numARows && m*TILE_WIDTH+tx < numAColumns){
            ds_A[ty][tx] = A[row*numAColumns + m*TILE_WIDTH+tx];
        } else {
            ds_A[ty][tx] = 0;
        }
        if(m*TILE_WIDTH+ty < numBRows && col < numBColumns){
            ds_B[ty][tx] = B[(m*TILE_WIDTH+ty)*numBColumns+col];
        } else {
            ds_B[ty][tx] = 0;
        }
        __syncthreads();
        if(row < numCRows && col < numCColumns){
            for (int k = 0; k < TILE_WIDTH; ++k)
                Cvalue += ds_A[ty][k] * ds_B[k][tx];
        }
        __syncthreads();
    }
    if(row < numCRows && col < numCColumns)
        C[row*numCColumns+col] = Cvalue;
}

その後、OpenCL バージョンで同じ上記のカーネル (若干の変更を加えたもの) を使用して、CUDA と OpenCL のパフォーマンスを一緒に比較しました。しかし、結果は私の予想をはるかに超えるものでした。OpenCL は CUDA よりも 6 ～ 7 倍高速でした。有効ですか？Nisght の出力は次のとおりです。

CUDA:

OpenCL:

アプリの起動とカーネルの実行の間に大きなギャップがあることがわかります。なぜそれが起こったのですか？

---

私の GPU は: GTX 580 | カーネル Ex 時間 (CUDA): 3.78 秒 | カーネル Ex 時間 (OpenCL): 0.53 秒 |

CUDA コード: http://pastebin.com/VQMp3Hba

OpenCL ホスト コード: http://pastebin.com/cjGYSLQf

OpenCL カーネル コード: http://pastebin.com/KKw3Ayz7

Answer 1

ツールからの出力を信頼する代わりに、明示的なタイマーをコードに挿入してみることができます。ツールが間違っている可能性があります。