cuda-tutorial

4.1 逐次処理

まずは逐次処理でベクトル計算をして、どれだけ時間がかかるのかを体感していただく。

4.1.1 タイマー

今回の計測では以下のタイマーを用いる。 sys/time.hはUnix, Linuxで動作するが、Windowsでは動作しない点に注意する。 Windowsユーザーは各自タイマーを作成していただきたい。

"timer.cuh"

#ifndef TIMER_CUH
#define TIMER_CUH
#include <math.h>
#include <sys/time.h>

double measureTime(void)
{
    static bool active = false;
    static time_t s;
    static suseconds_t us;
    double ms;
    struct timeval tv;
    struct timezone tz;

    if(active){
        active = false;
        gettimeofday(&tv,&tz);
        ms = 1.0e+3 * (tv.tv_sec - s) + 1.0e-3 * (tv.tv_usec - us);
    }
    else{
        active = true;
        ms = 0.0;
        gettimeofday(&tv,&tz);
        s = tv.tv_sec;
        us = tv.tv_usec;
    }
    return ms;
}

#endif

measureTime() を2回呼ぶ間の時間を計算し、msで返すようにしている。

4.1.2 CPUを用いたベクトルの加算

さて実際に1000000次元のベクトルの加算をCPUで行ってみる。

"add_vec_serial_cpu.cu"

#include <iostream>
#include <cuda.h>
#include "timer.cuh"

using namespace std;

const int Nv = 1000000;

void setup_vec(int *vec, int a){
    for(int i = 0; i < Nv; i++){
        vec[i] = i * a;
    }
    return;
}

void add_vec(int *c, int *a, int* b){
    for(int i = 0; i < Nv; i++){
        c[i] = a[i] + b[i];
    }
    return;
}

int main(){
    int *a, *b, *c;

    //Allocation
    a = (int*)malloc(Nv * sizeof(int));
    b = (int*)malloc(Nv * sizeof(int));
    c = (int*)malloc(Nv * sizeof(int));

    //Setup input vecs
    setup_vec(a, 1);
    setup_vec(b, 2);

    //Launch add_vec()
    double ms;
    measureTime();
    for(int i = 0; i < 1000; i++){
        add_vec(c, a, b);
    }
    ms = measureTime();
    cout << "Time: " << ms/1000. << "ms" << endl;

    //Free
    free(a);
    free(b);
    free(c);

    return 0;
}

今回は1000回計算したときの平均を取ることにした。 私の環境での実行結果は以下の通り。

$ nvcc add_vec_serial_cpu.cu
$ ./a.out
Time: 3.62592ms

4.1.3 GPUを用いたベクトルの加算(逐次計算)

続いて、GPUを用いて全く同じことをしてみる。

"add_vec_serial_gpu.cu"

...

__global__ void add_vec(int *c, int *a, int* b){
    for(int i = 0; i < Nv; i++){
        c[i] = a[i] + b[i];
    }
}

int main(){
    int *a, *b, *c;
    int *a_dev, *b_dev, *c_dev;

    //Allocation
    a = (int*)malloc(Nv * sizeof(int));
    b = (int*)malloc(Nv * sizeof(int));
    c = (int*)malloc(Nv * sizeof(int));

    cudaMalloc((void**)&a_dev, Nv * sizeof(int));
    cudaMalloc((void**)&b_dev, Nv * sizeof(int));
    cudaMalloc((void**)&c_dev, Nv * sizeof(int));

    //Setup input vecs
    setup_vec(a, 1);
    setup_vec(b, 2);

    //Transfer input vecs to device
    cudaMemcpy(a_dev, a, Nv * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(b_dev, b, Nv * sizeof(int), cudaMemcpyHostToDevice);

    //Launch add_vec() on the device
    double ms;
    measureTime();
    for(int i = 0; i < 1000; i++){
        add_vec<<<1, 1>>>(c_dev, a_dev, b_dev);
    }
    cudaMemcpy(c, c_dev, sizeof(int), cudaMemcpyDeviceToHost);
    ms = measureTime();
    cout << "Time: " << ms/1000. << "ms" << endl;

    //Free
    free(a);
    free(b);
    free(c);
    cudaFree(a_dev);
    cudaFree(b_dev);
    cudaFree(c_dev);

    return 0;
}

CPUはGPU側に処理を投げたあとは、GPUとメモリのやり取りをするまで続きの処理を行うようになっている。 そこで、タイマーを切る寸前にcudaMemcpy()を挟むことでGPU側の処理が完全に終わるまで強制的に待機させている点に注意する。 この際、先頭のアドレスのみやり取りさせているので、計測時間と比べて極めて短く、転送時間はほぼ無視できる。 私の環境での実行結果は以下の通り。

$ nvcc add_vec_serial_gpu.cu
$ ./a.out
Time: 120.14ms

なんとCPUでの場合に比べて極めて遅くなってしまった。 しかしこれはよく考えてみると当然の結果だ。 なぜなら、今回のプログラムではCUDAコアを1つしか用いない逐次処理をGPUで行っていたのであって、コアを沢山使う並列処理をしているわけではないためだ。

GPUでの高速な処理は並列化させるのが前提だ。 次回以降、いよいよ並列化による高速化を行う。

Links

• 次へ
• 前へ
• 目次