CPU/GPU高速化セミナー～性能モデルの理論と実践　理論編～（2022/01/28）

6.3K Views

January 28, 22

#cpu #gpu #cpu高速化 #gpu高速化 #高速化シリーズ #CPU高速化 #GPU高速化 #性能モデル #ルーフラインモデル #ソフトウェア高速化

スライド概要

二部構成のCPU/GPU高速化セミナーの前編として、ソフトウェア性能の理論的根拠となる性能モデルについて解説します。
ソフトウェアを高速に動かすことに興味がある・コンピュータアーキテクチャに基づきソフトウェア性能の根拠を説明できるようになりたいという方にオススメです。

＜講演内容＞
・理論編のおさらい
・AMD GPU のアーキテクチャ
・AMD GPU における行列積の実装

株式会社フィックスターズ

@fixstars

スライド一覧

フィックスターズは、コンピュータの性能を最大限に引き出すソフトウェア開発のスペシャリストです。車載、産業機器、金融、医療など、幅広い分野での開発経験があります。また、ディープラーニングや機械学習などの最先端技術にも力を入れています。並列化や最適化技術を駆使して、マルチコアCPU、GPU、FPGA、量子アニーリングマシンなど、さまざまなハードウェアでソフトウェアを高速化するサービスを提供しています。さらに、長年の経験から培ったハードウェアの知識と最適化ノウハウを活かし、高精度で高性能なアルゴリズムの開発も行っています。　　　　　　・開催セミナー一覧：https://www.fixstars.com/ja/seminar　　・技術ブログ：https://proc-cpuinfo.fixstars.com/

またはPlayer版

埋め込む »CMSなどでJSが使えない場合

ダウンロード

ダウンロード(pdf - 2.55MB)

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

kaggle 画像処理機械学習深層学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 63.6K

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

arm neon

株式会社フィックスターズ 38.6K

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

ros2 gpu ロボット自律走行 ros2シリーズ高速化シリーズ

株式会社フィックスターズ 30.6K

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

kaggle 画像処理機械学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 30K

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

cuda gpu cuda高速化高速化シリーズ

株式会社フィックスターズ 28.9K

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

gpu cuda cuda高速化高速化シリーズ cuda高速化シリーズ

株式会社フィックスターズ 27K

各ページのテキスト

Fixstars Group www.fixstars.com 発表者紹介 • 冨田明彦（とみたあきひこ） • 秋山茂樹（あきやましげき）ソリューションカンパニー営業企画執行役ソリューション第一事業部リードエンジニア 2008年に入社。金融、医療業界において、ソフトウェア高速化業務に携わる。その後、新規事業企画、半導体業界の事業を担当し、現職。 2016年に入社。主に画像処理・機械学習ソフトウェアについて x86-64 CPU や NVIDIA/AMD GPU, InfiniBand を用いた高速化業務を担当。 3 Copyright © Fixstars Group

Fixstars Group www.fixstars.com ソフトウェア高速化サービス (概要) お客様のソースコードをご提供いただき、最適化やアルゴリズムの改良を施して高速化してお返ししますオリジナルソースコードのご提供当社コンサルティング高速化したソースコード高速化お客様サポート要件分析アルゴリズムの改良・開発実製品への組込み支援先行研究等の調査ハードウェアへの最適化レポートやコードへのQ&A 7 Copyright © Fixstars Group

Fixstars Group www.fixstars.com ソフトウェア高速化サービス様々な領域でソフトウェア高速化サービスを提供しています大量データの高速処理は、お客様の製品競争力の源泉となっています Semiconductor Industrial ・NAND型フラッシュメモリ向けファー・Smart Factory化支援ムウェア開発・マシンビジョンシステムの高速化・次世代AIチップ向け開発環境基盤開発 Mobility Life Science ・自動運転の高性能化、実用化・ゲノム解析の高速化・次世代パーソナルモビリティの研究開発・医用画像処理の高速化・AI画像診断システムの研究開発 Finance ・デリバティブシステムの高速化・HFT(アルゴリズムトレード)の高速化 8 Copyright © Fixstars Group

10.

11.

12.

13.

14.

15.

16.

17.

