在GPU中查找最低速度比CPU慢

Question

我已经实现了此代码：http://www.cuvilib.com/Reduction.pdf以便计算矩阵元素的总和。

但是，在GPU中，它的运行速度比在CPU中慢得多。

我得到了i7处理器和NVIDIA GT 540M显卡。

它应该是这样或其他什么？

编辑：我在Ubuntu 13.04中使用上述代码的第3版，并使用Eclipse Nsight进行编译。矩阵的大小是2097152个元素。它在3.6 ms内执行，而CPU版本在1.0 ms左右。下面是整个代码：

#include <stdio.h> 
#include <stdlib.h> 
#include <thrust/sort.h> 
#include <sys/time.h> 
#include <omp.h> 
#include <iostream> 
#include <algorithm> 

#define MIN(a,b) (((a)<(b))?(a):(b)) 



static const int WORK_SIZE = 2097152; 



int find_min(int *a,int length){ 
    int min = a[0]; 
    for (int i=1;i<length;i++) 
      if (a[i]<min) 
     min=a[i]; 
    return min; 
} 


__global__ static void red_min(int *g_idata,int *g_odata) { 
    extern __shared__ int sdata[]; 
    unsigned int tid = threadIdx.x; 
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; 
    sdata[tid]= g_idata[i]; 
    __syncthreads(); 

    for(unsigned int s=blockDim.x/2; s > 0; s >>= 1) { 
     if (tid<s) { 
      sdata[tid] = MIN(sdata[tid],sdata[tid + s]); 
     } 
     __syncthreads(); 
    } 
    if (tid == 0) 
     g_odata[blockIdx.x] = sdata[0]; 
} 





int main(void) { 
    int *d1,*d2; 
    int i,*result; 
    int *idata,*fdata; 
    srand (time(NULL)); 
    result = (int *)malloc(sizeof(int)); 
    idata = (int *)malloc(WORK_SIZE*sizeof(int)); 
    fdata = (int *)malloc(WORK_SIZE*sizeof(int)); 
    cudaMalloc((int**)&d1,WORK_SIZE*sizeof(int)); 
    cudaMalloc((int**)&d2,WORK_SIZE*sizeof(int)); 


    for (i = 0; i < WORK_SIZE; i++){ 
     idata[i] = rand(); 
     fdata[i] = i; 
    } 
    struct timeval begin, end; 
    gettimeofday(&begin, NULL); 
    *result = find_min(idata,WORK_SIZE); 
    printf("Minimum Element CPU: %d \n", *result); 
    gettimeofday(&end, NULL); 
    int time = (end.tv_sec * (unsigned int)1e6 + end.tv_usec) - (begin.tv_sec * (unsigned int)1e6 + begin.tv_usec); 
    printf("Microseconds elapsed CPU: %d\n", time); 

    cudaMemcpy(d1,idata,WORK_SIZE*sizeof(int),cudaMemcpyHostToDevice); 



    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start,0); 
    int num_blocks = 16384; 
    bool flag = true; 
    while (num_blocks>0){ 
     if (flag) { 
      red_min<<<num_blocks,128,128*sizeof(int)>>>(d1,d2); 
     } 
     else { 
      red_min<<<num_blocks,128,128*sizeof(int)>>>(d2,d1); 
     } 
     num_blocks /= 128; 
     flag = !flag; 
} 

Answer 1

GT540M是移动GPU，所以我假设你在笔记本电脑上运行，而且您可能会被托管在540M GPU的X显示。

我建你的代码的完整版本：

#include <stdio.h> 
#include <stdlib.h> 
#include <thrust/sort.h> 
#include <sys/time.h> 
#include <omp.h> 
#include <iostream> 
#include <algorithm> 

#define MIN(a,b) (((a)<(b))?(a):(b)) 



static const int WORK_SIZE = 2097152; 



int find_min(int *a,int length){ 
    int min = a[0]; 
    for (int i=1;i<length;i++) 
      if (a[i]<min) 
     min=a[i]; 
    return min; 
} 


