はじめてのCUDAプログラミング
サポート情報
最終更新2011年01月19日
| はじめてのCUDAプログラミング サポート情報 最終更新2011年01月19日 |
■サンプル・プログラムのダウンロード(2009年11月19日）
｢はじめてのCUDAプログラミング｣の中で解説しているサンプル・プログラムです｡
Linux/Windows/MacOSに対応してコンパイル､実行できるようになっています｡
サンプル・プログラムを実際に動かしながら本書を読むことをお勧めします｡
| Windows 版のサンプル・プログラムの利用ついては、以下の補足説明(PDF)を参照してください。 [Windows 版を利用する場合のインストールについての補足説明(pdf)] |
| 第8.1 節 | サンプル・プログラム１ …メモリへのアクセス Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 第11.2節〜第11.7節 | サンプル・プログラム２ …粒子計算１ Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 第11.8節 | サンプル・プログラム３ …粒子計算２ Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 第12.2節 | サンプル・プログラム４ …拡散方程式(シェアード・メモリの利用無) Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 第12.3節 | サンプル・プログラム５ …拡散方程式(シェアード・メモリの利用) Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 第12.4節 | サンプル・プログラム６ …拡散方程式(シェアード・メモリの節約) Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 第12.5節 | サンプル・プログラム７ …拡散方程式(レジスタの利用) Windows Linux Mac(CUDA2.2用) Mac(CUDA2.3用) |
| 誤 | set CUDA_PROFILE=prof.conf |
| 正 | set CUDA_PROFILE_CONFIG=prof.conf |
・p.4 解説図および解説文を以下のように変更いたします｡
| 旧 | レイノルズ数が｢3000｣の計算を､｢768×384×384｣という… |
| 新 | レイノルズ数が｢13000｣の計算を､｢2000×1000×1000｣という… |
| 誤 | 58GPUを使うと7.2TFlopsにも達する｡スパコンの7000CPUで計算… |
| 正 | 60GPUを使うと10TFlopsにも達する｡東京工業大学・学術国際情報センターのスパコン､｢TSUBAME1.2｣のOpteron CPU 10000 Coreで計算… |
| 誤 | 消費電力を｢1/100｣以下に… |
| 正 | 消費電力を｢1/50｣以下に… |
| 誤 |
※2命令同時実行の場合 ※※3命令同時実行の場合 |
| 正 |
※積和演算命令のみを実行の場合 ※※積和演算命令 + 乗算命令を実行の場合 |
| 誤 | 配列｢f_d[n]｣ |
| 正 | 配列｢f_h[n]｣ |
| 誤 | このことにより､このことにより… |
| 正 | このことにより､… |
| 誤 | Elapsed Time= 2.956e+00 [sec] |
| 正 | Elapsed Time= 2.275e-01 [sec] |
| 誤 | Elapsed Time= 5.923e+02 [sec] |
| 正 | Elapsed Time= 3.716e+01 [sec] |
| 誤 | ｢CPU｣より200倍も… |
| 正 | ｢CPU｣より163倍も… |
| 誤 | 何と1557倍も速く… |
| 正 | 何と976倍も速く… |
| 誤 | int ny, /* grid number in the x-direction */ |
| 正 | int ny, /* grid number in the y-direction */ |
| 誤 |
|
| 正 |
|
| 誤 |
|
| 正 |
|
| 誤 | int ny, /* grid number in the x-direction */ |
| 正 | int ny, /* grid number in the y-direction */ |
| 誤 | int ny, /* grid number in the x-direction */ |
| 正 | int ny, /* grid number in the y-direction */ |
| 誤 | blockIdx.y == gridDim.y - 1 |
| 正 |
blockDim_y*blockIdx.y+jy == ny - 1 または、 blockIdx.y == gridDim.y - 1 && jy == blockDim_y - 1 |
| 誤 |
__syncthreads(); fn[j-nx] = c0*(fs[j1][js - 1] + fs[j1][js + 1]) + c1*(fs[j0][js] + fs[j2][js]) + c2*fs[j1][js]; j += nx; |
| 正 |
__syncthreads(); fn[j-nx] = c0*(fs[j1][js - 1] + fs[j1][js + 1]) + c1*(fs[j0][js] + fs[j2][js]) + c2*fs[j1][js]; __syncthreads(); j += nx; |
| 誤 |
また､｢シェアード・メモリ｣に書き込みを行なったら､読み出す前に必ず｢__syncthreads( );｣を入れる鉄則も忘れてはいけません｡ |
| 正 |
｢シェアード･メモリ｣を介してのスレッド間でデータのやり取りを確実に行なうためには､スレッド間の同期をとることが大切です｡ ｢__syncthreads();｣を入れることで､すべてのスレッドが｢シェアード･メモリ｣にデータを書き込み終わっていることを確認します｡ また､すべてのスレッドが｢シェアード･メモリ｣からデータを読み終わってから､次に｢シェアード･メモリ｣に書き込みを行なうまでの間にも､｢__ syncthreads();｣が必要です｡ ここで同期をとらないと､｢シェアード･メモリ｣のデータを読み終わる前に､次のデータで書き換えられてしまう可能性があるからです｡ |
| 誤 | int ny, /* grid number in the x-direction */ |
| 正 | int ny, /* grid number in the y-direction */ |
| 誤 | blockIdx.y == gridDim.y - 1 |
| 正 |
blockDim_y*blockIdx.y+jy == ny - 1 または、 blockIdx.y == gridDim.y - 1 && jy == blockDim_y - 1 |
| 誤 |
a.y = 10: b.y = 8: |
| 正 |
a.y = 10; b.y = 8; |
| 誤 |
a.y = 1: b.y = 1: |
| 正 |
a.y = 1; b.y = 1; |
| 誤 |
a.y = 10: b.y = 8: |
| 正 |
a.y = 10; b.y = 8; |
| 誤 |
a.y = 10: b.y = 8: |
| 正 |
a.y = 10; b.y = 8; |
| 誤 | Db(16, 16) |
| 正 | Db(16, 16); |
| 誤 | CUDA Visual Profile |
| 正 | CUDA Visual Profiler |
| 誤 | CUDA Visual Profile |
| 正 | CUDA Visual Profiler |
| 誤 | f[j-1] f[j+nx]… |
| 正 | f[j-1], f[j+nx]… |
| 誤 | f[j] f[j+nx+1]… |
| 正 | f[j], f[j+nx+1]… |
（株）工学社 第二I/O編集部
本書の内容に関するご質問／お問い合わせは、次の方法で工学社編集部宛にお願いします。
なお、電話によるお問い合わせはご遠慮ください。
※ご質問／お問い合わせの際、お客様の使用・動作環境などに添えて、具体的な症状をできるだけ詳しくお知らせください。
※FAXでのご質問／お問い合わせの場合は、必ずお客様のFAX番号を明記してください。