Movatterモバイル変換

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

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

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

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

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

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

Fixstars Group www.fixstars.comCopyright © Fixstars Group性能モデルの原始的な例 (1/3)• シンプルな性能モデル• 命令の種類ごとに実行回数を数えて重み付け• 命令数• 浮動小数点演算: 2N• メモリアクセス数: 3N• 実行時間: T = F * 2N + M * 3N• F: 浮動小数点演算1命令あたり実行時間• M: メモリアクセス1命令あたり実行時間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];}F, M は理論演算性能,メモリ帯域から計算24

Fixstars Group www.fixstars.comCopyright © Fixstars Group性能モデルの原始的な例 (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 sec25

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

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

Fixstars Group www.fixstars.comCopyright © Fixstars Groupルーフラインモデル*1• 概要• プログラムが達成可能な演算性能 [FLOPS] を見積もるための性能モデル• 考慮する要素• プログラムにおける演算量・メモリアクセス量• プロセッサの理論演算性能・メモリ帯域• 実行モデル29プロセッサメモリ浮動小数点演算性能:π [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)キャッシュは考慮しない

Fixstars Group www.fixstars.comCopyright © 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]30*1: 整数演算が重要な場合は整数演算数で考える*2: 動的周波数制御 (Intel Turbo Boost 等) も考慮する必要がある

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

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

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

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

Fixstars Group www.fixstars.comCopyright © Fixstars Group例: ナイーブな行列積の性能見積り (1/2)• 問題: 以下の行列積コードの得られる性能の上限は？• プロセッサの性能• 浮動小数点演算性能: 4000 GFLOPS• メモリ帯域: 200 GB/sec• M = N = K = 1000, キャッシュは考慮しないものとする36float 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;}}

Fixstars Group www.fixstars.comCopyright © Fixstars Group例: ナイーブな行列積の性能見積り (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]37どれだけ命令レベル高速化を頑張っても理論演算性能比 1.25% の性能しか得られない*2→ データ局所性を活用して演算強度を上げる必要がある (実践編へ)知見: ナイーブな行列積はメモリ律速*1*1 プロセッサの演算性能を 4000 GFLOPS, メモリ帯域を 200 GB/s とした場合*2 実際にはキャッシュ等の影響でこれ以上の性能となりうる

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

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

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

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

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

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

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

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

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

Fixstars Group www.fixstars.comCopyright © Fixstars Group演算性能の計測• ベンチマーク内容• 1命令あたり8加算可能な vaddps 命令を大量に実行し1秒あたりの演算回数 [FLOPS] を計測する*1• 1コアのみ使用48__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; ++i) {vz0 = _mm256_add_ps(vz0, vy); vz1 = _mm256_add_ps(vz1, vy);vz2 = _mm256_add_ps(vz2, vy); vz3 = _mm256_add_ps(vz3, vy);vz4 = _mm256_add_ps(vz4, vy); vz5 = _mm256_add_ps(vz5, vy);vz6 = _mm256_add_ps(vz6, vy); vz7 = _mm256_add_ps(vz7, vy);}// Avoid dead-code eliminationauto 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);依存関係のない8個の vaddps 命令を繰り返し実行(パイプラインハザードが発生しないようにする)*1: 理論ピーク演算性能を計測する場合、 FMA 命令 (vfmaddps) を使用する必要があるが、今回は vaddps 命令を用いる。

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

Fixstars Group www.fixstars.comCopyright © Fixstars Groupメモリ帯域の計測• ベンチマーク内容• キャッシュに収まらないサイズの配列に対して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 * 8);auto vx2 = _mm256_load_ps(sp0 + 2 * 8), vx3 = _mm256_load_ps(sp0 + 3 * 8);auto vy0 = _mm256_load_ps(sp1 + 0 * 8), vy1 = _mm256_load_ps(sp1 + 1 * 8);auto vy2 = _mm256_load_ps(sp1 + 2 * 8), vy3 = _mm256_load_ps(sp1 + 3 * 8);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;} 50

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

Fixstars Group www.fixstars.comCopyright © Fixstars Group演算・メモリ複合実行性能の計測• ベンチマーク内容• 配列の各要素に対してロード・N回加算・ストアを行う• read:write:add=1:1:N• 演算強度: N / ((1 + 1) * sizeof(float)) = N/8• 1コアのみ使用52const 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;}

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

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

Fixstars Group www.fixstars.comCopyright © Fixstars Group高速化結果• 最大 6351 GFLOPS (ピーク性能比 73.8%)• 単純な実装から11倍の高速化を達成02000400060008000100001024 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 optimizations57

Movatterモバイル変換

Change Language

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

Embed presentation

Recommended

More Related Content

What's hot

Similar to CPU / GPU高速化セミナー！性能モデルの理論と実践：理論編

More from Fixstars Corporation

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

Editor's Notes