cuda-tutorial

4.2 並列計算

4.2.1 スレッドの活用

まず、スレッドによって並列化してみる。前節から変化のある部分を中心に書き記した。 指定できるスレッド数の上限は1024であることに気をつける。

"add_vec_thread.cu"

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

using namespace std;

const int Nv = 1000000;
const int NT = 512; //The number of threads

...

__global__ void add_vec(int *c, int *a, int* b){
    //store thread ID
    int i_global = threadIdx.x;

    for(int i = i_global; i < Nv; i += NT){
        c[i] = a[i] + b[i];
    }
}

int main(){
    
    ...

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

    ...

    return 0;
}

__global__ add_vec()内で、複数の配列がある特定のアドレスにアクセスすることがないように、forループの始点をthreadが持つIDにし、NT毎にアクセスする配列をずらすようにしている点に注意する。 実際に実行してみると、次のようになった。

$ nvcc add_vec_thread.cu
$ ./a.out
Time: 0.865521ms

前節での結果と比較して、明らかに高速化出来ていることがわかる。

4.2.2 ブロックの活用

続いて、ブロックによって並列化してみたものが以下である。

"add_vec_block.cu"

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

using namespace std;

const int Nv = 1000000;
const int NB = Nv;

...

__global__ void add_vec(int *c, int *a, int* b){
    //store block ID
    int i_global = blockIdx.x;

    if(i_global < Nv){
        c[i_global] = a[i_global] + b[i_global];
    }
}

int main(){

    ...

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

    ...

    return 0;
}

__global__ add_vec()にてブロック数の上限値を並列数が上回る場合(めったに無い)はスレッドの時と同じようにforループを回すとよい。 実行結果は以下の通り。

$ nvcc add_vec_thread.cu
$ ./a.out
Time: 1.70792ms

CPUでの処理と比べるとやや速度は落ちたものの、GPUの逐次処理の場合と比較した場合、70倍ほどの高速化が出来ている。 このことからもブロックによる並列化はGPUで言うところのSMによる並列化と対応しており、1.5節のAutomatic ScalabilityによってGPUの詳細な構造がCUDAによって隠蔽されているということが伺える。

4.2.3 大規模な並列計算

先程までは、スレッドのみ、もしくはブロックのみによる並列化を行ったが、実際は以下のように両方を駆使して大規模な並列計算を実現する。

"add_vec_parallel.cu"

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

using namespace std;

const int Nv = 1000000;
const int NT = 512;
const int NB = (Nv + NT - 1)/NT;

...

__global__ void add_vec(int *c, int *a, int* b){
    //Culculate and store ID
    int i_global = blockIdx.x * blockDim.x + threadIdx.x;

    if(i_global < Nv){
        c[i_global] = a[i_global] + b[i_global];
    }
}

int main(){
    
    ...

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

    ...

    return 0;
}

ここで新たに出た来たblockDim.xはadd_vec<<<NB, "Here!">>>に格納される数字と対応している。 実際に実行してみると以下の通りになった。

$ nvcc add_vec_parallel.cu
$ ./a.out
Time: 0.025694ms

これを見て明らかなように、GPUの逐次処理と比較して4桁、CPUの逐次処理と比較しても3桁の高速化が実現できている。

このように考慮するべきことをきちんと考慮することで、ようやくマシンパワーを生かしたプログラムを組むことが出来る。

引用、紹介