プログラムがうまく動かない！　―CUDA のバグの見つけ方―

プログラム

が

う

ま

く

動

かない

！

―CUDA

の

バグ

の

見

つけ

方―

Agenda

- 1 of 3

• デバッグ

の

ノウハウ

を

紹介

します

。

• 商用ソフトウェア開発

での

実例

をとりあげます。

– Particleworks

の

K20 対応

• (注)

今回

は

触

れません

– NVIDIA Parallel Nsight

– CUDA-GDB

– CUDA-MEMCHECK

– Etc.

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

Agenda

- 3 of 3

• デバッグに役立ったこと

1.

詳細

な

実行ログ

を

出力

できるようにしておこう

•

意外

と

役立

つ

2.

CUDA カーネル

に

対応

した

HOST コード

を

用意

しよう

•

単体テスト

ができるように

3.

HOST コード

に置き換えて

実行

できるようにしておこう

•

結合テスト

ができるように

4.

HOST

と

DEVICE

の

計算結果

を

比較

できるようにしておこう

•

単体テスト

と

結合テスト

の

両方

で

5.

CUDA

の

しくみ

に

詳

しくなろう

Company Information

■会社名

プロメテック・ソフトウェア株式会社

2004年10月29日

■設立年月日

201,610千円

■資本金

■主要株主

株式会社構造計画研究所

三菱 UFJ キャピタル株式会社

大和企業投資株式会社

SMBC ベンチャーキャピタル株式会社

安田企業投資株式会社

りそなキャピタル株式会社

プロメテック・ソフトウェア協力研究者持株会

プロメテック・ソフトウェア従業員持株会

■役員

岡本伸一 藤澤智光 越塚誠一

角家強志 島田憲成 花田孔明

Access

■所在地

〒113-0033 東京都文京区本郷 7-3-1

東京大学アントレプレナープラザ 3 階

Timeline

CUDA Toolkit v5.0

Particleworks v4.5

Tesla K20

2012-10

2013-02

2012-11

2012-12

2013-01

Tesla K20 Early Access Program

作業期間 (約2ヶ月)

NVIDIA Manufacturing Day 2013

Tasks

• Software Testing

– すべて

の

例題

(Particleworks に付録)

– いくつか

の

顧客事例

• Performance Measurements

– いくつか

の

例題

– いくつか

の

顧客事例

• Performance Tunings

• (Software Debugs)

dam-break gearbox

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

Performance Turnings

for Kepler

• Read Only Cache Memory

• Warp Shuffle Operations

• Grid and Block size Optimization

cf. NVIDIA Manufacturing Day 2013, Particleworks

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

1.47x

(C2075

/

K20c)

Status

• 動作確認済

み

– GeForce GTX 640 (K10)

– Early Access Program (K20)

Logs

• 実行ログ

を

チェック

– ど

ういう

状況

で

停止

しているか

確認

できる

• printf debug

• Logger

– 出力

の

詳細度

を

変更

できるようにしておく

• 変数

の

ウォッチ

• コールツリー

1. 詳細

な

実行ログ

を

出力

できるようにしておこう

Bugs

• 特定

の

テストケース

で

解析

が

発散

する

– し

かも Tesla

K20c

で

実行

したときだけ

• クーラン条件

を

満

たせなくなり

解析

が

停止

する

– 安定

した

解析

のための

条件

– 粒子

の

速度

が

大

きくなりすぎること

– ど

こかの

カーネル

の

計算

がおかしい

• 止

まるときと

止

まらないときがある

– 粒子数

が

多

い

解析

(20 万以上)

で

止

まりやすい

– 並列計算

に

関

する

バグ？

CFL condition

– Courant-Friedrichs-Lewy

条件

を

満

たしている

条件

を

満

たしていない

Time Step: t

Time Step: t+1

Unit Testing

(Tesla K20)

[spmv.cpp]

void

hst_spmv(

…

)

{

…

}

[spmv.cu]

__global__

void

spmv_kernel(

…

)

{

…

}

void

dev_spmv(

…

)

{

spmv_kernel<<<

…

>>>(

…

);

}

[spmv.h]

void

hst_spmv(

…

);

void

dev_spmv(

…

);

