Slide 1
Slide 1 text
OpenCL Programming for
FPGA
株式会社フィックスターズ
@iitaku
Slide 2
Slide 2 text
フィックスターズという会社で
ソフトウェアエンジニアやってます
Slide 3
Slide 3 text
弊社の
お仕事
Slide 4
Slide 4 text
Cell/B.E.でマイクロベンチマーク
してソフトウェア最適化する簡単なお仕事
Slide 5
Slide 5 text
GPUでマイクロベンチマーク
してソフトウェア最適化する簡単なお仕事
Slide 6
Slide 6 text
x86でマイクロベンチマーク
してソフトウェア最適化する簡単なお仕事
Slide 7
Slide 7 text
POWERで (ry
Slide 8
Slide 8 text
ARMで (ry
Slide 9
Slide 9 text
MICで
Slide 10
Slide 10 text
Tileraで
Slide 11
Slide 11 text
Parallellaで
Slide 12
Slide 12 text
FPGAで
Slide 13
Slide 13 text
FPGAで
Slide 14
Slide 14 text
FPGA
Slide 15
Slide 15 text
Ͱɺ
)%-͚ͩɺ
ઈରʹɺࢮΜͰɺΠϠ
Slide 16
Slide 16 text
殺伐としたTLに
並列コンピューティングフレームワークが！
＿人人人人人＿
＞ OpenCL ＜
￣Y^Y^Y^Y^Y￣
Slide 17
Slide 17 text
至上命題
Performance
Slide 18
Slide 18 text
下位レイヤ意識して
プログラミングしてますか？
Slide 19
Slide 19 text
アセンブリ意識して
プログラミングしてますか？
Slide 20
Slide 20 text
マイクロアーキテクチャ意識して
プログラミングしてますか？
Slide 21
Slide 21 text
Question
OpenCL for FPGAの下位レイヤとは？
Slide 22
Slide 22 text
Answer
論理回路である
Slide 23
Slide 23 text
本日の内容
Slide 24
Slide 24 text
OpenCLアーキテクチャ
OpenCL for FPGA実装
SHA-256 実装例
Slide 25
Slide 25 text
OpenCLアーキテクチャ
OpenCL for FPGA実装
SHA-256 実装例
Slide 26
Slide 26 text
OpenCLホスト
制御用の汎用CPU
OpenCLデバイス
演算用のCPU、GPU、FPGA
ヘテロジニアス・システム
＋
Slide 27
Slide 27 text
OpenCL C Languageで
演算部分を実装
OpenCL C APIで
デバイスを制御
Slide 28
Slide 28 text
OpenCL C Language
Slide 29
Slide 29 text
C言語ベース
シンタックス
組み込み関数
スレッドモデル メモリ空間の規定
ベクトル型
ベンダ
拡張
Slide 30
Slide 30 text
スレッドモデル
Slide 31
Slide 31 text
OpenCL Kernel
Workgroup
0
Workitem
0
Workitem
N-1
階層化されたスレッドモデル
実行コンテキスト
Workgroup
M-1
Workitem
0
Workitem
N-1
Slide 32
Slide 32 text
__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を取得
Slide 33
Slide 33 text
同期機構
Slide 34
Slide 34 text
Workgroup内のWorkitemのバリア同期＋メモリオーダリングの保証
void barrier(cl_mem_fence_flags flags)
barrier
Workgroup
Workitem
0
Workitem
N-1
Slide 35
Slide 35 text
ベクトル型
Slide 36
Slide 36 text
N は 2, 3, 4, 8, 16のいずれか
int2 add2(int2 x, int2 y)
{
return x + y;
};
charN
ucharN
shortN
ushortN
intN
uintN
longN
ulongN
floatN
Slide 37
Slide 37 text
メモリ空間
Slide 38
Slide 38 text
Workgroup
Local
Workitem
Private
Workitem
Private
Global
Constant
Workgroup
Local
Workitem
Private
Workitem
Private
Private：単一Workitemから読み書きできる
Local：Workgroup内のWorkitemから読み書きできる
Constant：全てのWorkitemから読める
Global：全てのWorkitemから読み書きできる
4つのメモリ空間
Slide 39
Slide 39 text
ベンダ拡張
Slide 40
Slide 40 text
#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
…
Slide 41
Slide 41 text
OpenCLアーキテクチャ
OpenCL for FPGA実装
SHA-256 実装例
Slide 42
Slide 42 text
Altera SDK for OpenCL
Xilinx SDAccel Environment
Slide 43
Slide 43 text
今日は（主に）
Alteraの話です
Slide 44
Slide 44 text
スレッドモデルの実装
Slide 45
Slide 45 text
__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]
Slide 46
Slide 46 text
__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パイプライニングによって
命令レベルで並列に動作する
Slide 47
Slide 47 text
多くのGPU実装と比べると
並列化モデルが逆転している
Workitemの展開方向 命令の展開方向
GPU 空間 時間
FPGA 時間 空間
Slide 48
Slide 48 text
タイミングアキュレートな実装はできない
Workitem 0
ld
ld
add st
Workitem 1
ld
ld
add st
Workitem2
ld
ld
add st
…
このレイテンシはコンパイラが決定し、プログラマは制御できない
Slide 49
Slide 49 text
同期機構の実装
Slide 50
Slide 50 text
パイプラインストールによる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
Slide 51
Slide 51 text
分岐の実装
Slide 52
Slide 52 text
__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
Slide 53
Slide 53 text
__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
Slide 54
Slide 54 text
ループの実装
Slide 55
Slide 55 text
フィードバックによるループの実装
ld
src[k]
add
st
dst[i]
accum
j
Slide 56
Slide 56 text
__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
Slide 57
Slide 57 text
__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
Slide 58
Slide 58 text
ベクトル型の実装
Slide 59
Slide 59 text
__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
Slide 60
Slide 60 text
__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
…
Slide 61
Slide 61 text
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];
}
}
Slide 62
Slide 62 text
メモリ空間の実装
Slide 63
Slide 63 text
Board
Nallatech PCIe385/395の場合
DRAM
Global
Constant
Stratix V
Local
Private
Constant
Cache
ボードベンダが配布する
Board Support Packageによってメモリ空間を定義
Slide 64
Slide 64 text
I/O
Slide 65
Slide 65 text
FPGA
DRAM
SFP+
DRAM
PCIe Root Complex
CPU
BRAM
OpenCLベンダ拡張による外部I/O Kernel内のGlobalメモリI/O
Kernel内のLocalメモリI/O
OpenCL APIによるデータ転送
PCIe Board
Slide 66
Slide 66 text
SoC
DRAM
CPU
FPGA
I/O Peripheral
BRAM
OpenCLベンダ拡張による外部I/O Kernel内のGlobalメモリI/O
Kernel内のLocalメモリI/O
OpenCL APIによるデータ転送
System on Chip
Slide 67
Slide 67 text
メモリI/OのCoalescing
ld 4byte 0x1000
ld 4byte 0x1004
ld 8byte 0x1000
連続アドレスのメモリトランザクションを
結合（Coalescing）することでバス効率を上げる
Slide 68
Slide 68 text
__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;
}
Slide 69
Slide 69 text
連続するワークアイテムが連続アドレスにアクセスすることで
プリフェッチ＆ストアキューのついた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と等価な書き方
Slide 70
Slide 70 text
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
Slide 71
Slide 71 text
PCIe背面のSFP+を経由して外部I/O
Slide 72
Slide 72 text
OpenCLアーキテクチャ
OpenCL for FPGA実装
SHA-256 実装例
Slide 73
Slide 73 text
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);
…
}
Slide 74
Slide 74 text
ロードストアは全て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;
…
}
Slide 75
Slide 75 text
スループット（MHash/s）
120
150
180
210
240
ハッシュ数（Kilo）
0 17500 35000 52500 70000
PCIe385で 230 MHash/s程度で安定
Slide 76
Slide 76 text
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で計算できている
Slide 77
Slide 77 text
まとめ
Slide 78
Slide 78 text
OpenCL知ってるソフトウェア屋なら
FPGAプログラミングは普通にできる
OpenCLも別に難しくない
コンパイル時間はもっと
短くなってほしい・・・
最適化するなら
VerilogとLLVM IRは
読めるほうがいい
Slide 79
Slide 79 text
タイミングアキュレートな
回路を作るのは難しい
ベンダ拡張でHDLと混ぜれるので
部分的に活用するとか
なんでもできるわけではない
Slide 80
Slide 80 text
単なるアクセラレータとしては
ハマる分野を選ぶ
外部I/Oやレイテンシ重視ならアリ
電力性能比やコスト性能比を
総合して考えるべき
スループット重視なら
GPUも検討すべき
Slide 81
Slide 81 text
Happy Programming !