プログラムがうまく動かない! CUDA のバグの見つけ方 北岡伸也 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 1
Agenda - 1 of 3 デバッグのノウハウを紹介します 商用ソフトウェア開発での実例をとりあげます Particleworks の K20 対応 ( 注 ) 今回は触れません NVIDIA Parallel Nsight CUDA-GDB CUDA-MEMCHECK Etc. 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 2
Agenda - 2 of 3 Particleworks - 商用流体解析ソフトウェア プロメテックソフトウェア株式会社で開発している製品 v4.5 から CUDA 5.0 & Tesla K20 に対応 Fundamental Algorithms MPS method - Moving Particle Simulation (Semi-implicit) DEM - Distinct Element Method Solver on GPUs コード行数 : ~150,000 カーネル数 : <150 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 3
Agenda - 3 of 3 デバッグに役立ったこと 1. 詳細な実行ログを出力できるようにしておこう 意外と役立つ 2. CUDA カーネルに対応した HOST コードを用意しよう 単体テストができるように 3. HOST コードに置き換えて実行できるようにしておこう 結合テストができるように 4. HOST と DEVICE の計算結果を比較できるようにしておこう 単体テストと結合テストの両方で 5. CUDA のしくみに詳しくなろう 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 4
Particleworks 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 5
Company Information 会社名 設立年月日 資本金 役員 主要株主 プロメテック ソフトウェア株式会社 2004 年 10 月 29 日 201,610 千円 岡本伸一藤澤智光越塚誠一角家強志島田憲成花田孔明 株式会社構造計画研究所三菱 UFJ キャピタル株式会社大和企業投資株式会社 SMBC ベンチャーキャピタル株式会社安田企業投資株式会社りそなキャピタル株式会社プロメテック ソフトウェア協力研究者持株会プロメテック ソフトウェア従業員持株会 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 6
Access 所在地 113-0033 東京都文京区本郷 7-3-1 東京大学アントレプレナープラザ 3 階 URL http://www.prometech.co.jp/ 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 7
Timeline 2012-10 2012-11 2012-12 CUDA Toolkit v5.0 Tesla K20 Early Access Program Tesla K20 K20c 2013-01 作業期間 ( 約 2 ヶ月 ) NVIDIA Manufacturing Day 2013 2013-02 Particleworks v4.5 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 8
Tasks Software Testing gearbox すべての例題 (Particleworks に付録 ) いくつかの顧客事例 Performance Measurements いくつかの例題 いくつかの顧客事例 dam-break Performance Tunings (Software Debugs) 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 9
Development Environment (Windows) Redmine Project Management Subversion (+ Git) Version Control System Microsoft Visual Studio 2010 (+ CUDA 5.0) IDE; Integrated Development Environment Google Test Testing Framework Jenkins CI; Continuous Integration 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 10
Performance Turnings for Kepler Read Only Cache Memory Warp Shuffle Operations Grid and Block size Optimization 1.47x (C2075 / K20c) Case 1 Case 2 Case 3 Case 4 Case 5 # of Particles 807,885 344,633 366,210 295,113 861,042 Pressure (Implicit) x x x x x Viscosity (Implicit) x x Surface Tension x x Turbulence x DEM x Performance Gain 1.44 1.57 1.43 1.45 1.49 cf. NVIDIA Manufacturing Day 2013, Particleworks 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 11
Status 動作確認済み GeForce GTX 640 (K10) Early Access Program (K20) すぐにテストをパスできるだろう 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 12
あれっ? 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 13
Logs 1. 詳細な実行ログを出力できるようにしておこう 実行ログをチェック どういう状況で停止しているか確認できる printf debug Logger 出力の詳細度を変更できるようにしておく 変数のウォッチ コールツリー 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 14
Bugs 特定のテストケースで解析が発散する しかも Tesla K20c で実行したときだけ クーラン条件を満たせなくなり解析が停止する 安定した解析のための条件 粒子の速度が大きくなりすぎること どこかのカーネルの計算がおかしい 止まるときと止まらないときがある 粒子数が多い解析 (20 万以上 ) で止まりやすい 並列計算に関するバグ? 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 15
CFL condition Courant-Friedrichs-Lewy 条件を満たしている 条件を満たしていない Time Step: t Time Step: t+1 衝突 すり抜ける 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 16
Unit Testing (Tesla K20) 2. CUDA カーネルに対応した HOST コードを用意しよう [spmv.h] void hst_spmv( ); void dev_spmv( ); [spmv.cpp] void hst_spmv( ) { } [spmv.cu] global void spmv_kernel( ) { } void dev_spmv( ) { spmv_kernel<<< >>>( ); } 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 17
Integration Testing (Tesla K20) - 1 of 2 3. HOST コードに置き換えて実行できるようにしておこう hst_spmv( y.get(hst_mode, write_mode), row_ptr.get(hst_mode), col_ind.get(hst_mode), a.get(hst_mode), x.get(hst_mode), n); dev_spmv( y.get(dev_mode, write_mode), row_ptr.get(dev_mode), col_ind.get(dev_mode), a.get(dev_mode), x.get(dev_mode), n); メモリバッファを抽象化しておく HOST と DEVICE を対応付けて管理する 変更を相互に反映させる 取得時に変更されていたらコピーする 動作 変数はすべて抽象化したバッファ 取得モードで返す生ポインタのアドレスを変更 読み込みモードで取得されたあと異なる取得モードでとりだされたら cudamemcpy する 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 18
Integration Testing (Tesla K20) - 2 of 2 4. HOST と DEVICE の計算結果を比較できるようにしておこう Logger の出力結果 : Prometech::NeighborSearchGPU::Calculate { Prometech::NeighborSearchGPU::calculate_distribution { pw::arraymanagermethod::exchange_distributed_buffer_all { } debug: exchange_distributed_buffer_all : end. @ pw::arraymanager::exec(2497) } Prometech::NeighborSearchGPU::calculate_particle { pw::arraymanagermethod::construct_neighbor_table_large { debug: buffer = particle.collide_hash.int2.1.2, size = 203334 @ pw::arraymanagermethod::set_array(526) debug: buffer = particle.collide_mibb_buf.double3.1.1, size = 64 @ pw::arraymanagermethod::set_array(526) debug: buffer = particle.collide_mabb_buf.double3.1.1, size = 64 @ pw::arraymanagermethod::set_array(526) debug: bbmin = -1.10553-0.685474-0.637836 @ pw::arraymanagermethod::set_collision_slice(304) debug: bbmax = 1.10555 1.10542 0.637635 @ pw::arraymanagermethod::set_collision_slice(305) debug: ngrid = 45 36 26 @ pw::arraymanagermethod::set_collision_slice_array(313) debug: buffer = particle.collide_slice_sum.int.1.1, size = 27 @ pw::arraymanagermethod::set_array(526) debug: buffer = particle.collide_slice.int4.1.1, size = 26 @ pw::arraymanagermethod::set_array(526) debug: buffer = particle.collide_slice_offset.int4.1.1, size = 26 @ pw::arraymanagermethod::set_array(526) 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 19
3 Bugs 1. Prefix Sum (scan) 2. Sort Thrust に置き換え 3. 粉体計算部 : 接触判定 + 摩擦力計算 原因不明 単体テストはパスする HOST も DEVICE も似ている 同じようなコード 計算結果が化けているようにみえる レジスタの値が変 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 20
ところで 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 21
NVCC - NVIDIA CUDA Compiler たまにコンパイラが落ちる 複雑なヘッダファイルを読ましていると字句解析でアサーションがでる 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 22
ん? 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 23
NVCC has bugs? これまで計算に問題はなかった Tesla K20 特有の問題? Tesla K20 で変わったこと : Compute Capability が 3.5 になった 利用できるレジスタ数が増えた バグがあるとしたらここ? ためしに launch_bounds (T, B) を調整してみる T: ブロックあたりの最大スレッド数 B: マルチプロセッサあたりの最小ブロック数 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 24
動いた! 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 25
PTX - Parallel Thread execution 変更前 : launch_bounds (384, 1).visible.entry _Z28dem_collision_pp_calc_... ).maxntid 384, 1, 1.minnctapersm 1 {.reg.pred %p<11>;.reg.s32 %r<59>;.reg.s64 %rd<59>;.reg.f64 %fd<222>; 変更後 : launch_bounds (1024, 1).visible.entry _Z28dem_collision_pp_calc_... ).maxntid 1024, 1, 1.minnctapersm 1 {.reg.pred %p<11>;.reg.s32 %r<59>;.reg.s64 %rd<59>;.reg.f64 %fd<222>; maxntid 以外はすべて同じ 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 26
LLVM - Low Level Virtual Machine 5. CUDA のしくみに詳しくなろう LLVM を基礎としている CUDA C/C++ (*.cu) CUDA C/C++ front-end NVVM IR (LLVM IR 互換 ) 独自拡張もできる CUDA Compiler SDK NVVM IR (libnvvm) LLVM optimizer PTX back-end PTX (*.ptx) CUDA driver があやしい CUDA driver (JIT compiler) CUDA binary 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 27
Bug Report 動作確認ができたのでソフトウェアは Fix (1 月末 ) CUDA Registered Developer Program https://developer.nvidia.com/rdp/cuda-registered-developer-program CUDA/GPU Bug Reporting https://developer.nvidia.com/rdp/bugs/cudagpu-bug-reporting Submissions https://developer.nvidia.com/node/233301/submissions NVIDIA 側でも不具合の再現 (2 月末 ) CUDA 5.5 RC の driver で修正 (5 月初 ) 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 28
プログラムがうまく動かない! CUDA のバグの見つけ方 北岡伸也 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 29
CUDA driver の バグを見つけました 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 30
Summary デバッグに役立ったこと 1. 詳細な実行ログを出力できるようにしておこう 意外と役立つ 2. CUDA カーネルに対応した HOST コードを用意しよう 単体テストができるように 3. HOST コードに置き換えて実行できるようにしておこう 結合テストができるように 4. HOST と DEVICE の計算結果を比較できるようにしておこう 単体テストと結合テストの両方で 5. CUDA のしくみに詳しくなろう 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 31
Event - Simulation Conference 2013 2013-09-12 ( 木 ) 10:00 @ 東京コンファレンスセンター 品川 参加費 : 無料 ( 要事前登録 ) 主催 プロメテック ソフトウェア株式会社日本 GPU コンピューティング有限責任事業組合 基調講演 http://www.prometech.co.jp/ 青木素直 ( 株式会社三菱総合研究所副理事長 ) 姫野龍太郎 ( 独立法人理化学研究所情報基盤センター長 ) 越塚誠一 ( 東京大学大学院工学系研究科教授 ) パネルディスカッション メニ コア新時代! ソフトウェア開発の現場から見えてきた課題と期待 協賛 NVIDIA Japan 株式会社構造計画研究所株式会社エルザジャパンサイバネットシステム株式会社 株式会社日立製作所 株式会社資生堂 積水エンジニアリング株式会社 株式会社キタック 株式会社トプコン 住友重機械工業株式会社大日本スクリーン製造株式会社 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 32 ほか 特別講演 / 事例講演
Job Offer Product Development Dept. Researches Mathematics Linear Algebra Mathematical Analysis Differential equations (Function Approximation) Physics Incompressible Fluid (Non-Newtonian Fluid) (Turbulence) (Surface tension) (Heat conduction/transfer) Powder / Rigid Body Numeric analysis MPS / SPH / DEM (LBM / FDM / FEM / BEM) Software Developments OSs Windows Linux / (Mac) Languages C++ (STL, Boost, 11/14) CUDA Java (Python) Techniques Algorithms & Data Structures OOP / (TMP) / Design Patterns SIMD (SPMD) / OpenMP / MPI (Concurrency Programming) HCI (UI / UX) 2013-07-30 GTC Japan 2013 @ Tokyo Midtown Hall & Conference 33