__global__ static void red_min(int *g_idata,int *g_odata) { 
    extern __shared__ int sdata[]; 
    unsigned int tid = threadIdx.x; 
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; 
    sdata[tid]= g_idata[i]; 
    __syncthreads(); 

    for(unsigned int s=blockDim.x/2; s > 0; s >>= 1) { 
     if (tid<s) { 
      sdata[tid] = MIN(sdata[tid],sdata[tid + s]); 
     } 
     __syncthreads(); 
    } 
    if (tid == 0) 
     g_odata[blockIdx.x] = sdata[0]; 
} 





int main(void) { 
    int *d1,*d2; 
    int i,*result; 
    int *idata,*fdata; 
    srand (time(NULL)); 
    result = (int *)malloc(sizeof(int)); 
    idata = (int *)malloc(WORK_SIZE*sizeof(int)); 
    fdata = (int *)malloc(WORK_SIZE*sizeof(int)); 
    cudaMalloc((int**)&d1,WORK_SIZE*sizeof(int)); 
    cudaMalloc((int**)&d2,WORK_SIZE*sizeof(int)); 


    for (i = 0; i < WORK_SIZE; i++){ 
     idata[i] = rand(); 
     fdata[i] = i; 
    } 
    struct timeval begin, end; 
    gettimeofday(&begin, NULL); 
    *result = find_min(idata,WORK_SIZE); 
    printf("Minimum Element CPU: %d \n", *result); 
    gettimeofday(&end, NULL); 
    int time = (end.tv_sec * (unsigned int)1e6 + end.tv_usec) - (begin.tv_sec * (unsigned int)1e6 + begin.tv_usec); 
    printf("Microseconds elapsed CPU: %d\n", time); 

    cudaMemcpy(d1,idata,WORK_SIZE*sizeof(int),cudaMemcpyHostToDevice); 



    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start,0); 
    int num_blocks = 16384; 
    bool flag = true; 
    int loops = 0; 
    while (num_blocks>0){ 
     if (flag) { 
      red_min<<<num_blocks,128,128*sizeof(int)>>>(d1,d2); 
     } 
     else { 
      red_min<<<num_blocks,128,128*sizeof(int)>>>(d2,d1); 
     } 
     num_blocks /= 128; 
     flag = !flag; 
     loops++; 
    } 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    float et = 0.0f; 
    cudaEventElapsedTime(&et, start, stop); 
    printf("GPU time: %fms, in %d loops\n", et, loops); 
    int gpuresult; 
    if (flag) 
     cudaMemcpy(&gpuresult, d1, sizeof(int), cudaMemcpyDeviceToHost); 
    else 
     cudaMemcpy(&gpuresult, d2, sizeof(int), cudaMemcpyDeviceToHost); 
    printf("GPU min: %d\n", gpuresult); 
    return 0; 
} 

编译它：

$ nvcc -O3 -arch=sm_20 -o t264 t264.cu 

并运行它在M2050 GPU，RHEL 5.5，CUDA 5.5，至强X5650 CPU

$ ./t264 
Minimum Element CPU: 288 
Microseconds elapsed CPU: 1217 
GPU time: 0.621408ms, in 3 loops 
GPU min: 288 
$ 

所以我的CPU结果与您的结果非常接近，但我的GPU结果大约快了5-6倍。如果我们将M2050与GT540M进行比较，我们可以看到M2050有14个SM，而GT540M有2个。更重要的是，M2050的内存带宽约为GT540M GPU的5倍（GT540M的理论峰值为28.8GB/s，而大约150GB/s峰值理论为M2050）

由于写得很好的并行减少是GPU上的内存带宽受限代码，所以GPU和我的GPU之间的速度差异是有道理的。

所以我会说你的结果可能是关于预期的，为了获得更好的结果，你可能需要更快的GPU。另外，如果您的GT540M还承载了X显示器，则显卡活动可能会损坏GPU时序。如果我们计算单个内核，这通常不是问题 - 内核执行会暂时中断显示处理。但是当我们连续对内核序列进行计时时，显示任务可能会跳入并在内核调用之间执行（当GPU被要求同时支持显示器并处理CUDA代码时，GPU是多任务的） 。因此，这可能会对您的情况产生影响。