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 | 内核执行时间 (CUDA)：3.78s | Kernel Ex time (OpenCL): 0.53s |

CUDA 代码：http://pastebin.com/VQMp3Hba

OpenCL 主机代码：http://pastebin.com/cjGYSLQf

OpenCL 内核代码：http://pastebin.com/KKw3Ayz7

Answer 1

您可以尝试在代码中插入显式计时器，而不是信任工具的输出。可能是工具有问题。