Fixstars Group www.fixstars.com 今回の話題 • CPU/GPU 高速化にあたって重要な「性能モデル」について紹介 • 場当たり的な高速化ではなく理論的な分析を通した高速化のための枠組み • 「性能モデル」を用いると... • プログラムの性能の上限を見積ることができる • 高速化余地がどれくらいあるかわかる • 高速なシステムの設計に役立つ • 性能ボトルネックが何かあらかじめわかる 17 Copyright © Fixstars Group

18.

19.

20.

Fixstars Group www.fixstars.com 典型的な高速化の流れ 1. 性能分析 • プロファイラを用いて以下を調査する • どの関数で時間がかかっているか • 関数のどの部分で時間がかかっているか • なぜ時間がかかっているか 2. 各種高速化テクニックを適用 • • • • アルゴリズム変更命令レベルの改善 (命令数削減, 近似命令の活用, SIMD化, etc.) メモリアクセスの改善 (キャッシュの活用, アクセスパターン改善, etc.) etc. 3. 以上を繰り返す 20 Copyright © Fixstars Group

21.

Fixstars Group www.fixstars.com 典型的な高速化の問題点 • 反復的な作業であるため、ゴールが見えない • 作業を始めるにあたって以下を明らかにしたい • 目標 • 高速化余地がどの程度あるか • 作業内容 • どういった高速化手法を適用すべきか、どの程度有効か • 工数 • 高速化作業にどの程度時間がかかるか特に受託開発ではこれらをうまく説明できることが顧客満足につながる • これらの疑問に答えるために性能モデルを活用できる 21 Copyright © Fixstars Group

22.

23.

Fixstars Group www.fixstars.com 性能モデルとは • 性能モデル • 対象とするコンピュータを簡略化してプログラムの性能 (実行時間等) を定式化したもの • 目的 • あるプログラムがどの程度の性能を達成しうるか、どのようにすればそれを達成できるかについて知見を得る • 厳密な性能予測を目的としたものではない 23 Copyright © Fixstars Group

24.

[beta]

Fixstars Group

www.fixstars.com

性能モデルの原始的な例 (1/3)
• シンプルな性能モデル
• 命令の種類ごとに実行回数を数えて重み付け
void SAXPY(int N, float a, const float *x, float *y) {
for (int i = 0; i < N; ++i)
y[i] = a * x[i] + y[i];
}

• 命令数
• 浮動小数点演算: 2N
• メモリアクセス数: 3N

• 実行時間: T = F * 2N + M * 3N
• F: 浮動小数点演算1命令あたり実行時間
• M: メモリアクセス1命令あたり実行時間
F, M は理論演算性能,
メモリ帯域から計算
Copyright © Fixstars Group

24

25.

Fixstars Group www.fixstars.com 性能モデルの原始的な例 (2/3) • N = 109 として計算してみる • プロセッサ: Core i7-4790 • クロック周波数: 3.6GHz • メモリ帯域: 25.6GB/s • 1コア1スレッドのみ使用 • 浮動小数点命令1回あたり実行時間 • F = 1 / (3.6*1e9) [sec] • メモリアクセス命令1回あたり実行時間 • M = 1 / (25.6*1e9 / 4) [sec] • プログラムの実行時間 • T = 2FN + 3MN = 1.02 [sec] • 実測した結果: 0.81 sec 25 Copyright © Fixstars Group

26.

Fixstars Group www.fixstars.com 性能モデルの原始的な例 (3/3) • 当然ながら、性能モデルによる予測と実測が合わない • 実行モデルと現実のプロセッサが乖離しているため • 性能モデルにより完璧に性能を予測できるわけではない • コンピュータのもつ多数の性質をモデル化するのは困難 • 「知見を得る」という目的に応じて考慮する性質を選択する 26 Copyright © Fixstars Group

27.

Fixstars Group www.fixstars.com よく知られている性能モデル (のようなもの) • Computational complexity • アルゴリズムのリソース使用量 (演算量, メモリ使用量等) を解析 • キャッシュミス回数に対する cache complexity などもある • アムダールの法則 • プログラムの並列化効率を「並列化可能な処理の割合」から定式化 • DAG Execution Model • タスク並列プログラムにおける並列化効率の定式化 • ルーフラインモデル • プロセッサの演算性能・メモリ帯域およびプログラムの演算数・メモリアクセス量を用いて得られる演算性能を定式化 27 Copyright © Fixstars Group

