GPGPUは難しいと主張するのをやめよう 〜GPGPUは難しくない 2015冬〜

今年のテーマは「何故HadoopとかMapReduce系は簡単だという風潮があるのにGPGPUは難しいという風潮になるのか？」だった。
別にテーマだったと言っても何かやったわけではないけど…


それで「GPGPUも簡単だと思う」みたいなことを書いてたら https://twitter.com/tanakmura/status/672681735990075393 なんか色々あって、今見たら「帰ったらちゃんと書きます」とかおととい書いてたのでちゃんと書こうと思った。

「GPGPUは難しい」と主張することのデメリット 〜Hadoop vs CUDA〜

MapReduce系とか登場当時は、「これで並列が簡単に使える」みたいなことが宣伝されてたよね。(最近あんま見かけない気がするけど)
一方、それに比べて、GPGPUとかマルチスレッドは「難しくて使えない」みたいなことが言われていた。

(以後、話を簡単にするために、Hadoop と CUDA だけの話にする)

  • Hadoop : 簡単
  • CUDA : むずかしい

みたいに言われがちである。ほんとうにそうか？

インストールから考えると、

Hadoop

  1. インストール
  2. Mapを書く
  3. Reduceを書く
  4. HDFSへ転送
  5. 実行
  6. 結果をHDFSから取得

CUDA

  1. インストール
  2. カウント処理を書く
  3. prefix sumを書く
  4. cudaMemcpy
  5. 実行
  6. cudaMemcpy

となる。さて、どちらが難しいだろうか？

僕はやったことないが、おそらく、Hadoopのほうが難しいのではないか？と思う。Hadoopはネットワークを越えて設定しないといけない。ネットワークを越えるのはいつでもトラブルを難しくさせると思う。

まあ、イントールの難しさは枝葉なので、誰か詳しい人にやってもらえば終わる問題だ。それでも、ここで忘れてはいけないことは、両者とも、

  1. アルゴリズムがMapとReduceに対応するように変換して
  2. 適切なMapとReduceを書いて
  3. データを適切に配置し
  4. 処理を実行し
  5. 終わったらデータを適切に取得する

という処理が必要で、本質的にその難しさに違いがない、という点だ。

つまり、僕の理解だと、

  • インストール : CUDAの方が簡単
  • 本質的に : そんな変わらない

というような感じだ。

まあ、JavaとC++どっちが慣れてるか、とか、デバッガの使いかたとか、そういう本質的ではないが、重要な違いがあって、人によって個人差(腹痛が痛い差)があるかもしれない。まあでもそれは個人差だし、CUDAのほうが簡単だという人も結構いるでしょ。


じゃあ、何故

  • Hadoop : 簡単
  • CUDA : むずかしい

というようになってしまったかというと、(色々要因はあると思うけど)雰囲気の問題で、Hadoopのほうは「Hadoopは簡単」という人が多かったのに対し、CUDAのほうは「CUDAは難しい」という人が多かったというのが問題なんではないかと思うんだよねー。


まあ、それも些細な問題、実際色々見た印象としては、HadoopもCUDAも必要な人はちゃんと勉強して使ってるし、必要でない人は使ってない。
HadoopがCUDAより広く使われているのは、最内ループの演算時間が減るのが嬉しい人よりも、大規模なクラスタの管理を簡単にしたいという人のほうが多かったというだけで、CUDAが簡単だという文化になったとしてもHadoopより普及するというようなことは起こらなかったのではないかな、と思う。

つまり、現状で構わない…構わないんだよ…

じゃあ、何が問題か？というと、まあ、自分のポジションをアピールするために、自分の作業を必要以上に難しくアピールする人間嫌いなので、「CUDAむずかしい」とか言われると、「そういうのはやめろ」と言いたくなるので、つまり俺の心の問題だ。
(「心の問題は身近な人に相談しましょう」というフレーズが、身近に相談できる人がいない人間の心を傷付けているということに何故気付かないのか？)


CUDAはむずかしくない

