GPGPUによる高速画像処理

Size: px
Start display at page:

Download "GPGPUによる高速画像処理"

Transcription

1 GPGPU による高速画像処理 ~ リアルタイム画像処理への挑戦 ~ 名古屋大学大学院情報科学研究科 出口大輔

2 リアルタイム画像処理 2

3 3 発表の流れ GPGPU を始める前に GPGPU の基礎知識 CUDA の使い方 CUDA を使う前に プログラミングの予備知識 CUDA を使って Hello World GPGPU にチャレンジ 行列積の計算 テンプレートマッチング ガウシアンフィルタ SIFT 特徴量の計算

4 ~GPGPU って何?~ 4

5 5 GPGPU って? GPGPU は何の略? General-Purpose computation on GPUs GPUを汎用計算に利用しようという試み 現在は GPUコンピューティング とも呼ばれる どうして GPGPU が注目されているのか?

6 6 CPU と GPU の性能比較 1600 [GFLOP/s] NVIDIA GPU 1400 Geforece GTX480 ピーク性能 Geforece 8800GTX Geforece GTX280 Core2 Duo 3.0GHz Quad Core Xeon 3.2GHz Intel CPU NVIDIA CUDA Programming Guide 4.0より引用

7 7 GPGPU って? GPGPU は何の略? General-Purpose computation on GPUs GPUを汎用計算に利用しようという試み 現在は GPUコンピューティング とも呼ばれる どうして GPGPU が注目されているのか? GPUの計算性能が飛躍的に向上 最新のGPUは 1.5 TFLOPS 以上の演算性能 (CPUの約 10 倍 ) GPUはCPUと比較して並列計算に優れている GeForce GTX580 では 512 コアによる並列計算が可能 手頃な価格で入手可能 GeForce GTX580 は約 5 万円で購入可能

8 8 GPGPU の活用例 動画像処理 CyberLink PowerDirector 7 ビデオエフェクトのレンダリングを高速化 TMPGEnc 4.0 XPress フィルター処理 デコード処理の高速化 画像処理 Adobe Photoshop CS4 / CS5,Adobe Premire CS5 各種フィルタ処理の高速化 OpenCV 各種画像処理の高速化 数値計算 MATLAB FFT の高速化

9 9 GPU の歴史 : 1999 年以前 1970 年 ~ 1990 年 初期開発の時代 ソフトウェアによるグラフィックス処理 プログラム可能な GPU に関する初期研究 Ikonas System [1], Pixel Machine [2] [1] J. N. England, A system for interactive modeling of physical curved surface objects, Proceedings of SIGGRAPH 78, pp [2] M. Potmesil and E. M. Hoffert, The Pixel Machine: A Parallel Image Computer, Proceedings of SIGGRAPH89, pp.69-78, 年 ~ 1999 年 GPU 技術の黎明期 3D グラフィックス アクセラレータの開発 グラフィックス向けライブラリの開発 OpenGL(1992 年 ~), DirectX(1995 年 ~)

10 10 GPU の歴史 : 1999 年 ~ GPU の誕生 NVIDIA 社の GeForce 256 の登場 ハードウェア T&L をサポート CPU 負荷を大幅に削減 グラフィックスパイプライン ハードウェア固定の処理 自由表現な表現は不可 頂点処理 ( ハードウェア固定 ) ジオメトリ処理 ( クリッピング等 ) ラスタライズ処理 ピクセル処理 ( ハードウェア固定 ) Transform & Lighting 画面出力

11 11 GPU の歴史 : 2003 年 ~ プログラマブルシェーダの登場 Vertex Shader 頂点座標の変換処理 Pixel Shader 画素の輝度計算処理 シェーダ言語の進化 アセンブラから高級言語へ Cg, HLSL, GLSL 柔軟な映像表現が可能に グラフィックス以外への応用 GPGPU への関心が高まる Vertex Shader ( プログラマブル ) ジオメトリ処理 ( クリッピング等 ) ラスタライズ処理 Pixel Shader ( プログラマブル ) 画面出力

12 12 GPU の歴史 : 2007 年 ~ GPGPU の開発環境の整備 CUDA (NVIDIA) ATI Stream (AMD) OpenCL DirectCompute GPGPU 時代の到来 物理シミュレーション 数値計算 信号解析 画像処理 認識 統合型シェーダ Vertex Shader ( プログラマブル ) Geometry Shader ( プログラマブル ) ラスタライズ処理 Pixel Shader ( プログラマブル ) 画面出力

13 13 GPU の歴史 : 2009 年 ~ Fermi アーキテクチャの登場 ビデオカードによって サポート状況は異なる 汎用計算向けのアーキテクチャ L1/L2キャッシュの搭載 複数カーネルの同時実行のサポート 倍精度浮動小数点演算の高速化 ECCメモリのサポート アトミックなメモリ操作の高速化 C++ のフルサポート GPGPU 関連のライブラリ & ツールの充実 CUBLAS, CUFFT, Thrust, NPP Parallel Nsight, Visual Profiler

14 ~CUDA を使う前に ~ 14

15 15 CUDA って何? CUDA(Compute Unified Device Architecture) 発音は クーダ もしくは キューダ NVIDIA 社が提供する GPU を利用するための統合開発環境 GeForce 8 以降のハードウェアで利用可能 グラフィックス処理 API の知識は不要 CPU でプログラムを実行する感覚 C/C++ を用いてプログラムの開発が可能 既存のアルゴリズムの移植も比較的容易

16 16 CUDA が利用可能な環境 CUDA 対応のグラフィックスカード GeForce GTX 580, GTX 480, GTX 260, 8800シリーズ, 他 Quadro Plex 2200 D2, FX 5800, FX 5600, 5000, 6000, 他 Tesla S2050, C2070, S1070, C1060, S870, D870, C870 TeslaはHPCに特化した製品であり映像出力を持たない OS の対応状況 (CUDA 4.0) Windows XP(32bit, 64bit) Windows Vista(32bit, 64bit) Windows Server 2008(32bit, 64bit) Linux(32bit, 64bit) Mac OS

17 17 CUDA を使うための準備 (Windows 編 ) CUDA ZONE から次をダウンロード NVIDIA Driver CUDA 対応のビデオドライバー CUDA Toolkit 4.0 コンパイラ (nvcc) CUBLASやCUFFT ライブラリ ドキュメント CUDA SDK 4.0 サンプル 2011 年 5 月 25 日にリリース

18 18 CUDA を動かしてみよう!! Volumetric Particle Shadows Image Denoising CUDA SDK 内のサンプルの実行結果

19 ~ プログラミングの予備知識 ~ 19

20 20 CUDA 対応の GPU(G80, GT200) ストリーミング マルチプロセッサ (SM) 8 個のスカラープロセッサ (SP) 16KB の共有メモリ スレッド間の同期機構 マルチプロセッサ内でのみ可能 マルチプロセッサ間での同期には CPU 処理が必要 GeForce GTX280 の場合 30 基のマルチプロセッサを搭載 1GByte 以上のグローバルメモリ マルチプロセッサ #1 SP SP SP SP SP SP SP SP マルチプロセッサ #2 SP SP SP SP SP SP SP SP

21 21 CUDA 対応の GPU(Fermi) ストリーミング マルチプロセッサ (SM) 32 個のスカラープロセッサ (SP) 48KB の共有メモリ スレッド間の同期機構 マルチプロセッサ内でのみ可能 マルチプロセッサ間での同期には CPU 処理が必要 GeForce GTX480 の場合 15 基のマルチプロセッサを搭載 1GByte 以上のグローバルメモリ SM #1 SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SM #2 SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP

22 22 CUDA のプログラミングモデル GPU の特徴 多数のスレッドが高い並列性をもって処理を実行 GPU のみでプログラムを実行できない CPUとの連携が不可欠 CUDA における計算の流れ GPUを並列演算可能なデバイスとして扱う 複数のスレッドを同時に実行できる外部 CPU 階層的にスレッドを管理 スレッドのまとまり = ブロック ブロックのまとまり = グリッド 問題を分割して計算する際に便利

23 23 CUDA におけるスレッド管理 スレッドを 3 次元的に配置 各スレッドの ID を X, Y, Z の3 要素で表現 スレッドのまとまりをブロックとして管理 グリッド ブロックブロック (0,0,0) (0,0,1) ブロックブロック (1,0,0) (1,0,1) スレッドスレッド (0,0,0) (0,0,1) スレッドスレッド (1,0,0) (0,1,0) (0,1,1) スレッドスレッド (2,0,0) (0,2,0) (0,2,1) スレッドスレッド (0,0,0) (0,0,1) スレッドスレッド (1,0,0) (0,1,0) (0,1,1) スレッドスレッド (2,0,0) (0,2,0) (0,2,1) スレッドスレッド (0,1,0) (1,0,0) (1,0,1) スレッドスレッド (1,1,0) (1,1,1) スレッドスレッド (2,1,0) (1,2,0) (1,2,1) スレッドスレッド (0,1,0) (1,0,0) (1,0,1) スレッドスレッド (1,1,0) (1,1,1) スレッドスレッド (2,1,0) (1,2,0) (1,2,1) ブロックブロック (0,1,0) (0,1,1) ブロックブロック (1,1,0) (1,1,1)

24 24 階層的スレッドの利用方法 一般的な画像処理で利用する場合 ブロック内のスレッド数を決定 = 256 スレッド 画像内にブロックを配置 ブロック内のスレッドが各画素を処理 計算範囲の求め方 ブロック ID スレッド ID ブロック内スレッド数 T 00 T 10 T 01 T 0M T N0 T NM ブロック #1 T 00 T 10 T N0 T 01 T 0M T NM

25 25 CUDA における計算の流れ グリッド 1 処理 1 ブロック (0,0,0) スレッド (0,0,0) スレッド (1,0,0) スレッド (2,0,0) ブロック (1,0,0) スレッド (0,0,0) スレッド (1,0,0) スレッド (2,0,0) スレッド (0,1,0) スレッド (1,1,0) スレッド (2,1,0) スレッド (0,1,0) スレッド (1,1,0) スレッド (2,1,0) ブロック (0,1,0) ブロック (1,1,0) 処理 2 グリッド 2 ブロック (0,0,0) ブロック (1,0,0)

26 26 CUDA のメモリモデル ブロック (0,0,0) 共有メモリ ブロック (1,0,0) 共有メモリ ローカルメモリ ローカルメモリ ローカルメモリ ローカルメモリ レジスタ レジスタ レジスタ レジスタ スレッド (0,0,0) スレッド (1,0,0) スレッド (0,0,0) スレッド (1,0,0) コンスタントメモリ テクスチャメモリ グローバルメモリ

27 27 CUDA で利用可能なメモリ グローバルメモリ 大量のメモリを利用可能 低速なメモリアクセス (400~600 クロック必要 ) Fermi アーキテクチャではキャッシュ機構 (L1/L2) を搭載 共有メモリ レジスタと同じ速度でアクセス可能 マルチプロセッサ 1 基あたり 16KB(Fermi は最大 48KB) Fermi の場合は一部を L1 キャッシュとして利用可能 テクスチャメモリ キャッシュ機構による高速なアクセスが可能 ハードウェア線形補間や正規化テクスチャ座標が利用可能 コンスタントメモリ キャッシュ機構による高速なアクセスが可能 マルチプロセッサ 1 基あたり 64KB

28 28 各 GPU で使用できる機能の違い 使用できる機能を Compute Capability で区別 現在 1.0 ~ 2.1 の GPU が存在 Compute Capability による機能の違い 1.0: 初期リリース GeForce 8800GTX, Quadro FX 5600, 他 1.1: アトミックなメモリ操作のサポート GeForce 9800 GTX, Quadro FX 3700, 他 1.3: 倍精度浮動小数点演算のサポート GeForce GTX 280, Quadro FX 5800, 他 2.x : グローバルメモリのキャッシュをサポート GeForce GTX480,Quadro 5000, 他

29 29 Compute Capability による機能の違い Compute Capability x 3 次元グリッド 倍精度浮動小数 32 ビット整数のアトミック演算 64 ビット整数のアトミック演算 単精度浮動小数のアトミック演算

30 30 CUDA における制限事項 Compute Capability x 1 ブロックあたりのスレッド数 1 マルチプロセッサあたりのスレッド数 1 マルチプロセッサあたりのレジスタ数 コンスタントメモリ KB 共有メモリ 16KB 48KB 2D テクスチャ

31 ~ CUDA を使って Hello World ~ 31

32 32 CUDA を使って Hello World!! Hello World!! を表示するサンプル (main.cu) global void hello( char *data ) { char *text = "Hello World!! n"; data[ threadidx.x ] = text[ threadidx.x ]; } int main( int argc, char *argv[] ) { char *ddata, hdata[ 15 ]; cudamalloc( ( void ** )&ddata, sizeof( char ) * 15 ); dim3 nthreads( 15, 1 ); dim3 nblocks( 1, 1 ); hello<<< nblocks, nthreads >>>( ddata ); cudamemcpy( hdata, ddata, sizeof( char ) * 15, cudamemcpydevicetohost ); printf( "%s", hdata ); cudafree( ddata ); } return( 0 );

33 33 CUDA を使って Hello World!! Hello World!! を表示するサンプル (main.cu) global void hello( char *data ) { char *text = "Hello World!! n"; data[ threadidx.x ] = text[ threadidx.x ]; } int main( int argc, char *argv[] ) { char *ddata, hdata[ 15 ]; global, threadidx CUDA で拡張された部分 cudamalloc( ( void ** )&ddata, sizeof( char ) * 15 ); dim3 nthreads( 15, 1 ); dim3 nblocks( 1, 1 ); hello<<< nblocks, nthreads >>>( ddata ); dim3, <<< >>> cudamemcpy( hdata, ddata, sizeof( char ) * 15, cudamemcpydevicetohost ); printf( "%s", hdata ); cudafree( ddata ); } return( 0 );

