Apache TVMは、CPU、GPU、そして様々な機械学習アクセラレーションチップに適した、エンドツーエンドのディープラーニング構築フレームワークです。中国語版のTVMドキュメントは、→ https://tvm.hyper.ai/ をご覧ください。
著者：Tianqi Chen

以下では、TVM (ニューラル ネットワークの一般的なパターン) で再帰計算を実行する方法について説明します。

 from __future__ import absolute_import, print_function import tvm import tvm.testing from tvm import te import numpy as np

TVMは線形演算子を用いて記号循環を記述します。以下の線形演算子は、列Xの累積和を計算します。

線形性はテンソルの最高次元で実行されます。`s_state` は線形遷移の状態を記述するプレースホルダーです。`s_init` は最初の k タイム ステップを初期化する方法を記述します。その最初の次元は 1 で、最初のタイム ステップの状態を初期化する方法を記述します。

`s_update` は、タイムステップ `t` における値の更新方法を記述します。更新された値は、状態プレースホルダを介して前のタイムステップの値から参照できます。ただし、現在のタイムステップまたはそれ以降のタイムステップで `s_state` を参照することは無効です。

線形ユニットには、状態プレースホルダ、初期値、および更新記述が含まれます。線形ユニットへの入力をリスト化することが推奨されます。線形ユニットの結果は、時間領域で更新されたテンソル（s_state）です。

 m = te.var("m") n = te.var("n") X = te.placeholder((m, n), name="X") s_state = te.placeholder((m, n)) s_init = te.compute((1, n), lambda _, i: X[0, i]) s_update = te.compute((m, n), lambda t, i: s_state[t - 1, i] + X[t, i]) s_scan = tvm.te.scan(s_init, s_update, s_state, inputs=[X])

線形ユニットのスケジュール

線形ボリュームは、更新部分と初期部分を別々にスケジュールすることでスケジュールされます。スケジュールされた更新部分の最初の反復次元は無効であることに注意してください。反復を時間的に分割するには、ユーザーはscan_op.scan_axisに基づいてスケジュールできます。

 s = te.create_schedule(s_scan.op) num_thread = 256 block_x = te.thread_axis("blockIdx.x") thread_x = te.thread_axis("threadIdx.x") xo, xi = s[s_init].split(s_init.op.axis[1], factor=num_thread) s[s_init].bind(xo, block_x) s[s_init].bind(xi, thread_x) xo, xi = s[s_update].split(s_update.op.axis[1], factor=num_thread) s[s_update].bind(xo, block_x) s[s_update].bind(xi, thread_x) print(tvm.lower(s, [X, s_scan], simple_mode=True))