28.

29.

Fixstars Group www.fixstars.com ルーフラインモデル*1 • 概要 • プログラムが達成可能な演算性能 [FLOPS] を見積もるための性能モデル • 考慮する要素 • プログラムにおける演算量・メモリアクセス量 • プロセッサの理論演算性能・メモリ帯域 • 実行モデルプロセッサ浮動小数点演算性能: π [GFLOPS] メモリ帯域: β [GB/sec] メモリキャッシュは考慮しない 29 *1: Samuel Williams, Andrew Waterman, and David Patterson. Roofline: an insightful visual performance model for multicore architectures. Copyright © Fixstars Group Commun. ACM 52, 4 (April 2009)

30.

Fixstars Group www.fixstars.com 補足1: プロセッサの演算性能とは • 1秒間に実行可能な浮動小数点演算数*1 • 単位: FLOPS (FLoating point number Operations per Seconds) • 例: Intel Core i7-4790 • クロック周波数: 3.6 GHz*2 (= clock/sec) • 1クロックあたり実行可能な浮動小数点演算数 (単精度の場合) • • • • CPUコア数: 4 CPUコアあたりSIMD演算器数: 2 SIMDレーン数: 8 (AVX) SIMDレーンあたり演算数: 2 (Fused Multiply-Add 命令) • 3.6 * 4 * 2 * 8 * 2 = 460.8 [GFLOPS] Copyright © Fixstars Group *1: 整数演算が重要な場合は整数演算数で考える 30 *2: 動的周波数制御 (Intel Turbo Boost 等) も考慮する必要がある

31.

Fixstars Group www.fixstars.com 補足2: メモリ帯域とは • 1秒間に読み書き可能なメモリアクセス量 • 単位: Byte/sec • 例: Intel Core i7-4790 • メモリ規格: DDR3-1600 (12.8 GB/s) • 最大メモリチャネル数: 2 • 積をとると 25.6 GB/s • あくまでスペック値なので実測した方がよい 31 Copyright © Fixstars Group

32.

Fixstars Group www.fixstars.com 演算強度と達成可能な性能 • 演算強度 (Operational Intensity, Arithmetic Intensity) • アプリにおける演算量とメモリアクセス量の比演算強度 𝐼 [Flop/Byte] = プロセッサに対して独立な指標演算量 𝑊 [Flop] メモリアクセス量 𝑄 [Byte] • 達成可能な性能 (Attainable Performance) • 理論的に達成可能な性能の上限理論演算性能 𝜋 [FLOPS] 達成可能な性能 𝑃 [FLOPS] = min ൝ メモリ帯域 𝛽[Byte/sec] × 演算強度 𝐼 [Flop/Byte] アプリの演算量・メモリアクセス量, プロセッサの演算性能・メモリ帯域から計算 32 Copyright © Fixstars Group

33.

Fixstars Group www.fixstars.com 達成可能な性能の導出 • 演算律速の場合: プロセッサメモリ性能 𝑃 = 𝜋 [FLOPS] 時間処理開始処理終了常に演算が行われメモリアクセスは断続的 • メモリ律速の場合: • メモリアクセスにかかる時間はメモリアクセス量 𝑄 [sec] メモリ帯域 𝛽 プロセッサメモリ • 実行時間 = データ転送時間なので演算量 𝑊 性能 𝑃 = = 𝛽 × 𝐼 [FLOPS] データ転送時間 𝑄𝛽 時間処理開始処理終了常にメモリアクセスが行われ演算は断続的 33 Copyright © Fixstars Group

34.

Fixstars Group www.fixstars.com ルーフラインモデルが成立する前提条件 1. 演算とデータ転送が常に並行して行われる (or どちらかが支配的である) • 対象プログラムがそのように実装されている必要がある • Out-of-Order プロセッサなら意識しなくてもある程度満たしている 2. メモリ階層が単一である • キャッシュがある場合は以下を考慮して拡張する必要がある • メモリ – キャッシュ間データ転送帯域 • キャッシュ – レジスタ間データ転送帯域 • 演算性能 34 Copyright © Fixstars Group

35.

[beta]

Fixstars Group

