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

CPU / GPU高速化セミナー!性能モデルの理論と実践:実践編

CPU / GPU高速化セミナー!性能モデルの理論と実践:実践編

2022年3月4日開催の「CPU / GPU高速化セミナー!性能モデルの理論と実践:実践編」セミナー資料です。

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

Other Decks in Programming

Transcript

  1. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars

    Group CPU / GPU 高速化セミナー 性能モデルの理論と実践:実践編
  2. Fixstars Group www.fixstars.com Copyright © Fixstars Group 発表者紹介 2 •

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

    • 性能に関する課題 • 高速化サービスと開発の流れ 性能モデルの活用 (60分) • 理論編のおさらい • AMD GPUアーキテクチャ • AMD GPUにおける行列積の高速化 Q&A / 告知 3
  4. Fixstars Group www.fixstars.com Copyright © Fixstars Group 性能に関する課題 5 生産効率の向上

    • より短時間で欠陥検出 • より安価なハードで 安全性の向上 • より精度の高い物体検出 • より低消費電力なハードで
  5. Fixstars Group www.fixstars.com Copyright © Fixstars Group ソフトウェア高速化サービス (概要) お客様のソースコードをご提供いただき、

    最適化やアルゴリズムの改良を施して高速化してお返しします 当社 お客様 オリジナルソースコードのご提供 高速化したソースコード コンサルティング 高速化 サポート 要件分析 先行研究等の調査 アルゴリズムの改良・開発 ハードウェアへの最適化 実製品への組込み支援 レポートやコードへのQ&A 6
  6. Fixstars Group www.fixstars.com Copyright © Fixstars Group ソフトウェア高速化サービス 様々な領域でソフトウェア高速化サービスを提供しています 大量データの高速処理は、お客様の製品競争力の源泉となっています

    ・NAND型フラッシュメモリ向けファー ムウェア開発 ・次世代AIチップ向け開発環境基盤開発 Semiconductor ・デリバティブシステムの高速化 ・HFT(アルゴリズムトレード)の高速化 Finance ・自動運転の高性能化、実用化 ・次世代パーソナルモビリティの研究開発 Mobility ・ゲノム解析の高速化 ・医用画像処理の高速化 ・AI画像診断システムの研究開発 Life Science ・Smart Factory化支援 ・マシンビジョンシステムの高速化 Industrial 7
  7. Fixstars Group www.fixstars.com Copyright © Fixstars Group 8 要件分析 研究調査・アルゴリズム実装

    高速化 / 評価 アルゴリズム改善 / 評価 品質確保 典型的な開発の流れ • 課題のヒアリング • ユースケース洗い出し • 時間、ハードウェア等への制約条件抽出
  8. Fixstars Group www.fixstars.com Copyright © Fixstars Group 9 要件分析 研究調査・アルゴリズム実装

    高速化 / 評価 アルゴリズム改善 / 評価 品質確保 典型的な開発の流れ • 論文等サーベイ • アルゴリズム候補の絞り込み • アルゴリズムの比較と評価 • アルゴリズムの決定と実装
  9. Fixstars Group www.fixstars.com Copyright © Fixstars Group 10 要件分析 研究調査・アルゴリズム実装

    高速化 / 評価 アルゴリズム改善 / 評価 品質確保 典型的な開発の流れ • 対象ハードウェアへの移植 / 評価 • 計算量 / 精度面からの改善案検討 • 改善案の実装
  10. Fixstars Group www.fixstars.com Copyright © Fixstars Group 11 要件分析 研究調査・アルゴリズム実装

    高速化 / 評価 アルゴリズム改善 / 評価 品質確保 典型的な開発の流れ • データ並列プログラミング • 専用アクセラレータ / 演算器の活用 • 処理時間の計測 / 評価
  11. Fixstars Group www.fixstars.com Copyright © Fixstars Group 12 要件分析 研究調査・アルゴリズム実装

    高速化 / 評価 アルゴリズム改善 / 評価 品質確保 典型的な開発の流れ • 異常系処理の実装 • 静的解析ツールによる評価 • テストケース検討 / 実施 • カバレッジ向上
  12. Fixstars Group www.fixstars.com Copyright © Fixstars Group 本ウェビナーの対象プロセス 13 ココ

    要件分析 研究調査・アルゴリズム実装 高速化 / 評価 アルゴリズム改善 / 評価 品質確保
  13. Fixstars Group www.fixstars.com Copyright © Fixstars Group • 社内大学 •

    プログラミングコンテスト • 勉強会 • 各種コンテストへの参加 • 勉強会 • 論文・学会発表 • 社内向け • 社外向け よりよいサービスのご提供を目指して 14 • 組込み開発 • アルゴリズム開発 • AI・深層学習 • 組合せ最適化 各種高速化サービス 技術力強化
  14. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars

    Group 性能モデルの活用 AMD GPU における行列積の実装
  15. Fixstars Group www.fixstars.com Copyright © Fixstars Group 今回の話題 • CPU/GPU

    高速化にあたって重要な「性能モデル」について紹介 • 場当たり的な高速化ではなく、理論的な分析を通した高速化のための枠組み • 理論編に引き続き、性能モデルの活用例を紹介 • AMD GPU における行列積の実装 16
  16. Fixstars Group www.fixstars.com Copyright © Fixstars Group 目次 • 理論編のおさらい

    • AMD GPU アーキテクチャ • AMD GPU における行列積の実装 • ルーフラインモデルによる性能見積り・実装方針の決定 • レジスタブロッキング • 共有メモリブロッキング • ソフトウェアパイプライニングによる 命令・メモリアクセスレイテンシ隠蔽 17
  17. Fixstars Group www.fixstars.com Copyright © Fixstars Group ルーフラインモデル *1 •

    概要 • プログラムが達成可能な演算性能 [FLOPS] を 見積もるための性能モデル • 考慮する要素 • プログラムにおける演算量・メモリアクセス量 • プロセッサの理論演算性能・メモリ帯域 • 実行モデル 19 プロセッサ メモリ 浮動小数点演算性能: π [GFLOPS] メモリ帯域: β [GB/sec] *1: Samuel Williams, Andrew Waterman, and David Patterson. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) キャッシュは考慮しない
  18. Fixstars Group www.fixstars.com Copyright © Fixstars Group 補足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] 20 *1: 整数演算が重要な場合は整数演算数で考える *2: 動的周波数制御 (Intel Turbo Boost 等) も 考慮する必要がある
  19. Fixstars Group www.fixstars.com Copyright © Fixstars Group 補足2: メモリ帯域とは •

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

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

    性能 𝑃 = 𝜋 [FLOPS] • メモリ律速の場合: • メモリアクセスにかかる時間は メモリアクセス量 𝑄 メモリ帯域 𝛽 [sec] • 実行時間 = データ転送時間 なので 性能 𝑃 = 演算量 𝑊 データ転送時間 𝑄 𝛽 = 𝛽 × 𝐼 [FLOPS] 23 プロセッサ メモリ プロセッサ メモリ 常に演算が行われ メモリアクセスは断続的 常にメモリアクセスが行われ 演算は断続的 時間 処理開始 処理終了 時間 処理開始 処理終了
  22. Fixstars Group www.fixstars.com Copyright © Fixstars Group ベンチマークプログラムによる検証 • おおむねルーフラインモデルに近い値が得られている

    0 20 40 60 80 0 2 4 6 8 性能 [GFLOPS] Core i7-3770 (IvyBridge) Loofline 実測値 0 20 40 60 80 0 2 4 6 8 Core i7-4790 (Haswell) Loofline 実測値 0 20 40 60 80 0 2 4 6 8 性能 [GFLOPS] 演算強度 Core i7-6500U (Skylake) Loofline 実測値 0 20 40 60 80 0 2 4 6 8 演算強度 Ryzen 7 3700X (Zen 2) Loofline 実測値 24
  23. Fixstars Group www.fixstars.com Copyright © Fixstars Group Graphics Processing Unit

    (GPU) • ASCII.jpデジタル用語辞典より引用: 26 グラフィクスだけでなく Computer Vision, AI, 数値計算,ビッグデータ処理 等にも利用されるようになっている “3Dグラフィックスを表示するために必要な計算を、 CPUに代わって処理するプロセッサー。” “GPUの高性能化にともない、最近では、 GPUにグラフィックス以外の計算処理を行わせ、 汎用的に活用する動きもある。”
  24. Fixstars Group www.fixstars.com Copyright © Fixstars Group Graphics Processing Unit

    (GPU) • 現代のデスクトップ・サーバ向け GPU アーキテクチャ • 演算器を多数並べた高並列プロセッサ • 多数のHWスレッドを実行できる 27 Block diagram of the AMD Instinct™ MI200 multi-chip module (AMD Instinct™ MI250/MI250X) 出典: INTRODUCING AMD CDNA™ 2 ARCHITECTURE 1つの GPU に 28,160 個 の 32bit float 演算器
  25. Fixstars Group www.fixstars.com Copyright © Fixstars Group GPU の利点・欠点 •

    利点: 高スループット • 大量の処理を高速に実行するのに向く • 欠点: 高レイテンシ • シングルスレッド性能が低い • 性能を最大限活用するためには十分な並列性が必要 28 Xeon Platinum 8180 Tesla V100 32bit float 演算性能 2.06 TFLOPS 15.7 TFLOPS メモリ帯域 127. 8GB/s 900 GB/s クロック周波数 2.50 GHz 1.53 GHz コア数 28 80 32bit float 演算器数 448 5,120 同時実行可能な HWスレッドの数 コアあたり 2スレッド コアあたり 64スレッド *1 CPU はシングルスレッド性能が高い (GPUには分岐予測・投機実行等の機構もない) コア数・演算器数は GPU の方が圧倒的に多い GPU は多数のHWスレッドによりスループット向上, レイテンシも増加 GPU は演算性能で 7.6 倍、 メモリ性能で 7.0 倍高速 Intel CPU と NVIDIA GPU のサーバ向けプロセッサ比較 *1: ここではスレッド数 = プログラムカウンタ数としている
  26. Fixstars Group www.fixstars.com Copyright © Fixstars Group AMD Graphics-Core-Next (GCN)

    Architecture • AMD GPU のマイクロアーキテクチャ・命令セット • 2012 - 2018年までの AMD GPU にて採用 *1 • SIMT 方式 *2 を採用 • SIMD + HW スレッド • NVIDIA GPU と似た構成 • プログラミング環境 • OpenCL • ROCm Languages (HIP, Heterogeneous Compute C++) • 各種グラフィックス API + シェーダ言語 29 *1: 最新モデルは RDNA/CDNA アーキテクチャに移行 *2: Single Instruction Multiple Thread (NVIDIA 用語) 多数の HW スレッドにより 演算・メモリアクセスレイテンシを隠蔽
  27. Fixstars Group www.fixstars.com Copyright © Fixstars Group AMD Graphics-Core-Next (GCN)

    Architecture • 多数の Compute Unit*1 から構成 30 出典: White Paper | AMD GRAPHICS CORES NEXT (GCN) ARCHITECTURE *1: CPU におけるコアのようなもの
  28. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit 概要

    • AMD GPU における「プロセッサコア」 • NVIDIA GPU における Streaming Multiprocessor に類似 • 特徴 • 4つの 16-lane SIMD ユニット • 64 個の FP/INT 演算器 • SIMD ユニットあたり 10 個のプログラムカウンタ • 最大 40 個の HW スレッドを同時に実行可能 • 7種類の実行ユニット • Vector ALU, Vector Memory, Scalar ALU/Memory, Branch, LDS, etc. • 64KB のスクラッチパッドメモリ (LDS: Local Data Share) • メインメモリとは別のアドレス空間を持つ高速なオンチップメモリ 31
  29. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit: ブロック図

    32 SIMD3 10 PC & 命令バッファ SIMD2 10 PC & 命令バッファ SIMD1 10 PC & 命令バッファ 命令フェッチアービトレーション SIMD0 10 PC & 命令バッファ Branch Vector Decode Vector Memory Decode Scalar Decode Local Data Share 64KB ... 命令フェッチ 命令デコード Scalar Unit Scalar ALU SGPR 4KB 命令実行 Vector Memory Unit Data Return Unit Address Unit Vector Unit VGPR 64KB VALU 16 lane SIMD0 VGPR 64KB VALU 16 lane SIMD1 VGPR 64KB VALU 16 lane SIMD2 VGPR 64KB VALU 16 lane SIMD3 L1 Data Cache 16KB L1 Inst Cache 16KB 命令アービトレーション LDS Decode ※ PC: Program Counter, GPR: General-Purpose Register
  30. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit: 命令フェッチ

    33 SIMD3 10 PC & 命令バッファ SIMD2 10 PC & 命令バッファ SIMD1 10 PC & 命令バッファ 命令フェッチアービトレーション SIMD0 10 PC & 命令バッファ Branch Vector Decode Vector Memory Decode Scalar Decode Local Data Share 64KB ... 命令フェッチ 命令デコード Scalar Unit Scalar ALU SGPR 4KB 命令実行 Vector Memory Unit Data Return Unit Address Unit Vector Unit VGPR 64KB VALU 16 lane SIMD0 VGPR 64KB VALU 16 lane SIMD1 VGPR 64KB VALU 16 lane SIMD2 VGPR 64KB VALU 16 lane SIMD3 L1 Data Cache 16KB L1 Inst Cache 16KB 命令アービトレーション LDS Decode ※ PC: Program Counter, GPR: General-Purpose Register 40 個の PC からひとつを選んで 8命令程度をフェッチ
  31. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit: 命令デコード

    34 SIMD3 10 PC & 命令バッファ SIMD2 10 PC & 命令バッファ SIMD1 10 PC & 命令バッファ 命令フェッチアービトレーション SIMD0 10 PC & 命令バッファ Branch Vector Decode Vector Memory Decode Scalar Decode Local Data Share 64KB ... 命令フェッチ 命令デコード Scalar Unit Scalar ALU SGPR 4KB 命令実行 Vector Memory Unit Data Return Unit Address Unit Vector Unit VGPR 64KB VALU 16 lane SIMD0 VGPR 64KB VALU 16 lane SIMD1 VGPR 64KB VALU 16 lane SIMD2 VGPR 64KB VALU 16 lane SIMD3 L1 Data Cache 16KB L1 Inst Cache 16KB 命令アービトレーション LDS Decode ※ PC: Program Counter, GPR: General-Purpose Register 4つのSIMDユニットからひとつを選んで 最大5種類の命令を同時にデコード・発行
  32. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit: 命令実行

    35 SIMD3 10 PC & 命令バッファ SIMD2 10 PC & 命令バッファ SIMD1 10 PC & 命令バッファ 命令フェッチアービトレーション SIMD0 10 PC & 命令バッファ Branch Vector Decode Vector Memory Decode Scalar Decode Local Data Share 64KB ... 命令フェッチ 命令デコード Scalar Unit Scalar ALU SGPR 4KB 命令実行 Vector Memory Unit Data Return Unit Address Unit Vector Unit VGPR 64KB VALU 16 lane SIMD0 VGPR 64KB VALU 16 lane SIMD1 VGPR 64KB VALU 16 lane SIMD2 VGPR 64KB VALU 16 lane SIMD3 L1 Data Cache 16KB L1 Inst Cache 16KB 命令アービトレーション LDS Decode ※ PC: Program Counter, GPR: General-Purpose Register 各 16-lane SIMD ユニットは 64-lane SIMD 命令を 4 サイクルかけて実行 (命令レイテンシが自動的に隠蔽される)
  33. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit: Local

    Data Share (LDS) 36 SIMD3 10 PC & 命令バッファ SIMD2 10 PC & 命令バッファ SIMD1 10 PC & 命令バッファ 命令フェッチアービトレーション SIMD0 10 PC & 命令バッファ Branch Vector Decode Vector Memory Decode Scalar Decode Local Data Share 64KB ... 命令フェッチ 命令デコード Scalar Unit Scalar ALU SGPR 4KB 命令実行 Vector Memory Unit Data Return Unit Address Unit Vector Unit VGPR 64KB VALU 16 lane SIMD0 VGPR 64KB VALU 16 lane SIMD1 VGPR 64KB VALU 16 lane SIMD2 VGPR 64KB VALU 16 lane SIMD3 L1 Data Cache 16KB L1 Inst Cache 16KB 命令アービトレーション LDS Decode *1: ハイエンドモデルの場合 スクラッチパッドメモリ 1サイクルあたり 128 byte ロード可能 *1
  34. Fixstars Group www.fixstars.com Copyright © Fixstars Group Compute Unit: レジスタ

    • 32bit Vector GPR • SIMD ユニットあたり 64KB • SIMD lane あたり最大 256 個 • 1スレッドあたり VGPR 数は可変 • レジスタ使用量を増やすと同時実行可能な HW スレッド数が減る • 32bit Scalar GPR • SIMD ユニットあたり 2KB • 16-lane SIMD あたり 104 個 (ユーザが利用可能な数) 37
  35. Fixstars Group www.fixstars.com Copyright © Fixstars Group AMD GPU のピーク性能

    • 例: Radeon R9 Fury X • 浮動小数点演算性能 • クロック周波数: 1.05 GHz • 積和命令: 2 ops/cycle • SIMD: 16 lane x 4 • Compute Unit 数: 64 → 8601.6 GFLOPS (floating operation per sec) • メモリ帯域 • 理論値: 512 GB/s • 実測値: 392 GB/s (read:write=2:1 の場合) 38
  36. Fixstars Group www.fixstars.com Copyright © Fixstars Group OpenCL を用いた並列プログラミング •

    1 work-item (≈ 1演算器, SIMD lane) で行う処理を記述 (SPMD, Single Program Multiple Data) • 例: ベクトル和 C = A + B 39 __kernel void vector_add_gpu_kernel( __global float *A, __global float *B, __global float *C, int N) { int i = get_global_id(0); if (i < N) C[i] = A[i] + B[i]; } ベクトル和の OpenCL コード ※OpenCL カーネル起動時に N 個 work-item を生成 1 work-item が行う処理を記述 get_global_id(n): work-item 番号を取得 自 work-item の担当する 要素の和を計算 void vector_add(float *A, float *B, float *C, int N) { for (int i = 0; i < N; ++i) C[i] = A[i] + B[i]; } ベクトル和の逐次コード
  37. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars

    Group AMD GPU における 行列積の高速化 40
  38. Fixstars Group www.fixstars.com Copyright © Fixstars Group 行列積のGPU実装: 並列性の抽出 •

    行列積 C = A * B では A の i 行目と B の j 列目について内積を計算し その結果を C(i, j) に書き込む • i, j についてデータ依存が存在せず並列化可能 • 内積計算はデータ依存が存在し、単純には並列化不可 *1 41 void matrix_multiply( float *A, float *B, float *C, int M, int N, int K) { for (int i = 0; i < M; ++i) { for (int j = 0; j < N; ++j) { float ab = 0.0f; for (int k = 0; k < K; ++k) { ab += A[i * K + k] * B[k * N + j]; } C[i * N + j] = ab; } } } 行列積 C = A * B の逐次コード j i k k C A B Cの各要素について 対応するAの行, Bの列の内積を計算 各 i, j について独立した要素に データを書き込んでいる (= 並列化可能) *1: リダクションや atomic add で並列化可能
  39. Fixstars Group www.fixstars.com Copyright © Fixstars Group 単純な GPU 実装

    • C の1要素を1 work-item に割り当て • 各 work-item は A の i 行目と B の j 列目の内積を計算し、C(i, j) に書き込む 42 j i k k C A B __kernel void matrix_multiply_gpu_kernel( __global float *A, __global float *B, __global float *C, int M, int N, int K) { int i = get_global_id(1); int j = get_global_id(0); if (i >= M || j >= N) return; float ab = 0.0f; for (int k = 0; k < K; ++k) { ab += A[i * K + k] * B[k * N + j]; } C[i * N + j] = ab; } 行列積 C = A * B の OpenCL コード MN 個の work-item を生成, 各 work-item で C の1要素を計算 逐次コードの i, j ループ内部を そのまま記述
  40. Fixstars Group www.fixstars.com Copyright © Fixstars Group 0 2000 4000

    6000 8000 10000 1024 2048 3072 4096 5120 6144 7168 8192 性能 (GFLOPS) 行列サイズ (M=N=K) Peak Performance (1) Naïve 単純な GPU 実装 • Radeon R9 Fury X にて性能を計測 43 最大 567 GFLOPS (ピーク性能比6.5%)
  41. Fixstars Group www.fixstars.com Copyright © Fixstars Group 単純な GPU 実装の演算強度

    • 1 work-item あたりの演算数とメモリアクセス量を調べる 44 __kernel void matrix_multiply_gpu_kernel( __global float *A, __global float *B, __global float *C, int M, int N, int K) { int i = get_global_id(0); int j = get_global_id(1); if (i >= M || j >= N) return; float ab = 0.0f; for (int k = 0; k < K; ++k) { ab += A[i * K + k] * B[k * N + j]; } C[i * N + j] = ab; } ロード 2K回, 浮動小数点演算: 2K回 (積+和), ストア: 1回 演算強度: 2K (2K + 1) ∗ sizeof(float) ≈ 0.25
  42. Fixstars Group www.fixstars.com Copyright © Fixstars Group GPU のピーク性能と性能見積り •

    例: AMD Radeon R9 Fury X • 浮動小数点演算性能: 8602 GFLOPS • 実効メモリバンド幅: 392 GB/s • 演算・メモリ帯域比: 21.94 • 1 byte 転送する間に 21.94 演算を実行可能 • ルーフラインモデルによる達成可能な性能 • min(8602, 392 * 0.25) = 98 GFLOPS 45 *1: すべてのメモリアクセスがキャッシュミスする場合。 実際はキャッシュの影響を受ける可能性がある。 単純な行列積コードの演算強度 ≈ 0.25 [flop/byte] に対して Fury X の演算・メモリ帯域比は 21.94 [flop/byte] と非常に大きい メモリ帯域が足りず、演算性能を 1.1%*1 しか活用できない
  43. Fixstars Group www.fixstars.com Copyright © Fixstars Group どのように高速化すればよいか • 問題点

    • 演算量に対してメモリアクセス量が多すぎる • 解決策 • 行列積のもつメモリアクセスの局所性を利用して 一度ロードしたデータを再利用する → 「ブロッキング」により演算強度を高める 46
  44. Fixstars Group www.fixstars.com Copyright © Fixstars Group 高速化1: 演算強度を上げるループ変形 •

    行列積のデータ局所性 • 隣り合う要素の計算に必要なデータは一部共通している 47 j i k k C A B C(i+0, j+0) の計算には A(i+0, *), B(*, j+0) が必要 C(i+0, j+1) の計算には A(i+0, *), B(*, j+1) が必要 C(i+1, j+0) の計算には A(i+1, *), B(*, j+0) が必要 C(i+1, j+1) の計算には A(i+1, *), B(*, j+1) が必要 1 work-item あたり複数個の要素を処理し, 1度ロードした結果を使い回せばロード回数を減らすことができる (レジスタブロッキング)
  45. Fixstars Group www.fixstars.com Copyright © Fixstars Group 高速化1: 演算強度を上げるループ変形 •

    レジスタブロッキング例: ブロックサイズ 2x2 の場合 48 __kernel void matrix_multiply_gpu_kernel( __global float *A, __global float *B, __global float *C, int M, int N, int K) { int gy = get_global_id(1), gx = get_global_id(0); int i0 = gy * 2 + 0, j0 = gx * 2 + 0; int i1 = gy * 2 + 1, j1 = gx * 2 + 1; float ab[2][2] = { 0.0f }; for (int k = 0; k < K; ++k) { float a0 = A[i0 * K + k], b0 = B[k * N + j0]; float a1 = A[i1 * K + k], b1 = B[k * N + j1]; ab[0][0] += a0 * a0; ab[0][1] += a0 * a1; ab[1][0] += a1 * a0; ab[1][1] += a1 * a1; } C[i0 * N + j0] = ab[0][0]; C[i0 * N + j1] = ab[0][1]; C[i1 * N + j0] = ab[1][0]; C[i1 * N + j1] = ab[1][1]; } ロード: 4K回 (16K bytes) 演算: 8K回 ストア: 4回 (16 bytes) 2x2 レジスタブロッキングにより 演算強度が 0.25 から 0.5 に増加
  46. Fixstars Group www.fixstars.com Copyright © Fixstars Group ブロックサイズと演算強度 • ブロックサイズを大きくすればするほど演算強度は上がる

    49 ブロックサイズ ロード量 演算数 演算強度 2 16 bytes 8 0.5 4 32 bytes 32 1 8 64 bytes 128 2 64 512 bytes 8192 16 128 1024 bytes 32768 32 n 8*n 2*n*n n/4 ただし、128x128 の場合のレジスタ要求量は少なくとも16384個 → レジスタが足りない 128x128 以上のブロッキングで Fury X の演算・メモリ帯域比 21.94 を上回る → 演算ネックとなる
  47. Fixstars Group www.fixstars.com Copyright © Fixstars Group 階層ブロッキング • 各メモリ階層上でブロッキング

    • レジスタブロッキング • LDS ブロッキング • 実行イメージ • 演算と各データ転送をオーバーラップ、パイプライニング 50 メモリ→LDS転送 LDS→レジスタ転送 0 1 2 3 4 5 0 1 2 3 4 5 0 1 2 3 4 5 演算 時間 (k ループ) 演算中にメモリ-LDS間, LDS-レジスタ間データ転送が 完了するようにブロッキングサイズを選択
  48. Fixstars Group www.fixstars.com Copyright © Fixstars Group メモリ階層とブロッキングサイズ選択 • メモリ階層

    • 階層ごとのブロッキングサイズの選択 51 メモリ種別 サイズ 実効帯域 演算・帯域比 デバイスメモリ 4GB (Fury X) 392 GB/s 21.94 LDS 64KB / CU 128 byte/cycle (load only) 1 ブロック サイズ ロード量 演算数 演算強度 4 32 bytes 32 1 8 64 bytes 128 2 64 512 bytes 8192 16 128 1024 bytes 32768 32 n 8*n 2*n*n n/4 LDSで128x128以上とすると 演算中にメモリ-LDS間データ転送が終わる レジスタで4x4以上とすると 演算中にLDS-レジスタ間データ転送が終わる
  49. Fixstars Group www.fixstars.com Copyright © Fixstars Group 高速化2: LDS ブロッキング

    • データの配置 • LDS: A 128 要素, B 128 要素を配置 • レジスタ: A 8要素, B 8要素, C 8x8 要素を配置 52 i k k C A B 1 work-item の処理領域 8x8 j __kernel void matrix_multiply_gpu_kernel( __global float *A, __global float *B, __global float *C, int M, int N, int K) { // snip float ab[8][8] = { 0.0f }; for (int k = 0; k < K; ++k) { (A 128 要素を LDS にロード) (B 128 要素を LDS にロード) (LDS からレジスタに A 8要素, B 8要素をロード) (8x8 の部分行列積を計算) } (結果を C 8x8 要素にストア) } 128 128 1 work-item あたり 8x8 要素, 16x16 work-item で 128x128 要素を処理 LDS へのロードは 16x16 work-item が協調して行う
  50. Fixstars Group www.fixstars.com Copyright © Fixstars Group 各ブロッキング実装の性能 • 最大

    2018 GFLOPS (ピーク性能比 23.4%) • LDS ブロッキングにより性能低下 53 0 2000 4000 6000 8000 10000 1024 2048 3072 4096 5120 6144 7168 8192 性能 (GFLOPS) 行列サイズ (M=N=K) Peak Performance (1) Naïve (2) Register blocking only (3) (2) + LDS blocking
  51. Fixstars Group www.fixstars.com Copyright © Fixstars Group LDSブロッキングによる性能低下の原因 • データを効率良くロードできていない

    • メモリアクセスの粒度が小さい • メモリアクセスが連続でない 54 i k k C A B 1 work-item の処理領域 8x8 j 128 128 256 work-item で 128要素のロード 半分の work-item が空く メモリアクセスが 連続でない (ストライド)
  52. Fixstars Group www.fixstars.com Copyright © Fixstars Group 高速化3: ループアンローリング •

    Kループをアンロール • 一度にロードするデータ量の増加, 連続アクセス幅増加 • インデックス計算の簡素化 55 i k k C A B j __kernel void matrix_multiply_gpu_kernel( __global float *A, __global float *B, __global float *C, int M, int N, int K) { // snip float ab[4][4] = { 0.0f }; for (int k = 0; k < K; k += 8) { (A 128x8 要素を LDS にロード) (B 8x128 要素を LDS にロード) (LDS からレジスタに A 8x8 要素, B 8x8 要素をロード) (8x8x8 の部分行列積を計算) } // snip } 8 8 LDS へのロードは 16x16 work-item が協調して行う メモリアクセス粒度の増加, 8要素連続アクセス
  53. Fixstars Group www.fixstars.com Copyright © Fixstars Group Kループアンロール実装の性能 • 大幅に性能向上

    • 最大 5141 GFLOPS (ピーク性能比 59.8%) 56 0 2000 4000 6000 8000 10000 1024 2048 3072 4096 5120 6144 7168 8192 性能 (GFLOPS) 行列サイズ (M=N=K) Peak Performance (1) Naïve (2) Register blocking only (3) (2) + LDS blocking (4) (3) + K loop unrolling
  54. Fixstars Group www.fixstars.com Copyright © Fixstars Group 高速化4: ソフトウェアパイプライニング •

    演算とデータ転送を同時に行う • LDS バッファを2つ用意し、片方のバッファを使って計算している間に もう片方のバッファにデータをロードする 57 __kernel void matrix_multiply_gpu_kernel( __global float *A, __global float *B, __global float *C, int M, int N, int K) { // snip int X = 0, Y = 1; (A, B 128x8 要素を LDS[X] にロード) float ab[4][4] = { 0.0f }; for (int k = 0; k < K; k += 8) { (次の A, B 128x8 要素をロード開始) (LDS[X] からレジスタに A 8要素, B 8要素をロード) (8x8x8 の部分行列積を計算) (A, B のロード完了を待ち、LDS[Y] にストア) X ^= 1, Y ^= 1; // バッファの入れ替え } // snip } k = i の計算中に k = i + 1 のロードを 先行して行う
  55. Fixstars Group www.fixstars.com Copyright © Fixstars Group 最終的な高速化結果 • 最大

    6351 GFLOPS (ピーク性能比 73.8%) • 単純な実装から11倍の高速化を達成 58 0 2000 4000 6000 8000 10000 1024 2048 3072 4096 5120 6144 7168 8192 性能 (GFLOPS) 行列サイズ (M=N=K) Peak Performance (1) Naïve (2) Register blocking only (3) (2) + LDS blocking (4) (3) + K loop unrolling (5) (4) + Other optimizations
  56. Fixstars Group www.fixstars.com Copyright © Fixstars Group さらなる高速化余地 • 命令数の削減・ループ外への移動

    • メインループに含まれる命令: 積和, ロード, インデックス計算, branch 等 • ベクタALUではできるだけ積和演算だけを行う状態にしたい • インデックス計算の割合が小さくなるようにループ外へ移動 • LDS-レジスタ転送と演算のパイプライニング • LDS ロードのレイテンシを隠蔽し LDS からのデータ転送と演算をオーバーラップさせる • 問題サイズが小さい場合への対応 • ブロッキングサイズを小さくする • 並列度が足りない場合には K 方向を並列化する • etc. 59
  57. Fixstars Group www.fixstars.com Copyright © Fixstars Group まとめ • ルーフラインモデル

    • 「達成可能な性能」を見積るための性能モデル • プログラムの演算量・メモリアクセス量, プロセッサの演算性能・メモリ帯域を考慮 • AMD GPU アーキテクチャ • SIMD + HW スレッドからなる • AMD GPU における行列積の実装 • GPU 行列積を題材にルーフラインモデルの活用法を紹介 • 「達成可能な性能」の見積り • 性能ボトルネックの分析 • レジスタ・LDSブロッキングサイズ選択 60
  58. Fixstars Group www.fixstars.com Copyright © Fixstars Group 参考文献 • Samuel

    Williams, Andrew Waterman, and David Patterson. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009) • White Paper | AMD GRAPHICS CORES NEXT (GCN) ARCHITECTURE • AMD Radeon™ HD 7970 with graphics core next (GCN) architecture. Hot Chips 24. https://ieeexplore.ieee.org/document/7476485 • AMD GRAPHIC CORE NEXT - AMD Developer Central http://developer.amd.com/wordpress/media/2013/06/2620_final.pdf • AMD GCN3 ISA Architecture Manual https://gpuopen.com/compute-product/amd-gcn3-isa-architecture-manual/ • AMDGPU Compute Application Binary Interface. https://github.com/ROCm-Developer-Tools/ROCm-ComputeABI- Doc/blob/master/AMDGPU-ABI.md • AMD Catalyst OpenCL 2.0 ABI description. https://github.com/CLRX/CLRX-mirror/wiki/AmdCl2Abi • Junjie Lai et al., “Performance Upper Bound Analysis and Optimization of SGEMM on Fermi and Kepler GPUs” • Scott Gray, “SGEMM · NervanaSystems/maxas Wiki · GitHub” https://github.com/NervanaSystems/maxas/wiki/SGEMM 61