Upgrade to Pro — share decks privately, control downloads, hide ads and more …

OpenCL Programming for FPGA

Avatar for Takuro IIZUKA Takuro IIZUKA
February 01, 2015

OpenCL Programming for FPGA

Avatar for Takuro IIZUKA

Takuro IIZUKA

February 01, 2015
Tweet

More Decks by Takuro IIZUKA

Other Decks in Technology

Transcript

  1. __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を取得
  2. N は 2, 3, 4, 8, 16のいずれか int2 add2(int2 x,

    int2 y) { return x + y; }; charN ucharN shortN ushortN intN uintN longN ulongN floatN
  3. Workgroup Local Workitem Private Workitem Private Global Constant Workgroup Local

    Workitem Private Workitem Private Private:単一Workitemから読み書きできる Local:Workgroup内のWorkitemから読み書きできる Constant:全てのWorkitemから読める Global:全てのWorkitemから読み書きできる 4つのメモリ空間
  4. __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]
  5. __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パイプライニングによって 命令レベルで並列に動作する
  6. タイミングアキュレートな実装はできない Workitem 0 ld ld add st Workitem 1 ld

    ld add st Workitem2 ld ld add st … このレイテンシはコンパイラが決定し、プログラマは制御できない
  7. パイプラインストールによる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
  8. __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
  9. __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
  10. フィードバックによるループの実装 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; }
  11. __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
  12. __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
  13. __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
  14. __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 …
  15. 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]; } }
  16. Board Nallatech PCIe385/395の場合 DRAM Global Constant Stratix V Local Private

    Constant Cache ボードベンダが配布する Board Support Packageによってメモリ空間を定義
  17. I/O

  18. FPGA DRAM SFP+ DRAM PCIe Root Complex CPU BRAM OpenCLベンダ拡張による外部I/O

    Kernel内のGlobalメモリI/O Kernel内のLocalメモリI/O OpenCL APIによるデータ転送 PCIe Board
  19. メモリI/OのCoalescing ld 4byte 0x1000 ld 4byte 0x1004 ld 8byte 0x1000

    連続アドレスのメモリトランザクションを 結合(Coalescing)することでバス効率を上げる
  20. __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; }
  21. 連続するワークアイテムが連続アドレスにアクセスすることで プリフェッチ&ストアキューのついた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と等価な書き方
  22. 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
  23. ロードストアは全て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; … }
  24. 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で計算できている