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

Size: px
Start display at page:

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

Transcription

1 GPU のメモリ階層 長岡技術科学大学電気電子情報工学専攻出川智啓

2 今回の内容 GPU のメモリ階層 グローバルメモリ 共有メモリ モザイク処理への適用 コンスタントメモリ 空間フィルタへの適用 577

3 GPU の主要部品 基盤 GPU( チップ )+ 冷却部品 画面出力端子 電源入力端子 メモリ 特性の把握が重要 電源入力端子 画面出力端子 メモリ チップ PCI Ex 端子 new/articles /introducing the geforce gtx 780 に公開されている写真を基に作成 578

4 CPU のメモリ階層 オフチップ (off chip) メモリ CPUのチップ外部に置かれたメモリ メインメモリ ( 主記憶 ) 利用可能なメモリの中で最低速, 最大容量 オンチップ (on chip) メモリ CPU のチップ内部に置かれたメモリ レベル 1(L1) キャッシュ レベル 2(L2) キャッシュ レベル 2(L3) キャッシュ 高速, 容量小 低速, 容量大 579

5 GPU のメモリ階層 オフチップメモリ PCI Exカードの基板上に実装 ビデオメモリ 利用可能なメモリの中で最低速, 最大容量 オンチップメモリ GPU のチップ内部に置かれたメモリ レベル1(L1) キャッシュ レベル2(L2) キャッシュ 高速, 容量小低速, 容量大 CPU の構造に類似 580

6 GPU メモリの独自の特徴 CPU とは異なるメモリを複数搭載 各メモリの特徴を知り, 適材適所で利用する事により高速化 GPU から読み書き可能か 処理を行うスレッドから読み書き可能か, 読み込みのみか 複数のスレッドでデータを共有できるか CPU から読み書き可能か C 言語の機能のみで直接読み書きは不可能 の専用関数 (API) を利用して読み書き 581

7 メモリの階層 CPU のメモリ階層 コアごとに L2 キャッシュ, 全体で L3 キャッシュを持つこともある チップ コア コア コア 演算器 演算器 演算器 演算器 演算器 演算器 L1 キャッシュ L1 キャッシュ L1 キャッシュ L2 キャッシュ メインメモリ 582

8 メモリの階層 GPU のメモリ階層 CPU にはない独自のメモリを複数持つ チップ L1 キャッシュ 共有メモリ L2 キャッシュ テクスチャキャッシュ コンスタントキャッシュ グローバルメモリ ローカルメモリ テクスチャメモリ コンスタントメモリ 583

9 メモリの種類 オンチップメモリ (GPU のチップ内部に置かれたメモリ ) 高速アクセス, 小容量 CPUからはアクセス不可 L1キャッシュと共有メモリは一定サイズを共用 L1 キャッシュ / 共有 ( シェアード ) メモリ 容量小小 速度高速高速 GPU からの読み書き CPU からのアクセス 読み書き可ブロック内の全スレッドが同じアドレスにアクセス ( データを共有 ) することが可能 * 読み書き不可 読み書き可各スレッドが異なるアドレスにアクセス 読み書き不可 * スレッドごとに異なるアドレスにアクセスすることも可能 584

10 メモリの種類 オフチップメモリ (GPU のチップ外部に置かれたメモリ ) 低速アクセス, 大容量 CPUから直接アクセス可能 ローカルメモリだけはアクセス不可 グローバルメモリ ローカルメモリ テクスチャメモリ コンスタントメモリ 容量 大 小 大 小 速度 低速 低速 高速 * 高速 * GPU からの読み書き 読み書き可全てのスレッドが同じアドレスにアクセス可能 ** 読み書き可各スレッドが異なるアドレスにアクセス 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** CPU からのアクセス 読み書き可読み書き不可読み書き可読み書き可 * キャッシュが効く場合 ** スレッドごとに異なるアドレスにアクセスすることも可能 585

11 メモリの種類 共有メモリと L1 キャッシュは一定サイズを共用 グローバルメモリへのアクセスは L2 キャッシュ経由 Fermi 世代以前の GPU はキャッシュ無し * GPU Chip SM L1 キャッシュ 共有メモリ SM L1 キャッシュ 共有メモリ オフチップメモリ オンチップメモリ ホストメモリ ローカルメモリ L2キャッシュコンスタントメモリテクスチャメモリ *Tesla 世代でもテクスチャキャッシュは存在. Read/Write 可能なキャッシュは Fermi 世代から. ローカルメモリ グローバルメモリ 586

12 メモリの種類と並列化階層の対応 オンチップメモリ ブロックまたはスレッドごとに異なる値を持てる が不足するとローカルメモリが使われる オフチップメモリ GPU 全体で共通の値を持てる Grid Block(0,0,0) L1 キャッシュ Thre ad 0 Thre ad 1 共有メモリ Thre ad 2 Thre ad 3 Block(1,0,0) L1 キャッシュ Thre ad 0 Thre ad 1 共有メモリ Thre ad 2 Thre ad 3 各 GPU(Grid) 内でデータを共有各ブロック内でデータを共有各スレッドが個別のデータを保有 ホストメモリ ローカルメモリ ローカルメモリ L2キャッシュコンスタントメモリテクスチャメモリグローバルメモリ 587

13 各スレッドが個別に利用 カーネル内で変数を宣言するとを利用 非常に高速 キャッシュとしても利用可能 少容量 本 * 32bit 利用可能分を超えるとローカルメモリへ追い出される スピル *Kepler からは 本 ホストメモリ GPU Chip SM L1 キャッシュ ローカルメモリ ローカルメモリ 共有メモリ SM L2 キャッシュ L1 キャッシュ コンスタントメモリ テクスチャメモリ グローバルメモリ 共有メモリ 588

14 グローバルメモリ ビデオメモリ ( 数 GB) CPU のメインメモリに相当 読み込みがある一定サイズでまとめて行われる レイテンシが大きい 読み出し命令を発効してからデータが得られるまでの時間 効率よくアクセスするための条件がある コアレスアクセス ( コアレッシング ) 589 ホストメモリ GPU Chip SM L1 キャッシュ ローカルメモリ ローカルメモリ 共有メモリ SM L2 キャッシュ L1 キャッシュ コンスタントメモリ テクスチャメモリ グローバルメモリ 共有メモリ

15 グローバルメモリへのアクセス * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. 16 スレッドが協調して同時にアクセス では 32 スレッドを Warp という単位で管理 Warp の半分 (Half Warp) が協調して読み書き コアレスアクセスか否かで読込速度が大幅に変化 新しい世代のGPUでは速度の落ち込みが緩和 コアレスアクセスはGPUプログラムで最重要 GPUの処理能力と比較するとデータ供給不足が発生 効率よくデータを供給するためにコアレスアクセスは必須 590

16 コアレスアクセスになる条件 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. データのサイズ 32bit, 64bit, 128bit(4 バイト, 8 バイト, 16 バイト ) アドレスの隣接 16 スレッドがアクセスするアドレスがスレッド番号順に隣接 アクセスする最初のアドレス 16 スレッドがアクセスするアドレスの先頭 ( スレッド 0 がアクセスするアドレス ) が,64 バイトまたは 128 バイト境界 アドレスが 64 の倍数で始まる 64 バイトの領域か, アドレスが 128 の倍数で始まる 128 バイトの領域 591

17 コアレスアクセスの例 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. データ型が 32bit=4 バイト T0 T1 A128 A132 データ型が 32bit=4 バイト T0 T1 A128 A132 各スレッドが連続して隣接アドレスにアクセス 先頭アドレスが 128 バイト境界 16 スレッド T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 A136 A140 A144 A148 A152 A156 A160 A164 A168 A172 A176 各スレッドが連続して隣接アドレスにアクセス 実際にデータを取得するかは無関係 先頭アドレスが 128 バイト境界 T2 T3 T4 T5 T6 T7 T8 T9 T10 T11 T12 A136 A140 A144 A148 A152 A156 A160 A164 A168 A172 A176 T13 A180 T13 A180 T14 A184 T14 A184 T15 A188 T15 A

18 コアレスアクセスにならない例 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. T0 A128 T0 A128 T0 A128 T0 A128 T1 A132 T1 A132 T1 A132 T1 T2 A136 T2 A136 T2 A136 T2 T3 A140 T3 A140 T3 A140 T3 A140 T4 A144 T4 A144 T4 A144 T4 T5 A148 T5 A148 T5 A148 T5 T6 A152 T6 A152 T6 A152 T6 A152 T7 T10 A156 各スレッドが番 T8 A160 号順にアクセス T9 A164 していない A168 T7 T10 A156 先頭が128バイト T8 A160 境界ではない T9 A164 ( 現在は緩和 ) A168 T7 A156 アドレスが連続し T8 A160 ていない T9 A164 T10 A168 T7 データが 32bit, 64bit, T8 128bit T9 ではないA164 ( 構造体など ) T10 T11 A172 T11 A172 T11 A172 T11 T12 A176 T12 A176 T12 A176 T12 A176 T13 A180 T13 A180 T13 A180 T13 T14 A184 T14 A184 T14 A184 T14 T15 A188 T15 A188 T15 A188 T15 A

