cuda-tutorial

3.2 GPUを使う

前節の内容ではまだGPUに触れてもいなかった。この節で実際にGPUを触っていく。

3.2.1 カーネルの導入

はじめに、何もしないカーネルを導入しよう。

"hello.cu"

...

__global__ void myKernel(){
}

int main(){
    cout << "Hello CUDA!" << endl;
    myKernel<<<1, 1>>>();
    return 0;
}

これで「GPUに何もしない」という命令を与えることができるようになった。ちなみに通常CUDAのカーネル内でprintf()を使うことはできないが、先輩の話や参考資料[1]によれば、 $ nvcc -arch=sm_60 なり $ nvcc -arch=sm_20 といったオプションを指定すれば、可能らしい。しかし、私の環境ではそれができなかったのであくまで紹介に留める。(誰か教えて下さい。)

3.2.2 GPUを用いた足し算

何もしないプログラムだけではなんとも味気ない。実際にGPUに少しだけでも働いてもらってこの章を終わりにする。

"add.cu"

#include <iostream>
#include <cuda.h>
using namespace std;

__global__ void add(int *c, int *a, int *b){
    *c = *a + *b;
}

int main(){
    int a, b, c; //values on the host
    int *a_dev, *b_dev, *c_dev; //values on the device

    //Allocate memories on the devices
    cudaMalloc((void**)&a_dev, sizeof(int));
    cudaMalloc((void**)&b_dev, sizeof(int));
    cudaMalloc((void**)&c_dev, sizeof(int));

    //Setup input values
    a = 2;
    b = 7;

    //Transfer the inputs to device
    cudaMemcpy(a_dev, &a, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(b_dev, &b, sizeof(int), cudaMemcpyHostToDevice);

    //Launch add() kernel on the device
    add<<<1, 1>>>(c_dev, a_dev, b_dev);

    //Transfer the output to host
    cudaMemcpy(&c, c_dev, sizeof(int), cudaMemcpyDeviceToHost);
    cout << "c: " << c << endl;

    //Free
    cudaFree(a_dev);
    cudaFree(b_dev);
    cudaFree(c_dev);
}

カーネルにてデバイスメモリ上の値を引数に取る時、全てポインタで指定する必要がある。またカーネルに返却値は存在しないので、引数に返却値のポインタを据える必要がある。ちなみに今回は関係ないが、デバイス上にメモリを確保した時、初期値として自動的に0が入る。

実際に上のプログラムをコンパイルし実行すると以下の通りになる。

$ nvcc add.cu
$ ./a.out
c: 9

引用、紹介

[1] GPUによる並列計算とCUDA C プログラミング概要

cuda-tutorial

3.2 GPUを使う

3.2.1 カーネルの導入

3.2.2 GPUを用いた足し算

Links

引用、紹介