我在CUDA中編寫了簡單的平鋪矩陣乘法。它是這樣的:OpenCL和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.78s |內核Ex時間(OpenCL):0.53s |
CUDA代碼:http://pastebin.com/VQMp3Hba
OpenCL的主機代碼:http://pastebin.com/cjGYSLQf
的OpenCL內核代碼:http://pastebin.com/KKw3Ayz7
我想你應該爲OpenCL實現添加代碼,那麼至少有機會找到差異的原因... – ppeterka
我已經添加了代碼。 OpenCL:http://pastebin.com/cjGYSLQf CUDA:http://pastebin.com/VQMp3Hba –
您的openCL代碼不包含內核源代碼... – talonmies