34 34 CUDA を使って Hello World!! Hello World!! を表示するサンプル (main.cu) global void hello( char *data ) { char *text = "Hello World!! n"; data[ threadidx.x ] = text[ threadidx.x ]; } int main( int argc, char *argv[] ) { char *ddata, hdata[ 15 ]; cudamalloc( ( void ** )&ddata, sizeof( char ) * 15 ); dim3 nthreads( 15, 1 ); dim3 nblocks( 1, 1 ); hello<<< nblocks, nthreads >>>( ddata ); cudamemcpy( hdata, ddata, sizeof( char ) * 15, cudamemcpydevicetohost ); printf( "%s", hdata ); cudafree( ddata ); } return( 0 );

35 35 CUDA の言語拡張 (1) 言語拡張により追加された修飾子 関数に対する修飾子 global device host CPU から呼び出され,GPU で実行される関数 GPU から呼び出され,GPU で実行される関数 CPU から呼び出され,CPU で実行される関数 変数に対する修飾子 constant shared device GPU 上のコンスタントメモリに存在する変数 GPU 上の共有メモリに存在する変数 GPU 上のメモリに存在する変数

36 36 CUDA を使って Hello World!! Hello World!! を表示するサンプル (main.cu) global void hello( char *data ) { char *text = "Hello World!! n"; data[ threadidx.x ] = text[ threadidx.x ]; } int main( int argc, char *argv[] ) { char *ddata, hdata[ 15 ]; cudamalloc( ( void ** )&ddata, sizeof( char ) * 15 ); dim3 nthreads( 15, 1 ); dim3 nblocks( 1, 1 ); hello<<< nblocks, nthreads >>>( ddata ); cudamemcpy( hdata, ddata, sizeof( char ) * 15, cudamemcpydevicetohost ); printf( "%s", hdata ); cudafree( ddata ); } return( 0 );

37 37 CUDA の言語拡張 (2) CUDA で利用可能な組み込み型 dim3 uchar2, int2, float2, uchar3, int3, float3, uchar4, int4, float4, 整数 x, y, z からなる 3 次元ベクトル ( スレッド数やブロック数の指定に利用 ) x, y からなる 2 次元ベクトル x, y, z からなる 3 次元ベクトル x, y, z, w からなる 4 次元ベクトル GPU 内のスレッドを識別する組み込み変数 griddim blockidx blockdim threadidx グリッドの次数スレッドが属するブロックのインデックススレッドが属するブロックの次数ブロック内のスレッドのインデックス

38 38 スレッドを識別する組み込み変数 グリッド blockidx ブロックブロック (0,0,0) (0,0,1) スレッドスレッドスレッドスレッドスレッド (0,0,1) スレッドスレッド (0,1,1) (0,2,1) スレッドスレッド (0,0,0) (0,0,0) (1,0,0) (0,1,0) (2,0,0) (0,2,0) スレッドスレッドスレッドスレッド (1,0,1) (1,1,1) スレッド (1,2,1) スレッドスレッドスレッドスレッド (0,1,0) (1,0,0) (1,1,0) (1,1,0) (2,1,0) (1,2,0) ブロックブロック (1,0,0) (1,0,1) blockdim = スレッドスレッドスレッド (0,0,1) (0,0,0) (0,0,0) スレッドスレッド (1,0,1) スレッド (0,1,0) (1,0,0) スレッドスレッドスレッド (0,1,1) (1,0,0) (0,1,0) スレッド (1,1,1) スレッドスレッド (1,1,0) (1,1,0) スレッド (0,2,1) スレッドスレッド (2,0,0) (0,2,0) スレッド (1,2,1) スレッドスレッド (2,1,0) (1,2,0) threadidx ブロックブロック (0,1,0) (0,1,1) ブロックブロック (1,1,0) (1,1,1) griddim = 2 2 2

39 39 C/C++ プログラミングとの相違点 CPU と GPU で実行するコードを明確に区別 CPUで実行する場合は host ( 省略可 ) を付与 GPUで実行する場合は global を付与 CPU と GPU で処理可能なメモリ空間の違い CPU と GPU 間でメモリ転送が必要 スレッド数を指定して GPU 上の関数を呼び出し 関数名 <<< ブロック数, スレッド数 >>>( ); ブロック数とスレッド数は実行時に指定可能 問題サイズに合わせて調整可能

40 40 CUDA を使って Hello World!! Hello World!! を表示するサンプル (main.cu) global void hello( char *data ) { char *text = "Hello World!! n"; data[ threadidx.x ] = text[ threadidx.x ]; } int main( int argc, char *argv[] ) { char *ddata, hdata[ 15 ]; cudamalloc( ( void ** )&ddata, sizeof( char ) * 15 ); dim3 nthreads( 15, 1 ); dim3 nblocks( 1, 1 ); hello<<< nblocks, nthreads >>>( ddata ); GPU 上で実行される関数 GPU 上にメモリを確保 cudamemcpy( hdata, ddata, sizeof( char ) * 15, cudamemcpydevicetohost ); ブロック内に配置するスレッド数 15 1 = 15 総ブロック数 1 1 = 1 printf( "%s", hdata ); cudafree( ddata ); CPU から GPU へメモリ転送 } return( 0 ); GPU 上のメモリを解放

41 41 スレッドのレジスタ数を調査 Visual Studio 2008 コマンドプロンプトを起動 コマンドプロンプト上で main.cu をコンパイル C: Your Path>nvcc main.cu --ptxas-options=-v --compile main.cu ptxas info : Compiling entry function '_Z5helloPc' for 'sm_10' ptxas info : Used 2 registers, 8+16 bytes smem, 15 bytes cmem[0] レジスタ数 合計レジスタ数をチェック 2 15 = 30 レジスタ数生成スレッド数合計

42 ~Thrust ライブラリの使い方 ~ 42

43 43 Thrust ライブラリって? CUDA で利用できるテンプレートライブラリ CUDAとOpenMPに対応したC++ STLの並列版 URL: Thrust ライブラリの特徴 コンテナ CPUとGPUのデータをコンテナとして管理 アルゴリズム コンテナに対してアルゴリズムを適用可能 直感的なインターフェース CUDAの複雑なAPIに関する知識は不要 メモリの確保 / 解放 / 転送がとても簡単

44 44 Thrust ライブラリの使い方 GPU 上で 1,2,3,,100 の総和を計算 #include <thrust/host_vector.h> #include <thrust/device_vector.h> #include <thrust/sequence.h> #include <thrust/reduce.h> int main( int argc, char *argv[] ) { // CPU 上のメモリを確保 thrust::host_vector< int > hvec( 100 ); // データを初期化 1, 2, 3,... thrust::sequence( hvec.begin( ), hvec.end( ), 1 ); // GPU 上のメモリを確保 thrust::device_vector< int > dvec( 100 ); // CPU -> GPU のメモリ転送 dvec = hvec; int val = thrust::reduce( dvec.begin( ), dvec.end( ), 0, thrust::plus< int >( ) ); printf( "%d n", val ); } return( 0 );

45 45 コンテナの使い方 (1) メモリの確保 thrust::host_vector< int > hvec( 20 ); CPU 上のメモリに要素数 20の int 配列を確保 thrust::device_vector< float > dvec( 100 ); GPU 上のメモリに要素数 100の float 配列を確保 メモリの解放 コンテナオブジェクトの消滅時に自動解放 CPUの場合 :free( ) GPUの場合 :cudafree( ) 明示的なメモリ解放 :dvec.clear( )

46 46 コンテナの使い方 (2) メモリの転送 代入操作でメモリ転送が可能 thrust::host_vector< float > hvec( 20 ); thrust::device_vector< float > dvec( 20 ); dvec = hvec; コンテナの各要素へのアクセス 配列と同様にアクセス可能 hvec[ 5 ] = 100.0f; dvec[ 1 ] = 23.4f; 内部で cudamemcpy が呼ばれるので注意

47 47 イテレータ (1) コンテナの要素を指すポインタ STLのイテレータと同じように使用可能 thrust::device_vector< int > dvec( 4 ); thrust::device_vector< int >::iterator ite = dvec.begin( ); dvec.begin( ) dvec.end( ) dvec dvec.begin( ) + 3

48 48 イテレータ (2) コンテナの各要素へのアクセス thrust::device_vector< int > dvec( 4 ); thrust::device_vector< int >::iterator ite = dvec.begin( ); *ite = 10; dvec[ 0 ] = 10 と等価 ++ite; *ite = 25; dvec[ 1 ] = 25 と等価

49 49 アルゴリズムと使用方法 CPU と GPU のデータに対して同じ様に適用可能 GPU の場合は GPU を使用して並列処理される アルゴリズムの使用例 コンテナのデータすべてに1を代入する thrust::fill( dvec.begin( ), dvec.end( ), 1 ); コンテナのデータに 1, 2, 3, の値で初期化する thrust::sequence( dvec.begin( ), dvec.end( ) ); N

50 50 Parallel Reduction(1) コンテナ内の要素を 1 つの値に集約する処理 GPUのスレッドが並列に処理を実行 例 : 総和計算の場合 GPU 上のデータ CUDA スレッドの同期 CUDA スレッドの同期

51 51 Parallel Reduction(2) Thrust ライブラリを用いた集約処理 総和の計算 int val = thrust::reduce( dvec.begin( ), dvec.end( ) ); 最大値の計算 int val = thrust::reduce( dvec.begin( ), dvec.end( ), , thrust::maximum< int >( ) );

52 52 CUDA と Thrust を組み合わせる Hello World!! を表示するサンプル #include <thrust/host_vector.h> #include <thrust/device_vector.h> global void hello( char *data ) { char *text = "Hello World!! n"; data[ threadidx.x ] = text[ threadidx.x ]; } int main( int argc, char *argv[] ) { thrust::host_vector< char > hvec( 15 ); thrust::device_vector< char > dvec( 15 ); dim3 nthreads( 15, 1 ); dim3 nblocks( 1, 1 ); hello<<< nblocks, nthreads >>>( thrust::raw_pointer_cast( &dvec[ 0 ] ) ); hvec = dvec; printf( "%s", &hvec[ 0 ] ); return( 0 ); CPUからGPUへメモリ転送 } GPU 上のメモリを指すポインタを取得

53 53

54 54 GPGPU にチャレンジ 行列積の計算 CUDA における基本的な実装方法 共有メモリ テンプレートマッチング 2 次元テクスチャメモリ ハードウェア線形補間 正規化テクスチャ座標 Parallel Reduction アルゴリズム ガウシアンフィルタ 1 次元テクスチャメモリ コンスタントメモリ SIFT 特徴量の計算 1 2 次元テクスチャメモリ コンスタントメモリ

55 ~ 行列積の計算 ~ 55

