--- title: いまさら聞けない！CUDA高速化入門 ～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29） tags: #cuda #gpu #cuda高速化 #高速化シリーズ author: [株式会社フィックスターズ](https://www.docswell.com/user/fixstars) site: [Docswell](https://www.docswell.com/) thumbnail: https://bcdn.docswell.com/page/2EVLWR8MJQ.jpg?width=480 description: GPUは幅広い産業分野において適用範囲が拡大し、近年では特に機械学習での活用が注目されています。 本ウェビナーでは、CUDAプログラミングモデルとハードウェアアーキテクチャについて基礎から解説を行い、高速化の実践例を紹介します。 published: October 29, 21 canonical: https://www.docswell.com/s/fixstars/KGXL6D-20211029 --- # Page. 1 ![Page Image](https://bcdn.docswell.com/page/2EVLWR8MJQ.jpg) Fixstars Corporation www.fixstars.com ＜開演前＞ 2021年10月29日(金) 15:00開始予定 マイク、カメラをOFFにしてしばらくお待ちください 質問ございましたら随時チャット欄へ書き込みください Copyright © Fixstars Corporation # Page. 2 ![Page Image](https://bcdn.docswell.com/page/57GXDM5XEL.jpg) Fixstars Corporation www.fixstars.com いまさら聞けない！ CUDA高速化入門 Copyright © Fixstars Corporation Copyright © Fixstars Corporation # Page. 3 ![Page Image](https://bcdn.docswell.com/page/4EQ81RZ5JP.jpg) Fixstars Corporation www.fixstars.com 本日のアジェンダ 15:00-15:10 • フィックスターズの紹介 15:10-16:20 • ＣUDA高速化入門 16:20-16:30 • • なぜGPUなのか？ • CUDAプログラミングモデル • ハードウェアアーキテクチャ • 高速化実践例 Q&A Copyright © Fixstars Corporation 3 # Page. 4 ![Page Image](https://bcdn.docswell.com/page/KJ42R83VE1.jpg) Fixstars Corporation www.fixstars.com フィックスターズのご紹介 Copyright © Fixstars Corporation Copyright © Fixstars Corporation # Page. 5 ![Page Image](https://bcdn.docswell.com/page/LE1MQ2147G.jpg) Fixstars Corporation www.fixstars.com 会社概要 グループ体制 株式会社フィックスターズ 会社名 株式会社フィックスターズ 本社所在地 東京都港区芝浦3-1-1 msb Tamachi 田町ステーションタワーN 28階 設立 2002年8月 株式会社Fixstars Autonomous Technologies 株式会社ネクスティ エレクトロニクスとのJV 自動運転向けソフトウェア開発に特化 Fixstars Solutions, Inc. 上場区分 東証一部（証券コード：3687） 代表取締役社 長 三木 聡 株式会社Sider 資本金 5億5,401万円（2021年3月現在） 当社完全子会社 ソースコードレビュー、監視・修正漏れ検知ツール開発 社員数（連 結） 253名（2020年9月現在） 当社完全子会社 米国での営業及び開発を担当 株式会社Smart Opinion プロディジーメディカル株式会社とのJV 乳がんAI画像診断支援事業を担当 キオクシア株式会社 主なお客様 株式会社日立製作所 オスカーテクノロジー株式会社 株式会社ネクスティ エレクトロニクス 連結子会社 ソフトウェア自動並列化サービスを提供 キヤノン株式会社 Copyright © Fixstars Corporation 5 # Page. 6 ![Page Image](https://bcdn.docswell.com/page/GEW3MR8ZE2.jpg) Fixstars Corporation www.fixstars.com フィックスターズの強み フィックスターズは、コンピュータの性能を最大限に引き出し大量データの高速処理を実現する、 高速化のエキスパート集団です。 低レイヤ ソフトウェア技術 アルゴリズム 実装力 Copyright © Fixstars Corporation 各産業・研究 分野の知見 6 # Page. 7 ![Page Image](https://bcdn.docswell.com/page/47ZVQR8LE3.jpg) Fixstars Corporation www.fixstars.com ソフトウェア高速化サービス概要 お客様のソースコードをご提供いただき、 ソフトウェアの最適化やアルゴリズムの改良を行い高速化したコードをお返しします。 オリジナルソースコードのご提供 当社 高速化したソースコード コンサルティング 高速化 お客様 サポート 性能評価 アルゴリズムの改良・開発 レポートやコードへのQ&A ボトルネックの特定 ハードウェアへの最適化 実製品への組込み支援 レポート作成 Copyright © Fixstars Corporation 7 # Page. 8 ![Page Image](https://bcdn.docswell.com/page/YJ656RPM7V.jpg) Fixstars Corporation www.fixstars.com ソフトウェア高速化サービス領域 大量データの高速処理がお客様の製品競争力の源泉となる、 様々な領域でソフトウェア開発・高速化サービスを提供しています。 Semiconductor Industrial ・NAND型フラッシュメモリ向けファー ・Smart Factory化支援 ムウェア開発 ・マシンビジョンシステムの高速化 ・次世代AIチップ向け開発環境基盤開発 Mobility Life Science ・自動運転の高性能化、実用化 ・ゲノム解析の高速化 ・次世代パーソナルモビリティの研究開発 ・医用画像処理の高速化 ・AI画像診断システムの研究開発 Finance ・デリバティブシステムの高速化 ・HFT(アルゴリズムトレード)の高速化 Copyright © Fixstars Corporation 8 # Page. 9 ![Page Image](https://bcdn.docswell.com/page/GJ5Y68KQJ4.jpg) Fixstars Corporation www.fixstars.com 画像処理・アルゴリズム開発サービス • お客様の課題 • • • 高度な画像処理、深層学習等のアルゴリズム開発を行える人材が社内に限られている 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件まで クリアできるか不安 研究開発の成果が製品化にうまく結びつかない • 弊社の支援内容 • • • 課題に応じたアルゴリズム調査 深層学習ネットワーク精度改善、推論高速化手法調査 論文調査、実装 https://www.cs.toronto.edu/~frossard/post/vgg16/ Copyright © Fixstars Corporation 9 # Page. 10 ![Page Image](https://bcdn.docswell.com/page/9E2GMKWWER.jpg) Fixstars Corporation www.fixstars.com AI・深層学習関連サービス • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績 ネットワーク設計 データの前処理、データ拡張 精度改善 分散処理による学習高速化 各種DLフレームワーク ターゲットデバイスへの ポーティング及び推論高速化 ■ Visconti, ARM, GPU, DSP ■ SIMD,NEON,CUDA,TensorRT モデル圧縮 - 量子化 - 枝刈り - 蒸留 クラウド・サーバ Copyright © Fixstars Corporation エッジ 10 # Page. 11 ![Page Image](https://bcdn.docswell.com/page/D7YVXNL97M.jpg) Fixstars Corporation www.fixstars.com GPU向け高速化サービス • お客様の課題 • • GPU 高速化の知見がない 自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • • • GPU 高速化に関するコンサルティング ボトルネック調査、GPU プログラムの高速化 CPU/GPU が混在するヘテロジニアス環境での最適化 Copyright © Fixstars Corporation 10～150 倍の 高速化事例あり 11 # Page. 12 ![Page Image](https://bcdn.docswell.com/page/VENM1X4DJ8.jpg) Fixstars Corporation www.fixstars.com GPU開発の課題と当社のサービス 製品企画 ハードウェア選定 • 原価はあまり上げたくないものの 必要な演算リソースは確保したい • GPUはどんな計算を速くできる？ • アルゴリズムをカスタムしても性能が 出るか？ アルゴリズム 設計・実装 • 前処理・後処理も速くできる？ • 機能を追加するため、もう少し処理を 速くしたい 性能・精度 チューニング 品質確保 • GPUで動かしてみたものの 期待した性能が出ない • 品質確保のため、精度を上げたく 演算量は増えるが性能は維持したい Copyright © Fixstars Corporation • GPU高速化に関するコンサルティング • ボトルネック解析 • アルゴリズムのGPU向け設計・実装 • GPUプログラム高速化 • CPU / GPU混在環境での高速化 • 精度向上 • CUDAプログラミングモデルの理解 • ハードウェアアーキテクチャの理解 12 # Page. 13 ![Page Image](https://bcdn.docswell.com/page/Y79N58Q873.jpg) Fixstars Corporation www.fixstars.com CUDA高速化入門 Copyright © Fixstars Corporation Copyright © Fixstars Corporation # Page. 14 ![Page Image](https://bcdn.docswell.com/page/G78KY8GZJD.jpg) Fixstars Corporation www.fixstars.com なぜGPUを使うのか • CPUと比べて • • ピーク性能の高さ 電力効率の良さ 浮動小数点数演算性能 CPU: AMD Ryzen 9 5950X GPU: NVIDIA GeForce RTX 3070 • メモリバンド幅 TDP 価格 2.25 [TFLOPS] 51.2 [GB/s] 105 [W] ¥90,000~ 20.31 [TFLOPS] 448.0 [GB/s] 220 [W] ¥93,000~ その他のアクセラレータと比べて • • 入手性・価格性能比の良さ プログラミングの容易さ Copyright © Fixstars Corporation 14 # Page. 15 ![Page Image](https://bcdn.docswell.com/page/L7LZ4RG1ER.jpg) Fixstars Corporation www.fixstars.com なぜGPUが速いのか • 並列計算に特化した構成 • 大量のコア・演算器 • • • • CPU: AMD EPYC 7763: 64 Cores, 32 FLOPs/Core/cycle GPU: NVIDIA A100: 108 SMs, 128 FLOPs/SM/cycle バス幅の広い広帯域メモリ もちろん弱点もある • • 並列に処理できない問題には弱い 最大メモリ容量が小さい Copyright © Fixstars Corporation 15 # Page. 16 ![Page Image](https://bcdn.docswell.com/page/4EMVKRQ5EW.jpg) Fixstars Corporation www.fixstars.com CUDAプログラミングモデル Copyright © Fixstars Corporation Copyright © Fixstars Corporation # Page. 17 ![Page Image](https://bcdn.docswell.com/page/PER2XR8Z79.jpg) Fixstars Corporation www.fixstars.com 例題: saxpy • Single-precision ax plus y • y←a×x+y • CPU向けの実装例: void saxpy(float *y, const float *x, float a, int n){ for(int i = 0; i < n; ++i){ y[i] = a * x[i] + y[i]; } } Copyright © Fixstars Corporation 17 # Page. 18 ![Page Image](https://bcdn.docswell.com/page/P7XD6R81JX.jpg) Fixstars Corporation www.fixstars.com CUDAを用いたプログラムの流れ • ホストメモリからデバイスメモリへデータを転送 • GPU上でカーネル（プログラム）を実行 • デバイスメモリからホストメモリへデータを転送 Copyright © Fixstars Corporation 18 # Page. 19 ![Page Image](https://bcdn.docswell.com/page/37KX1RKMJD.jpg) Fixstars Corporation www.fixstars.com ホストメモリとデバイスメモリ • CPUとGPUはそれぞれがメモリを持っている • • 目的に応じて適切なほうを利用する 必要に応じて片方から他方へデータをコピーする ～200 GB/s ホストメモリ (DDR) CPU ～20 GB/s ～2000 GB/s GPU デバイスメモリ (GDDR/HBM) Copyright © Fixstars Corporation 19 # Page. 20 ![Page Image](https://bcdn.docswell.com/page/LJ35Q8Z275.jpg) Fixstars Corporation www.fixstars.com ホストメモリからデバイスメモリへデータを転送 • cudaMalloc • • • デバイスメモリ上の領域を確保 標準Cにおけるmallocに対応 cudaMemcpy • • デバイスメモリに関係するメモリコピー 第4引数で転送の方向を指定 (HostToDevice, DeviceToHost など) float *d_y, d_x; // デバイスメモリの確保 cudaMalloc(&d_x, sizeof(float) * n); cudaMalloc(&d_y, sizeof(float) * n); // ホストメモリ (h_x, h_y) から sizeof(float) * n バイト転送 cudaMemcpy(d_x, h_x, sizeof(float) * n, cudaMemcpyHostToDevice); cudaMemcpy(d_y, h_y, sizeof(float) * n, cudaMemcpyHostToDevice); Copyright © Fixstars Corporation 20 # Page. 21 ![Page Image](https://bcdn.docswell.com/page/8JDD9MR6JG.jpg) Fixstars Corporation www.fixstars.com GPU上でカーネル（プログラム）を実行 • カーネルの呼び出し • • • スレッド数を指定する スレッドブロック数×ブロックあたりのスレッド数で表現 ここではループ1回を1スレッドで処理する const int bdim = 128; const int gdim = (n + bdim – 1) / bdim; kernel<<<gdim, bdim>>>(d_y, d_x, a, n); Copyright © Fixstars Corporation // 切り上げ 21 # Page. 22 ![Page Image](https://bcdn.docswell.com/page/VEPXZRMZJ8.jpg) Fixstars Corporation www.fixstars.com GPU上で動くカーネルの実装 • __global__ 修飾された関数として定義 • 定義済み変数から自身のインデックスを取得 • • • blockDim: 現在のカーネル実行におけるブロックサイズ blockIdx: 自身の属するスレッドブロックのインデックス threadIdx: 自身のスレッドブロック内におけるインデックス __global__ void kernel(float *y, const float *x, float a, int n){ const int i = threadIdx.x + blockIdx.x * blockDim.x; if(i < n) y[i] = a * x[i] + y[i]; } Copyright © Fixstars Corporation 22 # Page. 23 ![Page Image](https://bcdn.docswell.com/page/27VLWR9MEQ.jpg) Fixstars Corporation www.fixstars.com デバイスメモリからホストメモリへデータを転送 • cudaMemcpyで逆方向にコピー // デバイスメモリ (d_y) から sizeof(float) * n バイト転送 cudaMemcpy(h_y, d_y, sizeof(float) * n, cudaMemcpyDeviceToHost); // デバイスメモリの解放 cudaFree(d_x); cudaFree(d_y); Copyright © Fixstars Corporation 23 # Page. 24 ![Page Image](https://bcdn.docswell.com/page/5JGXDMZX7L.jpg) Fixstars Corporation www.fixstars.com スレッドの階層構造 • CUDAではスレッド間に階層構造がある • 近いスレッド同士はより密に通信・同期を行うことができる Grid Thread Block (～1024T) Warp (32T) … … Copyright © Fixstars Corporation … 24 # Page. 25 ![Page Image](https://bcdn.docswell.com/page/47Q81RL5EP.jpg) Fixstars Corporation www.fixstars.com スレッドの階層構造 • CUDAではスレッド間に階層構造がある • Warp: 同時に命令が発行されるスレッドをまとめたもの • • Thread Block: いくつかのスレッドをまとめたもの • • • 現行アーキテクチャでは32スレッド 現行アーキテクチャでは1ブロックあたり最大1024スレッド 同一ワープに属するスレッドは必ず同一スレッドブロックに属する Grid: いくつかのスレッドブロックをまとめたもの • カーネル呼び出しは1つのグリッドで処理される Copyright © Fixstars Corporation 25 # Page. 26 ![Page Image](https://bcdn.docswell.com/page/KE42R8DV71.jpg) Fixstars Corporation www.fixstars.com スレッドの階層構造 • 階層構造上で近いスレッド同士はより密に同期や通信を行うことができる • 同一グリッド • • 同一スレッドブロック • • • カーネル起動・終了時の同期のみ 同じブロックに属するスレッド同士での同期 シェアードメモリを用いたデータ共有 同一ワープ • • スレッド同士でのより軽量な同期 ワープシャッフルによるデータ共有 Copyright © Fixstars Corporation 26 # Page. 27 ![Page Image](https://bcdn.docswell.com/page/L71MQ2Z4JG.jpg) Fixstars Corporation www.fixstars.com メモリの階層構造 • メモリにも階層構造がある • おおむねスレッドの階層構造と対応 Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 27 # Page. 28 ![Page Image](https://bcdn.docswell.com/page/G7W3MR9Z72.jpg) Fixstars Corporation www.fixstars.com メモリの階層構造: レジスタ • プログラム中の自動変数に対応 • 各種演算命令に直接渡すことができる • 他のスレッドとは共有されない Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 28 # Page. 29 ![Page Image](https://bcdn.docswell.com/page/4JZVQR9L73.jpg) Fixstars Corporation www.fixstars.com メモリの階層構造: ローカルメモリ • プログラム中の自動変数に対応 • 何らかの理由でレジスタに乗せられないときに使用される • 演算命令に渡す際はいったんレジスタにロードする必要がある • 他のスレッドとは共有されない Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 29 # Page. 30 ![Page Image](https://bcdn.docswell.com/page/YE656RKMJV.jpg) Fixstars Corporation www.fixstars.com メモリの階層構造: シェアードメモリ • __shared__ 修飾された変数に対応 • 同一スレッドブロック内の全スレッドで共有される Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 30 # Page. 31 ![Page Image](https://bcdn.docswell.com/page/GE5Y68PQE4.jpg) Fixstars Corporation www.fixstars.com メモリの階層構造: グローバルメモリ • cudaMalloc などで確保された領域に対応 • デバイス全体で共有される • カーネル停止後も値が保持される Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 31 # Page. 32 ![Page Image](https://bcdn.docswell.com/page/972GMK6W7R.jpg) Fixstars Corporation www.fixstars.com メモリの階層構造: コンスタントメモリ • __constant__ 修飾された変数に対応 • デバイス全体で共有される • カーネルから値を書き換えることができない Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 32 # Page. 33 ![Page Image](https://bcdn.docswell.com/page/DJYVXN99JM.jpg) Fixstars Corporation www.fixstars.com ホストとデバイス間の同期 • カーネル呼び出しやデータ転送は基本的に非同期実行 • 明示的もしくは暗黙的に同期を挿入する必要がある • cudaMemcpy など一部のAPIは自動的に同期を挿入する kernel<<<1, 1>>>(); // この時点では kernel() はまだ実行されていないかもしれない foo(); cudaDeviceSynchronize(); // この時点では kernel() の処理は確実に完了している CPU GPU foo() cudaDeviceSynchronize() kernel() Copyright © Fixstars Corporation 33 # Page. 34 ![Page Image](https://bcdn.docswell.com/page/V7NM1XLDE8.jpg) Fixstars Corporation www.fixstars.com ストリーム • デバイスで実行される処理のキュー • • • 投入した順に処理される 同じストリームに投入された処理同士はオーバーラップしない 指定されなかった場合はデフォルトストリームが使用される kernel1<<<1, 1>>>(); kernel2<<<1, 1>>>(); cudaDeviceSynchronize(); CPU GPU cudaDeviceSynchronize() kernel1() Copyright © Fixstars Corporation kernel2() 34 # Page. 35 ![Page Image](https://bcdn.docswell.com/page/YJ9N5848J3.jpg) Fixstars Corporation www.fixstars.com ストリーム • ストリームは複数作成することができる • 別ストリームに投入された処理同士は並行するかもしれない kernel1<<<1, 1, 0, stream1>>>(); kernel2<<<1, 1, 0, stream2>>>(); cudaDeviceSynchronize(); CPU Synchronize GPU kernel1() kernel2() Copyright © Fixstars Corporation 35 # Page. 36 ![Page Image](https://bcdn.docswell.com/page/GJ8KY8QZED.jpg) Fixstars Corporation www.fixstars.com プログラミングモデルまとめ • 大量のスレッドの間には階層関係がある • • • メモリにも階層関係がある • • • ワープ・スレッドブロック・グリッド 距離に応じて同期や通信の制約が変化する レジスタ・ローカルメモリ・シェアードメモリ・グローバルメモリ 速度や共有する必要があるスレッド数など要求に応じて適切な領域を使い分ける デバイス上で動く処理は基本的に非同期実行となる • 細かい同期周りの制御にはストリームを活用する Copyright © Fixstars Corporation 36 # Page. 37 ![Page Image](https://bcdn.docswell.com/page/LJLZ4RX17R.jpg) Fixstars Corporation www.fixstars.com ハードウェア Copyright © Fixstars Corporation Copyright © Fixstars Corporation # Page. 38 ![Page Image](https://bcdn.docswell.com/page/47MVKRL57W.jpg) Fixstars Corporation www.fixstars.com カーネルが遅い原因と対応策 要求されている演算量が多すぎる • アルゴリズムを改善して演算量を減らす メモリトラフィックが多すぎる • アルゴリズムを改善してメモリアクセスを減らす • キャッシュなどのハードウェア機能を活用する リソースを有効活用できていない • ハードウェアの制約を理解して演算器やバスなどの稼働率を上げる Copyright © Fixstars Corporation 38 # Page. 39 ![Page Image](https://bcdn.docswell.com/page/P7R2XRKZJ9.jpg) Fixstars Corporation www.fixstars.com カーネルが遅い原因と対応策 要求されている演算量が多すぎる • アルゴリズムを改善して演算量を減らす メモリトラフィックが多すぎる • アルゴリズムを改善してメモリアクセスを減らす • キャッシュなどのハードウェア機能を活用する リソースを有効活用できていない • ハードウェアの制約を理解して演算器やバスなどの稼働率を上げる Copyright © Fixstars Corporation 39 # Page. 40 ![Page Image](https://bcdn.docswell.com/page/PJXD6RL1EX.jpg) Fixstars Corporation www.fixstars.com NVIDIA A100 Block Diagram • CC 8.0 • 108 SMs/Chip • 6912 FP32 CUDA Cores • コアを活用できるだけの並行実行可能なタスク (=スレッド) を投入する必要がある • スレッド数が足りないならタスクを分割することも視野に入れる https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf Copyright © Fixstars Corporation 41 # Page. 41 ![Page Image](https://bcdn.docswell.com/page/3JKX1RLMED.jpg) Fixstars Corporation www.fixstars.com Streaming Multiprocessor (SM) • スレッドブロックに対応する • • いくつかのスレッドブロックを並行して処理 以下の要素を束ねたもの • • • • • • • CUDA Core Tensor Core LD/ST Unit SFU Register File Cache/Shared Memory Scheduler, Dispatcher Copyright © Fixstars Corporation 42 # Page. 42 ![Page Image](https://bcdn.docswell.com/page/LE35Q832J5.jpg) Fixstars Corporation www.fixstars.com Streaming Multiprocessor (SM) • スレッドブロックに対応する • 以下の要素を束ねたもの • • • • • • • CUDA Core Tensor Core LD/ST Unit SFU Register File Cache/Shared Memory Scheduler, Dispatcher 演算器 メモリ Copyright © Fixstars Corporation 43 # Page. 43 ![Page Image](https://bcdn.docswell.com/page/8EDD9M46EG.jpg) Fixstars Corporation www.fixstars.com Processing Block • ワープに対応する • • いくつかのワープを並行して処理 SMからワープをまたがない要素を分割したもの • • • 各種演算器 レジスタファイル スケジューラ・ディスパッチャ Copyright © Fixstars Corporation 44 # Page. 44 ![Page Image](https://bcdn.docswell.com/page/V7PXZRVZE8.jpg) Fixstars Corporation www.fixstars.com CUDA Core • スレッドに対応する • 何らかの演算を行う • • • レジスタファイルから値を読んで 演算を行って レジスタファイルに書き出す FP32/INT32 • Volta以降でINTコアが分離された • 整数演算と浮動小数点数演算を同時に実行できる Copyright © Fixstars Corporation 45 # Page. 45 ![Page Image](https://bcdn.docswell.com/page/2JVLWRZM7Q.jpg) Fixstars Corporation www.fixstars.com Tensor Core • 深層学習向けのアクセラレータ • ワープ単位で協調して小さい行列積を効率よく行う • 世代によって対応する精度・サイズが異なる Copyright © Fixstars Corporation 46 # Page. 46 ![Page Image](https://bcdn.docswell.com/page/5EGXDM4XJL.jpg) Fixstars Corporation www.fixstars.com その他のユニット • LD/ST (Load/Store) • • メモリアクセスを行う SFU (Special Function Unit) • • 特殊関数 (指数関数・三角関数など) の処理を行う 演算器が少ない分スループットも落ちる Copyright © Fixstars Corporation 47 # Page. 47 ![Page Image](https://bcdn.docswell.com/page/4JQ81RG57P.jpg) Fixstars Corporation www.fixstars.com SIMTとWarp • ディスパッチャはワープに対して一つの命令を一度に発行する • • • SIMT: Single Instruction, Multiple Threads スレッドごとに異なる命令を発行することはできない 条件分岐の取り扱い • • 分岐によって実行の必要がなくなった命令も発行されうる そのような場合はその命令が無視される Copyright © Fixstars Corporation 49 # Page. 48 ![Page Image](https://bcdn.docswell.com/page/K742R85VJ1.jpg) Fixstars Corporation www.fixstars.com Warp Divergence • 条件分岐によって有効な演算を行わないスレッド (=コア) が発生する • ワープ内での異なる方向への分岐は性能劣化につながる • • Warp Divergence と呼ぶ 下の例では B(), C() の処理中にコアが半分遊んでいる A() A(); if(threadIdx.x % 2 == 0){ B(); }else{ C(); } B() C() Warp 0 Warp 1 Copyright © Fixstars Corporation 50 # Page. 49 ![Page Image](https://bcdn.docswell.com/page/LJ1MQ2W4EG.jpg) Fixstars Corporation www.fixstars.com Warp Divergence • できるだけ同じワープのスレッドが同じように動くことで効率を改善できる • • 連続するスレッドが同じ方向に分岐するようにする 下の例では B(), C() におけるコア稼働率が改善している A() A(); if(threadIdx.x < 4){ B(); }else{ C(); } B() C() Warp 0 Warp 1 Copyright © Fixstars Corporation 51 # Page. 50 ![Page Image](https://bcdn.docswell.com/page/GJW3MR6ZJ2.jpg) Fixstars Corporation www.fixstars.com レイテンシの隠蔽 • 命令のパイプライニングはGPUでも有効 • • 依存性のない命令同士を並列実行する 依存性のない命令の組をどう見つけるか • • 近くにある命令との依存性を解析する 別のスレッドの命令と組み合わせる Copyright © Fixstars Corporation 52 # Page. 51 ![Page Image](https://bcdn.docswell.com/page/4EZVQRYLJ3.jpg) Fixstars Corporation www.fixstars.com ワープスケジューリング • Processing Block はいくつかの実行中ワープの状態を保持している • • サイクルごとに実行可能なワープをその中から選択して命令を発行する • • 可能であれば物理コア数より多くのスレッドの状態を保持する 実行可能: 次に発行される命令が依存している処理がすべて完了している 実行可能なワープを絶やさないことが効率改善につながる • • 命令のレイテンシを考慮したプログラムを記述する 実行可能なワープの候補 (=状態を保持しているスレッド数) を増やす Copyright © Fixstars Corporation 53 # Page. 52 ![Page Image](https://bcdn.docswell.com/page/Y7656RDMEV.jpg) Fixstars Corporation www.fixstars.com レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • • FADDのレイテンシは4とする 並行実行しているワープ数が1の場合: 4 ops / 16 cycles 0x00: 0x01: 0x02: 0x03: FADD FADD FADD FADD R1, R1, R1, R1, R2 R3 R4 R5 Warp 0 R1 += R2 R1 += R3 Copyright © Fixstars Corporation R1 += R4 R1 += R5 54 # Page. 53 ![Page Image](https://bcdn.docswell.com/page/G75Y683Q74.jpg) Fixstars Corporation www.fixstars.com レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • FADDのレイテンシは4とする • 並行実行しているワープ数が1の場合: 4 ops / 16 cycles • 並行実行しているワープ数が4の場合: 16 ops / 19 cycles 0x00: 0x01: 0x02: 0x03: FADD FADD FADD FADD R1, R1, R1, R1, R2 R3 R4 R5 Warp 0 Warp 1 Warp 2 Warp 3 R1 += R2 R1 += R3 R1 += R2 R1 += R4 R1 += R3 R1 += R2 R1 += R4 R1 += R3 R1 += R2 Copyright © Fixstars Corporation R1 += R5 R1 += R3 R1 += R5 R1 += R4 R1 += R4 R1 += R5 R1 += R5 55 # Page. 54 ![Page Image](https://bcdn.docswell.com/page/9J2GMKZWJR.jpg) Fixstars Corporation www.fixstars.com Occupancy • SMがいくつのワープを並行実行できるかを表す指標 • • ブロックサイズ・消費レジスタ数・シェアードメモリサイズから求める • • • • 高ければ高いほどレイテンシを隠蔽しやすい ブロックサイズ: SMあたりの並行実行可能なブロック数 消費レジスタ数: SMあたりのレジスタファイル数 シェアードメモリサイズ: SMあたりのシェアードメモリサイズ プロファイラ・CUDA Toolkit 付属のExcelシートなどで求められる Copyright © Fixstars Corporation 56 # Page. 55 ![Page Image](https://bcdn.docswell.com/page/DEYVXNR9EM.jpg) Fixstars Corporation www.fixstars.com メモリ階層 • 上に行くほど高速な代わりに共有される範囲が狭まる SM Processing Block Register File L1 Cache Shared Memory L2 Cache Device Memory Copyright © Fixstars Corporation 57 # Page. 56 ![Page Image](https://bcdn.docswell.com/page/VJNM1XDD78.jpg) Fixstars Corporation www.fixstars.com デバイスメモリ • グローバルメモリ・ローカルメモリに対応 • GPUのスペックに書かれている容量はこの領域のもの • アクセスパターンによって大きく性能が変わる Copyright © Fixstars Corporation 58 # Page. 57 ![Page Image](https://bcdn.docswell.com/page/YE9N58G8E3.jpg) Fixstars Corporation www.fixstars.com Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 1トランザクション / 1アクセス スレッド メモリ 0 1 2 3 4 5 … 31 … Copyright © Fixstars Corporation 59 # Page. 58 ![Page Image](https://bcdn.docswell.com/page/GE8KY8VZ7D.jpg) Fixstars Corporation www.fixstars.com Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 2トランザクション / 1アクセス スレッド メモリ 0 1 2 3 4 5 … 31 … Copyright © Fixstars Corporation 60 # Page. 59 ![Page Image](https://bcdn.docswell.com/page/LELZ4R51JR.jpg) Fixstars Corporation www.fixstars.com Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 32トランザクション / 1アクセス スレッド メモリ 0 1 … 2 3 … 4 … … 5 … 31 … … Copyright © Fixstars Corporation … 61 # Page. 60 ![Page Image](https://bcdn.docswell.com/page/4JMVKR35JW.jpg) Fixstars Corporation www.fixstars.com L2 キャッシュ • デバイス中の全SMで共有されている • デバイスメモリへのアクセス時には常に使用される Copyright © Fixstars Corporation 62 # Page. 61 ![Page Image](https://bcdn.docswell.com/page/PJR2XRQZE9.jpg) Fixstars Corporation www.fixstars.com L1キャッシュ • SMごとに用意されている • 明示的に指定したものか読み取り専用のデータへのアクセスに対して使用される 読み取り専用かどうかの判定 • コンパイラが判定する • ポインタを const __restrict__ 修飾すると読み取り専用であることを明示できる 明示的なL1キャッシュの利用 • 組み込み関数 __ldg() を使用する • *ptr → __ldg(ptr) Copyright © Fixstars Corporation 63 # Page. 62 ![Page Image](https://bcdn.docswell.com/page/PEXD6R517X.jpg) Fixstars Corporation www.fixstars.com シェアードメモリ • SMごとに用意された領域 • L1キャッシュとシェアードメモリの割合は設定で変更可能 • • シェアードメモリとして使えるのは 16-96 [KB/SM] 程度 残りはL1キャッシュとして使用される Copyright © Fixstars Corporation 64 # Page. 63 ![Page Image](https://bcdn.docswell.com/page/3EKX1R6M7D.jpg) Fixstars Corporation www.fixstars.com メモリバンク • シェアードメモリはバンクを用いてに管理されている • • バンクは4バイトごとに切り替わる 同じバンクの異なる領域へのアクセスはまとめて処理できない: バンクコンフリクト Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 … 0x00000100 Copyright © Fixstars Corporation 65 # Page. 64 ![Page Image](https://bcdn.docswell.com/page/L735Q862E5.jpg) Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できるアクセスの例 • 素直なシーケンシャルアクセス Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 66 # Page. 65 ![Page Image](https://bcdn.docswell.com/page/87DD9MV67G.jpg) Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できるアクセスの例 • バンクが重複しないランダムアクセス Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 67 # Page. 66 ![Page Image](https://bcdn.docswell.com/page/VJPXZRRZ78.jpg) Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できるアクセスの例 • ブロードキャスト: バンクが重なっても同じアドレスなら問題ない Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 68 # Page. 67 ![Page Image](https://bcdn.docswell.com/page/2EVLWRRMJQ.jpg) Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できないアクセスの例 • ストライドアクセス: この場合は2回に分割される Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 69 # Page. 68 ![Page Image](https://bcdn.docswell.com/page/57GXDMMXEL.jpg) Fixstars Corporation www.fixstars.com レジスタファイル • プロセッシングブロックごとに用意された領域 • レジスタ幅は1要素あたり32bit • long, double, ポインタなどの64bit型には2つ使われる 自動変数に対する領域割り当て • 自動変数は可能ならレジスタに割り当てられる • 特定のケースで低速なローカルメモリに割り当てられる • • 自動変数がレジスタに収まりきらない場合（レジスタスピル） インデックスアクセスが必要な場合 Copyright © Fixstars Corporation 70 # Page. 69 ![Page Image](https://bcdn.docswell.com/page/4EQ81RR5JP.jpg) Fixstars Corporation www.fixstars.com ハードウェアまとめ 演算器 • 演算器を使い切るためには注意が必要なことがある • • 分岐によって何もしないコアが発生することがある レイテンシを埋めるだけの命令供給が必要 メモリ • アクセスパターン次第で効率が落ちることがある • • • グローバルメモリ: Coalescing, キャッシュ利用 シェアードメモリ: バンクコンフリクト ローカルメモリの利用にも注意する • コンパイラの出力を確認すると確実 Copyright © Fixstars Corporation 71 # Page. 70 ![Page Image](https://bcdn.docswell.com/page/KJ42R88VE1.jpg) Fixstars Corporation www.fixstars.com 実践例 Copyright © Fixstars Corporation Copyright © Fixstars Corporation # Page. 71 ![Page Image](https://bcdn.docswell.com/page/LE1MQ2247G.jpg) Fixstars Corporation www.fixstars.com 問題の概要 • 画像のステレオマッチング: Semi-Global Matching (SGM) • ステレオ画像の視差を計算するアルゴリズム • • • 視差: 片方の画像のある画素が他方の画像で何ピクセルずれたところにあるか 近くの物体ほど視差が大きくなること利用して距離を計算できる ターゲット環境: Pascal 世代のGPU (GeForce GTX 10xx など) Copyright © Fixstars Corporation 73 # Page. 72 ![Page Image](https://bcdn.docswell.com/page/GEW3MRRZE2.jpg) Fixstars Corporation www.fixstars.com チューニングする部分 • 動的計画法 (DP) である画素における視差が d [px] としたときのスコアを求める • • 対応する画素同士の特徴ベクトルの距離が大きいとコストが大きい 隣接画素に対して急激な視差の変化があるとコストが大きい 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 𝑠𝑐𝑜𝑠𝑡 𝑥, 𝑑 = 𝑙𝑐𝑜𝑠𝑡 𝑥, 𝑑 + min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃 1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 𝑖 𝑑 − min 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑖) 𝑖 𝑥 Copyright © Fixstars Corporation 74 # Page. 73 ![Page Image](https://bcdn.docswell.com/page/47ZVQRRLE3.jpg) Fixstars Corporation www.fixstars.com チューニングする部分 • 動的計画法である画素における視差が d [px] としたときのスコアを求める • • • 対応する画素同士の特徴ベクトルの距離が大きいとコストが大きい 隣接画素に対して急激な視差の変化があるとコストが大きい これを縦横斜めの8方向それぞれについてラインごとに計算する Copyright © Fixstars Corporation 75 # Page. 74 ![Page Image](https://bcdn.docswell.com/page/YJ656RRM7V.jpg) Fixstars Corporation www.fixstars.com アルゴリズムの概略 左から右方向のスキャン for(int y = 0; y < H; ++y){ int prev_min = 0; for(int x = 0; x < W; ++x){ int cur_min = INT_MAX; for(int d = 0; d < D; ++d){ X方向のループは依存性がある int cost = min({ P2, scost[y][x-1][d-1] - prev_min + P1, 主要な計算は O(HWD) 回行われる scost[y][x-1][d+1] - prev_min + P1, 計算処理はかなり軽い scost[y][x-1][d] - prev_min }); scost[y][x][d] = キャッシュヒットが期待できない cost + dist(left[y][x], right[y][x-d]); メモリアクセスも O(HWD) 回 cur_min = min(prev_min, cost); } prev_min = cur_min; } Y方向のループは完全に独立している } • Copyright © Fixstars Corporation 76 # Page. 75 ![Page Image](https://bcdn.docswell.com/page/GJ5Y688QJ4.jpg) Fixstars Corporation www.fixstars.com 並列化方針の検討 y方向ループ1回を1スレッドで担当する • 並列度が足りない: 数百スレッド程度しか利用できない y方向ループ1回を複数スレッドで担当する • x方向のループは分割できない: 前のループに対する依存性があるため • d方向のループは分割可能 • ただしx方向のループ1回ごとに同期が必要になる Copyright © Fixstars Corporation 77 # Page. 76 ![Page Image](https://bcdn.docswell.com/page/LE35Q88ZJ5.jpg) Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • • d方向の最小値を求める 端の値を隣のスレッドに渡す scostの計算 𝑑 Thread 1 Thread 0 𝑥 Copyright © Fixstars Corporation 78 # Page. 77 ![Page Image](https://bcdn.docswell.com/page/8EDD9MM4EG.jpg) Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • • d方向の最小値を求める 端の値を隣のスレッドに渡す 最小値の計算・共有 𝑑 Thread 1 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃 1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 𝑖 Thread 0 𝑥 Copyright © Fixstars Corporation 79 # Page. 78 ![Page Image](https://bcdn.docswell.com/page/V7PXZRRVE8.jpg) Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • • d方向の最小値を求める 端の値を隣のスレッドに渡す 端の値の共有 𝑑 Thread 1 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃 1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 𝑖 Thread 0 𝑥 Copyright © Fixstars Corporation 80 # Page. 79 ![Page Image](https://bcdn.docswell.com/page/2JVLWRRR7Q.jpg) Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 細かく分割する場合のメリット • スレッド数を増やすことによる Occupancy の向上 • スレッドあたりのレジスタ量の削減 粗く分割する場合のメリット • スレッド間通信などの非本質的な処理の占める割合の減少 • • 1ワープ以内になるとより軽量な通信が利用できる グローバルメモリへのアクセス効率の向上 • • スレッドあたりのメモリアクセス量が多くなる 1回のアクセスで4要素までアクセスできる Copyright © Fixstars Corporation 81 # Page. 80 ![Page Image](https://bcdn.docswell.com/page/5EGXDMM6JL.jpg) Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 細かく分割する場合のメリット • スレッド数を増やすことによる Occupancy の向上 • スレッドあたりのレジスタ量の削減 粗く分割する場合のメリット • スレッド間通信などの非本質的な処理の占める割合の減少 • • 1ワープ以内になるとより軽量な通信が利用できる グローバルメモリへのアクセス効率の向上 1ワープを境に実装が大きく変化する • スレッドあたりのメモリアクセス量が多くなる ⇒ 1ワープ以下の範囲で値を変えつつ試せるように実装する • 1回のアクセスで4要素までアクセスできる Copyright © Fixstars Corporation 82 # Page. 81 ![Page Image](https://bcdn.docswell.com/page/4JQ81RR27P.jpg) Fixstars Corporation www.fixstars.com ハイエンドGPU対策 • ハイエンドGPUだとd方向の分割を入れてもスレッド数が足りない • 8方向それぞれが独立なことを利用して複数カーネルを並行させる Copyright © Fixstars Corporation 83 # Page. 82 ![Page Image](https://bcdn.docswell.com/page/K742R88PJ1.jpg) Fixstars Corporation www.fixstars.com アルゴリズムの検討:局所特徴同士の距離の計算 • 局所特徴についての情報 • • • 局所特徴の表現: 64 bit のビット列 局所特徴の距離: 互いに異なるビットの数 特徴ベクトルの距離は同じ組の距離が何度も使われる • • 8方向すべての処理で同じ計算を行う 既存実装では事前計算してテーブル化されていた • table[y][x][d] = distance(left[y][x], right[y][x - d]) Copyright © Fixstars Corporation 84 # Page. 83 ![Page Image](https://bcdn.docswell.com/page/LJ1MQ22XEG.jpg) Fixstars Corporation www.fixstars.com 理論性能で比べる • テーブル引きと計算どちらが速い? • • • • popcounts/s: 763 [Gops/s] • • • 32 [ops/s/SM] × 28 [SM/s] × 1.481 [GHz] = 1326 [Gops/s] 1要素あたり2回必要なのでその半分 Bytes/s: 484.4 [GB/s] • • テーブル化した場合1要素当たり 1 [byte] 特徴同士の距離は popcount 命令2回で求められる GeForce GTX 1080 Ti (sm_61) を例に試算してみる 実測値だとおよそ 340 [GB/s] くらい 毎回計算するほうが速そう!! Copyright © Fixstars Corporation 85 # Page. 84 ![Page Image](https://bcdn.docswell.com/page/GJW3MRRKJ2.jpg) Fixstars Corporation www.fixstars.com プロファイル結果 (1) • 横方向の処理のプロファイル結果 • 演算器の稼働率80%弱: うまくリソースを活用できてそう Copyright © Fixstars Corporation 86 # Page. 85 ![Page Image](https://bcdn.docswell.com/page/4EZVQRRNJ3.jpg) Fixstars Corporation www.fixstars.com プロファイル結果 (2) • 本当に距離をテーブル化しないほうが速かったのか? • • 実効メモリ帯域で評価する テーブル引きする場合はメモリトラフィックが Reads = Writes になる • • テーブルサイズが結果バッファのサイズと等しいため 52.049×2 = 104.098 [GB/s] 出せなければテーブル化のほうが遅い • • bandwidthTest での帯域が 92.7 [GB/s] 程度 テーブル化する方針では勝てないだろうと考えられる Copyright © Fixstars Corporation 87 # Page. 86 ![Page Image](https://bcdn.docswell.com/page/Y7656RR9EV.jpg) Fixstars Corporation www.fixstars.com プロファイル結果 (3) • 縦方向・斜め方向でも同様の傾向 Copyright © Fixstars Corporation 88 # Page. 87 ![Page Image](https://bcdn.docswell.com/page/G75Y688D74.jpg) Fixstars Corporation www.fixstars.com 全体の評価 • 既存実装との性能比較 • • • 比較対象: Embedded real-time stereo estimation via Semi-Global Matching on the GPU, D. Hernandez-Juarez et al, ICCS 2016. https://github.com/dhernandez0/sgm 実際にはもう一つ大きいカーネルがあるのですがそちらの詳細は省略しています Copyright © Fixstars Corporation 89 # Page. 88 ![Page Image](https://bcdn.docswell.com/page/9J2GMKVMJR.jpg) Fixstars Corporation www.fixstars.com 評価結果 • 2.3-12.5% 程度の高速化 • 演算性能に対してメモリ帯域の細いチップで特に強い フレームレート 300 250 261 232 200 150 今回の実装 100 Hernandez+ 69.7 68.1 50 50.2 45.8 0 GTX 1080 Ti GTX 1050 Ti Copyright © Fixstars Corporation DRIVE PX2 90 # Page. 89 ![Page Image](https://bcdn.docswell.com/page/DEYVXN3PEM.jpg) Fixstars Corporation www.fixstars.com まとめ • 演算とメモリアクセスどちらが重要か見極める • • • 理論性能から見積もり 実測で裏付け コアあたりの効率とチップあたりの効率 • • 並列度を下げると演算量は減らしやすい 一方でリソースが余りやすくなるのでうまくバランスをとる Copyright © Fixstars Corporation 91 # Page. 90 ![Page Image](https://bcdn.docswell.com/page/VJNM1XVM78.jpg) Fixstars Corporation www.fixstars.com 全体のまとめ • パフォーマンスチューニングにおいてはハードウェアの知識も重要 • • • カーネルのチューニングにおいては特に演算器とメモリに気を配る • • • 使い方を誤ると数倍の性能劣化なども起こりうる もちろんアルゴリズムも重要で両方からのアプローチが必要 演算器を余らせない 不得意なアクセスパターンによる性能劣化を防ぐ 理論をもとに仮説を立てて実装したものを評価する • • プロファイラによる評価 理論ピークと実性能の差を読み取る Copyright © Fixstars Corporation 92 # Page. 91 ![Page Image](https://bcdn.docswell.com/page/YE9N586WE3.jpg) Fixstars Corporation www.fixstars.com Thank You お問い合わせ窓口 : contact@fixstars.com Copyright © Fixstars Corporation