在 GPU 中找到比 CPU 慢的最小值

Question

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

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

我有 i7 处理器和 NVIDIA GT 540M 显卡。

应该是这样还是其他？

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

#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，因此我假设您在笔记本电脑上运行，而且您可能将 X 显示器托管在 540M GPU 上。

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

#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、Xeon 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 代码时，它就是多任务处理） .因此，这也可能对您的情况造成性能影响。