この記事はC3 Advent Calendar 2022 5日目の記事です。
前書き
まず簡単な自己紹介から。どうも、Prokumaです。九州工業大学情報工学部を2021年度に卒業し、奈良のどこかで大学院生しています。C3には2020年の秋ごろから卒業するまでいました。
中学時代からプログラミングに興味を持ち、高校時代に進路を悩んだ結果情報系大学に進学、今は情報系の大学院にいます。
地元は大韓民国仁川広域市です。大きな空港のある街で、実家からも車で30分程度で遠くはありません。母語として言語を習得可能な時期を過ぎてから日本語を覚えたので日本語ネイティブではなく、時々おかしい日本語があるかもしれませんがご愛嬌で。
ご存知かもしれませんが、韓国ではネット検閲が行われています。政権に対する悪口はいくら投稿してもいいのですが、性的快楽を求めて名前は出せないようなウェブサイトに接続しようとすると接続が切れたり接続が制限されたと表示されたりします。好奇心溢れる僕としては頑張って接続したいと思い、迂回手段を探し、試したのがパソコンと仲良くなるきっかけではなかったかなと思っています。
自己紹介はここまで。ちなみに後輩にねだって(?)投稿枠を取りました、もしC3のアドカレに投稿したい人がいたらごめんなさい...
いまだにMastodonを眺めた感じ自作ハードウェアに取り組むような人はいなかったようなので、「自作ハードウェアをやって欲しい!」「もっと性能を追求してほしい!」という願望を込めて書いていきたいと思います。
以下、多少上から目線の口調になるかもしれませんが多少強めの口調で言わないと永遠とC3の中で低レイヤーが広まらなさそうなので...ご了承ください。
ハードウェアについて知ることの大切さ
ハードウェアをフル活用することにより実行時間の削減
C3の皆さんは大学の講義でC言語を習うと思います。講義で習ったことを覚えて上手く活用すれば、実装しようとするアルゴリズムはすぐ実装できると思います。
例えば線形代数で習う行列積を実装するとしましょう。多くの大学生はこういう課題が出されたら下のようなコードを書くと思います。
int A[N*N];
int B[N*N];
int C[N*N];
//...
for (int i=0; i<N; i++) {
for (int j=0; j<N; j++) {
int sum = 0;
for (int k=0; k<N; k++) {
sum += A[N*i+k] + B[N*k+j];
}
C[N*i+j] = sum;
}
}
アルゴリズム的に$ O(N^3) $になることがわかります。まぁ、順当なやり方ではあるし、決して間違ってはいません。
しかし、お気づきの方はいらっしゃると思います。同時にやっても計算結果変わらない部分があるということを...
そうなのです。C[N*i+j]は他のCの値に影響されません。つまり、外側二つのループは前のループへの依存性が全くありません。
また、最近のCPUだと最低でも2コア以上積んでいることが多く、コアごとに分割してやれば速くなるのではないかと思えます。
実際OpenMPを使って並列化をしているのと単に三重ループで計算しているコードの性能比較をM1 ProのMacbook Proで行いました。N=2048と設定しました。計測結果がこちら。
Time(Single): 46.949812 sec.
Time(Multi): 8.051264 sec.
実行時間が異なることがわかります。さらなる最適化手法はありますが今回は割愛。同じアルゴリズムですが並列化によって全然違う結果になり、ハードウェア的な知見を活用することによって時間短縮を図れます。
様々な専用ハードウェアへの対応
GPUは普通にコードを書いても動かない
普通にコードを書いても動くのはCPUの上です。しかし、CPUではあまりにも効率が悪い演算というのは存在します。上に挙げた行列積でもそうです。
機械学習に少し興味がある方ならCUDAをご存知かもしれません。NVIDIAのGPUで動くGP(General-propose)GPU用のプラットフォーム
ですね。CUDA以外にもOpenCLなども存在します。
しかし、普通に三重ループを書いてもGPU上で動かないので(正確には動きはしますがGPUの良さを全く生かせないしCPUより遅くなるだけ)構造を考えた上でプログラムを書く必要があります。
__global__
void matrixMul(int *C, int *A, int *B) {
int row = blockIdx.y*blockDim.y + threadIdx.y;
int col = blockIdx.x*blockDim.x + threadIdx.x;
if (row < N && col < N) {
int acc = 0;
for (int k = 0; k < N; k++) {
acc += A[row*N+k] * B[k*N+col];
}
C[row*N+col] = acc;
}
}
行列積を求めるCUDAのカーネルです。外側の二重ループがなくなっていることがわかります。代わりにrowとcolに謎の式が代入されていることがわかります。これは一つのCUDAコアで外側の二重ループに相当する値を代入し、並列で実行しているからです。
しかし、カーネルを書いただけではもちろん動きません。まずメインメモリのデータをGPUのメモリにコピーして、さらにデータを格納するGPUのメモリの割当を行う必要があります。そしてカーネルの呼び出しもmatrixMul(C, A, B)だけでは終わりません。関数の実行はこういう風に行います。
dim3 block(BLOCK_SIZE, BLOCK_SIZE);
dim3 grid((N+block.x-1)/block.x, (N+block.y-1)/block.y);
matrixMul<<< grid, block >>>(C, A, B);
謎の<<< grid, block >>>というコードが見えますね。そうです。どういう風に分割するかを指定しなければなりません。ただカーネルを書いたら終わり、ではないのです。実際は、GPUで演算させるためにはもっとコードを書く必要がありますが、この記事を書く目的はCUDAの使い方ではないのでここで割愛。興味ある方は調べてみてください。
さらなる専用ハードウェア
専用ハードウェアはCUDAに限りません。普通のCPUまたはGPUでのプログラミング全く異なるようなプログラムを書く必要があるハードウェアはたくさん存在し、また新しく登場するかもしれません。これがハードウェアを知ることの大切さです。
GPUは並列化できる演算であればほぼなんでも効率よく計算できます。GP(General-Propose)に名乗っているだけあります。しかし、特定の演算であればGPUよりもっと効率のいいハードウェアを用いることもできます。
「RTX4090とかのハイエンドGPUをn枚積めば何とかなる」と思っているかもしれません。まぁ、あながち間違いではないのですが、それをスマホ上で動かしたい、となったらこの方法は使えません。どこぞの議員みたいに「時代はクラウド」と思うかもしれませんが、ネットワーク環境が弱い、またはネットワーク環境がないところでは使用困難となります。
最近のスマホにはNPU(Neural Processing Unit)が搭載されていることが多く、ここではニューラルネットワーク関係の演算を行います。ビット削減や行列積に最適化された回路構成などでスマホなどのエッジデバイスでもそれなりに計算量を要するニューラルネットワークを実用的な時間内に低消費電力での計算を実現しています。
もちろんこのような専用ハードウェアも、「ぶっ挿したら速度が速くなる魔法の代物」ではありません。これで何かを速くしたい、となったらそのハードウェアに合わせてプログラムを書く必要があります。
新しいハードウェアを理解する練習
ここでは、その一例として、CGRAの一種のIMAX2を挙げます。参考資料はこちら。TPUとかのもっと知られたものもありますが、ここではあえてあまり知られていないものを例として挙げて「ハードウェアの構造を理解し、APIからの操作だけではなくハードウェアの構造に寄り添うようなプログラムを書くことの大切さ」を感じてほしいところです。
IMAX(In-Memory Accelerator eXtention)2の紹介
まずCGRAの紹介から。CGRAは、粗粒度再構成可能アーキテクチャ(Coarse-Grained Reconfigurable Architecture)の略です。後で紹介するFPGAは論理ゲートレベルで再構成して回路をプログラムしますが、CGRAはそれより粗い粒度で再構成を行います。主に演算器単位で構成されます。CGRAのいいところは、データを途中で組み合わせてまるで流れるようにデータを処理できる点にあります。
IMAX2は、日本発のCGRAで、自分の所属研究室で研究開発が行われています。多くのCGRAはFPGAのように平面的にPE(Process Element)を並べているものが多く、FPGAより空間効率は改善されているがコンパイル時間はFPGAのように長くなり、最近用いるようになったテスト駆動開発(TDD)などの開発手法をこのようなデバイスで行うのは極めて困難です。
IMAXでは、線形的にPEを配置することにより、コンパイル時間の短縮を図り、メモリをPEの周辺に配置することによるメモリアクセスのコスト削減、また複雑な制御機構の削減により、省電力・小面積を達成しています。
残念ながらまだFPGA上での実装のみなので、実際のチップがあるわけではありません。「それなら意味ないのでは？」と思うかもしれませんが、ハードウェアを考案し、設計することには意義があり、またアイデアと設計のない状態でいきなりチップを作れるわけでもありません。
商業的に流行るのは難しいかもしれません。主流になってくるのはCGRAではなく真新しい別のハードウェアかもしれません。しかし、間違いなく新しい構造のハードウェアは流行り、それは既存のプログラミング方法ではプログラミングできないものになるはずです。
今回はその練習。IMAX2のプログラミングの仕方について見ていきましょう。
IMAX2のプログラミング
ホームページに載っているマニュアルは指導教員に対する批判にはなりますが読む気をなくすのでわかりやすい例を抜粋して一部だけ載せます。また、わかりやすいコードを重視する皆さんの思想に反するコードがマニュアルに平然と載っていたりしますので、できるだけわかりやすい等価なコードに書き換えています。
行列積のプログラムはここで紹介するには複雑すぎるし、トンカーブを計算するプログラムを紹介します。まず普通のC言語で書かれたプログラムから。
for (row=0; row<HT; row++) {
for (col=0; col<WD; col++) {
unsigned int pix = hin[row*WD+col];
hout0[row*WD+col] =
((ht)[pix >> 24]) << 24 |
(ht[256 + ((pix >> 16) & 255)]) << 16 |
(ht[512 + ((pix >> 8) & 255)]) << 8;
}
}
なんかわからないビット演算が並んでいてなんなのかよくわからない方のための解説。アルゴリズム的に等価なコードを示します。
// input[row][col][ch] = hin[row*WD+col] >> (ch+1)*8;
// output[row][col][ch] = hout0[row*WD+col] >> (ch+1)*8;
// nht[ch][color] = ht[ch*256 + color];
for (row=0; row<HT; row++) {
for (col=0; col<WD; col++) {
unsigned int *pix = input[row][col];
output[row][col][2] = nht[2][pix[2]];
output[row][col][1] = nht[1][pix[1]];
output[row][col][0] = nht[0][pix[0]];
}
}
効率のためにビット演算で表しているだけでやっていることは非常に単純です。ここで、IMAX2の実装を見ていきましょう。
#define Ull unsinged long long;
void tone_curve(unsigned int *r, unsigned int *d, unsigned char *t) {
Ull t1 = t;
Ull t2 = t+256;
Ull t3 = t+512;
Ull BR[16][4][4]; /* output registers in each unit */
Ull r0, r1, r2, r3, r4, r5, r6, r7 r8;
Ull r9, r10, r11, r12, r13, r14, r15, r16;
Ull r17, r18, r19, r20, r21, r22, r23, r24;
Ull r25, r26, r27, r28, r29, r30, r31;
int loop=WD;
//EMAX5A begin tone_curve mapdist=0
while (loop--) {
// rからデータを持ってくる
mop(
// 1word(4byte)をロード
OP_LDWR, 1,
// 格納先レジスタBR[0][1][1], ベースアドレスr++, オフセット0
// MSK_D0: オフセットの全てのビットを使用
&BR[0][1][1], (Ull)(r++), 0LL, MSK_D0,
// 主記憶領域の先頭アドレスr, 長さ320(単位:4byte)
(Ull)r, 320, 0,
// 前回の先頭アドレスおよび長さが同一であればLMMをそのまま利用
0,
// バースト演算中に行うDMAの先頭アドレスNULL, 長さ320(単位:4byte)
(Ull)NULL, 320
);/* stage#0 */
// 最初にロードしたデータをオフセットとし、t1+データをロード
mop(
// 1byteをロード
OP_LDBR, 1,
// 格納先レジスタBR[1][1][1], ベースアドレスt1, オフセットBR[0][1][1]
// MSK_B3: オフセットの31-24番目ビットを利用
&BR[1][1][1], (Ull)t1, BR[0][1][1], MSK_B3,
// 主記憶領域の先頭アドレスt1, 長さ64(単位:4byte)
(Ull)t1, 64, 0,
// 前回の先頭アドレスおよび長さが同一であればLMMをそのまま利用
0,
// バースト演算中に行うDMAの先頭アドレスNULL, 長さ64(単位:4byte)
(Ull)NULL, 64
); /* stage#1 */
// 最初にロードしたデータをオフセットとし、t1+データをロード
mop(
// 1byteをロード
OP_LDBR, 1,
// 格納先レジスタBR[1][2][1], ベースアドレスt1, オフセットBR[0][1][1]
// MSK_B2: オフセットの23-16番目ビットを利用
&BR[1][2][1], (Ull)t2, BR[0][1][1], MSK_B2,
// 主記憶領域の先頭アドレスt2, 長さ64(単位:4byte)
(Ull)t2, 64, 0,
// 前回の先頭アドレスおよび長さが同一であればLMMをそのまま利用
0,
// バースト演算中に行うDMAの先頭アドレスNULL, 長さ64(単位:4byte)
(Ull)NULL, 64
); /* stage#1 */
// 最初にロードしたデータをオフセットとし、t3+データをロード
mop(
OP_LDBR, 1,
// 格納先レジスタBR[1][3][1], ベースアドレスt1, オフセットBR[0][1][1]
// MSK_B1: オフセットの15-8番目ビットを利用
&BR[1][3][1], (Ull)t3, BR[0][1][1], MSK_B1,
// 主記憶領域の先頭アドレスt3, 長さ64(単位:4byte)
(Ull)t3, 64, 0,
// 前回の先頭アドレスおよび長さが同一であればLMMをそのまま利用
0,
// バースト演算中に行うDMAの先頭アドレスNULL, 長さ64(単位:4byte)
(Ull)NULL, 64
); /* stage#1 */
exe(
// 1番目の命令、MMRG演算, 演算結果格納先r1
// 上段のデータを合体
OP_MMRG, &r1,
// 先ほど格納したデータを入力として渡す
// EXP_H3210: 入力データを加工しない
BR[1][1][1], EXP_H3210,
BR[1][2][1], EXP_H3210,
BR[1][3][1], EXP_H3210,
// 2番目の命令、なし
OP_NOP, 0,
// 3番目の命令、なし
OP_NOP, 0
);
// データを保存
mop(
// 無条件ストア命令
OP_STWR, 3,
// 格納先アドレス&r1, メモリアドレスのベース(d++), オフセット0
&r1, (Ull)(d++), 0LL, MSK_D0,
// ホストメモリの先頭アドレスd, 長さ320(単位:4byte)
(Ull)d, 320, 0,
// バースト演算後IMAX2->ホスト
0,
// バースト演算中に行うDMAの先頭アドレスNULL, 長さ320(単位:4byte)
(Ull)NULL, 320
);/* stage#2 */
}
//EMAX
「何これ？」って思う方が多いと思います。しかし、勘のいい方なら気づくかもしれません。
あれ？これアセンブリに似てね？
というか、そうとしか思えん
そうなんです。これはIMAX2の演算器で実行される命令なのです。ここで、実際どういうふうにIMAX2でマッピングされるか見ていきましょう。
BR[16][4][4]の正体をここで知ることができます。これは、繋がっているユニットのどこにデータを置くかの話になります。
ハードウェア上でどういうふうに実装すればいいかを考えながらコードを書かないとならないことがわかります。この例では、IMAX2上でどういうふうにマッピングされ、データの流れがどうなるかを考えなければなりません。このようなCPU/GPUとは思想が違うハードウェアは擬似コードをほぼコピペしたようなコードではプログラミングできません。考えてみれば当たり前です。
IMAX2のいいところ
「プログラミングも難しいし普及してないハードウェアのことを理解して何がいいの？」と思うかもしれません。まぁ、当然な話です。大きなデメリットがあるにも関わらず使われている理由、圧倒的に電力効率がいいからです。
性能比較表となります。実行時間だけ見たらRTX3090にボロ負けですが、見て欲しいのはEDPです。EDPで大きな差でRTX3090に勝っています。
FPGA上の実装なのでバンド幅も狭いし、クロックも低いですが、結構高い性能を叩き出しています。左の一昔前のJetsonには普通に勝っています。
これでいいじゃん、と思わないでほしい
確かにIMAX2は特定の分野において（特に最近流行りのニューラルネットワーク）効率的な演算ができますが、すべてにおいて適しているとは限りません。自分もIMAX2を用いた研究を行っていますが、いい結果が出るとは限りません（いい結果が出ると思ってやってはいますが）。
IMAX2は5年以上前からその構想が存在し、実装されていました。しかし、さらなるハードウェアの改良を重ね、今に至ったのです。そこらへんで止まっていたわけではありません。また、IMAX2の2という数字が表しているように、これ以前のハードウェアも存在していました。つまり、これでいいと思ってないから誕生したわけです。
自作ハードウェアのススメ
序論が長くなりました。「何故ハードウェアを理解し、時には自分で設計する必要があるのか」を力説するためには本一冊が必要なくらい、一般の方々に理解してもらうのは大変です。
「外で買ってきた方が安い」と思うかもしれません。実際短期的観点から見たらそうだし、陳腐化した大きな組織の運営においては真っ当なやり方です。しかし、「外で買ってくる」はリスクの高いやり方なのを知っておくべきです。小さい企業だったらお金で会社ごと買っちゃえばいい話ですが、IntelやAMDなどの大企業を会社ごと買収するのは不可能です。そしてその「外」が自分の都合のいいハードウェアを作ってくれるかも予想できません。
別に安保の観点だけではなく、国際情勢により為替が変動したりして「外」のものが高くなったりします。今はある程度戻りましたが、一時期1ドル150円まで円安が進み、今でも決して昔の水準に戻ったとは言えません。高くて買えない、そもそも売ってくれないなど、いつも買えるとは限りません。
企業や国を挙げての製品開発だけではなく、個人の趣味領域においても、自作ハードウェアは役に立ちます。いつも企業が都合のいいハードウェアを出してくれるとは限りません。GPUの並列性能は一般レベルにおいて十二分に優れていますが、消費電力は大きく、常時稼働してAIを動かすのは「逸般の誤家庭」でも結構な負担になるくらいです。
与太話が長くなりました。本論に入りましょう。
チップ製造は無理、じゃどうする？
個人がTSMCとかに委託してチップを作るのは孫正義のような富豪にしかできません。最近は無料または低価格でチップ焼いてくれるようなプロジェクトも登場していますが、好きな時に好きなチップを作ってくれるわけではありません。また、製造プロセスは完成まで長い期間を要し、すぐポンとできるわけではありません。
「無理ならどうする？諦める？」と思うかもしれません。しかし、個人レベルでも自作ハードウェアを動かせる方法ではあります。上で説明したIMAX2でも用いてる方法、FPGAです。
FPGAで自作ハードウェア
そもそもFPGAとは
FPGAはField-Programmable Gate Arrayの略です。「プログラミング可能な論理ゲートの集合」と思ってもらったらいいです。アナログ的な信号を扱わない場合、ゲートが足りるのであればほとんどの回路は実装できます。物理的な限界などが指摘され、また後述する問題から一般向けデバイスでは採用されていませんが...
どうやって実装すればいいか
FPGAで論理回路を実装したい場合、HDL(Hardware Description Language, 日本語では「ハードウェア記述言語」)を用います。XilinxもIntelもVerilogとVHDLは対応しているため、この二つの中で選べばいいと思います。
VerilogもVHDLもそうですが、RTL(Register Transfer Level, 日本語では「レジスタ転送レベル」)で記述します。ゲートレベルよりは抽象化されているため、直接ANDゲートから設計する必要はほとんどありません。VerilogやVHDLの文法紹介は割愛しますが、頭に入れておいた方がいい設計思想については言及しておきます。
論理回路の講義を受講した情報系大学生なら知っていると思いますが、AND/OR/XORのような論理ゲートと前の値を保持できるフリップフロップがあります。この値を保持できるフリップフロップのようなレジスタを中心にその振る舞いを記述していくのがRTLです。プログラミングではなく、回路設計として考えてやればいいと思います。
どこで買えるか
FPGAはいろんなメーカーが出していますが、メジャーなのはXilinxとIntelです。シェア的にXilinx一択な気もしますが、ツールの使いやすさ・ドキュメントのわかりやすさなどを考えてIntelを選ぶのもいい選択です。
また、FPGA関係のツールチェーンは非常に大きいので、それなりに大きいストレージを用意した方がいいです。VivadoとVitisを入れるだけで100GB超え...
最近は、KV260、Ultra96など、学生にしては高いけど買えない程度ではないFPGA評価ボードが増えています。入手困難になっているものも多いですが、SoC FPGAでなければ２万未満で手に入るものもあります。それで何かを高速化するのは無理かもしれませんが、自作CPUで遊ぶことは十二分にできます。また、IoTデバイス自作において、小さいFPGAを用いることによりセンサーのデータ処理の高速化、モーターのリアルタイム制御もできるようになるので、ここでも自作ハードウェアを作る価値を再確認することができます。
自作ハードウェアはPython書くよりは難しいけど普通の人でもできる
もちろんハードウェア設計及びそれを実際動かすようなソフトウェアの開発はあるものを利用するよりは難しいです。しかし、普通の人でも自作ハードウェアはできます。僕自身も、2年前まではVerilogの存在すら知らず、なんとなく「自作CPU作れる板があるぞ」と知っていただけでした。また、学科の知能情報工学科だったため、授業でFPGAを触る機会もなく、ハードウェア設計に関する知識を講義で身につけたわけではありません。しかし、その一年後、初の独自設計CPU「Ayumu」、その次は初の自作RISC-V「Kasumi」を完成するなど、欠陥は多いけど自作ハードウェアを設計できるようになりました。間が空きすぎて「開発にめっちゃ時間かかるし、諦めるか」と思うかもしれませんが、僕も自作CPUだけやっていたわけではないので、実際開発に使った時間は一週間よりも短いです。
ツールチェーンを整備するのはかなり時間のかかることですが、自分で動かすくらいだったらツールチェーンは後回しにし、動かせる最低限のソフトウェアだけを用意するのも方法で、そんなに工数を要する作業でもありません。
Pythonのようなユーザビリティをよく考えて作った言語よりは断然難しいのですが、Pythonでプログラムを高速に動かすのは相当な工数を要します。むしろ専用ハードウェアを設計するより難しいかもしれません。安易に「誰かが高速なライブラリ作ってくれるっしょ」と思ってはいけません。自分が作ろうとするものは既存のライブラリでは実装できないかもしれませんし、ゼロから自分で作るしかありません。そうなるととても簡単だったPythonが姿を変えてとても難解な言語に化けます。
「俺には無理」と思うかもしれません。まずVerilogを書いて、何かしら動かしてみてからそう判断したのであればこれ以上はいいませんが、まず試してみてほしいです。資金的に無理な方は、2500円程度のFPGAも存在するので、まずこれで小さなハードウェアを実装してみるのはどうでしょうか。九工大のある福岡県の最低賃金は900円です。最低賃金もらっていても3時間働けば普通に買えます。
まずやってみろ！失敗してもそんなに損しない！
終わりに
長い与太話を挟みながら自作ハードウェアをやることについての大切さについて説明しました。単に動くだけではアイデア以上の価値は生み出せません。また、PyTorchやTensorflowなどの便利なライブラリがどこでも使えるわけではありません。
また、単に速度の側面だけではなく、セキュリティ、省電力など、目的に合わせて必要なハードウェアを設計し、実装することもできます。汎用性は売る側としてはいいのですが、使う側からしたらいつもいいとは言えません。
遊び心で自作ハードウェアをやっても楽しいです。できないと思い込んでいたものが可能になったり、大手メーカーが設計したアーキテクチャを利用せずとも自らの力で同じ仕組みを再現したという達成感を得ることもできます。
以上、今年３月に九工大を卒業した老害からの戯言は終わりです。みなさんも自作ハードウェア、やりませんか？
明日はCatくんの「Azureでマイブログ作ろうぜ👀」です！お楽しみに！
参考文献
[1] 藤澤誠、「CUDAで行列積」、筑波大学物理ベースコンピュータグラフィックス研究室PukiWiki、http://www.slis.tsukuba.ac.jp/~fujisawa.makoto.fu/cgi-bin/wiki/index.php?CUDA
[2] 中島康彦、「EMAX6/ZYNQ64 (IMAX2) Architecture Handbook」、奈良先端科学技術大学院大学コンピューティング・アーキテクチャ研究室、http://archlab.naist.jp/proj-arm64/doc/emax6/emax6j.pdf、2022.11.1更新