出力結果:

 @main = primfn(X_1: handle, scan_1: handle) -> () attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True} buffers = {X: Buffer(X_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"), scan: Buffer(scan_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")} buffer_map = {X_1: X, scan_1: scan} preflattened_buffer_map = {X_1: X_3: Buffer(X_2, float32, [m, n: int32], [stride, stride_2: int32], type="auto"), scan_1: scan_3: Buffer(scan_2, float32, [m, n], [stride_1, stride_3: int32], type="auto")} { attr [IterVar(blockIdx.x: int32, (nullptr), "ThreadIndex", "blockIdx.x")] "thread_extent" = floordiv((n + 255), 256); attr [IterVar(threadIdx.x: int32, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 256; if @tir.likely((((blockIdx.x*256) + threadIdx.x) < n), dtype=bool) { scan[(((blockIdx.x*256) + threadIdx.x)*stride_3)] = X[(((blockIdx.x*256) + threadIdx.x)*stride_2)] } for (scan.idx: int32, 0, (m - 1)) { attr [IterVar(blockIdx.x, (nullptr), "ThreadIndex", "blockIdx.x")] "thread_extent" = floordiv((n + 255), 256); attr [IterVar(threadIdx.x, (nullptr), "ThreadIndex", "threadIdx.x")] "thread_extent" = 256; if @tir.likely((((blockIdx.x*256) + threadIdx.x) < n), dtype=bool) { let cse_var_1: int32 = (scan.idx + 1) scan[((cse_var_1*stride_1) + (((blockIdx.x*256) + threadIdx.x)*stride_3))] = (scan[((scan.idx*stride_1) + (((blockIdx.x*256) + threadIdx.x)*stride_3))] + X[((cse_var_1*stride) + (((blockIdx.x*256) + threadIdx.x)*stride_2))]) } } }

構築と検証

線形カーネルは他の TVM カーネルと同様に構築できます。ここでは、結果の正確性を検証するために NumPy が使用されます。

 fscan = tvm.build(s, [X, s_scan], "cuda", name="myscan") dev = tvm.cuda(0) n = 1024 m = 10 a_np = np.random.uniform(size=(m, n)).astype(s_scan.dtype) a = tvm.nd.array(a_np, dev) b = tvm.nd.array(np.zeros((m, n), dtype=s_scan.dtype), dev) fscan(a, b) tvm.testing.assert_allclose(b.numpy(), np.cumsum(a_np, axis=0))

多段リニアユニット

上記の例は、s_update のテンソル計算ステージを使用する線形ユニットについて説明しており、線形ユニットでは複数のテンソル レベルを使用できます。

次のコードは、2 段階の操作を含む線形ユニット内の線形プロセスを示しています。

 m = te.var("m") n = te.var("n") X = te.placeholder((m, n), name="X") s_state = te.placeholder((m, n)) s_init = te.compute((1, n), lambda _, i: X[0, i]) s_update_s1 = te.compute((m, n), lambda t, i: s_state[t - 1, i] * 2, name="s1") s_update_s2 = te.compute((m, n), lambda t, i: s_update_s1[t, i] + X[t, i], name="s2") s_scan = tvm.te.scan(s_init, s_update_s2, s_state, inputs=[X])

これらの中間テンソルは通常通りスケジュールできます。正確性を保証するために、TVMは線形ループ外のcompute_at位置で線形ボリュームを無効にするグループ制約を作成します。

 s = te.create_schedule(s_scan.op) xo, xi = s[s_update_s2].split(s_update_s2.op.axis[1], factor=32) s[s_update_s1].compute_at(s[s_update_s2], xo)

出力結果:

 print(tvm.lower(s, [X, s_scan], simple_mode=True)) @main = primfn(X_1: handle, scan_1: handle) -> () attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True} buffers = {X: Buffer(X_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"), scan: Buffer(scan_2: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto")} buffer_map = {X_1: X, scan_1: scan} preflattened_buffer_map = {X_1: X_3: Buffer(X_2, float32, [m, n: int32], [stride, stride_2: int32], type="auto"), scan_1: scan_3: Buffer(scan_2, float32, [m, n], [stride_1, stride_3: int32], type="auto")} { allocate(s1: Pointer(global float32), float32, [32]), storage_scope = global { for (i: int32, 0, n) { scan[(i*stride_3)] = X[(i*stride_2)] } for (scan.idx: int32, 0, (m - 1)) { for (i.outer: int32, 0, floordiv((n + 31), 32)) { for (i_1: int32, 0, 32) { if @tir.likely((((i.outer*32) + i_1) < n), dtype=bool) { s1_1: Buffer(s1, float32, [32], [])[i_1] = (scan[((scan.idx*stride_1) + (((i.outer*32) + i_1)*stride_3))]*2f32) } } for (i.inner: int32, 0, 32) { if @tir.likely((((i.outer*32) + i.inner) < n), dtype=bool) { let cse_var_2: int32 = (scan.idx + 1) let cse_var_1: int32 = ((i.outer*32) + i.inner) scan[((cse_var_2*stride_1) + (cse_var_1*stride_3))] = (s1_1[i.inner] + X[((cse_var_2*stride) + (cse_var_1*stride_2))]) } } } } } }

複数の状態

RNNのような複雑なアプリケーションでは、複数の再帰状態が必要です。Linearは複数の再帰状態をサポートしています。次の例は、2つの状態を持つ再帰を構築する方法を示しています。

 m = te.var("m") n = te.var("n") l = te.var("l") X = te.placeholder((m, n), name="X") s_state1 = te.placeholder((m, n)) s_state2 = te.placeholder((m, l)) s_init1 = te.compute((1, n), lambda _, i: X[0, i]) s_init2 = te.compute((1, l), lambda _, i: 0.0) s_update1 = te.compute((m, n), lambda t, i: s_state1[t - 1, i] + X[t, i]) s_update2 = te.compute((m, l), lambda t, i: s_state2[t - 1, i] + s_state1[t - 1, 0]) s_scan1, s_scan2 = tvm.te.scan( [s_init1, s_init2], [s_update1, s_update2], [s_state1, s_state2], inputs=[X] ) s = te.create_schedule(s_scan1.op) print(tvm.lower(s, [X, s_scan1, s_scan2], simple_mode=True))

出力結果:

 @main = primfn(X_1: handle, scan_2: handle, scan_3: handle) -> () attr = {"from_legacy_te_schedule": True, "global_symbol": "main", "tir.noalias": True} buffers = {X: Buffer(X_2: Pointer(float32), float32, [(stride: int32*m: int32)], [], type="auto"), scan: Buffer(scan_4: Pointer(float32), float32, [(stride_1: int32*m)], [], type="auto"), scan_1: Buffer(scan_5: Pointer(float32), float32, [(stride_2: int32*m)], [], type="auto")} buffer_map = {X_1: X, scan_2: scan, scan_3: scan_1} preflattened_buffer_map = {X_1: X_3: Buffer(X_2, float32, [m, n: int32], [stride, stride_3: int32], type="auto"), scan_2: scan_6: Buffer(scan_4, float32, [m, n], [stride_1, stride_4: int32], type="auto"), scan_3: scan_7: Buffer(scan_5, float32, [m, l: int32], [stride_2, stride_5: int32], type="auto")} { for (i: int32, 0, n) { scan[(i*stride_4)] = X[(i*stride_3)] } for (i_1: int32, 0, l) { scan_1[(i_1*stride_5)] = 0f32 } for (scan.idx: int32, 0, (m - 1)) { for (i_2: int32, 0, n) { let cse_var_1: int32 = (scan.idx + 1) scan[((cse_var_1*stride_1) + (i_2*stride_4))] = (scan[((scan.idx*stride_1) + (i_2*stride_4))] + X[((cse_var_1*stride) + (i_2*stride_3))]) } for (i_3: int32, 0, l) { scan_1[(((scan.idx + 1)*stride_2) + (i_3*stride_5))] = (scan_1[((scan.idx*stride_2) + (i_3*stride_5))] + scan[(scan.idx*stride_1)]) } } }

要約

このチュートリアルでは、線形プリミティブの使用方法を説明します。

• 線形性は init と update を使用して記述されます。
• リニアユニットは通常のスケジュールとしてスケジュールされます。
• 複雑なワークロードの場合、複数の状態とステップが線形単位で使用されます。

Pythonソースコードをダウンロードする: scan.py

Jupyter Notebook をダウンロード: scan.ipynb