www.fixstars.com

例: ナイーブな行列積の性能見積り (1/2)
• 問題: 以下の行列積コードの得られる性能の上限は？
• プロセッサの性能
• 浮動小数点演算性能: 4000 GFLOPS
• メモリ帯域: 200 GB/sec

• M = N = K = 1000, キャッシュは考慮しないものとする
float A[M * K], B[K * N], C[M * N];
for (int i = 0; i < M; ++i) {
for (int j = 0; j < N; ++j) {
float value = 0.0f;
for (int k = 0; k < K; ++k)
value += A[i * K + k] * B[k * N + j];
C[i * N + j] = value;
}
}

36
Copyright © Fixstars Group

36.

Fixstars Group www.fixstars.com 例: ナイーブな行列積の性能見積り (2/2) • 答え • • • • 演算量 W = 2MNK = 2*10^9 [flop] メモリアクセス量 Q = 8MNK + 4MN = 8*10^9 + 4*10^6 [byte] 演算強度 I = W / Q ≈ 0.25 [flop/byte] 達成可能な性能 P ≈ min(4000, 200 * 0.25) = 50 [GFLOPS] 知見: ナイーブな行列積はメモリ律速*1 どれだけ命令レベル高速化を頑張っても理論演算性能比 1.25% の性能しか得られない*2 → データ局所性を活用して演算強度を上げる必要がある (実践編へ) *1 プロセッサの演算性能を 4000 GFLOPS, メモリ帯域を 200 GB/s とした場合 *2 実際にはキャッシュ等の影響でこれ以上の性能となりうる Copyright © Fixstars Group 37

37.

Fixstars Group www.fixstars.com ルーフラインモデルの可視化 • 演算強度を変化させたときの達成可能な性能をプロット達成可能な性能 • 例: Opteron X2 (17.6 GFLOPS, 15 GB/s) プロセッサのピーク演算性能 Opteron X2 のルーフライン出典: Samuel Williams, et al. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) アプリの演算強度メモリ律速 Copyright © Fixstars演算律速 Group 38

38.

Fixstars Group www.fixstars.com ルーフラインモデルの可視化 • 複数のプロセッサをプロット • 例: Opteron X2 (17.6 GFLOPS, 15 GB/s), Opteron X4 (73.6 GFLOPS, 15 GB/s) Opteron X4 Opteron X2 演算強度が 1 以上のアプリで初めて X4 の方が高速ありがちな誤謬: あるプロセッサでメモリ律速のアプリは別のプロセッサでもメモリ律速である X2でメモリ律速 X2で演算律速, X4でメモリ律速 X4 でメモリ律速 Copyright © Fixstars Group 出典: Samuel Williams, et al. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) 39

39.

Fixstars Group www.fixstars.com Computational Ceilings • ここまで考えてきた演算性能はピーク演算性能 • 実際には命令の並び次第で上限が決まる Fused Multiply-Add 命令を使用しない場合の性能上限 SIMD命令不使用の場合の性能上限出典: Samuel Williams, et al. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) Copyright © Fixstars Group 40

40.

Fixstars Group www.fixstars.com Bandwidth Ceilings • メモリ帯域についても同様 • 連続アクセス, Memory affininty, Prefetch, etc. プリフェッチを使用しない場合の性能上限 Memory affinity を考慮しない場合の性能上限出典: Samuel Williams, et al. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) Copyright © Fixstars Group 41

41.

Fixstars Group www.fixstars.com Optimization Regions • Computational ceilings と Bandwidth ceilings を組み合わせると演算強度ごとに、どの種の高速化技法が必要か判断できる緑色: 命令・メモリ両方の高速化が必要な領域青色: 命令レベルの高速化が必要な領域黄色: メモリ周りの高速化が必要な領域出典: Samuel Williams, et al. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) Kernel1 は命令・メモリの高速化の両方が必要 Kernel2 は命令レベル高速化のみでよい Copyright © Fixstars Group 42

42.

