Hatena::ブログ(Diary)

hogeなlog

プロフィール

hogelog

hogelog

小室 直(こむろ すなお)。電気通信大学2003年入学。2010年修士卒業。プログラミングとかしてます。

カレンダー
1984 | 01 |
2006 | 07 | 08 | 09 | 10 | 11 | 12 |
2007 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2008 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2009 | 01 | 02 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 |
2010 | 01 | 06 | 08 | 09 | 10 | 11 | 12 |
2011 | 01 | 02 | 03 | 05 | 08 | 09 | 10 | 12 |
2012 | 01 | 04 | 06 |

August 31(Tue), 2010

[][] 一定時間後にこの処理を実行、などとしたい時はandroid.os.HandlerからpostDelay

一定時間後にこの処理を実行、などということをJavaでやるにはTimerTaskとTimerなどを使う。Androidでもそれはできるけど、TimerでやるとGUIを動かしてるメインスレッドと別のスレッドで動くことになるから結局Handler経由で操作しなきゃいけないし、毎回TimerTaskをnewする必要なんかもあったりするのはAndroidみたいな環境では嬉しくない。

android.os.Handlerを使いましょう。

public class AndroidTest extends Activity {
    private int count = 0;
    private final Context context = this;
    private final Handler handler = new Handler();
    private final Runnable showMessageTask = new Runnable() {
        @Override
        public void run() {
            Toast.makeText(context, String.valueOf(++count), Toast.LENGTH_SHORT).show();
        }
    };
    @Override
    public void onCreate(Bundle savedInstanceState) {
        super.onCreate(savedInstanceState);
        setContentView(R.layout.main);
        handler.postDelayed(showMessageTask, 1000);
    }
}

起動からおよそ1000ミリ秒後にToastでメッセージを表示するだけ。


参考

Handler  |  Android Developers

Page not found

Timer(Task) = bad! Do it the Android way: Use a Handler :) Living in different worlds

[][] android.widget.RatingBarがあんまり使い心地良くないので改変したStarBar

ratingBar.setColor(color)とかしたかったのでそういうことができるように改変してみた。onDrawだけ自前適当実装。android.graphics.Pathで星の形を描いたんだけど、sinとcosのどっちがどっちかしばらく思い出せなくて数学とか頑張ってた時期は遥か遠くだなーと思った。

android.graphics以下のCanvas, Color, Paint, Pathなんかを使ってonDrawの中身実装する感じはだいたいJava2Dみたいなもの。

f:id:hogelog:20100831092526p:image

とりあえず適当に実装したものgistに置いておく。

続きを読む

トラックバック - http://d.hatena.ne.jp/hogelog/20100831

August 30(Mon), 2010

[][] android.widget.Toast便利

ステータスを表示するときなんかに使う。

public class AndroidTest extends Activity {
    @Override
    public void onCreate(Bundle savedInstanceState) {
        super.onCreate(savedInstanceState);
        setContentView(R.layout.main);
        Toast.makeText(this, "Show Text", Toast.LENGTH_LONG).show();
    }
}

makeTextの結果をshow()するのを忘れないようにしましょう。Androidのステータスメッセージっぽくふわっと出てきてそのうち消える。

f:id:hogelog:20100830234621p:image


文字列以外の任意のViewを表示することができる。

        AnalogClock	clock = new AnalogClock(this);
        Toast toast = new Toast(this);
        toast.setView(clock);
        toast.show();

f:id:hogelog:20100830235159p:image

参考

Toast  |  Android Developers

Y.A.M の 雑記帳: Android Toast で画像を表示する

トラックバック - http://d.hatena.ne.jp/hogelog/20100830

August 18(Wed), 2010

[][] CUDA memo (3) aobenchをCUDAで実装。

f:id:hogelog:20100818005003p:image

aobenchという上記画像を吐くプログラムがあります。オリジナルのC言語ソースコードは

http://code.google.com/p/aobenchcl/source/browse/#svn/trunk/ao/c_reference

などに。で、このプログラムが色々な言語で実装したりして速度を競うのが流行った時期がありました。似た様なプログラムにsmallptというものもあります。これらのプログラムはどちらも出力対象のデータ(画像)の要素(ピクセル)ごとの並列性が高い(とあるピクセルの色を求めるために、前のピクセルの値などが必要ない)ので、CUDAなどの並列性の高いアーキテクチャに圧倒的に向いています。

というわけでaobenchをCUDAに移植してみました。以下の256x256の画像がその出力結果です。

f:id:hogelog:20100818005002p:image

元のC実装は約1.42秒、CUDA実装は約0.11秒。画像サイズを2x2から256x256から変化させた場合の実行時間の変化もグラフにしてみました。

f:id:hogelog:20100818005004p:image

縦軸は秒、横軸はピクセルの数です。なんともまあわかりやすい。512以降はなんかCUDA側をちゃんと動かせなかったので256まで。ソースコードは全部githubに。

http://github.com/hogelog/cuda_learn

環境は以下のとおり。

  • Device: GeForce 9600 GT
  • PCI: PCI Express 2.0
  • Toolkit: CUDA Toolkit for Ubuntu Linux 9.10 release 3.1 (June 2010)
  • OS: Ubuntu 10.04.1 LTS, 2.6.32-24 i686
  • CPU: Intel(R) Core(TM) i7 CPU 960 @ 3.20GHz
  • CPU Cache: 8192 KB
  • Memory: 2517 MB