あわせて、「そんなにむずかしくない」という話も書いておこうと思う。


「CUDAが難しいか？」というのは、「Cのポインタは難しいか？」というのと、同クラスの問題であると思う。つまり、その理解はバイナリ値をとり、わかる人間にはわかるし、わからない人間にはわからない。

あなたは、「Cのポインタ難しい」という人間を見たら、どう思うだろうか？
「まあ俺も最初はむずかしかった」「一回理解すれば簡単」「あんなの機械語理解すれば一発」「コンクリートな機械語と抽象的なポインタを一緒にする人間乙、それはポインタの本質的な理解ではない」「そもそもポインタとは何か？6.2.5 20には〜」「いまどき機械語とかいう老害乙」「なんで引数は配列で宣言してもポイン(略)

僕から見たら、「CUDA難しい」というのは、そういう感じ、そういう感じなんだ…「まあ難しいという人が一定いるのはわかるけど、一回理解できれば難しくない」とか、そういう。


世の中もっと難しい問題はいくらでもある。最近はコンピュータで正しい答のわからない問題を解くことも一般的になってきた。そういう問題について勉強するときは、勉強すればするほど、勉強しないといけない量が発散していく。それに比べたら、CUDA、SIMD、マルチスレッド、キャッシュなんて、一回勉強すれば終わりだ。勉強量は収束していく。本当に難しい問題に比べたら、CUDAの難しさなんて、全然、全然大したことないよ…


あと、純粋にCUDAが実装としてよくできてると思うのでそういう話も書いておく。


並列フレームワークみたいなのは色々世の中にあるけど、CUDAはOpenMPの次ぐらいに書きやすいと思う。

CUDAで、配列加算は、

#define BLOCK_SIZE 256
#include <stdio.h>

__global__ void array_add(int *a, int *b, int *c) {
    int id = threadIdx.x + blockIdx.x * BLOCK_SIZE;
    a[id] = b[id]+c[id];
}

#define N 1024
int a[N];
int b[N]={1};
int c[N]={2};

int main() {
    int nblk = N/BLOCK_SIZE;
    void *a2, *b2, *c2;
    b[1023] = 9;
    c[1023] = 8;
    cudaMalloc(&a2, sizeof(int)*N);
    cudaMalloc(&b2, sizeof(int)*N);
    cudaMalloc(&c2, sizeof(int)*N);
    cudaMemcpy(b2, b, sizeof(int)*N, cudaMemcpyHostToDevice);
    cudaMemcpy(c2, c, sizeof(int)*N, cudaMemcpyHostToDevice);
    array_add<<<nblk, BLOCK_SIZE>>> ((int*)a2, (int*)b2, (int*)c2);
    cudaMemcpy(a, a2, sizeof(int)*N, cudaMemcpyDeviceToHost);
    printf("%d %d\n", a[0], a[1023]);
}

このぐらいに書ける。(これは速くないという人はHadoopのwordcountサンプルにも同じように主張してどうぞ)

cudaMalloc とか cudaMemcpy とか呼び出しがキモいという人の気持ちもわかる。ただ、これは、

  • メモリがGPUとCPUで分割されてると性能上かなり有利(倍速のDRAMを一個用意するより、普通のDRAMを二個並べたほうがコストが下げられる)
  • 並列数は、人間が与えたほうが、色々な問題が簡単になる

という、本質的な問題から来ていて、これを除去するのは難しい…難しいんだよ…つまり、僕の理解では、CUDAは、本質的な問題以外に余計なことを書かなくてよい、という点でよくできてる、と思う。(例えば、OpenCLと比べるとCUDAがよくできてるのがわかると思う)



そして、CUDAのポイントとして、性能を限界まで出そうと思えば出すことができるというのがある。つまり、

  • 簡単な問題は簡単に書けて
  • 性能を出したい人は頑張れば出せる

というのは、並列フレームワークまわりではよく言われることなんだけど、これを一番よく実現してるのがCUDAだと思うんだよね。