Fixstars Group www.fixstars.com ルーフラインモデルの可視化活用例 • 深層学習ベンチマーク (MLperf, DAWNBench, DeepBench) の性能分析 • ベンチマークプログラムが演算 or メモリ律速か, どれくらいチューニングされているかすべてのアプリが V100 SP 上でメモリ律速, MLperf はおおむね効率が良い, Deep_Conv_Cu は効率が悪い出典: Snehil Verma, et al. Demystifying the MLPerf Benchmark Suite. https://arxiv.org/abs/1908.09207 Copyright © Fixstars Group 43

https://arxiv.org/abs/1908.09207

43.

Fixstars Group www.fixstars.com ピーク性能を達成するためには • 命令レベルの高速化 • SIMD, FMA 命令の活用 • 命令レベル並列性 (ILP) の改善 • パイプラインハザードを減らす (命令レイテンシ隠蔽, 分岐予測改善, etc.) • スーパースカラ (複数の実行ユニット) を活用する • スレッドレベル並列性 (TLP) の改善 • 演算器を使い切れるだけの並列性を供給する • 浮動小数点命令の割合を増やす (FP命令とそれ以外の命令が同じ演算器で実行される場合) • メモリアクセスの高速化 • 適切な粒度, 量, alignment でメモリアクセス • メモリアクセスレイテンシの隠蔽 (w/ ILP, TLP, Prefetch) • プロセッサ・メモリトポロジの考慮 • 演算強度の向上 • レジスタ, キャッシュ等のメモリ階層を活用 • その他 (ルーフラインモデルの範疇外) • 負荷分散の改善, 同期の削減, etc. 44 Copyright © Fixstars Group

44.

Fixstars Group www.fixstars.com なぜルーフラインモデルが重要か • プログラムの性能を (ある程度) 見積もることができる • 高速化余地がどれくらいあるか • どのリソースがボトルネックか • 別のプロセッサに移植した場合にどの程度の性能となるか (→ 機種選定に役立つ) • 高速なソフトウェアの設計に役立つ • アルゴリズム選定, cache-aware algorithms で必要なキャッシュサイズ • 数ある高速化手法のうち、不要なものを事前に枝刈りできる • 試行錯誤を減らせるかも 45 Copyright © Fixstars Group

45.

46.

Fixstars Group www.fixstars.com ルーフラインモデルの検証 • ルーフラインモデルがどの程度正確かベンチマークプログラムを用いて検証する • 演算性能の計測 • メモリ帯域の計測 • 演算・メモリ複合実行性能の計測 • 実験環境 • x86-64 CPU • • • • Intel Intel Intel AMD Core i7-3770 3.40GHz (IvyBridge), Ubuntu 18.04 Core i7-4790 3.60GHz (Haswell), WSL1 Core i7-6500U 2.50GHz (Skylake), WSL1 Ryzen 7 3700X 3.60GHz (Zen 2), Ubuntu 18.04 47 Copyright © Fixstars Group

47.

[beta]

Fixstars Group

www.fixstars.com

演算性能の計測
• ベンチマーク内容
• 1命令あたり8加算可能な vaddps 命令を大量に実行し
1秒あたりの演算回数 [FLOPS] を計測する*1
• 1コアのみ使用
__m256 vx = ..., vy = ...;
auto vz0 = vx, vz1 = vx, vz2 = vx, vz3 = vx;
auto vz4 = vx, vz5 = vx, vz6 = vx, vz7 = vx;
for (size_t i = 0; i < n_times;
vz0 = _mm256_add_ps(vz0, vy);
vz2 = _mm256_add_ps(vz2, vy);
vz4 = _mm256_add_ps(vz4, vy);
vz6 = _mm256_add_ps(vz6, vy);
}