56 56 GPU による行列積の計算 行列積計算の特徴 各要素はすべて独立に計算可能 GPUによる並列計算が可能 同じメモリ領域へ頻繁にアクセス メモリアクセス速度がボトルネック GPU 上での実装方法 1. 行列積の計算式に忠実な実装 (GPU#1) 2. 共有メモリを利用した高速化 (GPU#2) ブロック単位 ( 部分行列 ) で行列積を計算共有メモリをキャッシュとして利用

57 57 行列積 C=A B の流れ 1. 行列 A, B, CのメモリをGPU 上に確保 2. 行列 A, BのデータをGPUに転送 3. 行列 Cの各要素 c mn を計算 N 次正方行列を仮定 ただし,N は16の倍数 c mn N k 1 mk b 4. 計算結果を CPU に転送 a kn m 行 a mk N 行列 A N n 列 行列 C b kn 行列 B c mn

58 58 処理の流れ CPU 側の処理 行列の初期化 GPU 側の処理 行列 C の各要素を計算 CPU GPU CPU GPU GPU GPU CPU 行列の初期化 メモリ確保 データ転送 行列積の計算 データ転送

59 59 行列積計算の CPU 処理 ( ソース ) int main( int argc, char *argv[] ) { int N = 512; // 行列 A, B, Cのサイズ float *ha, *hb, *hc; // CPU(host) 側で利用するメモリへのポインタ float *da, *db, *dc; // GPU(device) 側で利用するメモリへのポインタ /* CPU 側のメモリを確保 */ /* GPU 側のメモリを確保 */ /* CPU 側のメモリをGPU 側へ転送 */ /* 実行するGPUのスレッド数, ブロック数を設定 */ /* GPUのカーネルを実行し,C=A B の結果を dc に格納 */ /* GPUの計算結果をCPU 側へ転送 */ /* CPUとGPUそれぞれのメモリを解放 */ } return( 0 );

60 60 行列積計算の CPU 処理 (1) 行列積の計算で使用する変数 int N = 512; float *ha, *hb, *hc; float *da, *db, *dc; // 行列 A, B, Cのサイズ // CPU(host) 側で利用するメモリへのポインタ // GPU(device) 側で利用するメモリへのポインタ CPU と GPU でメモリを確保 CPU:malloc, new, cudamallochost でメモリを確保 GPU:cudaMalloc によりグローバルメモリを確保 /* CPU 側のメモリを確保 */ ha = ( float * )malloc( N * N * sizeof( float ) ); hb = ( float * )malloc( N * N * sizeof( float ) ); hc = ( float * )malloc( N * N * sizeof( float ) ); /* GPU 側のメモリを確保 */ cudamalloc( ( void ** )&da, N * N * sizeof( float ) ); cudamalloc( ( void ** )&db, N * N * sizeof( float ) ); cudamalloc( ( void ** )&dc, N * N * sizeof( float ) );

61 61 行列積計算の CPU 処理 (2) CPU と GPU 間のメモリ転送 CPU GPU 転送元 転送方向 cudamemcpy( dst, src, size, cudamemcpyhosttodevice ); 転送先 転送バイト数 GPU CPU cudamemcpy( dst, src, size, cudamemcpydevicetohost ); この部分を変更 CPU と GPU で確保したメモリの解放 CPU:free, delete, cudafreehost でメモリを解放 GPU:cudaFree でグローバルメモリを解放

62 62 行列積計算の CPU 処理 (3) GPU 上で処理を実行 計算範囲の設定 N 次正方行列を仮定 ただし N は 16 の倍数 行列 Cをサイズ16 16のブロックに分割 ブロック内の各要素を各スレッドが計算 合計 M M 個のブロックを配置 各スレッドが c mn を計算 n = threadidx.x + blockidx.x blockdim.x; m = threadidx.y + blockidx.y blockdim.y; GPU 上で関数を実行 multiply<<< ブロック数, スレッド数 >>>( ); (threadidx.y, threadidx.x) 行列 C (blockidx.y, blockidx.x) 行方向のブロック数 (=16)

63 63 行列積計算の CPU 処理 ( まとめ ) int main( int argc, char *argv[] ) { int N = 512; // 行列 A, B, Cのサイズ float *ha, *hb, *hc; // CPU(host) 側で利用するメモリへのポインタ float *da, *db, *dc; // GPU(device) 側で利用するメモリへのポインタ /* CPU 側のメモリを確保 */ /* GPU 側のメモリを確保 */ /* CPU 側のメモリを GPU 側へ転送 */ GPU 上で処理を実行 /* 実行する GPU のスレッド数, ブロック数を設定 */ dim3 nthreads( 16, 16 ); dim3 nblocks( N / nthreads.x, N / nthreads.y ); /* GPU のカーネルを実行し,C=A B の結果を dc に格納 */ multiply<<< nblocks, nthreads >>>( da, db, dc, N ); /* GPU の計算結果を CPU 側へ転送 */ /* CPU と GPU それぞれのメモリを解放 */ GPU 上のメモリを指定 } return( 0 );

64 64 実装方法 (GPU#1) n 列 行列積 C=A B の各要素 c mn を計算 N b kn c mn N k 1 a mk b kn m 行 a mk N 行列 B c mn 要素 c mn を GPU 上で並列計算 行列 A 行列 C global void multiply1( float *da, float *db, float *dc, int N ) { int n = threadidx.x + blockidx.x * blockdim.x; int m = threadidx.y + blockidx.y * blockdim.y; float sum = 0.0f; for( int k = 0 ; k < N ; k++ ) { sum += da[ m + k * N ] * db[ k + n * N ]; } } dc[ m + n * N ] = sum;

65 65 実装方法 (GPU#2) 各要素 c mn を部分行列 ( ブロック ) の積で計算 ブロックサイズが の場合 c mn N 16 t 1 16( t 1) k 16t ブロック内で同じメモリを参照 mk 高速な共有メモリを利用 a グローバルメモリへのアクセスを削減 2 N 16 2 回 16 分の1 2 N 16 回 b kn N a mk N 行列 A b kn 16 c mn 16 行列 B 行列 C

66 66 実装方法 (GPU#2) global void multiply2( float *da, float *db, float *dc, int N ) { int n = threadidx.x + blockidx.x * blockdim.x; 共有メモリの宣言 int m = threadidx.y + blockidx.y * blockdim.y; float sum = 0.0f; for( int k = 0 ; k < N ; k += 16 ) { shared float ta[ 16 ][ 16 ]; shared float tb[ 16 ][ 16 ]; 各スレッドが独立してメモリアクセス ta[ threadidx.y ][ threadidx.x ] = da[ m + ( k + threadidx.x ) * N ]; tb[ threadidx.x ][ threadidx.y ] = db[ ( k + threadidx.y ) + n * N ]; syncthreads( ); ブロック内のスレッドを同期 16 行列 B for( int t = 0 ; t < 16 ; t++ ) { sum += ta[ threadidx.y ][ t ] * tb[ threadidx.x ][ t ]; } N b kn 16 } } syncthreads( ); dc[ m + n * N ] = sum; ブロック内のスレッドを同期 各スレッドが独立して行列積を計算 N a mk 行列 A c mn 行列 C

67 67 計算時間の比較 (1) CPU, GPU#1, GPU#2 の計算時間を比較 使用計算機 CPU: Intel Core i7-x980(3.33ghz) OpenMPを使用して6スレッドで並列計算 GPU: NVIDIA Geforce GTX480 マルチプロセッサ数 : 15(SP 数 = = 480) Memory:1.5 GB OS:Windows 7 SP1

68 計算時間の比較 (2) [msec.] 計算時間 (N = 560 の場合 ) CPU#1: 91.7 msec. GPU#1: 18.3 msec. ( 5.0) GPU#2: 06.7 msec. ( 13.7) N = 560 CPU GPU# 行列のサイズ (N 次正方行列 ) GPU#2

69 69 まとめ ( 行列積 ) GPU 上での実装方法 1. 行列積の計算式に忠実な実装 (GPU#1) 2. 共有メモリを利用した高速化 (GPU#2) 大きさ16 16の部分行列積に分割共有メモリをキャッシュとして利用 CPU と GPU それぞれでの計算速度を比較 CPU < GPU#1 < GPU# メモリアクセスの工夫により約 34 倍の高速化

70 ~ テンプレートマッチング ~ 70

71 71 テンプレートマッチング (1) 画像処理の分野で広く用いられている手法 基板の品質検査 画像中の特定物体 ( 人物など ) の検出 入力画像中からテンプレートに類似する部分を探索 入力画像中に窓を設定 窓を移動させながら類似度を評価 比較 テンプレート 窓

72 72 テンプレートマッチング (2) 基本的なテンプレートマッチングの戦略 入力画像中の部分画像とテンプレートの類似度を評価 窓を移動させながら類似度計算 位置 (X 軸方向, Y 軸方向 ) 拡大 / 縮小 ( スケール変化 ) 回転 類似度の例 SAD(Sum of Absolute Difference) SSD(Sum of Squared Difference) NCC(Normalized Cross Correlation) 窓の大きさに比例して計算コストが増加 膨大な回数の類似度評価が必要 計算コスト大

73 デモ ( テンプレートマッチング ) 73

74 74 テンプレートマッチングの実装方法 テンプレートマッチングの特徴 類似度は各位置で独立に計算可能 GPU による並列計算が可能 頻繁に画像メモリへアクセス テクスチャメモリのキャッシュ機能を利用 複数スケールでテンプレートと入力画像を比較 正規化テクスチャ座標とハードウェア線形補間の利用 解像度に依存しないメモリアクセス スケール

75 75 処理の流れ CPU 側の処理 入力画像とテンプレートの読み込み GPUスレッドの同期 GPU 側の処理 類似度計算 スケールの変更 CPU GPU CPU GPU GPU CPU GPU CPU 画像読み込み メモリ確保 データ転送 類似度計算 スレッド同期 データ転送

76 76 テクスチャメモリとは? GPU 上の読み取り専用の特殊なメモリ領域 キャッシュを利用した高速なアクセスが可能 2 次元アクセスに対して効率的 ハードウェアを利用した高速な線形補間が可能 テクスチャメモリの定義 texture<datatype, Type, ReadMode> texref; DataType: データの型 (int,float 等 ) Type: テクスチャの種類 (1 次元テクスチャなど ) cudatexturetype1d,cudatexturetype2d, ReadMode: 値の範囲 cudareadmodeelementtype: 各データ型の値を使用 cudareadmodenormalizedfloat: 0 ~ 1 に正規化 ( 符号付きは -1~1)

77 77 メモリ確保と転送 (1) 入力画像のメモリを GPU 上に確保 テクスチャの定義 ( 画素表現 :RGBA) texture< uchar4, cudatexturetype2d, cudareadmodeelementtype > imgtex; cudaarray を利用して 2 次元テクスチャを確保 cudaarray *imgarray; cudachannelformatdesc c1 = cudacreatechanneldesc< uchar4 >( ); cudamallocarray( &imgarray, &c1, width, height ); 画素表現の指定 CPU から GPU へメモリ転送 転送データ量 cudamemcpytoarray( imgarray, 0, 0, psrc, nbytes, cudamemcpyhosttodevice ); 転送元ポインタ テクスチャメモリへの対応付け cudabindtexturetoarray( imgtex, imgarray, c1 ); CPU から GPU へ転送

78 78 メモリ確保と転送 (2) テンプレートのメモリを GPU 上に確保 テクスチャの定義 ハードウェア線形補間の有効化 ハードウェア線形補間の利用 (1) texture< uchar4, cudatexturetype2d, cudareadmodenormalizedfloat > reftex; reftex.filtermode = cudafiltermodelinear; ハードウェア線形補間の利用 (2) 正規化テクスチャ座標の有効化 reftex.normalized = 1; (0, 0) (W-1, 0) (0, 0) (1, 0) (0, 0) (1, 0) (0, 0) (1, 0) (0, 1) (0, H-1) (W-1, H-1) (0, 1) (0, 1) (1, 1)

79 79 GPU 上で類似度計算 global void kernel( float *error, float *scale, int imgw, int imgh, int areaw, int areah, int maskw, int maskh, float s ) { int i = threadidx.x + blockdim.x * blockidx.x; int j = threadidx.y + blockdim.y * blockidx.y; float err = 0.0f; float _1_w = 1.0f / maskw; // テンプレートの幅に対するスケーリング係数 float _1_h = 1.0f / maskh; // テンプレートの高さに対するスケーリング係数 for( int n = 0 ; n < maskh ; n++ ) { for( int m = 0 ; m < maskw ; m++ ) { uchar4 p1 = tex2d( imgtex, i + m, j + n ); float4 p2 = tex2d( reftex, _1_w * m, _1_h * n ) * 255.0f; テクスチャからデータを読み取り } } } err += ( p1.x - p2.x ) * ( p1.x - p2.x ); err += ( p1.y - p2.y ) * ( p1.y - p2.y ); err += ( p1.z - p2.z ) * ( p1.z - p2.z ); err *= _1_w * _1_h; if( error[ i + j * imgw ] > err ) { error[ i + j * imgw ] = err; scale[ i + j * imgw ] = s; } SSD が最小のものを記録 SSD を計算

80 80 CPU による類似度計算の制御 // スケール & 誤差を保持する GPU 側のメモリ領域 thrust::device_vector< float > derror( img.size( ), 1.0e10f ); thrust::device_vector< float > dscale( img.size( ), 0 ); スケールを変更して再探索 float s = MIN_SCALE; while( s <= MSCALE ) { // 実行する GPU のスレッド数を指定 int W = ( int )img.width( ), H = ( int )img.height( ); int w = ( int )( ref.width( ) * s ), h = ( int )( ref.height( ) * s ); int threadnumx = 16, threadnumy = 16; int blocknumx = ( W - w ) / threadnumx + ( ( ( W - w ) % threadnumx ) == 0? 0 : 1 ); int blocknumy = ( H - h ) / threadnumy + ( ( ( H - h ) % threadnumy ) == 0? 0 : 1 ); dim3 nthreads( threadnumx, threadnumy, 1 ); dim3 nblocks( blocknumx, blocknumy, 1 ); // GPU 側で類似度を計算する kernel<<< nblocks, nthreads >>>( thrust::raw_pointer_cast( &derror[ 0 ] ), thrust::raw_pointer_cast( &dscale[ 0 ] ), W, H, W - w, H - h, w, h, s ); cudadevicesynchronize( ); // 処理が完了するまで待機する } s *= SCALE_FACTOR; // スケールを変更する // GPU の処理結果を転送する thrust::host_vector< float > herror = derror; thrust::host_vector< float > hscale = dscale; GPU を用いて類似度を計算

81 81 計算時間の比較 (1) CPUとGPUで計算時間を評価 使用計算機 CPU: Intel Core i7-x980(3.33ghz) OpenMPを使用して6スレッドで並列計算 GPU: NVIDIA GeForce GTX480 マルチプロセッサ数 : 15(SP 数 = = 480) Memory:1.5 GB OS:Windows 7 SP1

82 82 計算時間の比較 (2) 実験パラメータ 入力画像 : 画素 テンプレート : 画素 スケール :0.3~1.9 倍 ( 拡大率 1.2) テンプレート 実験結果 CPU: 77.3 sec. GPU: 2.2 sec. 約 35 倍の高速化 結果画像

83 83 まとめ ( テンプレートマッチング ) テクスチャメモリを利用した高速化 キャッシュを利用した効率的なメモリアクセス 解像度に依存しないメモリアクセス ハードウェア線形補間 正規化テクスチャ座標 CPU と GPU で計算時間を比較 CPU: 77.3 sec. GPU: 2.2 sec. 約 35 倍の高速化を実現

84 ~ テンプレートマッチングを高速化 ~ 84

85 85 処理の流れ 類似度計算の打ち切り処理を導入 各スケールで計算した類似度の最大値をGPUで計算 Parallel Reduction アルゴリズムを利用 スケールの変更 CPU GPU CPU GPU GPU CPU GPU GPU CPU 画像読み込み メモリ確保 データ転送 類似度計算 スレッド同期 最大類似度計算 データ転送

86 86 GPU 上で類似度計算 global void kernel( float *error, float *scale, int imgw, int imgh, int areaw, int areah, int maskw, int maskh, float s, float maxerr ) { int i = threadidx.x + blockdim.x * blockidx.x; int j = threadidx.y + blockdim.y * blockidx.y; float err = 0.0f; float _1_w = 1.0f / maskw; // テンプレートの幅に対するスケーリング係数 float _1_h = 1.0f / maskh; // テンプレートの高さに対するスケーリング係数 for( int n = 0 ; n < maskh ; n++ ) { for( int m = 0 ; m < maskw ; m++ ) { uchar4 p1 = tex2d( imgtex, i + m, j + n ); float4 p2 = tex2d( reftex, _1_w * m, _1_h * n ) * 255.0f; } err += ( p1.x - p2.x ) * ( p1.x - p2.x ); err += ( p1.y - p2.y ) * ( p1.y - p2.y ); err += ( p1.z - p2.z ) * ( p1.z - p2.z ); } } if( maxerr < err *_1_w * _1_h ) { break; }... 以下同じ... 前スケールの探索時における SSD が最小値より大きい場合は計算を打ち切り

87 87 Parallel Reduction による最小値探索 GPU 上のメモリから最小値を計算 GPUのスレッドが並列に値の比較処理を実行 GPU 上のデータ 同期同期同期

88 88 Parallel Reduction による最小値探索 GPU 上のメモリから最小値を計算 GPU のスレッドが並列に値の比較処理を実行 Thrust ライブラリの reduce アルゴリズムを利用 thrust::reduce( 先頭, 末尾, 初期値, thrust::minimum<t>( ) ); 末尾を指すイテレータ 先頭を指すイテレータ GPU 上のデータ 同期同期同期

89 89 CPU による類似度計算の制御 // スケール & 誤差を保持する GPU 側のメモリ領域 thrust::device_vector< float > derror( img.size( ), 1.0e10f ); thrust::device_vector< float > dscale( img.size( ), 0 ); float s = MIN_SCALE, maxerr = 1.0e10f; while( s <= MSCALE ) { // 実行する GPU のスレッド数を指定 int W = ( int )img.width( ), H = ( int )img.height( ); int w = ( int )( ref.width( ) * s ), h = ( int )( ref.height( ) * s ); int threadnumx = 16, threadnumy = 16; int blocknumx = ( W - w ) / threadnumx + ( ( ( W - w ) % threadnumx ) == 0? 0 : 1 ); int blocknumy = ( H - h ) / threadnumy + ( ( ( H - h ) % threadnumy ) == 0? 0 : 1 ); dim3 nthreads( threadnumx, threadnumy, 1 ); dim3 nblocks( blocknumx, blocknumy, 1 ); // GPU 側で類似度を計算する kernel<<< nblocks, nthreads >>>( thrust::raw_pointer_cast( &derror[ 0 ] ), thrust::raw_pointer_cast( &dscale[ 0 ] ), W, H, W - w, H - h, w, h, s, maxerr ); cudadevicesynchronize( ); // 処理が完了するまで待機する // 誤差の最大値を取得する maxerr = thrust::reduce( derror.begin( ), derror.end( ), 1.0e10f, thrust::minimum< float >( ) ); } s *= SCALE_FACTOR; // スケールを変更する // GPU の処理結果を転送する... Trust ライブラリを用いて GPU を活用して誤差の最小値を計算

90 90 計算時間の比較 (1) CPUとGPUで計算時間を評価 使用計算機 CPU: Intel Core i7-x980(3.33ghz) OpenMPを使用して6スレッドで並列計算 GPU: NVIDIA GeForce GTX480 マルチプロセッサ数 : 15(SP 数 = = 480) Memory:1.5 GB OS:Windows 7 SP1

91 91 計算時間の比較 (2) 実験パラメータ 入力画像 : 画素 テンプレート : 画素 スケール :0.3~1.9 倍 ( 拡大率 1.2) テンプレート 実験結果 CPU: sec. GPU: sec. 約 33.7 倍の高速化 結果画像

92 92 まとめ ( テンプレートマッチング ) 類似度計算の打ち切り処理により高速化 Parallel Reduction アルゴリズムを利用 CPU と GPU で計算時間を比較 類似度の打ち切り なし あり CPU 77.3 秒 8.7 秒 GPU 2.2 秒 秒

93 ~ ガウシアンフィルタ ~ 93

94 ガウシアンフィルタに挑戦 94

95 95 空間フィルタリング ( 線形フィルタ ) 線形フィルタとは? i, j h x, y f i x j y 画像処理の基本的な処理 ガウシアンフィルタ LoGフィルタ ガボールフィルタ 他 h w g, y h x w 出力画像フィルタ 入力画像 フィルタ h(i,j) 積和 出力 g(i,j) 周辺領域 注目画素 f(i,j) 入力画像

96 ガウシアンフィルタ フィルタ係数に 2 次元ガウス分布を利用 問題点 フィルタ半径が大きくなるにつれ計算コスト大 フィルタを 1 次元ガウス分布の積に分解 exp 2 1, y x y x h exp 2 exp exp 2 1, y x y x y x h h h y w w x y j x i f x y j i g, 2 exp 2 exp 2 1, X 軸方向への 1 次元ガウシアンフィルタ Y 軸方向への 1 次元ガウシアンフィルタ

97 97 ガウシアンフィルタの特徴 フィルタ出力は各画素で独立に計算可能 GPU による並列計算が可能 フィルタ係数の算出 フィルタ適用前に事前計算が可能 フィルタ適用中は常に同じ値を参照 コンスタントメモリを利用 注目画素の周辺領域 (1 次元 ) にアクセス 1 次元テクスチャのキャッシュ機能を利用

98 98 コンスタントメモリとは? GPU 上に実装されている特殊なメモリ領域 読み取り専用 マルチプロセッサ1 基あたり 64 KB キャッシュを利用した高速なアクセスが可能 レジスタとほぼ同じ速度でアクセス可能 コンスタントメモリの定義 constant float coeff[ 512 ]; C 言語の配列のようにアクセス可能

99 99 処理の流れ GPU 上に 2 つの 1 次元配列を確保 X,Y 軸方向の処理時に入力と出力を入れ替え GPU 内でのメモリ転送コストを削減 CPU GPU CPU GPU GPU CPU GPU GPU CPU 画像読み込み メモリ確保 データ転送 ガウシ ( ア X ン軸フ ) ィルタ スレッド同期 ガウシ ( ア Y ン軸フ ) ィルタ データ転送 入力画像 メモリ 1 メモリ 2 メモリ 2

100 100 メモリ確保と転送 (1) 入力画像のメモリを GPU 上に確保 テクスチャの定義 texture< float4, 1 > imgtexx; texture< float4, 1 > imgtexy; GPU 上に 1 次元配列を確保 cudamalloc( ( void ** )&idata, nbytes ); cudamalloc( ( void ** )&odata, nbytes ); CPU から GPU へメモリ転送 cudamemcpy( odata, psrc, nbytes, cudamemcpyhosttodevice ); 1 次元テクスチャにマッピング cudabindtexture( 0, imgtexx, odata ); cudabindtexture( 0, imgtexy, idata );

101 101 メモリ確保と転送 (2) フィルタ係数をコンスタントメモリに確保 コンスタントメモリの定義 constant float coeff[ 512 ]; サイズ指定が必要コンスタントメモリを表す修飾子 CPU から GPU へメモリ転送 cudamemcpytosymbol( coeff, psrc, 512 * sizeof( float ) );

102 102 計算時間の比較 (1) CPU と GPU で計算時間を評価 使用計算機 CPU: Intel Core i7-x980(3.33ghz) OpenMP を使用して 6 スレッドで並列計算 GPU: NVIDIA Geforce GTX480 マルチプロセッサ数 : 15(SP 数 = = 480) Memory:1.5 GB OS:Windows 7 SP1 実験パラメータ 画像サイズ :100 2 ~ 画素 σ:5.0,10.0

103 103 計算時間の比較 (2) [msec.] 計算時間 (N 2 = の場合 ) CPU(σ=5) : msec. GPU(σ=5) : 20.9 msec. ( 5.0) CPU(σ=10) : msec. GPU(σ=10) : 30.7 msec. ( 5.5) CPU(σ=10) CPU(σ=5) GPU(σ=10) 画像サイズ N 2 GPU(σ=5)

104 104 まとめ ( ガウシアンフィルタ ) 2 次元ガウス分布を 1 次元ガウス分布の積で表現 ガウシアンフィルタを 1 次元フィルタに分解 コンスタントメモリを利用した高速化 キャッシュを利用したメモリアクセスの高速化 同じ値をスレッド間で共有する場合に有効 CPU と GPU で計算時間を比較 画像サイズが の場合 (σ=5) CPU: msec. GPU: 20.9 msec. 約 5 倍の高速化を実現

105 ~SIFT 特徴量の計算 ~ 105

106 106 SIFT(Scale Invariant Feature Transform) 回転 スケール変化等に頑健な特徴点の検出 画像間のマッチングや物体認識 検出に利用 1,2 1. David G. Lowe, Distinctive image features from scale-invariant keypoints, International Journal of Computer Vision, 60, 2, pp , 藤吉弘亘. "Gradient ベースの特徴抽出 - SIFT と HOG - ", 情報処理学会研究報告 CVIM 160, pp , 2007.

107 107 SIFT のアルゴリズム (1) DoG 画像の作成 (Difference-of-Gaussian) 特徴点の検出 エッジ上の点を削除 サブピクセル位置推定 コントラストの小さい点を削除 3 k 2 k 異なるスケールの平滑化画像の差分 (DoG) を計算 k 周辺 26 画素に対して極値をとる位置を特徴点として検出 DoG3 DoG2 DoG1 DoG3 DoG2 DoG1

108 108 SIFT のアルゴリズム (2) DoG 画像の作成 (Difference-of-Gaussian) 特徴点の検出 エッジ上の点を削除 サブピクセル位置推定 コントラストの小さい点を削除

109 109 SIFT のアルゴリズム (3) オリエンテーションの算出 特徴ベクトルを算出 周辺領域の勾配方向と強度からオリエンテーションを算出

110 デモ ( SIFT ) 110

111 111 SIFT の実装方法 さまざまなスケールにおける DoG 計算 高速なガウシアンフィルタを利用 キーポイント検出 各画素で独立に判定可能 GPU による並列計算が可能 判定に 26 近傍の画素値が必要 3 次元テクスチャを利用できるか? テクスチャサイズの制限から利用は困難 2 次元テクスチャで代用 オリエンテーションの算出 各キーポイントで独立に計算可能

112 処理の流れ 112 画像読み込み CPU D o G 計算 GPU スレッド同期 CPU データ転送 CPU GPU データ転送 GPU CPU キーポイント検出 GPU キーポイント DoG 画像オリエンテーション入力画像キーポイントリスト作成 CPU オリエンテーション算出 GPU データ転送 GPU CPU

113 113 SIFT 計算におけるメモリ配置 複数スケールの DoG 画像をテクスチャ 1 枚に配置 ミップマップ を構築 DoG 画像の出力先をソフトウェア的に調整 CUDA は未サポート

114 114 計算時間の比較 (1) CPU と GPU で計算時間を評価 SIFT 特徴量としてオリエンテーションを計算 使用計算機 CPU: Intel Core2 Quad Q9550(2.83 GHz) 4 スレッドで並列計算 GPU: NVIDIA GeForce GTX280 マルチプロセッサ数 : 30(SP 数 = 30 8 = 240) Memory:1.0 GB OS:Windows Vista SP1 実験パラメータ 画像サイズ :50 2 ~ 画素 σ 0 :1.6, 分割数 :3

115 115 計算時間の比較 (2) [msec.] 計算時間 (N 2 = の場合 ) CPU : msec. GPU : 10.2 msec. ( 44.9) 画像サイズ N 2 GPU CPU

116 デモ (SIFT) 116

117 117 まとめ (SIFT) テクスチャメモリとコンスタントメモリの利用 キャッシュを利用したメモリアクセスの高速化 ミップマップを構築 複数スケールの DoG 画像を 1 枚のテクスチャに格納 CPUとGPUで計算時間を比較 N 2 = の場合 CPU: msec. GPU: 10.2 msec. 約 44 倍の高速化を実現

118 ~GPGPU による高速画像処理に挑戦して ~ 118

119 119 GPGPU への挑戦を終えて CUDA を利用することで容易に GPGPU が可能 スレッドプログラミングの経験があれば非常に簡単 既存プログラムの移植も比較的容易 Thrust ライブラリによる簡単な GPGPU GPU を意識せずにプログラミングが可能 GPU を使うと 10 倍以上の高速化が可能? 多くの画像処理アルゴリズムは高速化が可能 空間フィルタリング, 局所特徴量計算, 他 逐次型の画像処理アルゴリズムは高速化が困難 ラベリング, 細線化, 他

120 120 GPGPU の問題点と今後の展望 複数 GPU の利用 各 GPU 上で別々に処理を実行 CUDA 4.0 は単一 CPUスレッドから複数 GPUを利用可能 CPU との連携 GPUの苦手な処理をCPUで計算 CPUとGPUの役割分担が重要 高性能な GPU が登場 リアルタイム画像処理 ( 大規模計算 ) への挑戦 GPGPU 開発環境の標準化 (OpenCL)

121 121 参考文献 [1] M. J. Harris, G. Coombe, T. Scheuermann, and A. Lastra, Physically- Based Visual Simulation on Graphics Hardware, Proceedings of SIGGRAPH 2002 / Eurographics Workshop on Graphics Hardware 2002, pp.1-10, 2002.(GPGPU の起源が書かれている論文 ) [2] J. D. Owens, D. Luebke, N. Govindaraju, M. Harris, J. Krüger, A. E. Lefohn, and T. J. Purcell, A Survey of General-Purpose Computation on Graphics Hardware, Computer Graphics Forum, Vol.26, No.1, pp , 2007.( 最近の GPGPU が詳しく述べられている論文 ) [3] GPGPU, [4] CUDA ZONE, [5] CUDA Programming Guide, [6] OpenCL,

122 122 MIST(Media Integration Standard Toolkit) 複数メディアを扱うためのライブラリ 音声 画像処理のアルゴリズムを多数実装 C/C++ を用いた高速な処理を実現 C++ のテンプレートを用いた汎用的な実装 複数のプラットフォームで動作 充実した日本語チュートリアルを用意 オープンソースとして公開中 BSDスタイルのライセンス 商用の製品開発でも利用可能

123 123 ~Visual Studio 2010 の詳細設定 ~ CUDA Toolkit 4.0 以降

124 124 Visual Studio 2010 の簡易設定 (1).cu ファイルの簡易コンパイル設定 プロジェクトメニューのビルドのカスタマイズを表示 CUDA 4.0 を選択.cu ファイルの追加前に行う 1. これを表示 2. これを選択

125 125 Visual Studio 2010 の簡易設定 (2).cu ファイルのコンパイルオプション.cu ファイルのプロパティを表示 NVCCのコンパイルオプションをGUIで調整可能 これを表示

126 126 Visual Studio 2010 の簡易設定 (3) プロジェクト プロパティ を選択 構成プロパティ リンカー を選択 追加の依存ファイル cudart.lib を指定 ここに入力

127 127 Visual Studio 2010 の簡易設定 (3) プログラムのコンパイルと実行 Visual Studio の ビルド を実行 実行ファイルが作成されることを確認 デバッグなしで開始 コマンドラインに Hello World!! が表示される Hello World!! 続行するには何かキーを押してください... 本講演のプログラムを実行した場合

128 ~ 補足資料 ~ 128

129 129 FLOPS FLoating point number Operations Per Second FLOPSやFLOP/s と表記される 1 秒あたりに実行可能な浮動小数点演算回数 スーパーコンピュータ等の性能を表す指標 代表的な CPU/GPU の FLOPS Core i GFLOPS GeforceGTX TFLOPS RadeonHD TFLOPS GPU は積和演算の性能

130 130 CUDA におけるエラー処理 API 関数 ( メモリ確保, 他 ) の場合 各 API 関数の戻り値を評価 cudasuccess 実行に成功 GPU で実行する関数の場合 cudathreadsynchronize( ) により同期 cudagetlasterror( ) の戻り値を評価 cudasuccess 実行に成功

131 ~OpenCV + CUDA~ 131

132 デモ ( 特徴点検出 & 対応付け ) 132

133 133 OpenCV をダウンロード Subversion 経由で最新版を入手 Cmake をダウンロード

134 134 OpenCV ライブラリのビルド CUDA を有効にする Cmake の設定で WITH_CUDA にチェック Configure を実行 ビルド設定を反映 Generate を実行 ビルドファイルを生成 3. ビルドファイルの生成 1. ここにチェック 2. 設定を反映 Cmake の設定画面

135 135 CPU/GPU のコードの違い SURF を用いた特徴点検出 & マッチング (CPU) #include <opencv2/opencv.hpp>... cv::mat_< float > desc2; cv::surffeaturedetector detector2( th2 ); // 特徴点を検出 detector2.detect( frame_gray, keys2 ); // 特徴量を計算 cv::surfdescriptorextractor extractor; extractor.compute( frame_gray, keys2, desc2 ); // 特徴点を対応付け cv::bruteforcematcher< cv::l2< float > > matcher; matcher.match( desc1, desc2, matches );

136 136 CPU/GPU のコードの違い SURF を用いた特徴点検出 & マッチング (GPU) #include <opencv2/gpu/gpu.hpp>... cv::gpu::gpumat desc_gpu2; cv::gpu::surf_gpu detector2; detector2.hessianthreshold = th2; // データを GPU へ転送 cv::gpu::gpumat frame_gpu; frame_gpu.upload( frame_gray ); // 特徴点検出 & 特徴量計算 detector2( frame_gpu, cv::gpu::gpumat(), keys2, desc_gpu2 ); // 特徴点を対応付け cv::gpu::bruteforcematcher_gpu< cv::l2< float > > matcher; matcher.match( desc_gpu1, desc_gpu2, matches );

137 137 OpenCV で GPU を利用できる機能 行列演算 テンプレートマッチング 歩行者検出 (HOG) 特徴点検出 & 特徴量 (SURF) 特徴点対応付け 画像フィルタ ラプラシアン, ソーベル, ガウシアン, 他 カメラキャリブレーション

07-二村幸孝・出口大輔.indd

07-二村幸孝・出口大輔.indd GPU Graphics Processing Units HPC High Performance Computing GPU GPGPU General-Purpose computation on GPU CPU GPU GPU *1 Intel Quad-Core Xeon E5472 3.0 GHz 2 6 MB L2 cache 1600 MHz FSB 80 GFlops 1 nvidia

More information

Microsoft Word - paper.docx

Microsoft Word - paper.docx による高速画像処理 名古屋大学大学院情報科学研究科出口大輔, 井手一郎, 村瀬洋 概要 : 本発表では, 近年注目を集めている GP(General Purpose computing on s) の技術に着目し,GP を利用するための開発環境の使い方やプログラミングのノウハウを分かりやすく解説する. GP は を汎用計算に利用しようという試みであり, 現在では物理シミュレーション, 数値計算, 信号解析,

More information

1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin

1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin Windows で始める CUDA 入門 GTC 2013 チュートリアル エヌビディアジャパン CUDA エンジニア森野慎也 1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境

More information

Slide 1

Slide 1 CUDA プログラミングの基本 パート II - カーネル CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パート II カーネルの起動 GPUコードの具体像 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください GPU 上でのコードの実行 カーネルは C 関数 + 多少の制約 ホストメモリはアクセスできない戻り値型は

More information

Microsoft PowerPoint - pr_12_template-bs.pptx

Microsoft PowerPoint - pr_12_template-bs.pptx 12 回パターン検出と画像特徴 テンプレートマッチング 領域分割 画像特徴 テンプレート マッチング 1 テンプレートマッチング ( 図形 画像などの ) 型照合 Template Matching テンプレートと呼ばれる小さな一部の画像領域と同じパターンが画像全体の中に存在するかどうかを調べる方法 画像内にある対象物体の位置検出 物体数のカウント 物体移動の検出などに使われる テンプレートマッチングの計算

More information

Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx

Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx GPU のプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU プログラミング環境 (CUDA) GPU プログラムの実行の流れ CUDA によるプログラムの記述 カーネル (GPU で処理する関数 ) の構造 記述方法とその理由 GPU 固有のパラメータの確認 405 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ

More information

熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date Type URL Presentation

熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date Type URL Presentation 熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date 2011-03-17 Type URL Presentation http://hdl.handle.net/2298/23539 Right GPGPU による高速演算について 榎本昌一 東京大学大学院工学系研究科システム創成学専攻

More information

CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン

CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン CUDA 画像処理入門 エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014 CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン RGB Y( 輝度 ) 変換 カラー画像から グレイスケールへの変換 Y = 0.299 R + 0.587

More information

1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU.....

1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU..... CPU GPU N Q07-065 2011 2 17 1 1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU...........................................

More information

Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments

Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments 計算機アーキテクチャ第 11 回 マルチプロセッサ 本資料は授業用です 無断で転載することを禁じます 名古屋大学 大学院情報科学研究科 准教授加藤真平 デスクトップ ジョブレベル並列性 スーパーコンピュータ 並列処理プログラム プログラムの並列化 for (i = 0; i < N; i++) { x[i] = a[i] + b[i]; } プログラムの並列化 x[0] = a[0] + b[0];

More information

GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓

GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU のアーキテクチャ CUDA CUDA によるプログラミング 58 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ チップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU と呼ぶことも多い

More information

TSUBAME2.0におけるGPUの 活用方法

TSUBAME2.0におけるGPUの 活用方法 GPU プログラミング 基礎編 東京工業大学学術国際情報センター 1. GPU コンピューティングと TSUBAME2.0 スーパーコンピュータ GPU コンピューティングとは グラフィックプロセッサ (GPU) は グラフィック ゲームの画像計算のために 進化を続けてきた 現在 CPU のコア数は 2~12 個に対し GPU 中には数百コア その GPU を一般アプリケーションの高速化に利用! GPGPU

More information

( CUDA CUDA CUDA CUDA ( NVIDIA CUDA I

(    CUDA CUDA CUDA CUDA (  NVIDIA CUDA I GPGPU (II) GPGPU CUDA 1 GPGPU CUDA(CUDA Unified Device Architecture) CUDA NVIDIA GPU *1 C/C++ (nvcc) CUDA NVIDIA GPU GPU CUDA CUDA 1 CUDA CUDA 2 CUDA NVIDIA GPU PC Windows Linux MaxOSX CUDA GPU CUDA NVIDIA

More information

Slide 1

Slide 1 CUDA プログラミングの基本 パート I - ソフトウェアスタックとメモリ管理 CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パートII カーネルの起動 GPUコードの具体項目 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください CUDA インストレーション CUDA インストレーションの構成

More information

NUMAの構成

NUMAの構成 GPU のプログラム 天野 アクセラレータとは? 特定の性質のプログラムを高速化するプロセッサ 典型的なアクセラレータ GPU(Graphic Processing Unit) Xeon Phi FPGA(Field Programmable Gate Array) 最近出て来た Deep Learning 用ニューロチップなど Domain Specific Architecture 1GPGPU:General

More information

Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx

Microsoft PowerPoint - GPGPU実践基礎工学(web).pptx GPU のメモリ階層 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU のメモリ階層 グローバルメモリ 共有メモリ モザイク処理への適用 コンスタントメモリ 空間フィルタへの適用 577 GPU の主要部品 基盤 GPU( チップ )+ 冷却部品 画面出力端子 電源入力端子 メモリ 特性の把握が重要 電源入力端子 画面出力端子 メモリ チップ PCI Ex 端子 http://www.geforce.com/whats

More information

GPGPUイントロダクション

GPGPUイントロダクション 大島聡史 ( 並列計算分科会主査 東京大学情報基盤センター助教 ) GPGPU イントロダクション 1 目的 昨今注目を集めている GPGPU(GPU コンピューティング ) について紹介する GPGPU とは何か? 成り立ち 特徴 用途 ( ソフトウェアや研究例の紹介 ) 使い方 ( ライブラリ 言語 ) CUDA GPGPU における課題 2 GPGPU とは何か? GPGPU General-Purpose

More information

main.dvi

main.dvi PC 1 1 [1][2] [3][4] ( ) GPU(Graphics Processing Unit) GPU PC GPU PC ( 2 GPU ) GPU Harris Corner Detector[5] CPU ( ) ( ) CPU GPU 2 3 GPU 4 5 6 7 1 toyohiro@isc.kyutech.ac.jp 45 2 ( ) CPU ( ) ( ) () 2.1

More information

3次多項式パラメタ推定計算の CUDAを用いた実装 (CUDAプログラミングの練習として) Implementation of the Estimation of the parameters of 3rd-order-Polynomial with CUDA

3次多項式パラメタ推定計算の CUDAを用いた実装 (CUDAプログラミングの練習として)  Implementation of the Estimation of the parameters of 3rd-order-Polynomial with CUDA 3 次多項式パラメタ推定計算の CUDA を用いた実装 (CUDA プログラミングの練習として ) Estimating the Parameters of 3rd-order-Polynomial with CUDA ISS 09/11/12 問題の選択 目的 CUDA プログラミングを経験 ( 試行錯誤と習得 ) 実際に CPU のみの場合と比べて高速化されることを体験 問題 ( インプリメントする内容

More information

1. マシンビジョンにおける GPU の活用

1. マシンビジョンにおける GPU の活用 CUDA 画像処理入門 GTC 213 チュートリアル エヌビディアジャパン CUDA エンジニア森野慎也 1. マシンビジョンにおける GPU の活用 1. 医用画像処理における GPU の活用 CT や MRI から画像を受信して三次元画像の構築をするシステム 2 次元スキャンデータから 3 次元 4 次元イメージの高速生成 CUDA 化により画像処理速度を約 2 倍に高速化 1. CUDA で画像処理

More information

TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日

TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日 TSUBAME2.0 における GPU の 活用方法 東京工業大学学術国際情報センター丸山直也第 10 回 GPU コンピューティング講習会 2011 年 9 月 28 日 目次 1. TSUBAMEのGPU 環境 2. プログラム作成 3. プログラム実行 4. 性能解析 デバッグ サンプルコードは /work0/gsic/seminars/gpu- 2011-09- 28 からコピー可能です 1.

More information

ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014

ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014 ストリームを用いたコンカレントカーネルプログラミングと最適化 エヌビディアジャパン CUDAエンジニア森野慎也 GTC Japan 2014 コンカレントな処理の実行 システム内部の複数の処理を 平行に実行する CPU GPU メモリ転送 カーネル実行 複数のカーネル間 ストリーム GPU 上の処理キュー カーネル実行 メモリ転送の並列性 実行順序 DEFAULT STREAM Stream : GPU

More information

Microsoft PowerPoint - GPU_computing_2013_01.pptx

Microsoft PowerPoint - GPU_computing_2013_01.pptx GPU コンピューティン No.1 導入 東京工業大学 学術国際情報センター 青木尊之 1 GPU とは 2 GPGPU (General-purpose computing on graphics processing units) GPU を画像処理以外の一般的計算に使う GPU の魅力 高性能 : ハイエンド GPU はピーク 4 TFLOPS 超 手軽さ : 普通の PC にも装着できる 低価格

More information

1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU

1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU GPGPU (I) GPU GPGPU 1 GPU(Graphics Processing Unit) GPU GPGPU(General-Purpose computing on GPUs) GPU GPGPU GPU ( PC ) PC PC GPU PC PC GPU GPU 2008 TSUBAME NVIDIA GPU(Tesla S1070) TOP500 29 [1] 2009 AMD

More information

memo

memo 数理情報工学演習第一 C プログラミング演習 ( 第 5 回 ) 2015/05/11 DEPARTMENT OF MATHEMATICAL INFORMATICS 1 今日の内容 : プロトタイプ宣言 ヘッダーファイル, プログラムの分割 課題 : 疎行列 2 プロトタイプ宣言 3 C 言語では, 関数や変数は使用する前 ( ソースの上のほう ) に定義されている必要がある. double sub(int

More information

PowerPoint プレゼンテーション

PowerPoint プレゼンテーション 講座準備 講座資料は次の URL から DL 可能 https://goo.gl/jnrfth 1 ポインタ講座 2017/01/06,09 fumi 2 はじめに ポインタはC 言語において理解が難しいとされる そのポインタを理解することを目的とする 講座は1 日で行うので 詳しいことは調べること 3 はじめに みなさん復習はしましたか? 4 & 演算子 & 演算子を使うと 変数のアドレスが得られる

More information

N08

N08 CPU のキモチ C.John 自己紹介 英語きらい 絵かけない 人の話を素直に信じない CPUにキモチなんてない お詫び 予告ではCとC# とありましたがやる気と時間の都合上 C++のみを対象とします 今日のネタ元 MSDN マガジン 2010 年 10 月号 http://msdn.microsoft.com/ja-jp/magazine/cc850829.aspx Windows と C++

More information

研究報告用MS-Wordテンプレートファイル

研究報告用MS-Wordテンプレートファイル マルチコアおよび GPGPU 環境における画像処理最適化 矢野勝久 高山征大 境隆二出宮健彦 スケーラを題材として, マルチコアおよび GPGPU 各々の HW 特性に適した画像処理の最適化を図る. マルチコア環境では, 数値演算処理の削減,SIMD 化など直列性能の最適化を行った後,OpenMP を利用して並列化を図る.GPGPU(CUDA) では, スレッド並列を優先して並列処理の設計を行いブロックサイズを決める.

More information

EnSightのご紹介

EnSightのご紹介 オープン CAE シンポジウム 2014 汎用ポストプロセッサー EnSight の大規模データ対応 CEI ソフトウェア株式会社代表取締役吉川慈人 http://www.ceisoftware.co.jp/ 内容 大規模データで時間のかかる処理 クライアント サーバー機能 マルチスレッドによる並列処理 サーバーの分散処理 クライアントの分散処理 ( 分散レンダリング ) EnSightのOpenFOAMインターフェース

More information

Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード]

Microsoft PowerPoint - GPUシンポジウム _d公開版.ppt [互換モード] 200/0/9 数値流体解析の並列効率とその GPU による高速化の試み 清水建設 ( 株 ) 技術研究所 PHAM VAN PHUC ( ファムバンフック ) 流体計算時間短縮と GPU の活用の試み 現 CPUとの比較によりGPU 活用の可能性 現 CPU の最大利用 ノード内の最大計算資源の利用 すべてCPUコアの利用 適切なアルゴリズムの利用 CPU コア性能の何倍? GPU の利用の試み

More information

untitled

untitled GPGPU NVIDACUDA Learn More about CUDA - NVIDIA http://www.nvidia.co.jp/object/cuda_education_jp.html NVIDIA CUDA programming Guide CUDA http://www.sintef.no/upload/ikt/9011/simoslo/evita/2008/seland.pdf

More information

CUDA 連携とライブラリの活用 2

CUDA 連携とライブラリの活用 2 1 09:30-10:00 受付 10:00-12:00 Reedbush-H ログイン GPU 入門 13:30-15:00 OpenACC 入門 15:15-16:45 OpenACC 最適化入門と演習 17:00-18:00 OpenACC の活用 (CUDA 連携とライブラリの活用 ) CUDA 連携とライブラリの活用 2 3 OpenACC 簡単にGPUプログラムが作成できる それなりの性能が得られる

More information

Microsoft PowerPoint - ICD2011UenoSlides.pptx

Microsoft PowerPoint - ICD2011UenoSlides.pptx 画像認識向け 3 次元積層 アクセラレータ アーキテクチャの検討 九州大学大学院システム情報科学府学院 * 九州大学大学院システム情報科学研究院 ** 上野伸也 * Gauthier Lovic Eric** 井上弘士 ** 村上和彰 ** 1 概要 画像認識技術 アクセラレータによる高性能 低消費エネルギー化 アプリケーション分析 アクセラレータ アーキテクチャ検討ア 性能 消費エネルギー評価 まとめ

More information

講習No.9

講習No.9 日本語は通常 2 バイトの文字コード.JIS コード, シフト JIS コード, Unicode (UTF-8) 等の様々な文字コードがある. アスキーコード表 (ASCII code) アスキーコード ( 値 ) 漢字変換無しでキーボードから直接入力できる半角文字 32 48 0 64 @ 80 P 96 ` 112 p 33! 49 1 65 A 81 Q 97 a 113 q 34 " 50

More information

Taro-ポインタ変数Ⅰ(公開版).j

Taro-ポインタ変数Ⅰ(公開版).j 0. 目次 1. ポインタ変数と変数 2. ポインタ変数と配列 3. ポインタ変数と構造体 4. ポインタ変数と線形リスト 5. 問題 問題 1 問題 2-1 - 1. ポインタ変数と変数 ポインタ変数には 記憶領域の番地が格納されている 通常の変数にはデータが格納されている 宣言 int *a; float *b; char *c; 意味ポインタ変数 aは 整数型データが保存されている番地を格納している

More information

プログラミング実習I

プログラミング実習I プログラミング実習 I 03 変数と式 人間システム工学科井村誠孝 m.imura@kwansei.ac.jp 3.1 変数と型 変数とは p.60 C 言語のプログラム中で, 入力あるいは計算された数や文字を保持するには, 変数を使用する. 名前がついていて値を入れられる箱, というイメージ. 変数定義 : 変数は変数定義 ( 宣言 ) してからでないと使うことはできない. 代入 : 変数には値を代入できる.

More information

PowerPoint プレゼンテーション

PowerPoint プレゼンテーション 講座を行う前に 自己紹介 僕と上回生について 1 年生同士で少しお話しよう! オリエンテーションの宿題 アルゴロジック http://home.jeita.or.jp/is/highschool/algo/index3.html どこまでできましたか? あまりできなかった人はこれから全部クリアしよう! 2016 年度 C 言語講座 第一回目 2016/6/11 fumi 今回の目標 プログラムを書いて実行するやり方を覚える

More information

コンピュータグラフィックス

コンピュータグラフィックス コンピュータグラフィックス 第 13 回 リアルタイム CG 理工学部 兼任講師藤堂英樹 CG 制作の主なワークフロー 3DCG ソフトウェアの場合 モデリング カメラ シーン アニメーション テクスチャ 質感 ライティング 画像生成 2015/12/21 コンピュータグラフィックス 2 リアルタイム CG CG をリアルタイムにする必要性 インタラクティブなユーザーとのやり取り 映像制作 モデリング,,

More information

スライド 1

スライド 1 知能制御システム学 画像処理の高速化 OpenCV による基礎的な例 東北大学大学院情報科学研究科鏡慎吾 swk(at)ic.is.tohoku.ac.jp 2007.07.03 リアルタイム処理と高速化 リアルタイム = 高速 ではない 目標となる時間制約が定められているのがリアルタイム処理である.34 ms かかった処理が 33 ms に縮んだだけでも, それによって与えられた時間制約が満たされるのであれば,

More information

画像ファイルを扱う これまでに学んだ条件分岐, 繰り返し, 配列, ファイル入出力を使って, 画像を扱うプログラムにチャレンジしてみよう

画像ファイルを扱う これまでに学んだ条件分岐, 繰り返し, 配列, ファイル入出力を使って, 画像を扱うプログラムにチャレンジしてみよう 第 14 回 応用 情報処理演習 ( テキスト : 第 10 章 ) 画像ファイルを扱う これまでに学んだ条件分岐, 繰り返し, 配列, ファイル入出力を使って, 画像を扱うプログラムにチャレンジしてみよう 特定色の画素の検出 ( テキスト 134 ページ ) 画像データが保存されているファイルを読み込んで, 特定色の画素の位置を検出するプログラムを作成しなさい 元画像生成画像 ( 結果の画像 )

More information

GPGPUクラスタの性能評価

GPGPUクラスタの性能評価 2008 年度理研 HPC シンポジウム第 3 世代 PC クラスタ GPGPU クラスタの性能評価 2009 年 3 月 12 日 富士通研究所成瀬彰 発表の概要 背景 GPGPU による高速化 CUDA の概要 GPU のメモリアクセス特性調査 姫野 BMT の高速化 GPGPU クラスタによる高速化 GPU Host 間のデータ転送 GPU-to-GPU の通信性能 GPGPU クラスタ上での姫野

More information

SICE東北支部研究集会資料(2013年)

SICE東北支部研究集会資料(2013年) 280 (2013.5.29) 280-4 SURF A Study of SURF Algorithm using Edge Image and Color Information Yoshihiro Sasaki, Syunichi Konno, Yoshitaka Tsunekawa * *Iwate University : SURF (Speeded Up Robust Features)

More information

Microsoft Word - 3new.doc

Microsoft Word - 3new.doc プログラミング演習 II 講義資料 3 ポインタ I - ポインタの基礎 1 ポインタとは ポインタとはポインタは, アドレス ( データが格納されている場所 ) を扱うデータ型です つまり, アドレスを通してデータを間接的に処理します ポインタを使用する場合の, 処理の手順は以下のようになります 1 ポインタ変数を宣言する 2 ポインタ変数へアドレスを割り当てる 3 ポインタ変数を用いて処理 (

More information

GPUコンピューティング講習会パート1

GPUコンピューティング講習会パート1 GPU コンピューティング (CUDA) 講習会 GPU と GPU を用いた計算の概要 丸山直也 スケジュール 13:20-13:50 GPU を用いた計算の概要 担当丸山 13:50-14:30 GPU コンピューティングによる HPC アプリケーションの高速化の事例紹介 担当青木 14:30-14:40 休憩 14:40-17:00 CUDA プログラミングの基礎 担当丸山 TSUBAME の

More information

MATLAB® における並列・分散コンピューティング ~ Parallel Computing Toolbox™ & MATLAB Distributed Computing Server™ ~

MATLAB® における並列・分散コンピューティング ~ Parallel Computing Toolbox™ & MATLAB Distributed Computing Server™ ~ MATLAB における並列 分散コンピューティング ~ Parallel Computing Toolbox & MATLAB Distributed Computing Server ~ MathWorks Japan Application Engineering Group Takashi Yoshida 2016 The MathWorks, Inc. 1 System Configuration

More information

C プログラミング演習 1( 再 ) 2 講義では C プログラミングの基本を学び 演習では やや実践的なプログラミングを通して学ぶ

C プログラミング演習 1( 再 ) 2 講義では C プログラミングの基本を学び 演習では やや実践的なプログラミングを通して学ぶ C プログラミング演習 1( 再 ) 2 講義では C プログラミングの基本を学び 演習では やや実践的なプログラミングを通して学ぶ 今回のプログラミングの課題 次のステップによって 徐々に難易度の高いプログラムを作成する ( 参照用の番号は よくわかる C 言語 のページ番号 ) 1. キーボード入力された整数 10 個の中から最大のものを答える 2. 整数を要素とする配列 (p.57-59) に初期値を与えておき

More information

Microsoft PowerPoint - exp2-02_intro.ppt [互換モード]

Microsoft PowerPoint - exp2-02_intro.ppt [互換モード] 情報工学実験 II 実験 2 アルゴリズム ( リスト構造とハッシュ ) 実験を始める前に... C 言語を復習しよう 0. プログラム書ける? 1. アドレスとポインタ 2. 構造体 3. 構造体とポインタ 0. プログラム書ける? 講義を聴いているだけで OK? 言語の要素技術を覚えれば OK? 目的のプログラム? 要素技術 データ型 配列 文字列 関数 オブジェクト クラス ポインタ 2 0.

More information

iphone GPGPU GPU OpenCL Mac OS X Snow LeopardOpenCL iphone OpenCL OpenCL NVIDIA GPU CUDA GPU GPU GPU 15 GPU GPU CPU GPU iii OpenMP MPI CPU OpenCL CUDA OpenCL CPU OpenCL GPU NVIDIA Fermi GPU Fermi GPU GPU

More information

スライド 1

スライド 1 GPU クラスタによる格子 QCD 計算 広大理尾崎裕介 石川健一 1.1 Introduction Graphic Processing Units 1 チップに数百個の演算器 多数の演算器による並列計算 ~TFLOPS ( 単精度 ) CPU 数十 GFLOPS バンド幅 ~100GB/s コストパフォーマンス ~$400 GPU の開発環境 NVIDIA CUDA http://www.nvidia.co.jp/object/cuda_home_new_jp.html

More information

Microsoft PowerPoint - comprog11.pptx

Microsoft PowerPoint - comprog11.pptx Outline プログラミング演習第 回エッジを検出する on 3..4 電気通信大学情報理工学部知能機械工学科長井隆行 画像の本質 輝度の境目に情報あり! 画像の微分と 階微分 エッジ検出 画像をぼかす 本日の課題 画像の本質 エッジ抽出 画像の情報は境目にあり! エッジ 輝度が大きく変化しているところ ( 境界 ) 画像の情報はエッジにあり 輝度 人間の視覚系でも特定のエッジの方向に発火するニューロンが見つかっている

More information

講習No.8

講習No.8 配列変数の要素 復習 int x[5]; x[0] x[1] x[2] x[3] x[4] 5 は配列の要素数 これらの変数をそれぞれ配列の要素と呼ぶ この数字を配列の添え字, またはインデックスと呼ぶ! 重要! インデックスの最大値 = 要素数ー 1 int x = 7; float aa[x]; int x = 7; float aa[7];! 重要! 配列宣言時の要素数は定数でなければならない

More information

バイオプログラミング第 1 榊原康文 佐藤健吾 慶應義塾大学理工学部生命情報学科

バイオプログラミング第 1 榊原康文 佐藤健吾 慶應義塾大学理工学部生命情報学科 バイオプログラミング第 1 榊原康文 佐藤健吾 慶應義塾大学理工学部生命情報学科 ポインタ変数の扱い方 1 ポインタ変数の宣言 int *p; double *q; 2 ポインタ変数へのアドレスの代入 int *p; と宣言した時,p がポインタ変数 int x; と普通に宣言した変数に対して, p = &x; は x のアドレスのポインタ変数 p への代入 ポインタ変数の扱い方 3 間接参照 (

More information

(MIRU2010) NTT Graphic Processor Unit GPU graphi

(MIRU2010) NTT Graphic Processor Unit GPU graphi (MIRU2010) 2010 7 889 2192 1-1 905 2171 905 NTT 243 0124 3-1 E-mail: ac094608@edu.okinawa-ct.ac.jp, akisato@ieee.org Graphic Processor Unit GPU graphic processor unit CUDA Fully automatic extraction of

More information

GPU GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1

GPU GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1 GPU 4 2010 8 28 1 GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1 Register & Shared Memory ( ) CPU CPU(Intel Core i7 965) GPU(Tesla

More information

Microsoft Word - no02.doc

Microsoft Word - no02.doc 使い方 1ソースプログラムの入力今回の講義では C++ 言語用の統合環境ソフトといわれるプログラムを利用します デスクトップにある CPad for C++ のアイコン ( 右参照 ) をダブルクリ ックしましょう ( 同じアイコンで Java_pad とかい エディタ部 てあるものもありますので気をつけてください ) これで 起 動します 統合環境を立ち上げると エディタ部とメッセージ部をもった画面が出てきます

More information

GPU 画像 動画処理用ハードウェア 低性能なプロセッサがたくさん詰まっている ピーク性能が非常に高い GPUを数値計算に用いるのがGPGPU Graphics Processing Unit General Purpose GPU TSUBAME2.0: GPUスパコン 本演習ではNVIDIA社の

GPU 画像 動画処理用ハードウェア 低性能なプロセッサがたくさん詰まっている ピーク性能が非常に高い GPUを数値計算に用いるのがGPGPU Graphics Processing Unit General Purpose GPU TSUBAME2.0: GPUスパコン 本演習ではNVIDIA社の 演習II (連続系アルゴリズム) 第2回: GPGPU 須田研究室 M1 本谷 徹 motoya@is.s.u-tokyo.ac.jp 2012/10/19 GPU 画像 動画処理用ハードウェア 低性能なプロセッサがたくさん詰まっている ピーク性能が非常に高い GPUを数値計算に用いるのがGPGPU Graphics Processing Unit General Purpose GPU TSUBAME2.0:

More information

3次元画像計測分野でのGPUを用いた高速化事例

3次元画像計測分野でのGPUを用いた高速化事例 3 次元画像計測分野での GPU を 用いた高速化事例 関口尚大, 三浦衛, 高橋徹, 青木孝文 東杜シーテック 東北大学大学院情報科学研究科 1 画像処理 GPGPU 関連事業製品組込みシステムソリューションカーナビ カーオー CUDA & OpenCL 最新画像処理技術ディオ 車載電装分ソリューション 画像処理の研究開発と 組込技術 野のソフトウェア開発その実用化 MATLAB/Simlinkを組み合わせで新たな価値の創造を目指しています

More information

tabaicho3mukunoki.pptx

tabaicho3mukunoki.pptx 1 2 はじめに n 目的 4倍精度演算より高速な3倍精度演算を実現する l 倍精度では足りないが4倍精度は必要ないケースに欲しい l 4倍精度に比べてデータサイズが小さい Ø 少なくともメモリ律速な計算では4倍精度よりデータ 転送時間を減らすことが可能 Ø PCIeやノード間通信がボトルネックとなりやすい GPUクラスタ環境に有効か n 研究概要 l DD型4倍精度演算 DD演算 に基づく3倍精度演算

More information

Microsoft PowerPoint - kougi7.ppt

Microsoft PowerPoint - kougi7.ppt C プログラミング演習 第 7 回メモリ内でのデータの配置 例題 1. 棒グラフを描く 整数の配列から, その棒グラフを表示する ループの入れ子で, 棒グラフの表示を行う ( 参考 : 第 6 回授業の例題 3) 棒グラフの1 本の棒を画面に表示する機能を持った関数を補助関数として作る #include "stdafx.h" #include void draw_bar( int

More information

GPUを用いたN体計算

GPUを用いたN体計算 単精度 190Tflops GPU クラスタ ( 長崎大 ) の紹介 長崎大学工学部超高速メニーコアコンピューティングセンターテニュアトラック助教濱田剛 1 概要 GPU (Graphics Processing Unit) について簡単に説明します. GPU クラスタが得意とする応用問題を議論し 長崎大学での GPU クラスタによる 取組方針 N 体計算の高速化に関する研究内容 を紹介します. まとめ

More information

System Requirements for Geomagic

System Requirements for Geomagic GEOMAGIC 動作環境 32-bit 版 64-bit 版 OS CPU RAM ハードディスクディスプレイ GPU - Windows XP (32-bitまたは64-bit SP2 以上 ) - Windows XP (64-bit SP2 以上 ) - Windows Vista (32-bitまたは64-bit SP1 - Windows Vista (64-bit SP1 以上 ) 以上

More information

02: 変数と標準入出力

02: 変数と標準入出力 C プログラミング入門 総機 1 ( 月 1) 11: 動的メモリ確保 Linux にログインし 以下の講義ページを開いておくこと http://www-it.sci.waseda.ac.jp/ teachers/w483692/cpr1/ 2015-06-22 1 まとめ : ポインタを使った処理 内容 説明 呼び出し元の変数を書き換える第 9 回 文字列を渡す 配列を渡す 第 10 回 ファイルポインタ

More information

概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran

概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran CUDA Fortran チュートリアル 2010 年 9 月 29 日 NEC 概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran を用いた Linux

More information

いまからはじめる組み込みGPU実装

いまからはじめる組み込みGPU実装 いまからはじめる組み込み GPU 実装 ~ コンピュータービジョン ディープラーニング編 ~ MathWorks Japan アプリケーションエンジニアリング部シニアアプリケーションエンジニア大塚慶太郎 2017 The MathWorks, Inc. 1 コンピュータービジョン ディープラーニングによる 様々な可能性 自動運転 ロボティクス 予知保全 ( 製造設備 ) セキュリティ 2 転移学習を使った画像分類

More information

Prog1_10th

Prog1_10th 2012 年 6 月 20 日 ( 木 ) 実施ポインタ変数と文字列前回は, ポインタ演算が用いられる典型的な例として, ポインタ変数が 1 次元配列を指す場合を挙げたが, 特に,char 型の配列に格納された文字列に対し, ポインタ変数に配列の 0 番の要素の先頭アドレスを代入して文字列を指すことで, 配列そのものを操作するよりも便利な利用法が存在する なお, 文字列リテラルは, その文字列が格納されている領域の先頭アドレスを表すので,

More information

02: 変数と標準入出力

02: 変数と標準入出力 C プログラミング入門 基幹 2 ( 月 4) 11: 動的メモリ確保 Linux にログインし 以下の講義ページを開いておくこと http://www-it.sci.waseda.ac.jp/ teachers/w483692/cpr1/ 2014-06-22 1 まとめ : ポインタを使った処理 内容 説明 呼び出し元の変数を書き換える第 9 回 文字列を渡す 配列を渡す 第 10 回 ファイルポインタ

More information

cp-7. 配列

cp-7. 配列 cp-7. 配列 (C プログラムの書き方を, パソコン演習で学ぶシリーズ ) https://www.kkaneko.jp/cc/adp/index.html 金子邦彦 1 本日の内容 例題 1. 月の日数配列とは. 配列の宣言. 配列の添え字. 例題 2. ベクトルの内積例題 3. 合計点と平均点例題 4. 棒グラフを描く配列と繰り返し計算の関係例題 5. 行列の和 2 次元配列 2 今日の到達目標

More information

GPGPU

GPGPU GPGPU 2013 1008 2015 1 23 Abstract In recent years, with the advance of microscope technology, the alive cells have been able to observe. On the other hand, from the standpoint of image processing, the

More information

Microsoft PowerPoint - 11.pptx

Microsoft PowerPoint - 11.pptx ポインタと配列 ポインタと配列 配列を関数に渡す 法 課題 : 配列によるスタックの実現 ポインタと配列 (1/2) a が配列であるとき, 変数の場合と同様に, &a[0] [] の値は配列要素 a[0] のアドレス. C 言語では, 配列は主記憶上の連続領域に割り当てられるようになっていて, 配列名 a はその配列に割り当てられた領域の先頭番地となる. したがって,&a[0] と a は同じ値.

More information

製品開発の現場では 各種のセンサーや測定環境を利用したデータ解析が行われ シミュレーションや動作検証等に役立てられています しかし 日々収集されるデータ量は増加し 解析も複雑化しており データ解析の負荷は徐々に重くなっています 例えば自動車の車両計測データを解析する場合 取得したデータをそのまま解析

製品開発の現場では 各種のセンサーや測定環境を利用したデータ解析が行われ シミュレーションや動作検証等に役立てられています しかし 日々収集されるデータ量は増加し 解析も複雑化しており データ解析の負荷は徐々に重くなっています 例えば自動車の車両計測データを解析する場合 取得したデータをそのまま解析 ホワイトペーパー Excel と MATLAB の連携がデータ解析の課題を解決 製品開発の現場では 各種のセンサーや測定環境を利用したデータ解析が行われ シミュレーションや動作検証等に役立てられています しかし 日々収集されるデータ量は増加し 解析も複雑化しており データ解析の負荷は徐々に重くなっています 例えば自動車の車両計測データを解析する場合 取得したデータをそのまま解析に使用することはできず

More information

プログラミング実習I

プログラミング実習I プログラミング実習 I 05 関数 (1) 人間システム工学科井村誠孝 m.imura@kwansei.ac.jp 関数とは p.162 数学的には入力に対して出力が決まるもの C 言語では入出力が定まったひとまとまりの処理 入力や出力はあるときもないときもある main() も関数の一種 何かの仕事をこなしてくれる魔法のブラックボックス 例 : printf() 関数中で行われている処理の詳細を使う側は知らないが,

More information

OpenGL GLSL References Kageyama (Kobe Univ.) Visualization / 58

OpenGL GLSL References Kageyama (Kobe Univ.) Visualization / 58 WebGL *1 2013.04.23 *1 X021 2013 LR301 Kageyama (Kobe Univ.) Visualization 2013.04.23 1 / 58 OpenGL GLSL References Kageyama (Kobe Univ.) Visualization 2013.04.23 2 / 58 Kageyama (Kobe Univ.) Visualization

More information

char int float double の変数型はそれぞれ 文字あるいは小さな整数 整数 実数 より精度の高い ( 数値のより大きい より小さい ) 実数 を扱う時に用いる 備考 : 基本型の説明に示した 浮動小数点 とは数値を指数表現で表す方法である 例えば は指数表現で 3 書く

char int float double の変数型はそれぞれ 文字あるいは小さな整数 整数 実数 より精度の高い ( 数値のより大きい より小さい ) 実数 を扱う時に用いる 備考 : 基本型の説明に示した 浮動小数点 とは数値を指数表現で表す方法である 例えば は指数表現で 3 書く 変数 入出力 演算子ここまでに C 言語プログラミングの様子を知ってもらうため printf 文 変数 scanf 文 if 文を使った簡単なプログラムを紹介した 今回は変数の詳細について習い それに併せて使い方が増える入出力処理の方法を習う また 演算子についての復習と供に新しい演算子を紹介する 変数の宣言プログラムでデータを取り扱う場合には対象となるデータを保存する必要がでてくる このデータを保存する場所のことを

More information

3/7 マイグレーション開発方針 顧客名 0 作成者 根岸正 < プログラム移行方針 > システム名称 A-VX システムマイグレーション作成日 2015/09/01 < COBOL 資産のプログラム移行 > COBOLソース ( メインとCOPYLIB) を入力としてSCC 言語変換ツールにてVB

3/7 マイグレーション開発方針 顧客名 0 作成者 根岸正 < プログラム移行方針 > システム名称 A-VX システムマイグレーション作成日 2015/09/01 < COBOL 資産のプログラム移行 > COBOLソース ( メインとCOPYLIB) を入力としてSCC 言語変換ツールにてVB 3/7 マイグレーション開発方針 顧客名 0 作成者 根岸正 < プログラム移行方針 > システム名称 A-VX システムマイグレーション作成日 2015/09/01 < COBOL 資産のプログラム移行 > COBOLソース ( メインとCOPYLIB) を入力としてSCC 言語変換ツールにてVB.netソリューションを作成します言語変換後にSDK( ソフトウェア開発キット ) にてデバッグおよびビルドにて実行可能アプリケーションを作成します

More information

GPGPU によるアクセラレーション環境について

GPGPU によるアクセラレーション環境について GPGPU によるアクセラレーション環境について 長屋貴量 自然科学研究機構分子科学研究所技術課計算科学技術班 概要 GPGPU とは 単純で画一的なデータを一度に大量に処理することに特化したグラフィックカードの演算資源を 画像処理以外の汎用的な目的に応用する技術の一つである 近年 その演算能力は CPU で通常言われるムーアの法則に則った場合とは異なり 飛躍的に向上しており その演算性能に魅力を感じた各分野での応用が広がってきている

More information

memo

memo 計数工学プログラミング演習 ( 第 4 回 ) 2016/05/10 DEPARTMENT OF MATHEMATICA INFORMATICS 1 内容 リスト 疎行列 2 連結リスト (inked ists) オブジェクトをある線形順序に並べて格納するデータ構造 単方向連結リスト (signly linked list) の要素 x キーフィールド key ポインタフィールド next x->next:

More information

Microsoft PowerPoint - Lec15 [互換モード]

Microsoft PowerPoint - Lec15 [互換モード] 情報デザイン専攻 画像情報処理論及び演習 II 周波数分解 FFT Gaussian フィルタと周波数分解 今日の授業内容 www.riken.jp/brict/yoshizawa/lectures/index.html www.riken.jp/brict/yoshizawa/lectures/lec5.pdf. 前回 前々回の復習 レポートの説明. 第 3, 回講義水曜日 限教室 68 吉澤信

More information

IPSJ SIG Technical Report Vol.2013-ARC-203 No /2/1 SMYLE OpenCL (NEDO) IT FPGA SMYLEref SMYLE OpenCL SMYLE OpenCL FPGA 1

IPSJ SIG Technical Report Vol.2013-ARC-203 No /2/1 SMYLE OpenCL (NEDO) IT FPGA SMYLEref SMYLE OpenCL SMYLE OpenCL FPGA 1 SMYLE OpenCL 128 1 1 1 1 1 2 2 3 3 3 (NEDO) IT FPGA SMYLEref SMYLE OpenCL SMYLE OpenCL FPGA 128 SMYLEref SMYLE OpenCL SMYLE OpenCL Implementation and Evaluations on 128 Cores Takuji Hieda 1 Noriko Etani

More information

CLEFIA_ISEC発表

CLEFIA_ISEC発表 128 ビットブロック暗号 CLEFIA 白井太三 渋谷香士 秋下徹 盛合志帆 岩田哲 ソニー株式会社 名古屋大学 目次 背景 アルゴリズム仕様 設計方針 安全性評価 実装性能評価 まとめ 2 背景 AES プロジェクト開始 (1997~) から 10 年 AES プロジェクト 攻撃法の進化 代数攻撃 関連鍵攻撃 新しい攻撃法への対策 暗号設計法の進化 IC カード, RFID などのアプリケーション拡大

More information

次に示す数値の並びを昇順にソートするものとする このソートでは配列の末尾側から操作を行っていく まず 末尾の数値 9 と 8 に着目する 昇順にソートするので この値を交換すると以下の数値の並びになる 次に末尾側から 2 番目と 3 番目の 1

次に示す数値の並びを昇順にソートするものとする このソートでは配列の末尾側から操作を行っていく まず 末尾の数値 9 と 8 に着目する 昇順にソートするので この値を交換すると以下の数値の並びになる 次に末尾側から 2 番目と 3 番目の 1 4. ソート ( 教科書 p.205-p.273) 整列すなわちソートは アプリケーションを作成する際には良く使われる基本的な操作であり 今までに数多くのソートのアルゴリズムが考えられてきた 今回はこれらソートのアルゴリズムについて学習していく ソートとはソートとは与えられたデータの集合をキーとなる項目の値の大小関係に基づき 一定の順序で並べ替える操作である ソートには図 1 に示すように キーの値の小さいデータを先頭に並べる

More information

GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 理化学研究所 共通コードプロジェクト

GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 理化学研究所 共通コードプロジェクト GPU チュートリアル :OpenACC 篇 Himeno benchmark を例題として 高エネルギー加速器研究機構 (KEK) 松古栄夫 (Hideo Matsufuru) 1 December 2018 HPC-Phys 勉強会 @ 理化学研究所 共通コードプロジェクト Contents Hands On 環境について Introduction to GPU computing Introduction

More information

ホワイトペーパー

ホワイトペーパー ホワイトペーパー NVIDIA の次世代 CUDA コンピュートアーキテクチャ : Fermi 目次 GPU コンピューティングの歴史 G80 アーキテクチャ計算処理とグラフィックスをカバーする NVIDIA の次世代アーキテクチャ CUDA ( 開発コード : Fermi ) CUDA の概要ハードウェア実行 Fermi アーキテクチャの概要第 3 世代のストリーミング マルチプロセッサ 512

More information

WebGL OpenGL GLSL Kageyama (Kobe Univ.) Visualization / 57

WebGL OpenGL GLSL Kageyama (Kobe Univ.) Visualization / 57 WebGL 2014.04.15 X021 2014 3 1F Kageyama (Kobe Univ.) Visualization 2014.04.15 1 / 57 WebGL OpenGL GLSL Kageyama (Kobe Univ.) Visualization 2014.04.15 2 / 57 WebGL Kageyama (Kobe Univ.) Visualization 2014.04.15

More information

Slide 1

Slide 1 電子情報通信学会研究会組込みシステム研究会 (IPSJ-EMB) 2010 年 1 月 28 日 超並列マルチコア GPU を用いた高速演算処理の実用化 NVIDIA Solution Architect 馬路徹 目次 なぜ今 GPU コンピューティングか? CPUの性能向上速度が減速 性能向上 = 並列処理 にGPUコンピューティングが応える CUDAシステムアーキテクチャによる超並列処理の実現

More information

untitled

untitled A = QΛQ T A n n Λ Q A = XΛX 1 A n n Λ X GPGPU A 3 T Q T AQ = T (Q: ) T u i = λ i u i T {λ i } {u i } QR MR 3 v i = Q u i A {v i } A n = 9000 Quad Core Xeon 2 LAPACK (4/3) n 3 O(n 2 ) O(n 3 ) A {v i }

More information

CCS HPCサマーセミナー 並列数値計算アルゴリズム

CCS HPCサマーセミナー 並列数値計算アルゴリズム 大規模系での高速フーリエ変換 2 高橋大介 daisuke@cs.tsukuba.ac.jp 筑波大学計算科学研究センター 2016/6/2 計算科学技術特論 B 1 講義内容 並列三次元 FFT における自動チューニング 二次元分割を用いた並列三次元 FFT アルゴリズム GPU クラスタにおける並列三次元 FFT 2016/6/2 計算科学技術特論 B 2 並列三次元 FFT における 自動チューニング

More information

<4D F736F F D20438CBE8CEA8D758DC F0939A82C282AB2E646F63>

<4D F736F F D20438CBE8CEA8D758DC F0939A82C282AB2E646F63> C 言語講座第 2 回 作成 : ハルト 前回の復習基本的に main () の中カッコの中にプログラムを書く また 変数 ( int, float ) はC 言語では main() の中カッコの先頭で宣言する 1 画面へ出力 printf() 2 キーボードから入力 scanf() printf / scanf で整数を表示 / 入力 %d 小数を表示 / 入力 %f 3 整数を扱う int 型を使う

More information

インテル(R) Visual Fortran コンパイラ 10.0

インテル(R) Visual Fortran コンパイラ 10.0 インテル (R) Visual Fortran コンパイラー 10.0 日本語版スペシャル エディション 入門ガイド 目次 概要インテル (R) Visual Fortran コンパイラーの設定はじめに検証用ソースファイル適切なインストールの確認コンパイラーの起動 ( コマンドライン ) コンパイル ( 最適化オプションなし ) 実行 / プログラムの検証コンパイル ( 最適化オプションあり ) 実行

More information

27_02.indd

27_02.indd GPGPU を用いたソフトウェア高速化手法 Technique to Speedup of the software by GPGPU 大田弘樹 馬場明子 下田雄一 安田隆洋 山本啓二 Hiroki Ota, Akiko Baba, Shimoda Yuichi, Takahiro Yasuta, Keiji Yamamoto PCやワークステーションにおいて画像処理に特化して使用されてきたGPUを

More information

PowerPoint Presentation

PowerPoint Presentation 工学部 6 7 8 9 10 組 ( 奇数学籍番号 ) 担当 : 長谷川英之 情報処理演習 第 7 回 2010 年 11 月 18 日 1 今回のテーマ 1: ポインタ 変数に値を代入 = 記憶プログラムの記憶領域として使用されるものがメモリ ( パソコンの仕様書における 512 MB RAM などの記述はこのメモリの量 ) RAM は多数のコンデンサの集合体 : 電荷がたまっている (1)/ いない

More information

Microsoft Word - HOKUSAI_system_overview_ja.docx

Microsoft Word - HOKUSAI_system_overview_ja.docx HOKUSAI システムの概要 1.1 システム構成 HOKUSAI システムは 超並列演算システム (GWMPC BWMPC) アプリケーション演算サーバ群 ( 大容量メモリ演算サーバ GPU 演算サーバ ) と システムの利用入口となるフロントエンドサーバ 用途の異なる 2 つのストレージ ( オンライン ストレージ 階層型ストレージ ) から構成されるシステムです 図 0-1 システム構成図

More information

4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司

4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司 4 倍精度基本線形代数ルーチン群 QPBLAS の紹介 [index] 1. Introduction 2. Double-double algorithm 3. QPBLAS 4. QPBLAS-GPU 5. Summary 佐々成正 1, 山田進 1, 町田昌彦 1, 今村俊幸 2, 奥田洋司 3 1 1 日本原子力研究開発機構システム計算科学センター 2 理科学研究所計算科学研究機構 3 東京大学新領域創成科学研究科

More information

DO 時間積分 START 反変速度の計算 contravariant_velocity 移流項の計算 advection_adams_bashforth_2nd DO implicit loop( 陰解法 ) 速度勾配, 温度勾配の計算 gradient_cell_center_surface 速

DO 時間積分 START 反変速度の計算 contravariant_velocity 移流項の計算 advection_adams_bashforth_2nd DO implicit loop( 陰解法 ) 速度勾配, 温度勾配の計算 gradient_cell_center_surface 速 1 1, 2 1, 2 3 2, 3 4 GP LES ASUCA LES NVIDIA CUDA LES 1. Graphics Processing Unit GP General-Purpose SIMT Single Instruction Multiple Threads 1 2 3 4 1),2) LES Large Eddy Simulation 3) ASUCA 4) LES LES

More information

hpc141_shirahata.pdf

hpc141_shirahata.pdf GPU アクセラレータと不揮発性メモリ を考慮した I/O 性能の予備評価 白幡晃一 1,2 佐藤仁 1,2 松岡聡 1 1: 東京工業大学 2: JST CREST 1 GPU と不揮発性メモリを用いた 大規模データ処理 大規模データ処理 センサーネットワーク 遺伝子情報 SNS など ペタ ヨッタバイト級 高速処理が必要 スーパーコンピュータ上での大規模データ処理 GPU 高性能 高バンド幅 例

More information

memo

memo 計数工学プログラミング演習 ( 第 3 回 ) 2016/04/26 DEPARTMENT OF MATHEMATICAL INFORMATICS 1 内容 ポインタ malloc 構造体 2 ポインタ あるメモリ領域 ( アドレス ) を代入できる変数 型は一致している必要がある 定義時には値は不定 ( 何も指していない ) 実際にはどこかのメモリを指しているので, #include

More information

VXPRO R1400® ご提案資料

VXPRO R1400® ご提案資料 Intel Core i7 プロセッサ 920 Preliminary Performance Report ノード性能評価 ノード性能の評価 NAS Parallel Benchmark Class B OpenMP 版での性能評価 実行スレッド数を 4 で固定 ( デュアルソケットでは各プロセッサに 2 スレッド ) 全て 2.66GHz のコアとなるため コアあたりのピーク性能は同じ 評価システム

More information

untitled

untitled A = QΛQ T A n n Λ Q A = XΛX 1 A n n Λ X GPGPU A 3 T Q T AQ = T (Q: ) T u i = λ i u i T {λ i } {u i } QR MR 3 v i = Q u i A {v i } A n = 9000 Quad Core Xeon 2 LAPACK (4/3) n 3 O(n 2 ) O(n 3 ) A {v i }

More information

高性能計算研究室の紹介 High Performance Computing Lab.

高性能計算研究室の紹介 High Performance Computing Lab. 高性能計算研究室 (HPC Lab) の紹介 High Performance Computing Lab. 静岡理工科大学総合情報学部コンピュータシステム学科 ( 兼 Web デザイン特別プログラム ) 幸谷智紀 543 研究室 幸谷研究室 @ 静岡 検索 概要 1. 幸谷智紀 個人の研究テーマ 2. 3 年生ゼミ ( 情報セミナー II) 3. 卒研テーマ 4. 過去の卒研 5. 今後について

More information

目次 目次... 1 はじめに... 3 概要... 4 サポート環境... 5 関数... 6 MEC_OpenDevice... 7 MECDevice_Release... 8 MECDevice_GetFirmVersion... 9 MECDevice_GetCoreTemperature

目次 目次... 1 はじめに... 3 概要... 4 サポート環境... 5 関数... 6 MEC_OpenDevice... 7 MECDevice_Release... 8 MECDevice_GetFirmVersion... 9 MECDevice_GetCoreTemperature MECodecAPI Reference Manual 2015 年 9 月 1 日 MEDIAEDGE 株式会社 目次 目次... 1 はじめに... 3 概要... 4 サポート環境... 5 関数... 6 MEC_OpenDevice... 7 MECDevice_Release... 8 MECDevice_GetFirmVersion... 9 MECDevice_GetCoreTemperature...

More information

Microsoft PowerPoint - 09.pptx

Microsoft PowerPoint - 09.pptx 情報処理 Ⅱ 第 9 回 2014 年 12 月 22 日 ( 月 ) 関数とは なぜ関数 関数の分類 自作関数 : 自分で定義する. ユーザ関数 ユーザ定義関数 などともいう. 本日のテーマ ライブラリ関数 : 出来合いのもの.printf など. なぜ関数を定義するのか? 処理を共通化 ( 一般化 ) する プログラムの見通しをよくする 機能分割 ( モジュール化, 再利用 ) 責任 ( あるいは不具合の発生源

More information