0

行列の要素の合計を計算するために、http ://www.cuvilib.com/Reduction.pdf というコードを実装しました。

ただし、GPU では、CPU よりもはるかに遅く実行されます。

私は i7 プロセッサと NVIDIA GT 540M グラフィックス カードを手に入れました。

それはそのようであるべきですか、それとも他の何かですか？

編集: Ubuntu 13.04 で上記のコードのバージョン 3 を使用し、Eclipse Nsight を使用してコンパイルします。行列のサイズは 2097152 要素です。CPUバージョンは約1.0ミリ秒で実行されるのに対し、3.6ミリ秒で実行されます。以下はコード全体です。

#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;
}
4

1 に答える 1

7

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、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 個の SM があることがわかります。さらに重要なことに、M2050 には GT540M GPU の約 5 倍のメモリ帯域幅があります (GT540M の理論上のピークは 28.8GB/s で、最大 150GB です)。 M2050 の理論上の /s ピーク)

適切に記述された並列リダクションは GPU 上のメモリ帯域幅に制約のあるコードであるため、あなたの GPU と私の GPU の速度の違いは理にかなっています。

したがって、結果はおそらく期待どおりのものであり、より良い結果を得るには、おそらくより高速な GPU が必要になるでしょう。

また、GT540M が X ディスプレイもホストしている場合、GPU タイミングがディスプレイ アクティビティによって破損している可能性があります。単一のカーネルのタイミングを計っている場合、これは通常は問題になりません。カーネルの実行によって表示処理が一時的に中断されます。しかし、カーネルのシーケンスを連続してタイミングを計っている場合、ディスプレイ タスクがカーネル呼び出しの合間に飛び込んで実行される可能性があります (GPU は、ディスプレイのサポートと CUDA コードの処理の両方を要求された場合にマルチタスクになります)。 . したがって、これはあなたの場合にもパフォーマンスに影響を与える可能性があります。

于 2013-11-14T23:12:11.007 に答える