++i) {
vz1 = _mm256_add_ps(vz1,
vz3 = _mm256_add_ps(vz3,
vz5 = _mm256_add_ps(vz5,
vz7 = _mm256_add_ps(vz7,

vy);
vy);
vy);
vy);

依存関係のない8個の vaddps 命令を繰り返し実行

// Avoid dead-code elimination
(パイプラインハザードが発生しないようにする)
auto vz = vz0;
vz = _mm256_add_ps(vz, vz1); vz = _mm256_add_ps(vz, vz2);
vz = _mm256_add_ps(vz, vz3); vz = _mm256_add_ps(vz, vz4);
vz = _mm256_add_ps(vz, vz5); vz = _mm256_add_ps(vz, vz6);
vz = _mm256_add_ps(vz, vz7);
_mm256_store_ps(tmp, vz);

48

Copyright
Fixstars(vfmaddps)
Group
*1: 理論ピーク演算性能を計測する場合、
FMA© 命令
を使用する必要があるが、今回は vaddps 命令を用いる。

48.

Fixstars Group www.fixstars.com 演算性能の計測結果 CPU CPU世代ベース / 計測時*1 クロック周波数理論演算性能*2 実測演算性能 Intel Core i7-3770 IvyBridge (VADD port x1) 3.4 / 3.9 GHz 31.2 GFLOPS 30.79 GFLOPS Intel Core i7-4790 Haswell (VADD port x1) 3.6 / 3.6 GHz 28.8 GFLOPS 28.63 GFLOPS Intel Core i7-6500U Skylake (VADD port x2) 2.5 / 2.7 GHz 43.2 GFLOPS 43.21 GFLOPS AMD Ryzen 7 3700X Zen 2 (VADD port x2) 3.6 / 4.4 GHz 70.4 GFLOPS 69.96 GFLOPS Copyright © Fixstars Group *1: Turbo Boost 有効時の推定値 *2: 1コア, 加算のみの場合の理論性能 49

49.

[beta]

Fixstars Group

www.fixstars.com

メモリ帯域の計測
• ベンチマーク内容
• キャッシュに収まらないサイズの配列に対して
4つのパターンでメモリアクセスし、スループットを計測する
• read only (read:write=1:0), write only (read:write=0:1),
copy (read:write=1:1), triad (read:write=2:1)

• 1コアのみ使用
float *sp0 = src0, *sp1 = src1, *dp = dst, *dst_end = dst + size;
while (dp < dst_end) {
// triad: dst[i] = alpha * src0[i] + src1[i];
auto vx0 = _mm256_load_ps(sp0 + 0 * 8), vx1 = _mm256_load_ps(sp0 + 1
auto vx2 = _mm256_load_ps(sp0 + 2 * 8), vx3 = _mm256_load_ps(sp0 + 3
auto vy0 = _mm256_load_ps(sp1 + 0 * 8), vy1 = _mm256_load_ps(sp1 + 1
auto vy2 = _mm256_load_ps(sp1 + 2 * 8), vy3 = _mm256_load_ps(sp1 + 3
auto vz0 = _mm256_fmadd_ps(valpha, vx0, vy0);
auto vz1 = _mm256_fmadd_ps(valpha, vx1, vy1);
auto vz2 = _mm256_fmadd_ps(valpha, vx2, vy2);
auto vz3 = _mm256_fmadd_ps(valpha, vx3, vy3);
_mm256_store_ps(dp + 0 * 8, vz0); _mm256_store_ps(dp + 1 * 8, vz1);
_mm256_store_ps(dp + 2 * 8, vz2); _mm256_store_ps(dp + 3 * 8, vz3);
sp0 += 4 * 8, sp1 += 4 * 8, dp += 4 * 8;
}
Copyright © Fixstars Group

*
*
*
*

8);
8);
8);
8);

50

50.

Fixstars Group www.fixstars.com メモリ帯域の計測結果 CPU CPU世代メモリ規格理論メモリ帯域 [GB/s] 実測メモリ帯域 [GB/s] read / write / copy / triad Intel Core i7-3770 IvyBridge DDR3-1600 x2 25.6 17.3 / 10.4 / 13.3 / 14.3 Intel Core i7-4790 Haswell DDR3-1600 x2 25.6 16.3 / 10.3 / 13.4 / 14.1 Intel Core i7-6500U Skylake LPDDR3-1866 x2 29.9 15.9 / 11.1 / 14.6 / 16.5 AMD Ryzen 7 3700X Zen 2 DDR4-2666 x2 42.6 22.8 / 12.1 / 17.6 / 20.5 • 1コアのみ使用する場合には最大でも理論値の60%程度に留まる • メモリアクセスパターンによって得られる帯域が異なる • プロセッサアーキテクチャに依存するが今回の対象ではおおむね read only > triad > copy > write only 51 Copyright © Fixstars Group

51.

[beta]

Fixstars Group

www.fixstars.com

演算・メモリ複合実行性能の計測
• ベンチマーク内容
• 配列の各要素に対してロード・N回加算・ストアを行う
• read:write:add=1:1:N
• 演算強度: N / ((1 + 1) * sizeof(float)) = N/8

• 1コアのみ使用
const auto va = ...;
const auto *sp0 = src0, *dp = dst, *dst_end = dst + size;
while (dp < dst_end) {
// dst[i] = src0[i] + (a + a + ... + a);
auto vz0 = _mm256_load_ps(sp0 + 0 * 8), vz1 = _mm256_load_ps(sp0 + 1 * 8);
auto vz2 = _mm256_load_ps(sp0 + 2 * 8), vz3 = _mm256_load_ps(sp0 + 3 * 8);
auto vz4 = _mm256_load_ps(sp0 + 4 * 8), vz5 = _mm256_load_ps(sp0 + 5 * 8);
auto vz6 = _mm256_load_ps(sp0 + 6 * 8), vz7 = _mm256_load_ps(sp0 + 7 * 8);
UNROLL for (int i = 0; i < N; ++i) { // N: 演算強度を変化させるためのパラメータ
vz0 = _mm256_add_ps(vz0, va); vz1 = _mm256_add_ps(vz1, va);
vz2 = _mm256_add_ps(vz2, va); vz3 = _mm256_add_ps(vz3, va);
vz4 = _mm256_add_ps(vz4, va); vz5 = _mm256_add_ps(vz5, va);
vz6 = _mm256_add_ps(vz6, va); vz7 = _mm256_add_ps(vz7, va);
}
_mm256_store_ps(dp + 0 * 8, vz0); _mm256_store_ps(dp + 1 * 8, vz1);
_mm256_store_ps(dp + 2 * 8, vz2); _mm256_store_ps(dp + 3 * 8, vz3);
_mm256_store_ps(dp + 4 * 8, vz4); _mm256_store_ps(dp + 5 * 8, vz5);
_mm256_store_ps(dp + 6 * 8, vz6); _mm256_store_ps(dp + 7 * 8, vz7);
sp0 += 8 * 8, dp += 8 * 8;
}
Copyright © Fixstars Group

52

52.

Fixstars Group www.fixstars.com 演算・メモリ複合実行性能の計測結果 Core i7-3770 (IvyBridge) 性能 [GFLOPS] 80 60 60 40 40 Loofline 実測値 20 Core i7-4790 (Haswell) 80 Loofline 実測値 20 0 0 0 2 4 6 8 Core i7-6500U (Skylake) 性能 [GFLOPS] 80 0 60 40 40 Loofline 実測値 20 4 20 0 6 8 Ryzen 7 3700X (Zen 2) 80 60 2 Loofline 実測値 0 0 2 4 演算強度 6 8 0 2 4 演算強度 • おおむねルーフラインモデルに近い値が得られている Copyright © Fixstars Group 6 8 53

53.

Fixstars Group www.fixstars.com まとめ • CPU/GPU 高速化において重要な「ルーフラインモデル」について解説 • • • • プログラムの性能上限を見積もることができるどのリソースがボトルネックか判定できる色んなプロセッサに移植した場合の性能を推定できる有効な高速化手法を選ぶのに役立つ • ベンチマークによる検証を実施 • ルーフラインモデルによる理論性能と実性能を比較し、おおむね理論性能に近い結果が得られることを確認 54 Copyright © Fixstars Group

54.

55.

56.

Fixstars Group www.fixstars.com 高速化結果 10000 性能 (GFLOPS) 8000 Peak Performance (1) Naïve (2) Register blocking only (3) (2) + LDS blocking (4) (3) + K loop unrolling (5) (4) + Other optimizations 6000 4000 2000 0 1024 2048 3072 4096 5120 6144 行列サイズ (M=N=K) 7168 8192 • 最大 6351 GFLOPS (ピーク性能比 73.8%) • 単純な実装から11倍の高速化を達成 Copyright © Fixstars Group 57

57.

CPU/GPU高速化セミナー ～性能モデルの理論と実践 理論編～（2022/01/28）

株式会社フィックスターズ

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

いまさら聞けない！CUDA高速化入門 ～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

各ページのテキスト

CPU/GPU高速化セミナー～性能モデルの理論と実践　理論編～（2022/01/28）

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）