OpenCL Programming for FPGA

Transcript

1. OpenCL Programming for FPGA 株式会社フィックスターズ @iitaku
2. フィックスターズという会社で ソフトウェアエンジニアやってます
3. 弊社の お仕事
4. Cell/B.E.でマイクロベンチマーク してソフトウェア最適化する簡単なお仕事
5. GPUでマイクロベンチマーク してソフトウェア最適化する簡単なお仕事
6. x86でマイクロベンチマーク してソフトウェア最適化する簡単なお仕事
7. POWERで (ry
8. ARMで (ry
9. MICで
10. Tileraで
11. Parallellaで
12. FPGAで
13. FPGAで
14. FPGA
15. Ͱ΋ɺ
    )%-͚ͩ͸ɺ
    ઈରʹɺࢮΜͰ΋ɺΠϠ
16. 殺伐としたTLに 並列コンピューティングフレームワークが！ ＿人人人人人＿ ＞　OpenCL　＜ ￣Y^Y^Y^Y^Y￣
17. 至上命題 Performance
18. 下位レイヤ意識して プログラミングしてますか？
19. アセンブリ意識して プログラミングしてますか？
20. マイクロアーキテクチャ意識して プログラミングしてますか？
21. Question OpenCL for FPGAの下位レイヤとは？
22. Answer 論理回路である
23. 本日の内容
24. OpenCLアーキテクチャ OpenCL for FPGA実装 SHA-256 実装例
25. OpenCLアーキテクチャ OpenCL for FPGA実装 SHA-256 実装例
26. OpenCLホスト 制御用の汎用CPU OpenCLデバイス 演算用のCPU、GPU、FPGA ヘテロジニアス・システム ＋
27. OpenCL C Languageで 演算部分を実装 OpenCL C APIで デバイスを制御
28. OpenCL C Language
29. C言語ベース シンタックス 組み込み関数 スレッドモデル メモリ空間の規定 ベクトル型 ベンダ 拡張
30. スレッドモデル
31. OpenCL Kernel Workgroup 0 Workitem 0 Workitem N-1 階層化されたスレッドモデル 実行コンテキスト Workgroup M-1 Workitem 0 Workitem N-1
32. __kernel void copy(__global int *src, __global int *dst) { const int i = get_global_id(0); dst[i] = src[i]; } src i:0 i:1 i:N-1 dst Workitemによる並行性 WorkitemのIDを取得
33. 同期機構
34. Workgroup内のWorkitemのバリア同期＋メモリオーダリングの保証 void barrier(cl_mem_fence_flags flags) barrier Workgroup Workitem 0 Workitem N-1
35. ベクトル型
36. N は 2, 3, 4, 8, 16のいずれか int2 add2(int2 x, int2 y) { return x + y; }; charN ucharN shortN ushortN intN uintN longN ulongN floatN
37. メモリ空間
38. Workgroup Local Workitem Private Workitem Private Global Constant Workgroup Local Workitem Private Workitem Private Private：単一Workitemから読み書きできる Local：Workgroup内のWorkitemから読み書きできる Constant：全てのWorkitemから読める Global：全てのWorkitemから読み書きできる 4つのメモリ空間
39. ベンダ拡張
40. #pragma OPENCL EXTENSION extension_name : enable 可搬性を破る禁断のおまじない cl_intel_accelerator cl_intel_motion_estimation … cl_amd_device_memory_flags cl_amd_media_ops cl_amd_svm … cl_altera_channels …
41. OpenCLアーキテクチャ OpenCL for FPGA実装 SHA-256 実装例
42. Altera SDK for OpenCL Xilinx SDAccel Environment
43. 今日は（主に） Alteraの話です
44. スレッドモデルの実装
45. __kernel void vecadd(__global int *a, __global int *b, __global int *c) { const int i = get_global_id(0); c[i] = a[i] + b[i]; } ASTがそのまま論理回路に変換される ld a[i] ld b[i] add st c[i]
46. __kernel void vecadd(__global int *a, __global int *b, __global int *c) { const int i = get_global_id(0); c[i] = a[i] + b[i]; } t Workitem 0 ld ld add st Workitem 1 ld ld add st Workitem2 ld ld add st … Workitemパイプライニングによって 命令レベルで並列に動作する
47. 多くのGPU実装と比べると 並列化モデルが逆転している Workitemの展開方向 命令の展開方向 GPU 空間 時間 FPGA 時間 空間
48. タイミングアキュレートな実装はできない Workitem 0 ld ld add st Workitem 1 ld ld add st Workitem2 ld ld add st … このレイテンシはコンパイラが決定し、プログラマは制御できない
49. 同期機構の実装
50. パイプラインストールによるWorkitem間同期 __kernel void vecadd(__global int *a, __global int *b, __global int *c) { const int i = get_global_id(0); const int av = a[i]; const int bv = b[i]; barrier(CLK_LOCAL_MEM_FENCE); c[i] = av + bv; } Workitem 0 ld ld add st Workitem 1 ld ld add st ・・・ Workitem N-1 ld ld add st barrier t
51. 分岐の実装
52. __kernel void branch(__global int *a, __global int *b, __global int *c) { const int i = get_global_id(0); if ((i % 2) == 0) { c[i] = a[i] + b[i]; } else { c[i] = a[i] * b[i]; } } マルチデータパスとセレクトによる分岐の実装 ld a[i] ld b[i] add st c[i] mul select
53. __kernel void branch(__global int *a, __global int *b, __global int *c) { const int i = get_global_id(0); if ((i % 2) == 0) { c[i] = a[i] + b[i]; } else { c[i] = a[i] * b[i]; } } パイプライン性能は遅い方の分岐ブロックに依存 Workitem 0 ld ld add mul select st Workitem1 ld ld add mul select st Workitem2 ld ld add mul select st … t
54. ループの実装
55. フィードバックによるループの実装 ld src[k] add st dst[i] accum j<N __kernel void loop(__global const int *src, __global int *dst) { const int i = get_global_id(0); int accum = 0; for (uint j=0; j<N; ++j) { accum += src[N * i + j]; } dst[i] = accum; }
56. __kernel void loop(__global const int *src, __global int *dst) { const int i = get_global_id(0); int accum = 0; for (uint j=0; j<N; ++j) { accum += src[N * i + j]; } dst[i] = accum; } ループ1回目 性能は1/ループ回数になる Workitem 0 ld add ld add st Workitem 1 ld add … ld add st Workitem 2 ld add ld add st … t
57. __kernel void loop(__global const int *src, __global int *dst) { const int i = get_global_id(0); int accum = 0; #pragma unroll 2 for (uint j=0; j<N; ++j) { accum += src[N * i + j]; } dst[i] = accum; } ループ1, 2回目 ループアンロールすることでスループット向上 Workitem 0 ld ld add add ld ld add add st Workitem 1 ld ld add add … ld ld add add st Workitem 2 ld ld add add ld ld add add st … t
58. ベクトル型の実装
59. __kernel void vecadd(__global int4 *a, __global int4 *b, __global int4 *c) { const int i = get_global_id(0); c[i] = a[i] + b[i]; } int4 ベクトル演算回路が生成され、リソース消費は増大 vld a[i] vld b[i] vadd vst c[i] int4 int4
60. __kernel void vecadd(__global int4 *a, __global int4 *b, __global int4 *c) { const int i = get_global_id(0); c[i] = a[i] + b[i]; } 性能はベクトル幅倍に増加 t Workitem 0 vld vld vadd vst Workitem 1 vld vld vadd vst Workitem 2 vld vld vadd vst …
61. Loop Carried Dependencyの無いループを フルアンロールするのと等価 __kernel void vecadd(__global int *a, __global int *b, __global int *c) { const int i = get_global_id(0); #pragma unroll for (int j=0; j<4; ++j) { c[4*i+j] = a[4*i+j] + b[4*i+j]; } }
62. メモリ空間の実装
63. Board Nallatech PCIe385/395の場合 DRAM Global Constant Stratix V Local Private Constant Cache ボードベンダが配布する Board Support Packageによってメモリ空間を定義
64. I/O
65. FPGA DRAM SFP+ DRAM PCIe Root Complex CPU BRAM OpenCLベンダ拡張による外部I/O Kernel内のGlobalメモリI/O Kernel内のLocalメモリI/O OpenCL APIによるデータ転送 PCIe Board
66. SoC DRAM CPU FPGA I/O Peripheral BRAM OpenCLベンダ拡張による外部I/O Kernel内のGlobalメモリI/O Kernel内のLocalメモリI/O OpenCL APIによるデータ転送 System on Chip
67. メモリI/OのCoalescing ld 4byte 0x1000 ld 4byte 0x1004 ld 8byte 0x1000 連続アドレスのメモリトランザクションを 結合（Coalescing）することでバス効率を上げる
68. __kernel void add_adj(__global const int *src, __global int *dst) { const int i = get_global_id(0); dst[i] = src[2*i+0] + src[2*i+1]; } ld src [2*i+0] ld src [2*i+1] add st dst[i] add st dst[i] vld src [2*i] メモリI/Oを明示的にベクトル化することでCoalescingする __kernel void add_adj(__global const int *src,__ global int *dst) { const int i = get_global_id(0); const int2 v = vload(&src[2*i]); dst[i] = v.x + v.y; }
69. 連続するワークアイテムが連続アドレスにアクセスすることで プリフェッチ＆ストアキューのついたLSUを作成 OpenCL Kernel Global Memory streaming ld streaming st __kernel void vecadd(__global int4 *a, __global int4 *b, __global int4 *c) { const int i = get_global_id(0); c[i] = a[i] + b[i]; } streaming ld 後続ワークアイテムのアクセスが連続する GPUにおけるCoalescingと等価な書き方
70. Alteraチャネル拡張を使用することで HDLモジュールのAvalon STインタフェースに接続できる #pragma OPENCL_EXTENSION cl_altera_channels : enable channel int input_ch __attribute__((io(“input_ch0”))); channel int output_ch __attribute__((io(“output_ch0”))); __kernel void echo(void) { int v = read_channel_altera(input_ch); write_channel_altera(output_ch, v); } FPGA I/O Peripheral OpenCL Kernel HDL I/O Module Avalon ST
71. PCIe背面のSFP+を経由して外部I/O
72. OpenCLアーキテクチャ OpenCL for FPGA実装 SHA-256 実装例
73. Workitemあたりの動作 1. Globalメモリから32バイト読む 2. SHA-256でハッシュを計算 3. Globalメモリに32バイト書きこむ __kernel __attribute__((reqd_work_group_size(1, 1, 1))) void sha256(const __global uchar * restrict src, __global uchar * restrict dst) { const uint i = get_global_id(0); … }
74. ロードストアは全てCoalescing #pragma unroll for (int j=0; j<32; ++j) { buffer[j] = src[i*32+j]; } #pragma unroll for (int j=0; j<8; ++j) { uchar4 v = scatter_in_be(hs[j]); vstore4(v, 0, &dst[(i*8+j)*4]); } Hashの計算はフルアンローリング #pragma unroll for (int j=0; j<64; ++j) { const uint s1 = rotate_right(e, 6) ^ rotate_right(e, 11) ^ rotate_right(e, 25); const uint ch = (e & f) ^ (~e & g); const uint temp1 = h + s1 + ch + keys[j] + ws[j]; const uint s0 = rotate_right(a, 2) ^ rotate_right(a, 13) ^ rotate_right(a, 22); const uint maj = (a & b) ^ (a & c) ^ (b & c); const uint temp2 = s0 + maj; … }
75. スループット（MHash/s） 120 150 180 210 240 ハッシュ数（Kilo） 0 17500 35000 52500 70000 PCIe385で 230 MHash/s程度で安定
76. ALUTs: 67265 Registers: 109,475 Logic utilization: 59,476 / 172,600 ( 34 % ) I/O pins: 384 / 664 ( 58 % ) DSP blocks: 0 / 1,590 ( 0 % ) Memory bits: 1,695,184 / 41,246,720 ( 4 % ) RAM blocks: 335 / 2,014 ( 17 % ) Actual clock freq: 268.599999201 Kernel fmax: 268.6 1x clock fmax: 268.6 2x clock fmax: 10000 Highest non-global fanout: 3167 ALMsは34%程度を使用 Kernel Clockは268 MHz Workitem(=1ハッシュ)あたり1.2cycleで計算できている
77. まとめ
78. OpenCL知ってるソフトウェア屋なら FPGAプログラミングは普通にできる OpenCLも別に難しくない コンパイル時間はもっと 短くなってほしい・・・ 最適化するなら VerilogとLLVM IRは 読めるほうがいい
79. タイミングアキュレートな 回路を作るのは難しい ベンダ拡張でHDLと混ぜれるので 部分的に活用するとか なんでもできるわけではない
80. 単なるアクセラレータとしては ハマる分野を選ぶ 外部I/Oやレイテンシ重視ならアリ 電力性能比やコスト性能比を 総合して考えるべき スループット重視なら GPUも検討すべき
81. Happy Programming !