CUDA移植する際に

  • CUDAにはdrand48などの擬似乱数生成器がないためthrustライブラリ(http://code.google.com/p/thrust/)を利用。
  • doubleを全てfloatに。
  • コンパイルオプション-use-fast-math利用
  • constant メモリ領域の利用(__constant__変数フラグ)
  • 出力画像形式をPFM 形式に変更(生のfloat 値のラスタ画像形式)

とかそんなことも一応変更。

GPU(CUDA)はすごい?

元のC実装は叩き台とするための簡単な実装。CPUで動作する実装と比較するというならばせめてOpenMPやTBBなんかで並列化されたプログラムと比較するべきだろうなと思う。やってない。CUDAに移植するより圧倒的に楽だと思う。

参考

CUDAでAO bench - imHo

既にid:mokeheheさんが通った道でもあります。擬似乱数生成器が無くて苦労されてるようですが、今はthrustライブラリにあるから楽チンですね。

トラックバック - http://d.hatena.ne.jp/hogelog/20100818

August 10(Tue), 2010

[][] CUDA memo (2)

f:id:hogelog:20100810222951p:image

まずCで左側から右側に向かって黒から白にグラデーションしてるPFM画像*1を吐くプログラムを記述。

#include <stdio.h>
#include <stdlib.h>

#define WIDTH  (256)
#define HEIGHT (256)

#define SCALE (1.0)
#define ENDIAN (-1)

void writeppm(FILE *output, const size_t size) {
    int x, y;
    float *ppmbuf = malloc(sizeof(float)*size);
    for (x=0;x<WIDTH;++x) {
        for (y=0;y<HEIGHT;++y) {
            ppmbuf[y*WIDTH+x] = (float)x / (float)WIDTH;
        }
    }
    fprintf(output, "Pf\n%d %d\n%f\n", WIDTH, HEIGHT, ENDIAN*SCALE);
    fwrite(ppmbuf, sizeof(float), size, output);
    free(ppmbuf);
}

int main() {
    const size_t ppmsize = WIDTH * HEIGHT;
    FILE *ppm = fopen("white2black.ppm", "w");
    if (ppm) {
        writeppm(ppm, ppmsize);
    }
    return 0;
}

それを適当にCUDAで書き直したもの。

#include <stdio.h>
#include <stdlib.h>

#define WIDTH  (256)
#define HEIGHT (256)

#define SCALE (1.0)
#define ENDIAN (-1)

__global__ void white2black(float *ppmbuf) {
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    float gray = (threadIdx.x / (float)WIDTH);
    ppmbuf[index] = gray;
}
void writeppm(FILE *output, const size_t size) {
    float *ppmbuf = (float*)malloc(sizeof(float)*size);
    float *d_ppmbuf = NULL;

    dim3 blockNum(HEIGHT);
    dim3 threadNum(WIDTH);

    cudaMalloc((void**)&d_ppmbuf, sizeof(float)*size);

    if (cudaConfigureCall(blockNum, threadNum)) {
        fprintf(stderr, "%s\n", cudaGetErrorString(cudaGetLastError()));
        exit(EXIT_FAILURE);
    }
    white2black<<<blockNum, threadNum>>>(d_ppmbuf);

    cudaMemcpy(ppmbuf, d_ppmbuf, sizeof(float)*size, cudaMemcpyDeviceToHost);
    cudaFree(d_ppmbuf);

    fprintf(output, "Pf\n%d %d\n%f\n", WIDTH, HEIGHT, ENDIAN*SCALE);
    fwrite(ppmbuf, sizeof(float), size, output);
    free(ppmbuf);
}

int main(int argc, char **argv) {
    const size_t size = WIDTH * HEIGHT;
    FILE *ppm = fopen("white2black.ppm", "w");
    if (ppm) {
        writeppm(ppm, size);
    }
    return 0;
}

特にどうというところはないのだけれども、スレッドが512個までしか作れないということを知らずに書いていて途中しばらくハマった。

*1:このフォーマットの画像を開けるプログラムがImageMagicぐらいしか見当たらない

トラックバック - http://d.hatena.ne.jp/hogelog/20100810

August 08(Sun), 2010

[][] Ubuntu 10.04 on USBメモリで CUDA 2.3 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なんだろか。

no titleそのままのプログラム。なにをしてるコードなのか詳細はリンク先参照。

$ 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

no titleを参考にして並列な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

no title

曰く、

s<<<1,COUNT>>>/<<<COUNT,1>>>/;

s/threadIdx/blockIdx/g

でマルチブロックなプログラムになるらしい。CUDAアーキテクチャのデバイスは複数のSM(Streaming Multiprocessor)を持ち、SMはそれぞれ複数のSP(Streaming Processor)を持ち、SMで動く処理単位をブロック、SPで動く処理単位をスレッドと呼ぶらしい。

参考文献

  • 本家のマニュアル (まだあんまり読んでない)
  • no title
  • no title

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

トラックバック - http://d.hatena.ne.jp/hogelog/20100808
最近のコメント
Connection: close