open-mncore2-sdk について、基本的な使い方をまとめます。
open-mncore2-sdk とは
open-mncore2-sdkSinby が開発している、MN-Core2 向けの非公式開発環境です。 PFN 公式配布のエミュレータ環境と組み合わせて使います。
とりあえず N-Body が動くところまで来ましたよ。
なお、いろいろ歴史的経緯でまぎららわしいですが、本家の mncl やmncore sdk とは一切関係ありません。
PFCP と Kubernetes
現状、実行環境は#(:a :href "https://pfcomputing.com/" ”PFCP")しかないと思います。Kubernetes(k8s とも) の環境なのでDocker しか使ったことがない私はちょっと苦労しました。
Kubernetes
Kubernetes は Docker のイメージを動作させることのできるコンテナ環境のようです。いまや Docker のイメージはもともとの Docker から離れて、いろんな動作環境があるようで(詳しくない)、Kubernetes もその１つです。
Kubernetes はその動作環境の再現性が高いのが特徴で、構成情報をコードとして定義する「宣言的設定（Declarative Configuration）」と、稼働中の環境を変更せず丸ごと差し替える「イミュータブル・インフラストラクチャ（Immutable Infrastructure）」という設計思想に基づいているためです。
最近は、言語もDBも”イミュータブル"という言葉をよく使うので特別違和感はないでしょう。簡単に言えば静的で、あとから書き換えができないということです。
私は Docker の使い方としてイメージを作って、そのなかに入ってbash で作業をするという使い方を良くします。Kubernetes もそのような使い方はできます。
ただ、Kubernetes 自体はイミュータブルにサービスを作ってそれを起動あるいはHPCならバッチ処理をさせる、という使い方をするようです。
Pod
Kubernetes という新しい(わたしにとって)トピックでとっつきにくいのがその中で使われる専用用語です。Pod もその一つで、簡単に言うと「実行時の実体（最小のデプロイ単位）」のことです。
なんで Pod なのかは不明ですが、Pod をつくって起動すれば Docker 的に使える！！ということを覚えておけば当面は十分でしょう。
イミュータブルであること
Kubernetes の基本がイミュータブルであるとうことは Docker とは運用が違ってくるという点に注意しなければなりません。具体的にはコンテナ内で /home や /root などにファイルをを作成してもPod が再起動したり、再作成されたりすると、保存したはずの情報はきれいさっぱりなくなってしまうのです。
そのため、奥の手の k8s で用意される PVC(PersitentVolumeClaim) という永続性のあるオブジェクトを使う必要があります。PVC がなにものかはあまり理解していませんが、File System として mount でき、Pod が消えてもデータは保存されます(永続的だけに)。Docker 的な使い方をする私としてはこの点が非常に重要です。
PVC を使う
PVC のもう一つの使い方として /opt などのツール提供です。今回の open-mncore2-sdk では当面、PVC にツール群をおいて、/opt にマウントし、さらに、/root に PVC もマウントしてます。
PVC をつくる
PVC をつくるの簡単かな？とおもったですが、いまそのシェルプロを見てみると結構複雑なことしてますね。そのうち gist とかで公開しましょう。
1. 既存の Deployment を削除（MN-Core2 枠を解放）
2. loader pod を起動（リトライ付き）
3. 古い rc を削除（PVC の容量確保。今デプロイする SUB_PATH は除外）
4. SDK をコピー（tar で中身だけ転送）
5. sdk.ninja を書き込み
6. activate を修正（sdk.ninja が既にあればスキップ）
7. loader pod を削除
基本的な使い方
PCV を２つ無事作って、/opt にツール群が /root にホームのようなものができたとしましょう。
まずおもむろに activate します。これは open-mncore2-sdk に含まれているbash 用のスクリプトです。Python の venv みたいなのをイメージしました。
text
source /opt/mncore2-sdk/bin/activate
以後、OpenCL てきな clang が使えるようになります。ここではサンプル環境をつくり、コンパイルしてみます。
text
$ /opt/mncore2-sdk/bin/sdk-examples mncl
$ cd mncl-0.3.3
$ lit -j1 --order=lexical lit/test-device
-- Testing: 19 tests, 1 workers --
PASS: MNCL examples :: test-device/01-vecadd.test (1 of 19)
PASS: MNCL examples :: test-device/02-vecadd-2d.test (2 of 19)
PASS: MNCL examples :: test-device/03-vecadd-3d.test (3 of 19)
PASS: MNCL examples :: test-device/04-vecadd-3d-batch.test (4 of 19)
PASS: MNCL examples :: test-device/05-stencil1d.test (5 of 19)
PASS: MNCL examples :: test-device/06-vecadd-f32.test (6 of 19)
PASS: MNCL examples :: test-device/07-reduce-add.test (7 of 19)
PASS: MNCL examples :: test-device/08-reduce-max.test (8 of 19)
PASS: MNCL examples :: test-device/09-abs.test (9 of 19)
PASS: MNCL examples :: test-device/10-odd-even-sort.test (10 of 19)
PASS: MNCL examples :: test-device/11-vecdiv-with-newton-raphson.test (11 of 19)
PASS: MNCL examples :: test-device/12-vecadd-i64.test (12 of 19)
PASS: MNCL examples :: test-device/13-odd-even-sort-gid.test (13 of 19)
PASS: MNCL examples :: test-device/14-odd-even-sort-full.test (14 of 19)
PASS: MNCL examples :: test-device/15-broadcast.test (15 of 19)
PASS: MNCL examples :: test-device/16-broadcast-f64.test (16 of 19)
PASS: MNCL examples :: test-device/17-nbody-2d.test (17 of 19)
PASS: MNCL examples :: test-device/18-broadcast-vec.test (18 of 19)
PASS: MNCL examples :: test-device/19-nbody-3d.test (19 of 19)
Testing Time: 0.33s
Total Discovered Tests: 19
Passed: 19 (100.00%)
$ ls
01-vecadd 09-abs 16-broadcast-f64
02-vecadd-2d 10-odd-even-sort 17-nbody-2d
03-vecadd-3d 11-vecdiv-with-newton-raphson 18-broadcast-vec
04-vecadd-3d-batch 12-odd-even-sort-full 19-nbody-3d
05-stencil1d 12-vecadd-i64 README.md
06-vecadd-f32 13-odd-even-sort-gid data
07-reduce-add 14-odd-even-sort-full lit
08-reduce-max 15-broadcast scripts
sdk-examples コマンドでサンプルをコピーしてくるので、cd mncl-0.3.3 して、サンプルのあるディレクトリに行きます。ここでは lit で全サンプルをテストしてしまいました。
ここではサンプル内容は説明しません。たぶん README.md や各ディレクトリの README.md に説明があるはず。
サンプル
あるはずだけだとなんなんで、一応、サンプルコード提示しておきます。説明はなしです。
odd-even-sort-gid-odd.cl
// 奇偶転置ソート 1 フェーズ (odd kernel、 get_global_id 版)。
//
// even と同じく、 左メンバ判定を get_global_id で行う。 例 10 の odd フェーズの flag は
// {0,1,0,1,...}（奇数 PE が左メンバ）なので、 「PE 番号が奇数 = 左メンバ」で置き換える。
// is_left = distribute(flag) -> is_left = ((id & 1) != 0) ? 1.0 : 0.0
// ping-pong の向き（data_in / data_out の割り当て）だけ even と逆にする（.param 参照）。
// 両端 (PE 0 と PE 4095) は例 10 と同じく boundary_flags の clamp で self 比較となり swap しない。
//
// data_in: in (このフェーズの入力配列。 distribute / neighbor が読む)
// data_out: out (1 フェーズ後の配列。 collect 先)
// swap_flags: out (PE ごとの swap 有無 0.0/1.0。 reduce kernel が読み直して集計する)
__kernel void odd_even_sort_gid_odd(__global double* data_in,
__global double* data_out,
__global double* swap_flags) {
long id = get_global_id(0);
double self = distribute(data_in);
double right = neighbor(data_in, 1);
double left = neighbor(data_in, -1);
double is_left = ((id & 1) != 0) ? 1.0 : 0.0; // 奇数 PE がペアの左メンバ
double left_swap;
if (self > right) { left_swap = 1.0; } else { left_swap = 0.0; }
double right_swap;
if (left > self) { right_swap = 1.0; } else { right_swap = 0.0; }
double do_left = is_left * left_swap;
double do_right = (1.0 - is_left) * right_swap;
double result;
if (do_left > 0.5) {
result = right;
} else {
result = self;
}
double result2;
if (do_right > 0.5) {
result2 = left;
} else {
result2 = result;
}
/* この PE が swap に関与したか (0 or 1)。 reduce kernel が全 PE を集計する */
double swap_indicator = do_left + do_right;
collect(swap_flags, swap_indicator);
collect(data_out, result2);
}
リンク集
open-mncore2-sdk として github で公開中 ただドキュメントがほぼない。
PFN が提供する Kubernetes ベースのコンテナ。 LLM うごいているのね。