Ubuntu 10.04 on USBメモリで CUDA 2.3 memo

cuda memo

nvccを使う。

CUDAプログラミングのやり方は色々あるようだが、一つの方法としてはnvccというコンパイラからcuda用のコードをコンパイルするという方法がある。らしい。基本的に普通のccコマンドみたいな感じ。詳しくはcudatoolkitインストールすると入る /usr/local/cuda/doc/nvcc_2.3.pdf で。

とりあえずはろわ
#include <stdio.h>

int main() {
    puts("Hello, world\n");
    return 0;
}

$ nvcc -o hello hello.cu
$ ./hello
Hello, world

わーちょうかんたんにGPUプログラミングができたよ！！！

まるでGPUを使っていないのでこれはGPUプログラミングではありません。でもとりあえずnvccが動くことを確認。

まがりなりにもGPUを使うはろわ

意味はないけど一応GPUを使っているハローワールド hello_cuda.cu。

#include <stdio.h>
__device__ void dev_strcpy(char *dst, const char *src) {
    while (*dst++ = *src++);
}
__global__ void gen_hello(char *buf) {
    dev_strcpy(buf, "Hello, World!");
}
int main(){
    char *hello_ptr;
    char hello_buf[128];

    cudaMalloc(&hello_ptr, 128);
    gen_hello<<<1,1>>>(hello_ptr);
    cudaMemcpy(hello_buf, hello_ptr, 128, cudaMemcpyDeviceToHost);
    cudaFree(hello_ptr);
    puts(hello_buf);
    return 0;
}

cudaMemcpyDeviceToHostは組み込み定数だけどCAMEL_CASEじゃない。CUDAの定数は全部camelCaseなんだろか。

簡単なプログラムを書いてみよう - CUDA Information Siteそのままのプログラム。なにをしてるコードなのか詳細はリンク先参照。

$ nvcc -o hello_cuda hello_cuda.cu
$ ./hello_cuda
Hello, World!

動いているようです。

vim設定

上では省いたが実際のコードは

#include <stdio.h>

int main() {
    puts("Hello, world\n");
    return 0;
}
// vim: set ft=cpp:

とvimのmodelineでc++として開くようにしている。cuda.vimとかを設定してもいいんだけど、あっちこっちで書いたりするのでとりあえずこの程度で。本腰据えてでかいの書く気分になったら設定する。

fizzbuzz

簡単な並列プログラムを書いてみよう - CUDA Information Siteを参考にして並列なFizzBuzzプログラムを書いてみる。

#include <stdio.h>
#include <assert.h>
#define COUNT 40

enum FizzBuzzEnum {
    FIZZBUZZ, FIZZ, BUZZ, NONE
};

__device__ FizzBuzzEnum to_enum(int num) {
    return (num%15)==0 ? FIZZBUZZ : (num%3)==0 ? FIZZ : (num%5)==0 ? BUZZ : NONE;
}
__global__ void dev_fizzbuzz(FizzBuzzEnum *d_fizzbuzz) {
    int i = threadIdx.x;
    d_fizzbuzz[i] = to_enum(i+1);
}

int main() {
    FizzBuzzEnum h_fizzbuzz[COUNT];
    FizzBuzzEnum *d_fizzbuzz;

    cudaMalloc(&d_fizzbuzz, sizeof(h_fizzbuzz));

    dev_fizzbuzz<<<1,COUNT>>>(d_fizzbuzz);

    cudaMemcpy(h_fizzbuzz, d_fizzbuzz, sizeof(h_fizzbuzz), cudaMemcpyDeviceToHost);

    for (int i=0;i<COUNT;++i) {
        switch(h_fizzbuzz[i]) {
            case FIZZBUZZ:
                puts("FIZZBUZZ");
                break;
            case FIZZ:
                puts("FIZZ");
                break;
            case BUZZ:
                puts("BUZZ");
                break;
            case NONE:
                printf("%d\n", i+1);
                break;
            default:
                assert(0);
        }
    }

    cudaFree(d_fizzbuzz);

    return 0;
}

$ nvcc -o fizzbuzz01 fizzbuzz01.cu
$ ./fizzbuzz01
1
2
Fizz
4
Buzz
Fizz
7
8
Fizz
Buzz
11
Fizz
13
14
FizzBuzz
16
17
Fizz
19
Buzz
Fizz
22
23
Fizz
Buzz
26
Fizz
28
29
FizzBuzz
31
32
Fizz
34
Buzz
Fizz
37
38
Fizz
Buzz

普通のマルチスレッドプログラミングよりも簡単な感じでマルチスレッドなプログラムが書けるしCUDAはすごいなあ！ 詳細はリンク先で。

fizzbuzz02

簡単な並列プログラムを書いてみよう2 - CUDA Information Site
曰く、
s<<<1,COUNT>>>/<<<COUNT,1>>>/;
s/threadIdx/blockIdx/g
でマルチブロックなプログラムになるらしい。CUDAアーキテクチャのデバイスは複数のSM(Streaming Multiprocessor)を持ち、SMはそれぞれ複数のSP(Streaming Processor)を持ち、SMで動く処理単位をブロック、SPで動く処理単位をスレッドと呼ぶらしい。

参考文献

*1:デバイスのあれこれが新しいせいかhttp://www.yasuoka.mech.keio.ac.jp/cuda/が動かなかった。あれこれいじって動かそうとしたが、Knoppixはカスタマイズが面倒すぎる

test