2. CUDA カーネル

に

対応

した

HOST コード

を

用意

しよう

Integration Testing

(Tesla K20) - 1 of 2

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);

メモリバッファ

を

抽象化しておく

• HOST

と

DEVICE

を

対応付

けて

管理

する

• 変更

を

相互

に

反映

させる

• 取得時

に

変更

されていたら

コピー

する

3. HOST コード

に

置

き

換

えて

実行

できるようにしておこう

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);

動作

• 変数

は

すべて抽象化

した

バッファ

• 取得モード

で

返

す

生ポインタ

の

アドレス

を

変更

• 読

み

込

み

モード

で

取得

されたあと

異

なる

取得モード

でとりだされたら

cudaMemcpy

する

Integration Testing

(Tesla K20) - 2 of 2

… 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) …

Logger

の

出力結果:

3 Bugs

1. Prefix Sum (scan)

2. Sort

3. 粉体

計算部

:

接触判定 + 摩擦力計算

– 原因不明

– 単体テスト

は

パス

する

– HOST

も

DEVICE

も

似

ている

• 同

じような

コード

– 計算結果

が

化

けているようにみえる

• レジスタ

の

値

が

変

Thrust

に

置

き

換

え

NVCC

- NVIDIA CUDA Compiler

• た

まに

コンパイラ

が

落

ちる

– 複雑

な

ヘッダファイル

を

読

ましていると

NVCC has bugs?

こ

れまで

計算

に

問題

はなかった

– Tesla K20 特有

の

問題？

Tesla K20

で

変

わったこと:

– Compute Capability

が

3.5

になった

– 利用

できる

レジスタ数

が

増

えた

• バグ

があるとしたら

こ

こ

？

た

めしに

__launch_bounds__(T, B)

を

調整

してみる

……

– T: ブロック

あたりの

最大スレッド数

– B: マルチプロセッサ

あたりの

最小ブロック数

PTX

- Parallel Thread eXecution

.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__(

384

, 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>; …

変更後: __launch_bounds__(

1024

, 1)

maxntid 以外

はすべて同じ

LLVM

- Low Level Virtual Machine

– LLVM

を

基礎

としている

• 独自拡張

もできる

– CUDA Compiler SDK

• NVVM IR

(libNVVM)

• CUDA driver

があやしい

CUDA C/C++ front-end

NVVM IR

(LLVM IR 互換)

LLVM optimizer

PTX back-end

CUDA C/C++

(*.cu)

PTX

(*.ptx)

CUDA binary

CUDA driver

(JIT compiler)

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

の

バグ

の

見

つけ

方

―

CUDA driver

の

Summary

• デバッグに役立ったこと

1.

詳細

な

実行ログ

を

出力

できるようにしておこう

•

意外

と

役立

つ

2.

CUDA カーネル

に

対応

した

HOST コード

を

用意

しよう

•

単体テスト

ができるように

3.

HOST コード

に置き換えて

実行

できるようにしておこう

•

結合テスト

ができるように

4.

HOST

と

DEVICE

の

計算結果

を

比較

できるようにしておこう

•

単体テスト

と

結合テスト

の

両方

で

5.

CUDA

のしくみに

詳

しくなろう

Event

- Simulation Conference 2013

2013-09-12 (木) 10:00 @ 東京コンファレンスセンター・品川

参加費: 無料 (要事前登録)

株式会社日立製作所 株式会社資生堂 積水エンジニアリング株式会社 株式会社キタック 株式会社トプコン 住友重機械工業株式会社

■主催

プロメテック・ソフトウェア株式会社 日本 GPU コンピューティング有限責任事業組合 NVIDIA Japan 株式会社構造計画研究所 株式会社エルザ ジャパン サイバネットシステム株式会社 ほか

■協賛

■基調講演

青木 素直 (株式会社三菱総合研究所 副理事長) 姫野龍太郎 (独立法人理化学研究所 情報基盤センター長) 越塚 誠一 (東京大学大学院工学系研究科 教授)

■パネルディスカッション

メニ―コア新時代！ ソフトウェア開発の現場から見えてきた課題と期待

■特別講演／事例講演

http://www.prometech.co.jp/

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)