19 コアレスアクセスにならない例 * * かなり古い情報のため要注意. Fermi 世代以降では状況が大きく異なる. 128 バイト境界からわずかにずれている場合 T0 T1 T2 A128 A132 A136 Tesla 世代以降は 64 バイトブロックと 32 バイトブロックに分けて読込 メモリアクセス要求は 2 回 コアレスアクセスの半分程度の性能は得られる T3 T4 T5 T6 T7 T8 T9 T10 T11 A140 A144 A148 A152 A156 A160 A164 A168 A バイトブロックでデータ読込 T12 A176 T13 A180 T14 A184 T15 A188 A バイトブロックでデータ読込 594

20 モザイク処理 前回授業で取り上げたモザイク処理 画像を小さな領域に分け, その領域を全て同じ色にする 領域内の全画素を, 領域内の画素の平均値に置き換える 原画像処理画像 595

21 モザイク処理 前回授業で取り上げたモザイク処理 高速化が今ひとつ 他の処理より一桁遅い 処理 処理時間 [ms] CPU GPU ネガティブ処理 水平反転 垂直反転 空間フィルタ モザイク処理

22 モザイク処理 ( 前回授業で作成 ) global void mosaic(unsigned char *p, unsigned char *filtered, int mosaicsize){ int i,j, isub,jsub; int average; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; if(threadidx.x == 0 && threadidx.y == 0){// ブロック内の1スレッドのみが処理 // 領域内の画素の平均値を計算 average = 0; for(jsub = 0; jsub<mosaicsize; jsub++){ for(isub = 0; isub<mosaicsize; isub++){ average += p[(i+isub) + WIDTH*(j+jsub)]; average /= (mosaicsize*mosaicsize); // 領域内の画素を計算した平均値で塗りつぶす for(jsub = 0; jsub<mosaicsize; jsub++){ for(isub = 0; isub<mosaicsize; isub++){ filtered[(i+isub) + WIDTH*(j+jsub)] = (unsigned char)average; imageproc.cu 597

23 モザイク処理 ( 前回授業で作成 ) i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; if(threadidx.x == 0 && threadidx.y == 0){//1 スレッドのみが処理 // 領域内の画素の平均値を計算 // 計算した平均値をグローバルメモリへ書き込む ブロックに分割し, ブロック内の 1 スレッドのみが動作 ブロックの数だけ並列に処理ブロックの中では1スレッドのみが処理 グローバルメモリから画素情報を読込 ブロック内の画素の平均値を計算 グローバルメモリに画素の平均値を書込 残りのスレッドは i,j の計算をするだけ p[] 598

24 モザイク処理が高速化されない原因 各ブロックの 1 スレッドのみが処理を実行 グローバルメモリから画素情報を読込 ブロック内の画素の平均値を計算 グローバルメモリに画素の平均値を書込 グローバルメモリへコアレスアクセスできていない 複数のスレッドが協調し, アドレスが隣接したメモリを読むと高速 ある 1 スレッドが 1 画素ずつメモリアドレスを変えながら読むのは最悪の処理 599

25 モザイク処理が高速化されない原因 せめてコアレスアクセスしたい 平均を並列に計算するのは中級レベルの処理 カーネル内で変数を宣言するとが使われる 下の例ではaverageとcacheはに確保 ブロック内で値を共有できない global void mosaic(unsigned char *p, unsigned char *filtered, int mosaicsize){ int i,j, isub,jsub; int cache, average; //cache と average はに確保 ( 各スレッドが異なる値を保持 ) i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; cache = p[(i+isub) + WIDTH*(j+jsub)];// グローバルメモリから読み込み ( コアレスアクセス ) if(threadidx.x == 0 && threadidx.y == 0){ // 領域内の画素の平均値を計算したいが, 他のスレッドが持つ cache の値を読む事は不可能 // グローバルメモリに書き出し ( コアレスアクセス ) filtered[(i+isub) + WIDTH*(j+jsub)] = (unsigned char)average; 600

26 メモリの種類 オンチップメモリ (GPU のチップ内部に置かれたメモリ ) 高速アクセス, 小容量 CPUからはアクセス不可 L1キャッシュと共有メモリは一定サイズを共用 L1 キャッシュ / 共有 ( シェアード ) メモリ 容量小小 速度高速高速 GPU からの読み書き CPU からのアクセス 読み書き可ブロック内の全スレッドが同じアドレスにアクセス ( データを共有する ) ことが可能 * 読み書き不可 読み書き可各スレッドが異なるにアクセス 読み書き不可 * スレッドごとに異なるアドレスにアクセスすることも可能 601

27 共有 ( シェアード ) メモリ ブロック内のスレッドが共通のデータにアクセス可能 1 回目のアクセスに時間がかかるが, それ以降は非常に高速にアクセス可能 Fermi 世代以前の GPU で管理可能なキャッシュとして利用 1 ブロックあたり 16kB~48kB ホストメモリ GPU Chip SM L1 キャッシュ ローカルメモリ ローカルメモリ 共有メモリ SM L2 キャッシュ L1 キャッシュ コンスタントメモリ テクスチャメモリ グローバルメモリ 共有メモリ 602

28 共有 ( シェアード ) メモリの宣言 カーネル内で修飾子 shared を付けて宣言 配列として宣言 配列サイズを静的 ( コンパイル時 ) に決定する場合 shared 型変数名 [ 要素数 ] 多次元配列も宣言可能 配列サイズを動的 ( カーネル実行時 ) に決定する場合 extern shared 型変数名 [] サイズはカーネル呼出時のパラメータで指定 <<< ブロック数, スレッド数, 共有メモリのバイト数 >>> 603

29 共有メモリを使ったモザイク処理の高速化 ブロック内の全スレッドで共有したい値 各スレッドがグローバルメモリから読んだ画素情報 ブロック内の画素の平均値 処理の流れ 1. ブロック内の全スレッドがグローバルメモリから画素の値を読み出し, 共有メモリに置く ( コアレスアクセス ) 2. ある1スレッドが共有メモリに置かれた画素の値を読み, 画素の平均値を計算し, 共有メモリに置く 3. ブロック内の全スレッドが共有メモリに置かれた画素の平均値を読み, グローバルメモリに書き込む ( コアレスアクセス ) 604

30 global void mosaic_shared(unsigned char *p,unsigned char *filtered, int mosaicsize){ int i,j, isub,jsub; shared int average; shared unsigned char cache[threadx][thready]; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; モザイク処理 ( 共有メモリを利用 ) cache[threadidx.x][threadidx.y] = p[i + WIDTH*j]; if(threadidx.x == 0 && threadidx.y == 0){ average = 0; for(jsub = 0; jsub<mosaicsize; jsub++){ for(isub = 0; isub<mosaicsize; isub++){ average += cache[isub][jsub]; average /= (mosaicsize*mosaicsize); filtered[i + WIDTH*j] = (unsigned char)average; 605

31 共有メモリの宣言 shared int average; shared unsigned char cache[threadx][thready]; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; cache[threadidx.x][threadidx.y] = p[i + WIDTH*j]; ブロック内のスレッド数分の共有メモリを確保 画素情報を共有するのは 1 ブロックの中だけなので, 配列サイズはブロック内のスレッド数分でよい 各スレッドがグローバルメモリからデータを読み, 共有メモリに書込む p[] cache[][] 606

32 共有メモリの宣言 if(threadidx.x == 0 && threadidx.y == 0){ average = 0; for(isub = 0; isub<mosaicsize; isub++){ for(jsub = 0; jsub<mosaicsize; jsub++){ average += cache[isub][jsub]; average /= (mosaicsize*mosaicsize); 1スレッドが共有メモリに置かれた画素の値を読み, 画素の平均値を計算 画素の平均値 averageも共有メモリに存在 ブロック内の全スレッドが average にアクセス可能 cache[][] average 607

33 共有メモリの宣言 filtered[i + WIDTH*j] = (unsigned char)average; 各スレッドが average を読み込み グローバルメモリの各位置に画素の平均値を書き込む filtered[] average 608

34 実行結果 正しく処理できていない ブロック内のスレッドの協調が不十分 原画像 処理画像 609

35 モザイク処理 ( 共有メモリを利用 ) int i,j, isub,jsub; shared int average; shared unsigned char cache[threadx][thready]; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; スレット 0 スレット 0 以外 cache[threadidx.x][threadidx.y] = p[i + WIDTH*j]; if(threadidx.x == 0 && threadidx.y == 0){ average = 0; for(jsub = 0; jsub<mosaicsize; jsub++){ for(isub = 0; isub<mosaicsize; isub++){ average += cache[isub][jsub]; average /= (mosaicsize*mosaicsize); filtered[i + WIDTH*j] = (unsigned char)average; 他のスレッドが共有メモリに書き込む前に cache[][] にアクセスする可能性がある 610

36 モザイク処理 ( 共有メモリを利用 ) int i,j, isub,jsub; shared int average; shared unsigned char cache[threadx][thready]; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; スレット 0 スレット 0 以外 cache[threadidx.x][threadidx.y] = p[i + WIDTH*j]; if(threadidx.x == 0 && threadidx.y == 0){ average = 0; for(jsub = 0; jsub<mosaicsize; jsub++){ for(isub = 0; isub<mosaicsize; isub++){ average += cache[isub][jsub]; average /= (mosaicsize*mosaicsize); filtered[i + WIDTH*j] = (unsigned char)average; スレッド0 以外は平均値を計算しないので, 直ちにこの行に到達し, averageの値が確定する前に値を読んでfilteredに書き込む可能性がある 611

37 ブロック内でのスレッドの同期 syncthreads(); カーネル実行中にスレッドの同期を取る syncthreads() が書かれた行にスレッドが到達すると, 同一ブロック内の他の全てのスレッドがその行に達するまで待機 異なるブロック間での同期は不可能 if の中に記述するとカーネルが終了しないことがある if( 条件 ){ syncthreads(); // 条件が真にならないスレッドは if の中に入らないため, // カーネルが永久に終わらない 612

38 モザイク処理 ( 共有メモリを利用 ) int i,j, isub,jsub; shared int average; shared unsigned char cache[threadx][thready]; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; スレット 0 スレット 0 以外 cache[threadidx.x][threadidx.y] = p[i + WIDTH*j]; syncthreads(); if(threadidx.x == 0 && threadidx.y == 0){ average = 0; for(isub = 0; isub<mosaicsize; isub++){ for(jsub = 0; jsub<mosaicsize; jsub++){ average += cache[isub][jsub]; average /= (mosaicsize*mosaicsize); syncthreads(); filtered[i + WIDTH*j] = (unsigned char)average; スレッド 0 以外は if 文を実行しないが, スレッド 0 が到達する ( 平均値を計算し終わる ) まで syncthreads() で停止 if 文の前でブロック内の全スレッドが同期しているので, cache[][] には画素情報が入っている ブロック内で同期 ブロック内で同期 imageproc_mem.cu 613

39 実行結果 正しく処理できている 実行時間 260 ms(cpu) 38.5 ms 18.7 ms 処理時間が約 1/2 に短縮原画像 処理画像 614

40 空間フィルタ ある画素とその周囲の画素を使って処理 処理の仕方を規定したカーネルを定義 カーネルは 1 次元配列で表現 原画像 フィルタ ( カーネル ) 輪郭抽出 a b c d e f g h i = b+d 4e+f+h 615

41 空間フィルタ カーネルは 1 次元配列で表現 ぼかし ( 平均フィルタ ) 1/9 1/9 1/9 1/9 1/9 1/9 1/9 1/9 1/9 float blur[9] ={1.0f/9.0f,1.0f/9.0f,1.0f/9.0f, 1.0f/9.0f,1.0f/9.0f,1.0f/9.0f, 1.0f/9.0f,1.0f/9.0f,1.0f/9.0f; 輪郭抽出 float laplacian[9] ={ 0.0f, 1.0f, 0.0f, 1.0f, 4.0f, 1.0f, 0.0f, 1.0f, 0.0f; 616

42 空間フィルタ ( 前回授業で作成 ) global void boxfilter(unsigned char *p,unsigned char *filtered, float *filter){ int i,j; int result = BLACK; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; if(result<black) result = result; // 数値が負になっていれば 1 をかける if(result>white) result = WHITE; // 数値が 255 を超えていれば 255 に収める filtered[i+width*j] = (unsigned char)result; imageproc.cu 617

43 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 618 p[]

44 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 619 j p[] i

45 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 620 j p[] i

46 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 621 j p[] i

47 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 622 j p[] i

48 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 623 j p[] i

49 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 624 j p[] i

50 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 625 j p[] i

51 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 626 j p[] i

52 空間フィルタ ( 前回授業で作成 ) if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = filter[0]*p[(i 1) + WIDTH*(j 1)] +filter[1]*p[(i ) + WIDTH*(j 1)] +filter[2]*p[(i+1) + WIDTH*(j 1)] +filter[3]*p[(i 1) + WIDTH*(j )] +filter[4]*p[(i ) + WIDTH*(j )] +filter[5]*p[(i+1) + WIDTH*(j )] +filter[6]*p[(i 1) + WIDTH*(j+1)] +filter[7]*p[(i ) + WIDTH*(j+1)] +filter[8]*p[(i+1) + WIDTH*(j+1)]; メモリアクセス filter[] 画像 ( 配列 p) へは複数スレッドが隣接したメモリアドレスにアクセス コアレスアクセス可能 フィルタ ( 配列 filter) へは複数スレッドが同じ要素にアクセス コアレスアクセス不可能 L2キャッシュに入る可能性はある 627 j p[] i

53 空間フィルタ処理の高速化 フィルタ ( 配列 filter) へのアクセス コアレスアクセスできていない L2 キャッシュに入る可能性は高いが, 有効活用されているかは不明 配列 filter へのアクセスの高速化 共有メモリを利用すると処理が冗長 1 ブロックから 9 スレッドを選び, グローバルメモリから共有メモリへコピーし, 同期をとる コンスタントメモリが活用できる 628

54 メモリの種類 オフチップメモリ (GPU のチップ外部に置かれたメモリ ) 低速アクセス, 大容量 CPUから直接アクセス可能 ローカルメモリだけはアクセス不可 グローバルメモリ ローカルメモリ テクスチャメモリ コンスタントメモリ 容量 大 小 大 小 速度 低速 低速 高速 * 高速 * GPU からの読み書き 読み書き可全てのスレッドが同じアドレスにアクセス可能 ** 読み書き可各スレッドが異なるアドレスにアクセス 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** 読み込み可全てのスレッドが同じアドレスにアクセス可能 ** CPU からのアクセス 読み書き可読み書き不可読み書き可読み書き可 * キャッシュが効く場合 ** スレッドごとに異なるアドレスにアクセスすることも可能 629

55 コンスタントメモリ GPU 全体で同じメモリにアクセス GPU Chip SM SM メモリを読み取り専用とすることで値をキャッシュし, 一度読んだ値を再利用 GPU 全体で 64kB L1 キャッシュ ローカルメモリ 共有メモリ L2 キャッシュ L1 キャッシュ 共有メモリ ホストメモリ ローカルメモリ コンスタントメモリ テクスチャメモリ グローバルメモリ 630

56 コンスタントメモリの宣言 グローバル領域で修飾子 constant を付けて宣言 配列サイズは静的に決定 constant 型変数名 ; constant 型変数名 [ 要素数 ]; 配列としても宣言可能 サイズはコンパイル時に確定している必要がある cudamalloc() や cudafree() は不要 グローバル変数として宣言し, 複数のカーネルからアクセスすることが多い 読込専用なので許される 書込可能なメモリでは厳禁 631

57 コンスタントメモリの利用 メモリは読込専用 CPUからは変更可能 専用のメモリ転送命令でコピー cudamemcpytosymbol CPU 上のメモリをコンスタントメモリにコピーする cudamemcpytosymbol( 転送先変数名, 転送元アドレス, バイト数, オフセット, 方向 ); オフセット, 方向は省略可 632

58 コンスタントメモリへのアクセス コンスタントメモリへ高速にアクセスできる要因 コンスタントメモリはオフチップメモリ コンスタントメモリへのアクセス自体は高速ではない 1. データの配分 32 スレッド (Warp) 単位でアクセスし,1 回の読込を 32 スレッドで共有できる 2. キャッシュによる値の再利用 他の Warp がキャッシュされたデータへアクセスすることで, コンスタントメモリから直接読むよりも高速化 633

59 コンスタントメモリを使った空間フィルタの高速化 空間フィルタに用いるカーネル 1 次元の配列として宣言,GPU( グローバルメモリ ) へ転送 ( 端を除く ) 全スレッドからアクセス 値は固定値で,GPUから書き換えない コンスタントメモリを利用 1 次元の配列として宣言,GPU( コンスタントメモリ ) へ転送 全スレッドがコンスタントメモリにアクセス コンスタントキャッシュが有効利用される 634

60 メイン関数 ( コンスタントメモリの宣言と転送 ) :( 省略 ) : constant float cfilter[9];// コンスタントメモリにフィルタのカーネル分のメモリを確保 //main の外で宣言しているので, どの関数からもアクセス可能 int main(void){ :( 省略 ) float laplacian[9] ={ 0.0f, 1.0f, 0.0f, 1.0f, 4.0f, 1.0f, 0.0f, 1.0f, 0.0f; // グローバルメモリに確保していたフィルタのカーネルは不要 //float *filter; //cudamalloc( (void **)&filter, sizeof(float)*9); //cudamemcpy(filter, laplacian, sizeof(float)*9, cudamemcpyhosttodevice); // コンスタントメモリにフィルタのカーネルを送る cudamemcpytosymbol(cfilter, laplacian, sizeof(float)*9); :( 省略 ) boxfilter_constant<<<block,thread>>>(dev_p,dev_filtered); :( 省略 ) imageproc_mem.cu 635

61 空間フィルタ ( コンスタントメモリ利用 ) global void boxfilter_constant(unsigned char *p, unsigned char *filtered){ int i,j; int result = BLACK; i = blockidx.x*blockdim.x + threadidx.x; j = blockidx.y*blockdim.y + threadidx.y; if(0<i && i<width 1 && 0<j && j<height 1){ // 端の画素は処理をしないため,ifで処理を分岐 result = cfilter[0]*p[(i 1) + WIDTH*(j 1)] +cfilter[1]*p[(i ) + WIDTH*(j 1)] +cfilter[2]*p[(i+1) + WIDTH*(j 1)] +cfilter[3]*p[(i 1) + WIDTH*(j )] +cfilter[4]*p[(i ) + WIDTH*(j )] +cfilter[5]*p[(i+1) + WIDTH*(j )] +cfilter[6]*p[(i 1) + WIDTH*(j+1)] +cfilter[7]*p[(i ) + WIDTH*(j+1)] +cfilter[8]*p[(i+1) + WIDTH*(j+1)]; if(result<black) result = result; // 数値が負になっていれば 1 をかける if(result>white) result = WHITE; // 数値が 255 を超えていれば 255 に収める filtered[i+width*j] = (unsigned char)result; imageproc_mem.cu 636

62 実行結果 実行時間 553 ms(cpu) 4.13 ms 3.38 ms CPU からかなり高速化されていたが, さらに 2 割短縮 原画像 処理画像 637

63 レポート課題 4( 提出期限は 2 学期末 ) ガウシアンフィルタ ( ガウスぼかし ) を実装せよ 実行の条件 画像の形状は自身で定め, プログラム中で生成せよ 既存の画像を読む機能を実装できる場合は, 既存の画像を用いてよい ガウシアンフィルタのカーネルサイズは 5 5 とする 小さいサイズの画像を用いて 1. 原画像が正しく生成できている事を確認せよ 2. フィルタが正しくかけられている事を確認せよ 大きいサイズの画像を用いて 1. フィルタのカーネルをグローバルメモリから読んだ場合とコンスタントメモリから読んだ場合の実行時間の違いを比較せよ 2.1ブロックあたりのスレッド数を変更し, 実行時間がどのように変化するかを考察せよ 638

64 レポートの書式 必ず表紙を付けること 授業名, 課題番号, 学籍番号, 氏名, 提出日に加えて課題に要した時間を書く 課題内容, プログラム, 実行結果, 考察で構成 プログラムを実行した tesla?? および GPU の番号も明記すること 利用する GPU を cudasetdevice 命令で選択すること pdf 形式に変換してメールで提出 宛先 degawa at vos.nagaokaut.ac.jp メール題目 課題 4( 氏名 ) 639

Slide 1

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

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

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

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

More information

スライド 1

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

More information

RX ファミリ用 C/C++ コンパイラ V.1.00 Release 02 ご使用上のお願い RX ファミリ用 C/C++ コンパイラの使用上の注意事項 4 件を連絡します #pragma option 使用時の 1 または 2 バイトの整数型の関数戻り値に関する注意事項 (RXC#012) 共用

RX ファミリ用 C/C++ コンパイラ V.1.00 Release 02 ご使用上のお願い RX ファミリ用 C/C++ コンパイラの使用上の注意事項 4 件を連絡します #pragma option 使用時の 1 または 2 バイトの整数型の関数戻り値に関する注意事項 (RXC#012) 共用 RX ファミリ用 C/C++ コンパイラ V.1.00 Release 02 ご使用上のお願い RX ファミリ用 C/C++ コンパイラの使用上の注意事項 4 件を連絡します #pragma option 使用時の 1 または 2 バイトの整数型の関数戻り値に関する注意事項 (RXC#012) 共用体型のローカル変数を文字列操作関数で操作する場合の注意事項 (RXC#013) 配列型構造体または共用体の配列型メンバから読み出した値を動的初期化に用いる場合の注意事項

More information

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

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

More information

スライド 1

スライド 1 知能制御システム学 画像処理の高速化 東北大学大学院情報科学研究科鏡慎吾 swk(at)ic.is.tohoku.ac.jp 2008.07.22 今日の内容 ビジュアルサーボのようなリアルタイム応用を考える場合, 画像処理を高速に実装することも重要となる いくつかの基本的な知識を押さえておかないと, 同じアルゴリズムを実行しているのに性能が上がらないということがしばしば生じる 今日は, あくまで普通の

More information

Microsoft PowerPoint - OS07.pptx

Microsoft PowerPoint - OS07.pptx この資料は 情報工学レクチャーシリーズ松尾啓志著 ( 森北出版株式会社 ) を用いて授業を行うために 名古屋工業大学松尾啓志 津邑公暁が作成しました 主記憶管理 主記憶管理基礎 パワーポイント 27 で最終版として保存しているため 変更はできませんが 授業でお使いなる場合は松尾 (matsuo@nitech.ac.jp) まで連絡いただければ 編集可能なバージョンをお渡しする事も可能です 復習 OS

More information

Microsoft PowerPoint - 09.pptx

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

More information

プログラミング実習I

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

More information

概要 プログラミング論 変数のスコープ, 記憶クラス. メモリ動的確保. 変数のスコープ 重要. おそらく簡単. 記憶クラス 自動変数 (auto) と静的変数 (static). スコープほどではないが重要.

概要 プログラミング論 変数のスコープ, 記憶クラス. メモリ動的確保. 変数のスコープ 重要. おそらく簡単. 記憶クラス 自動変数 (auto) と静的変数 (static). スコープほどではないが重要. 概要 プログラミング論 変数のスコープ, 記憶クラス. メモリ動的確保. 変数のスコープ 重要. おそらく簡単. 記憶クラス 自動変数 (auto) と静的変数 (static). スコープほどではないが重要. http://www.ns.kogakuin.ac.jp/~ct13140/progc/ C-2 ブロック 変数のスコープ C 言語では, から をブロックという. for( ) if( )

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

熊本大学学術リポジトリ 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

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

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

More information

memo

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

More information

ex04_2012.ppt

ex04_2012.ppt 2012 年度計算機システム演習第 4 回 2012.05.07 第 2 回課題の補足 } TSUBAMEへのログイン } TSUBAMEは学内からのログインはパスワードで可能 } } } } しかし 演習室ではパスワードでログインできない設定 } 公開鍵認証でログイン 公開鍵, 秘密鍵の生成 } ターミナルを開く } $ ssh-keygen } Enter file in which to save

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

Microsoft PowerPoint - kougi7.ppt

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

More information

7 ポインタ (P.61) ポインタを使うと, メモリ上のデータを直接操作することができる. 例えばデータの変更 やコピーなどが簡単にできる. また処理が高速になる. 7.1 ポインタの概念 変数を次のように宣言すると, int num; メモリにその領域が確保される. 仮にその開始のアドレスを 1

7 ポインタ (P.61) ポインタを使うと, メモリ上のデータを直接操作することができる. 例えばデータの変更 やコピーなどが簡単にできる. また処理が高速になる. 7.1 ポインタの概念 変数を次のように宣言すると, int num; メモリにその領域が確保される. 仮にその開始のアドレスを 1 7 ポインタ (P.61) ポインタを使うと, メモリ上のデータを直接操作することができる. 例えばデータの変更 やコピーなどが簡単にできる. また処理が高速になる. 7.1 ポインタの概念 変数を次のように宣言すると, int num; メモリにその領域が確保される. 仮にその開始のアドレスを 10001 番地とすると, そこから int 型のサイズ, つまり 4 バイト分の領域が確保される.1

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

gengo1-8

gengo1-8 問題提起その 1 一文字ずつ文字 ( 数字 ) を読み込み それぞれの文字が何回入力されたかを数えて出力するプログラム int code, count_0=0, count_1=0, count_2=0, count_3=0,..., count_9=0; while( (code=getchar())!= EOF ){ } switch(code){ case 0 : count_0++; break;

More information

Microsoft PowerPoint - ARCEMB08HayashiSlides.ppt [互換モード]

Microsoft PowerPoint - ARCEMB08HayashiSlides.ppt [互換モード] 演算 / メモリ性能バランスを考慮した CMP 向けオンチップ メモリ貸与法の提案 九州大学 林徹生今里賢一井上弘士村上和彰 1 発表手順 背景 目的 演算 / メモリ性能バランシング 概要 アクセスレイテンシの削減とオーバーヘッド 提案手法の実現方法 着目する命令 (Cell プロセッサへの ) 実装 性能評価 姫野ベンチマーク Susan@MiBench おわりに 2 チップマルチプロセッサ (CMP)

More information

N 体問題 長岡技術科学大学電気電子情報工学専攻出川智啓

N 体問題 長岡技術科学大学電気電子情報工学専攻出川智啓 N 体問題 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 天体の運動方程式 天体運動の GPU 実装 最適化による性能変化 #pragma unroll 855 計算の種類 画像処理, 差分法 空間に固定された観測点を配置 観測点 ( 固定 ) 観測点上で物理量がどのように変化するかを追跡 Euler 型 多粒子の運動 観測点を配置せず, 観測点が粒子と共に移動 Lagrange 型 観測点

More information

第1回 プログラミング演習3 センサーアプリケーション

第1回 プログラミング演習3 センサーアプリケーション C プログラミング - ポインタなんて恐くない! - 藤田悟 fujita_s@hosei.ac.jp 目標 C 言語プログラムとメモリ ポインタの関係を深く理解する C 言語プログラムは メモリを素のまま利用できます これが原因のエラーが多く発生します メモリマップをよく頭にいれて ポインタの動きを理解できれば C 言語もこわくありません 1. ポインタ入門編 ディレクトリの作成と移動 mkdir

More information

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

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

More information

Microsoft PowerPoint - 計算機言語 第7回.ppt

Microsoft PowerPoint - 計算機言語 第7回.ppt 計算機言語第 7 回 長宗高樹 目的 関数について理解する. 入力 X 関数 f 出力 Y Y=f(X) 関数の例 関数の型 #include int tasu(int a, int b); main(void) int x1, x2, y; x1 = 2; x2 = 3; y = tasu(x1,x2); 実引数 printf( %d + %d = %d, x1, x2, y);

More information

PowerPoint Presentation

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

More information

memo

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

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 - os ppt [互換モード]

Microsoft PowerPoint - os ppt [互換モード] 4. メモリ管理 (1) 概要メモリ管理の必要性静的メモリ管理と動的メモリ管理スワッピング, 仮想記憶ページングとセグメンテーション 2008/5/ 20 メモリ管理 (1) 1 メモリはコンピュータの 5 大構成要素 装置 ( キーボード, マウス ) CPU ( 中央演算装置 ) 出 装置 ( モニタ, プリンタ ) 主記憶装置 ( メインメモリ ) 外部記憶装置 (HDD) 2008/5/ 20

More information

本書は INpMac v2.20(intime 5.2 INplc 3 Windows7/8/8.1に対応 ) の内容を元に記載しています Microsoft Windows Visual Studio は 米国 Microsoft Corporation の米国及びその他の国における登録商標です

本書は INpMac v2.20(intime 5.2 INplc 3 Windows7/8/8.1に対応 ) の内容を元に記載しています Microsoft Windows Visual Studio は 米国 Microsoft Corporation の米国及びその他の国における登録商標です ACTIVE TOUCH 拡張部品取扱説明書 - 共有メモリアクセスコンポーネント - 1. はじめに 1 (1) 概要... 1 (2) INpMac のインストール... 1 2. Windows アプリケーションとの連携 2 (1) コントロール ( 部品 ) の登録... 2 (2) データの関連付け... 3 3. INtime アプリケーションとの連携 4 (1) 部品 ( コンポーネント

More information

出 アーキテクチャ 誰が 出 装置を制御するのか 1

出 アーキテクチャ 誰が 出 装置を制御するのか 1 出 アーキテクチャ 誰が 出 装置を制御するのか 1 が 出 装置を制御する メモリ ( 主記憶 ) 命令データ 出 装置 2 が 出 装置を制御する 命令 実 入出力装置を制御する命令を実行する メモリ ( 主記憶 ) 命令データ 制御 出 装置 3 が 出 装置を制御する メモリ ( 主記憶 ) 命令 実 制御 命令データ データを出力せよ 出 装置 4 が 出 装置を制御する メモリ ( 主記憶

More information

Microsoft PowerPoint - OpenMP入門.pptx

Microsoft PowerPoint - OpenMP入門.pptx OpenMP 入門 須田礼仁 2009/10/30 初版 OpenMP 共有メモリ並列処理の標準化 API http://openmp.org/ 最新版は 30 3.0 バージョンによる違いはあまり大きくない サポートしているバージョンはともかく csp で動きます gcc も対応しています やっぱり SPMD Single Program Multiple Data プログラム #pragma omp

More information

GPU CUDA CUDA 2010/06/28 1

GPU CUDA CUDA 2010/06/28 1 GPU CUDA CUDA 2010/06/28 1 GPU NVIDIA Mark Harris, Optimizing Parallel Reduction in CUDA http://developer.download.nvidia.com/ compute/cuda/1_1/website/data- Parallel_Algorithms.html#reduction CUDA SDK

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

Prog1_6th

Prog1_6th 2012 年 5 月 24 日 ( 木 ) 実施 多分岐のプログラム 前回は多段階の 2 分岐を組み合わせて 3 種類以上の場合分けを実現したが, 式の値の評価によって, 一度に多種類の場合分けを行う多分岐の利用によって見通しのよいプログラムを作成できる場合がある ( 流れ図は右図 ) 式の評価 : 値 1 : 値 2 : 値 n : 該当値無し 処理 1 処理 2 処理 n 既定の処理 switch

More information

-2 外からみたプロセッサ GND VCC CLK A0 A1 A2 A3 A4 A A6 A7 A8 A9 A10 A11 A12 A13 A14 A1 A16 A17 A18 A19 D0 D1 D2 D3 D4 D D6 D7 D8 D9 D10 D11 D12 D13 D14 D1 MEMR

-2 外からみたプロセッサ GND VCC CLK A0 A1 A2 A3 A4 A A6 A7 A8 A9 A10 A11 A12 A13 A14 A1 A16 A17 A18 A19 D0 D1 D2 D3 D4 D D6 D7 D8 D9 D10 D11 D12 D13 D14 D1 MEMR 第 回マイクロプロセッサのしくみ マイクロプロセッサの基本的なしくみについて解説する. -1 マイクロプロセッサと周辺回路の接続 制御バス プロセッサ データ バス アドレス バス メモリ 周辺インタフェース バスの基本構成 Fig.-1 バスによる相互接続は, 現在のコンピュータシステムのハードウェアを特徴づけている. バス (Bus): 複数のユニットで共有される信号線システム内の データの通り道

More information

Prog1_10th

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

More information

(1) プログラムの開始場所はいつでも main( ) メソッドから始まる 順番に実行され add( a,b) が実行される これは メソッドを呼び出す ともいう (2)add( ) メソッドに実行が移る この際 add( ) メソッド呼び出し時の a と b の値がそれぞれ add( ) メソッド

(1) プログラムの開始場所はいつでも main( ) メソッドから始まる 順番に実行され add( a,b) が実行される これは メソッドを呼び出す ともいう (2)add( ) メソッドに実行が移る この際 add( ) メソッド呼び出し時の a と b の値がそれぞれ add( ) メソッド メソッド ( 教科書第 7 章 p.221~p.239) ここまでには文字列を表示する System.out.print() やキーボードから整数を入力する stdin.nextint() などを用いてプログラムを作成してきた これらはメソッドと呼ばれるプログラムを構成する部品である メソッドとは Java や C++ などのオブジェクト指向プログラミング言語で利用されている概念であり 他の言語での関数やサブルーチンに相当するが

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

スライド 1

スライド 1 東北大学工学部機械知能 航空工学科 2015 年度 5 セメスター クラス D 計算機工学 6. MIPS の命令と動作 演算 ロード ストア ( 教科書 6.3 節,6.4 節 ) 大学院情報科学研究科鏡慎吾 http://www.ic.is.tohoku.ac.jp/~swk/lecture/ レジスタ間の演算命令 (C 言語 ) c = a + b; ( 疑似的な MIPS アセンブリ言語 )

More information

memo

memo 計数工学プログラミング演習 ( 第 3 回 ) 2017/04/25 DEPARTMENT OF MATHEMATICAL INFORMATICS 1 内容 ポインタの続き 引数の値渡しと参照渡し 構造体 2 ポインタで指されるメモリへのアクセス double **R; 型 R[i] と *(R+i) は同じ意味 意味 R double ** ポインタの配列 ( の先頭 ) へのポインタ R[i]

More information

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

PowerPoint プレゼンテーション

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

More information

Microsoft PowerPoint - 11Web.pptx

Microsoft PowerPoint - 11Web.pptx 計算機システムの基礎 ( 第 10 回配布 ) 第 7 章 2 節コンピュータの性能の推移 (1) コンピュータの歴史 (2) コンピュータの性能 (3) 集積回路の進歩 (4) アーキテクチャ 第 4 章プロセッサ (1) プロセッサの基本機能 (2) プロセッサの構成回路 (3) コンピュータアーキテクチャ 第 5 章メモリアーキテクチャ 1. コンピュータの世代 計算する機械 解析機関 by

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

Microsoft PowerPoint - sp ppt [互換モード]

Microsoft PowerPoint - sp ppt [互換モード] システムプログラム概論 メモリ管理 (1) 第 x 講 : 平成 20 年 10 月 15 日 ( 水 ) 2 限 S1 教室 今日の講義概要 メモリ管理の必要性 静的メモリ管理と動的メモリ管理 スワッピング, 仮想記憶 ページングとセグメンテーション 中村嘉隆 ( なかむらよしたか ) 奈良先端科学技術大学院大学助教 y-nakamr@is.naist.jp http://narayama.naist.jp/~y-nakamr/

More information

Microsoft PowerPoint - ca ppt [互換モード]

Microsoft PowerPoint - ca ppt [互換モード] 大阪電気通信大学情報通信工学部光システム工学科 2 年次配当科目 コンピュータアルゴリズム 良いアルゴリズムとは 第 2 講 : 平成 20 年 10 月 10 日 ( 金 ) 4 限 E252 教室 中村嘉隆 ( なかむらよしたか ) 奈良先端科学技術大学院大学助教 y-nakamr@is.naist.jp http://narayama.naist.jp/~y-nakamr/ 第 1 講の復習

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

(2) 構造体変数の宣言 文法は次のとおり. struct 構造体タグ名構造体変数名 ; (1) と (2) は同時に行える. struct 構造体タグ名 { データ型変数 1; データ型変数 2;... 構造体変数名 ; 例 : struct STUDENT{ stdata; int id; do

(2) 構造体変数の宣言 文法は次のとおり. struct 構造体タグ名構造体変数名 ; (1) と (2) は同時に行える. struct 構造体タグ名 { データ型変数 1; データ型変数 2;... 構造体変数名 ; 例 : struct STUDENT{ stdata; int id; do 8 構造体と供用体 ( 教科書 P.71) 構造体は様々なデータ型,int 型,float 型や char 型などが混在したデータを一つのまとまり, 単位として扱える.( 配列は一つのデータ型しか扱えない.) 構造体は柔軟なデータ構造を扱えるので, プログラムを効率よく開発できる. つまり構造体を使用すると, コード量を抑え, バグを少なくし, 開発時間を短くし, 簡潔なプログラムが作れる. 共用体は,

More information

情報処理演習 B8クラス

情報処理演習 B8クラス 予定スケジュール ( 全 15 回 ) 1 1. 終了 プログラミング言語の基礎 2. 終了 演算と型 3. 終了 プログラムの流れの分岐 (if 文,switch 文など ) 4. 終了 プログラムの流れの繰返し (do, while, for 文など ) 5. 終了 中間レポート1 6. 終了 配列 7. 終了 関数 8. 終了 文字列 ( 文字列の配列, 文字列の操作 ) 9. 終了 ポインタ

More information

Microsoft PowerPoint - kougi6.ppt

Microsoft PowerPoint - kougi6.ppt C プログラミング演習 第 6 回ファイル処理と配列 1 ファイル処理 2 ファイル読み込み ファイル プログラム ファイルの中身は変わらない 3 ファイル書き出し ファイル プログラム ファイルの中身が変わる ファイルは伸び縮みすることがある 4 例題 1. テキストファイル形式の ファイルからのデータ読み込み 次のような名簿ファイル ( テキストファイル形式 ) を読み込んで,1 列目の氏名と,3

More information

計算機プログラミング

計算機プログラミング プログラミング言語 C 第 8 講 システム標準関数 ( 入出力関数 ) システム標準関数 システムに備え付けの関数 例 ) printf( ); scanf( ); 標準出力関数 標準入力関数 A. 入出力用の関数 高水準入出力関数 高水準言語 (OS に依存しない ) 低水準入出力関数 機械語レベル (OS に依存 ) B. それ以外の関数 引数と関数の型 ( 戻り値 ) に注目しましょう 例

More information

program7app.ppt

program7app.ppt プログラム理論と言語第 7 回 ポインタと配列, 高階関数, まとめ 有村博紀 吉岡真治 公開スライド PDF( 情報知識ネットワーク研 HP/ 授業 ) http://www-ikn.ist.hokudai.ac.jp/~arim/pub/proriron/ 本スライドは,2015 北海道大学吉岡真治 プログラム理論と言語, に基づいて, 現著者の承諾のもとに, 改訂者 ( 有村 ) が加筆修正しています.

More information

本文ALL.indd

本文ALL.indd Intel Xeon プロセッサにおける Cache Coherency 時間の性能測定方法河辺峻田口成美古谷英祐 Intel Xeon プロセッサにおける Cache Coherency 時間の性能測定方法 Performance Measurement Method of Cache Coherency Effects on an Intel Xeon Processor System 河辺峻田口成美古谷英祐

More information

02: 変数と標準入出力

02: 変数と標準入出力 C プログラミング入門 基幹 2 ( 月 4) 09: ポインタ 文字列 Linux にログインし 以下の講義ページを開いておくこと http://www-it.sci.waseda.ac.jp/ teachers/w483692/cpr1/ 2014-06-09 1 関数できなかったこと 配列を引数として渡す, 戻り値として返す 文字列を扱う 呼び出し元の変数を直接書き換える 例 : 2 つの変数の値を入れ替える関数

More information

Slide 1

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

More information

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

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

More information

プログラミング実習I

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

More information

GR-SAKURA-SAのサンプルソフト説明

GR-SAKURA-SAのサンプルソフト説明 フルカラーシリアル LED テープ (1m) を GR-KURUMI で使ってみる 2014/2/25 がじぇっとるねさす鈴木 Rev. 1.00 フルカラーシリアル LED の特徴 http://www.switch-science.com/catalog/1399/ 3570 円 1m で 60 個の LED がついている 電源と信号線 1 本で制御する x 24 この信号を 24 個送信して

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

Microsoft Word - Cプログラミング演習(10)

Microsoft Word - Cプログラミング演習(10) 第 10 回 (6/25) 3. ファイルとその応用 (3) ファイルの更新 シーケンシャルファイルの更新 シーケンシャルファイルでは, 各レコードが可変長で連続して格納されており, その中の特定のレコードを変更することができない そこで一般的には, マスタファイルからデータを取り出し, 更新処理を行ったあとに新マスタファイルに書き込む 注 ) マスタファイル : 主ファイル, 基本ファイルと呼ばれるファイルで内容は比較的固定的であり,

More information

Microsoft PowerPoint - ARC2009HashiguchiSlides.pptx

Microsoft PowerPoint - ARC2009HashiguchiSlides.pptx 3 次元 DRAM プロセッサ積層実装を 対象としたオンチップ メモリ アーキテクチャの提案と評価 橋口慎哉 小野貴継 ( 現 ) 井上弘士 村上和彰 九州大学大学院システム情報科学府 九州大学大学院システム情報科学研究院 発表手順 研究背景 研究目的 ハイブリッド キャッシュ アーキテクチャ 評価実験 まとめと今後の課題 2 3 次元実装技術 研究背景 グローバル配線長の削減 チップ面積縮小 異なるプロセスを経て製造されたダイ同士の積層

More information

PowerPoint プレゼンテーション

PowerPoint プレゼンテーション 部内向けスキルアップ研修 組込み OS 自作入門 2014 年 2 月 10st ステップ担当 : 中村 目次 はじめに OSの役割 メモリ管理 メモリ管理実装 プログラムの実行 まとめ はじめに 前回やったこと OS の原型を作成 今回やること 9th ステップでは CPU 時間 という資源管理 本ステップでは メモリ という資源管理 10.1 OS の役割 10.1.1 コンピュータの 3 大要素

More information

MMUなしプロセッサ用Linuxの共有ライブラリ機構

MMUなしプロセッサ用Linuxの共有ライブラリ機構 MMU なしプロセッサ用 Linux の共有ライブラリ機構 大谷浩司 高岡正 近藤政雄 臼田尚志株式会社アックス はじめに μclinux には 仮想メモリ機構がないので共有ライブラリ機構が使えない でもメモリ消費抑制 ストレージ消費抑制 保守性の向上のためには 欲しい 幾つかの実装があるが CPU ライセンス 機能の制限のためにそのまま利用できない RidgeRun 社 (Cadenux 社 )

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

スライド 1

スライド 1 プログラミング 第 3 週 静岡大学工学部機械工学科知能 材料コースロボット 計測情報分野臼杵深光電 精密コース光ナノバイオ分野居波渉 講義の前に 講義資料や演習課題 LiveCampusよりダウンロード可能 成績評価期末試験および課題により行う. 評価の配分は, おおむね試験 90%, 課題 10% である. 再試験期末試験で40 点以上 60 点未満の場合, 再試験となる. 2 月 26 日 (

More information

Microsoft PowerPoint - CproNt02.ppt [互換モード]

Microsoft PowerPoint - CproNt02.ppt [互換モード] 第 2 章 C プログラムの書き方 CPro:02-01 概要 C プログラムの構成要素は関数 ( プログラム = 関数の集まり ) 関数は, ヘッダと本体からなる 使用する関数は, プログラムの先頭 ( 厳密には, 使用場所より前 ) で型宣言 ( プロトタイプ宣言 ) する 関数は仮引数を用いることができる ( なくてもよい ) 関数には戻り値がある ( なくてもよい void 型 ) コメント

More information

/*Source.cpp*/ #include<stdio.h> //printf はここでインクルードして初めて使えるようになる // ここで関数 average を定義 3 つの整数の平均値を返す double 型の関数です double average(int a,int b,int c){

/*Source.cpp*/ #include<stdio.h> //printf はここでインクルードして初めて使えるようになる // ここで関数 average を定義 3 つの整数の平均値を返す double 型の関数です double average(int a,int b,int c){ ソフトゼミ A 第 6 回 関数 プログラムは関数の組み合わせでできています 今までのゼミAでも printf や scanf など様々な関数を使ってきましたが なんと関数は自分で作ることもできるのです!! 今日は自作関数を中心に扱っていきます ゲーム制作でも自作関数は避けては通れないので頑張りましょう そもそもまず 関数とは 基本的には 受け取った値に関数によって定められた操作をして その結果の値を返す

More information

arduino プログラミング課題集 ( Ver /06/01 ) arduino と各種ボードを組み合わせ 制御するためのプログラミングを学 ぼう! 1 入出力ポートの設定と利用方法 (1) 制御( コントロール ) する とは 外部装置( ペリフェラル ) が必要とする信号をマイ

arduino プログラミング課題集 ( Ver /06/01 ) arduino と各種ボードを組み合わせ 制御するためのプログラミングを学 ぼう! 1 入出力ポートの設定と利用方法 (1) 制御( コントロール ) する とは 外部装置( ペリフェラル ) が必要とする信号をマイ arduino プログラミング課題集 ( Ver.5.0 2017/06/01 ) arduino と各種ボードを組み合わせ 制御するためのプログラミングを学 ぼう! 1 入出力ポートの設定と利用方法 (1) 制御( コントロール ) する とは 外部装置( ペリフェラル ) が必要とする信号をマイコンから伝える 外部装置の状態をマイコンで確認する 信号の授受は 入出力ポート 経由で行う (2) 入出力ポートとは?

More information

Microsoft PowerPoint - Lec24 [互換モード]

Microsoft PowerPoint - Lec24 [互換モード] 第 11 回講義水曜日 1 限教室 618 情報デザイン専攻 画像情報処理論及び演習 II - 動画像処理 - Video Styliztion 吉澤信 shin@riken.jp, 非常勤講師 大妻女子大学社会情報学部 今日の授業内容 www.riken.jp/brict/yoshizw/lectures/index.html www.riken.jp/brict/yoshizw/lectures/lec4.pdf

More information

情報処理Ⅰ

情報処理Ⅰ Java フローチャート -1- フローチャート ( 流れ図 ) プログラムの処理手順 ( アルゴリズム ) を図示したもの 記号の種類は下記のとおり 端子記号 ( 開始 終了 ) 処理記号計算, 代入等 条件の判定 条件 No ループ処理 LOOP start Yes データの入力 出力 print など 定義済み処理処理名 end サンプルグログラム ( 大文字 小文字変換 ) 大文字を入力して下さい

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

講習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

Microsoft PowerPoint - 5Chap15.ppt

Microsoft PowerPoint - 5Chap15.ppt 第 15 章文字列処理 今日のポイント 15.1 文字列処理の基本 strcpy strcat strlen strchr などの使い方をマスターする strcpy はなんて読むの? 普通はストリングコピー C のキーワードの読み方に悩んだら下記サイトを参考 ( 前回紹介とは別サイト ) http://www.okakogi.go.jp/people/miwa/program/c_lang/c_furoku.html

More information

Microsoft PowerPoint - 第3回目.ppt [互換モード]

Microsoft PowerPoint - 第3回目.ppt [互換モード] 第 3 回プログラミング応用 目的ファイル入出力 1. ファイルの概念 2. ファイルの読み込み 3. ファイルの書き込み CPU 演算 判断 ファイルの概念 内部記憶装置 OS 機械語プログラム 入力装置 キーボード 出力装置 ディスプレイ ファイル 外部記憶装置ハードディスク CD-ROM CPU が外部とデータをやり取りするための媒介 printf 関数や scanf 関数でもうすでにファイルのやり取りの基本は学んでいる

More information

JavaプログラミングⅠ

JavaプログラミングⅠ Java プログラミング Ⅰ 12 回目クラス 今日の講義で学ぶ内容 クラスとは クラスの宣言と利用 クラスの応用 クラス クラスとは 異なる複数の型の変数を内部にもつ型です 直観的に表現すると int 型や double 型は 1 1 つの値を管理できます int 型の変数 配列型は 2 5 8 6 3 7 同じ型の複数の変数を管理できます 配列型の変数 ( 配列変数 ) クラスは double

More information

ガイダンス

ガイダンス 情報科学 B 第 2 回変数 1 今日やること Java プログラムの書き方 変数とは何か? 2 Java プログラムの書き方 3 作業手順 Java 言語を用いてソースコードを記述する (Cpad エディタを使用 ) コンパイル (Cpad エディタを使用 ) 実行 (Cpad エディタを使用 ) エラーが出たらどうしたらよいか??? 4 書き方 これから作成する Hello.java 命令文 メソッドブロック

More information

Microsoft PowerPoint - ARC-SWoPP2011OkaSlides.pptx

Microsoft PowerPoint - ARC-SWoPP2011OkaSlides.pptx データ値の局所性を利用した ライン共有キャッシュの提案 九州大学大学院 岡慶太郎 福本尚人 井上弘士 村上和彰 1 キャッシュメモリの大容量化 マルチコア プロセッサが主流 メモリウォール問題の深刻化 メモリアクセス要求増加 IOピンの制限 大容量の LL(Last Level) キャッシュを搭載 8MB の L3 キャッシュを搭載 Core i7 のチップ写真 * * http://www.atmarkit.co.jp/fsys/zunouhoudan/102zunou/corei7.html

More information

Microsoft PowerPoint - suda.pptx

Microsoft PowerPoint - suda.pptx GPU の HWアーキテクチャと高性能化手法 須田礼仁 ( 東京大学 ) 2011/03/22 GPU 高性能プログラミング GPU のハードウェアを理解する CUDA のソフトウェアを理解する CUDA でプログラムを書くのは難しくないが, CUDA で高速なプログラムを書くのは難しい どうすれば遅くなるかを理解する! 効果が大きいものから順に説明します 1 高性能プログラミングの手順 1. 現在のコードの,

More information

今回のプログラミングの課題 ( 前回の課題で取り上げた )data.txt の要素をソートして sorted.txt というファイルに書出す ソート (sort) とは : 数の場合 小さいものから大きなもの ( 昇順 ) もしくは 大きなものから小さなもの ( 降順 ) になるよう 並び替えること

今回のプログラミングの課題 ( 前回の課題で取り上げた )data.txt の要素をソートして sorted.txt というファイルに書出す ソート (sort) とは : 数の場合 小さいものから大きなもの ( 昇順 ) もしくは 大きなものから小さなもの ( 降順 ) になるよう 並び替えること C プログラミング演習 1( 再 ) 4 講義では C プログラミングの基本を学び 演習では やや実践的なプログラミングを通して学ぶ 今回のプログラミングの課題 ( 前回の課題で取り上げた )data.txt の要素をソートして sorted.txt というファイルに書出す ソート (sort) とは : 数の場合 小さいものから大きなもの ( 昇順 ) もしくは 大きなものから小さなもの ( 降順

More information

演算増幅器

演算増幅器 構造体 ここまでに char int doulbe などの基本的なデータ型に加えて 同じデータ型が連続している 配列についてのデータ構造について習った これ以外にも もっと複雑なデータ型をユーザが定義 することが可能である それが構造体と呼ばれるもので 異なる型のデータをひとかたまりのデー タとして扱うことができる 異なるデータをまとめて扱いたい時とはどんな場合だろうか 例えば 住民データを管理したい

More information

2006年10月5日(木)実施

2006年10月5日(木)実施 2010 年 7 月 2 日 ( 金 ) 実施 ファイル処理ファイルとはファイル (file) は日常用語では紙などを綴じたものを表すが, コンピュータ用語ではデータの集合体を指す言葉である ファイルは例えば, 文書ファイルやプログラムファイルのように, 用途によって分類されることもあれば, また, テキストファイルやバイナリファイルのように, ファイルの作り方によって分類されることもある なお,

More information

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

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

More information

Microsoft PowerPoint - 14th.ppt [互換モード]

Microsoft PowerPoint - 14th.ppt [互換モード] 工学部 6 7 8 9 10 組 ( 奇数学籍番号 ) 担当 : 長谷川英之 情報処理演習 第 14 回 2011 年 1 月 20 日 1 今日のテーマ ファイル入出力 ですが, キーボード入力などもおさらいします 2 標準入力 キーボードで入力 : 標準入力という例 )scanf( %d,&i) 前回までの講義でファイルからデータを読み込む場合に使用した関数 : fscanf 例 )fscanf(fin,

More information

Microsoft PowerPoint - No6note.ppt

Microsoft PowerPoint - No6note.ppt 前回 : 管理 管理の目的 : の効率的利用 ( 固定区画方式 可変区画方式 ) しかし, いかに効率よく使ったとしても, 実行可能なプログラムサイズや同時に実行できるプロセス数は実装されているの大きさ ( 容量 ) に制限される 256kB の上で,28kB のプロセスを同時に 4 個実行させることはできないか? 2 256kB の上で,52kB のプロセスを実行させることはできないか? 方策 :

More information

Microsoft PowerPoint - 12.ppt [互換モード]

Microsoft PowerPoint - 12.ppt [互換モード] 第 12 回構造体 1 今回の目標 構造体を理解する 構造体の定義の仕方を理解する 構造体型を理解する 構造体型の変数 引数 戻り値を理解する 複素数同士を足し算する関数を作成し その関数を利用するプログラムを作成する 2 複素数の足し算 複素数は実部と虚部の2つの実数で 表現される 表現される z = a+ bi 2 つの複素数 z 1 = a 1+ bi 1 と z2 = a2 + b2i の和

More information

プログラミング基礎

プログラミング基礎 C プログラミング Ⅰ 授業ガイダンス C 言語の概要プログラム作成 実行方法 授業内容について 授業目的 C 言語によるプログラミングの基礎を学ぶこと 学習内容 C 言語の基礎的な文法 入出力, 変数, 演算, 条件分岐, 繰り返し, 配列,( 関数 ) C 言語による簡単な計算処理プログラムの開発 到達目標 C 言語の基礎的な文法を理解する 簡単な計算処理プログラムを作成できるようにする 授業ガイダンス

More information

Microsoft PowerPoint - 11.pptx

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

More information

SuperH RISC engineファミリ用 C/C++コンパイラパッケージ V.7~V.9 ご使用上のお願い

SuperH RISC engineファミリ用 C/C++コンパイラパッケージ V.7~V.9 ご使用上のお願い ツールニュース RENESAS TOOL NEWS 2014 年 02 月 01 日 : 140201/tn1 SuperH RISC engine ファミリ用 C/C++ コンパイラパッケージ V.7~V.9 ご使用上のお願い SuperH RISC engine ファミリ用 C/C++ コンパイラパッケージ V.7~V.9の使用上の注意事項 4 件を連絡します 同一ループ内の異なる配列要素に 同一の添え字を使用した場合の注意事項

More information

Microsoft PowerPoint - prog04.ppt

Microsoft PowerPoint - prog04.ppt プログラミング言語 2 第 04 回 (2007 年 05 月 14 日 ) 今日の配布物 片面の用紙 1 枚 今日の課題が書かれています 本日の出欠を兼ねています 1 今日やること http://www.tnlab.ice.uec.ac.jp/~s-okubo/class/language/ にアクセスすると 教材があります 2007 年 05 月 14 日分と書いてある部分が 本日の教材です 本日の内容

More information

Microsoft Word - Cプログラミング演習(9)

Microsoft Word - Cプログラミング演習(9) 第 9 回 (6/18) 3. ファイルとその応用 外部記憶装置に記録されたプログラムやデータを, ファイルと呼ぶ シーケンシャルファイルやランダムファイルへのデータの記録や読み出し, 更新の手順について学習する (1) ファイルとレコードファイル複数の関連したデータを一つに集めたり プログラムを外部記憶装置に保存したものレコードファイルを構成する一塊のデータ ex. 個人カードフィールドレコードを構成する個別の要素

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

02: 変数と標準入出力

02: 変数と標準入出力 C プログラミング入門 基幹 7 ( 水 5) 09: ポインタ 文字列 Linux にログインし 以下の講義ページを開いておくこと http://www-it.sci.waseda.ac.jp/ teachers/w483692/cpr1/ 2016-06-08 1 関数できなかったこと 配列を引数として渡す, 戻り値として返す 文字列を扱う 呼び出し元の変数を直接書き換える 例 : 2 つの変数の値を入れ替える関数

More information

gengo1-11

gengo1-11 関数の再帰定義 自然数 n の階乗 n! を計算する関数を定義してみる 引数は整数 返却値も整数 n! = 1*2*3*... * (n 1)*n である ただし 0! = 1 とする int factorial(int n) int i, tmp=1; if( n>0 ) for(i=1; i

More information

Operating System 仮想記憶

Operating System 仮想記憶 Operating System 仮想記憶 2018-12 記憶階層 高速 & 小容量 ( 高価 ) レジスタ アクセスタイム 数ナノ秒 容量 ~1KB CPU 内キャッシュ (SRAM) 数ナノ秒 1MB 程度 ランダムアクセス 主記憶 (DRAM) 数十ナノ秒 数 GB 程度 ランダムアクセス フラッシュメモリ (SSD) 約 100 万倍 シーケンシャルアクセス 磁気ディスク (HDD) 数十ミリ秒

More information

3,, となって欲しいのだが 実際の出力結果を確認すると両方の配列とも 10, 2, 3,, となってしまっている この結果は代入後の配列 a と b は同じものになっていることを示している つまり 代入演算子 = によるの代入は全要素のコピーではなく 先をコピーする ため 代入後の a と b は

3,, となって欲しいのだが 実際の出力結果を確認すると両方の配列とも 10, 2, 3,, となってしまっている この結果は代入後の配列 a と b は同じものになっていることを示している つまり 代入演算子 = によるの代入は全要素のコピーではなく 先をコピーする ため 代入後の a と b は 配列 2 前回には 配列の基本的な使い方と拡張 for 文について学んだ 本日は配列に付いての追加の説明として 配列のコピー 文字列配列 ガーベジコレクション 多次元配列について学んでいく 配列のコピー配列を用意し その全ての要素を別の配列にコピーすることを考える まず 以下に間違った例を示していく プログラム例 1 public class Prog07_01 int[] a = 1, 2, 3,,

More information

Microsoft PowerPoint - handout07.ppt [互換モード]

Microsoft PowerPoint - handout07.ppt [互換モード] Outline プログラミング演習第 7 回構造体 on 2012.12.06 電気通信大学情報理工学部知能機械工学科長井隆行 今日の主眼 構造体 構造体の配列 構造体とポインタ 演習課題 2 今日の主眼 配列を使うと 複数の ( 異なる型を含む ) データを扱いたい 例えば 成績データの管理 複数のデータを扱う 配列を使う! 名前学籍番号点数 ( 英語 ) 点数 ( 数学 ) Aomori 1 59.4

More information

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

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

More information