CUDA技術を利用したGPUコンピューティングの実際（前編） ―― グラフィックス分野で磨かれた並列処理技術を汎用数値計算に応用

　CUDAにはブロック番号を格納したblockIdxという構造体が用意されており，blockIdx.xがグリッド内のブロックの水平方向番号を，blockIdx.yが垂直方向番号を示します．さらに，ブロック内のスレッド番号に対してthreadIdxという構造体が用意されており，threadIdx.xがブロック内のスレッド水平方向番号をthreadIdx.yが垂直方向番号を示します．また，ブロック内の水平・垂直方向のスレッド数をblockDimという構造体から取得できます．

　これらのCUDA組み込み変数を使用することで，配列からスレッドに対応したデータを取得できるようになります．例えば，配列Aの水平方向座標mはm＝blockDim.x ×blockIdx.x＋threadIdx.xで算出でき，垂直方向座標nはn＝blockDim.y×blockIdx.y＋threadIdx.yで算出できます．配列AはN×Nの要素を持つので，配列中のある1要素へはA［n×N+m］でアクセスできます．

4）メモリのスコープ

　表2のスコープ欄に，グリッド，ブロック，スレッドがどのメモリへアクセス可能かをまとめます．グローバル・メモリ，テクスチャ・メモリ，コンスタント・メモリのスコープはグリッド全体なので，すべてのブロック，スレッドで内容が共有化され，アクセスが可能となっています．

　シェアード・メモリのスコープはブロックなので，ブロック内の各スレッドは，そのブロックのシェアード・メモリの内容が共有化されアクセスが可能です．ただし，異なるブロック間のシェアード・メモリは互いにアクセスすることができません．

　レジスタのスコープはスレッドです．同一ブロックにあるスレッド同士でも，そのスレッドのレジスタの内容へはアクセスできません．

● ブロック数とスレッド数には上限がある

　以下に，マルチプロセッサで実行できる最大のブロック数とスレッド数などについてまとめておきます．

• ブロック当たりの最大スレッド数：512スレッド
• グリッドの各次元に指定できる最大ブロック数：65535ブロック
• ワープ：32スレッド
• マルチプロセッサで実行できる最大ブロック数：8個
• マルチプロセッサで実行できる最大ワープ数：24個
• マルチプロセッサで実行できる最大スレッド数：768スレッド（24ワープ×32スレッド/ワープ）

3．CUDA環境のインストールとプロジェクト作成

　NVIDIA社のWebサイトから下記の三つのファイルをダウンロードします（筆者が使用したCUDAのバージョンは1.1）．本稿では，Windows XPにおける環境構築について解説します．CUDAはLinuxにも対応しており，このWebサイトからLinux用のファイルもダウンロードできるようになっています．

• NVIDIA Driver for Microsoft Windows XP with CUDA Support（169.21）
• CUDA Toolkit version 1.1 for Windows XP
• CUDA SDK version 1.1 for Windows XP

　三つのファイルを上から順にインストールすれば，CUDAの環境が整います．

● コンパイラがGPU用コードとホスト用コードを分離

　CUDAはC言語を拡張した文法を採用しているので，C言語で開発経験のある方ならほとんど違和感なくプログラミングを進められると思います．CUDAによるプログラミングは，ホスト側でGPUを制御するプログラムと，GPU側で計算を行わせるプログラム（カーネル）の両方を作成する必要があります．このように書くと何か特殊な作業を行うように感じられるかもしれませんが，GPU側のプログラミングでも，計算したい内容をC言語に似た文法で記述できます．

　ホスト側のプログラムとカーネルのプログラムを別ファイルに記述し，コンパイル時に統合することもできますが，本稿では，同一のファイル（cuファイル）にホストとカーネルのプログラムを記述していくことにします．

　CUDAのコンパイラnvccは，まずcuファイルをGPU側のコードとホスト側のコードに分離します．分離したホスト側のコードは，Visual Studio付属のコンパイラでコンパイルされます．また，GPU側のプログラムはnvccによりコンパイルされ，GPU用のバイナリ・コードへと変換されます．