CUDA 8.0 が出たようだ

CUDA 8.0 Downloads | NVIDIA Developer:
https://developer.nvidia.com/cuda-downloads

CUDA 8 Features Revealed | Parallel Forall:
https://devblogs.nvidia.com/parallelforall/cuda-8-features-revealed/

Microsoft Visual Studio 2015 (updates 2 and 3) をサポートしたとか、NVCCコンパイラの速度が上がったとか。

CUDA 8 RC を入れてみた

CUDA 8 RC (8.0.27) が出ていたので、インストールしてみた。

Visual Studio 2015 に対応していたので、Visual Studio も更新した。

日本語の CUDA に関するページ

Coding/CUDA - ClockAhead 記憶の欠片
http://wiki.clockahead.com/index.php?Coding%2FCUDA

CUDA入門・サンプル集
http://cudasample.net/

トータル・ディスクロージャ・サイト
http://topsecret.hpc.co.jp/wiki/index.php/%E3%83%A1%E3%82%A4%E3%83%B3%E3%83%9A%E3%83%BC%E3%82%B8

CUDA Information Site
http://gpu.fixstars.com/index.php/%E3%83%A1%E3%82%A4%E3%83%B3%E3%83%9A%E3%83%BC%E3%82%B8

GPGPUをもふもふする会 - xhl Wiki*
http://wikiwiki.jp/xhl/?GPGPU%A4%F2%A4%E2%A4%D5%A4%E2%A4%D5%A4%B9%A4%EB%B2%F1

tips : tips/02.プログラミングなど/GPGPU/CUDAメモ.txt
http://homepage2.nifty.com/takaaki024/tips/programs/gpgpu/cuda.html

良いもの。悪いもの。: CUDAで作成した分子動力学計算プログラムを書き直してみた
http://handasse.blogspot.com/2009/12/cuda.html

CUDA技術を利用したGPUコンピューティングの実際（後編） ―― FFTを利用した光波の伝播（フレネル回折）をGPUで高速計算｜Tech Village （テックビレッジ）　／　CQ出版株式会社
http://www.kumikomi.net/archives/2008/10/22gpu2.php?page=9

CUDA でカーネルからカーネルを呼び出す (Dynamic Parallelism)

Visual Studio 2013 の CUDA で、カーネルからカーネルを呼び出す (Dynamic Parallelism) には、次のページに書かれている設定を行う必要がある。

Compiling CUDA Projects with Dynamic Parallelism (VS 2012/13) | Viral F#:
http://viralfsharp.com/2014/08/17/compiling-cuda-projects-with-dynamic-parallelism-vs-201213/

Geforce GTX 970 で CUDA カーネル呼び出しにかかるオーバーヘッドを計測した。

(1) Dynamic Parallelism なし： ホストからカーネルを呼び出し
(2) Dynamic Parallelism あり： ホストからカーネルを呼び出し
(3) Dynamic Parallelism あり： カーネルからカーネルを呼び出し

結果は次のとおり。

(1) 3.80 µsec (マイクロ秒)
(2) 7.56 µsec
(3) 7.19 µsec

Dynamic Parallelism を使用にすると、コンパイルオプションの [Generate Relocatable Device Code] の影響なのか分からないが、ホストからカーネルを呼び出す場合もオーバーヘッドが増加している。

CPU から GPU の呼び出しを GPU から GPU の呼び出しに単純に変更しても、速度は速くならないようだ。

[参考]
CUDA の Kernel 呼び出しのオーバーヘッドと引数の数について ( 周辺機器 ) - 正統納豆天国ブログ - Yahoo!ブログ:
http://blogs.yahoo.co.jp/natto_heaven/33360615.html

CUDA の Kernel 呼び出しオーバーヘッド: Dynamic Parallelism 編 ( 周辺機器 ) - 正統納豆天国ブログ - Yahoo!ブログ:
http://blogs.yahoo.co.jp/natto_heaven/33395337.html

CUDA カーネルでの malloc とメモリ使用量

CUDA カーネルで malloc に大きいサイズを指定するとヒープ不足で失敗する。
事前に下記の API を呼び出して余裕を持たせたヒープサイズを指定する必要がある。

cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

Geforce GTX 970 メモリ 4GB で、cudaDeviceSetLimit(cudaLimitMallocHeapSize) に 2GB を指定する。そして、カーネルから malloc(1) を呼び出して 1バイトのメモリを確保する。その後、cudaMemGetInfo API が返す free は 2GB へ減少する。malloc されるたびにメモリが確保されるのではなく、一度に最大ヒープサイズのメモリが確保されるようだ。

Thrust: CUDA の C++ template ライブラリ

CUDA SDK に Thrust という CUDA の C++ template ライブラリが
標準で入っているのにさっき気付いた。

知らずに host_vector とかを自分で実装してしまった…。

Thrust :: CUDA Toolkit Documentation:
http://docs.nvidia.com/cuda/thrust/

C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v6.5/include/thrust

にあるライブラリのバージョンは 1.7.2 のようだ。

GitHub の Thrust project page の最新版は 1.8.0 となっている。

CUDA で分岐中の __syncthreads の動作

CUDA で分岐中に __syncthreads() を実行させた場合にデッドロックが起きると
インターネット上で見かけた文書に書かれていた。

そこで実際に GTX 970 で実験してみたが、デッドロックは起きなかった。

__global__
void my_kernel()
{
    __shared__ int shared[2];
    shared[0] = -1;
    shared[1] = -1;

    int val;
    if (threadIdx.x == 0)
    {
        shared[1 - threadIdx.x] = threadIdx.x;
        __syncthreads();
        val = shared[threadIdx.x];
    }
    else
    {
        shared[1 - threadIdx.x] = threadIdx.x;
        __syncthreads();
        val = shared[threadIdx.x];
    }

    printf("threadIdx.x=%d, val=%d.\n", threadIdx.x, val);
}

my_kernel<<<1, 2>>>(); でカーネル関数を呼び出す。

実行結果は次のとおり。

threadIdx.x=0, val=1.
threadIdx.x=1, val=-1.

デッドロックはしなかったが、予期しない結果になった。

また、上記の関数で分岐の片方の __syncthreads() をコメントアウトしても
デッドロックはしなかった。

分岐中での __syncthreads() の動作は未定義ということだが、
これをテストした環境では何もしないという動作のようだ。

なお、上記のコードから if 文を取り除けば、
次のように正しい結果になる。

threadIdx.x=0, val=1.
threadIdx.x=1, val=0.

Visual Studio で CUDA のソースコード編集時の赤い波線を消す

[注意]
この方法は副作用があるようなので、おすすめはしない。

Visual Studio で CUDA のソースコードを編集していると、
構文間違いではないのに赤い波線のエラーが表示される。

例えば、

__syncthreads();

や

my_cuda_kernel_func<<<1, 1>>>();

などで赤い波線が表示される。

これを解消するには、ソースコードの先頭に

#ifndef __CUDACC__
#define __CUDACC__
#endif

を挿入する。

[参考]
visual studio 2010 - CUDA __syncthreads() compiles fine but is underlined with red - Stack Overflow:
http://stackoverflow.com/questions/13893919/cuda-syncthreads-compiles-fine-but-is-underlined-with-red

[追記]
これをすると、IntelliSense が効かなくなるようだ。

CUDA の block と warp、threadIdx と thread ID について

1つの block が 31*31 のスレッドで構成される場合、

warpSize = 32
31 * 31 / warpSize = 30
31 * 31 % warpSize = 1

となり、block は 31個の warp に分割される。

また、thread ID は、

(thread ID) = threadIdx.x + threadIdx.y * blockDim.x

[参考]
cuda - How is the 2D thread blocks padded for warp scheduling? - Stack Overflow:
http://stackoverflow.com/questions/15044671/how-is-the-2d-thread-blocks-padded-for-warp-scheduling

CUDA で高速に配列の合計値を計算する方法

GPU を使って配列の合計値を計算する（parallel reductions）には、共有メモリとスレッド間の同期をとるためのバリアを使う方法が一般的だ。

CUDA には、warp という32個のスレッドのまとまりがある。warp 内のスレッドは常に同期しているので、バリアが不要になり、その分だけ高速化できる。

Kepler 以降（__CUDA_ARCH__ >= 300）の場合は __shfl_down という命令を使ってさらに高速化できる。この命令は、あるスレッドが同じ warp 内の別のスレッドのレジスタを直接参照できるので、共有メモリを使わずに warp 内のレジスタの合計を計算することができる。

詳しくは下記のページを参照されたい。

Faster Parallel Reductions on Kepler | Parallel Forall:
http://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/

また、CUDA SDK に含まれるサンプルソースコード reduction_kernel.cu も参考になる。CUDA SDK 6.5 の場合は下記のパスにある。
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v6.5\6_Advanced\reduction

OpenCL のソースコード

CUDA の SDK に含まれている OpenCL のバージョンは 1.1 のようだが、
cl.hpp ファイルが見つからない。

なお、cl.hpp のソースコードは次のページにある。

Khronos OpenCL Registry:
https://www.khronos.org/registry/cl/

CUDA の Stream を使って多重データ転送

下記のページに Stream を使って多重データ転送する方法と
ベンチマークが載っている。

How to Overlap Data Transfers in CUDA C/C++ | Parallel Forall:
http://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/

最初に全てのストリームでデータを転送してカーネルを実行する方法と
データを転送してカーネルを実行するのをストリーム数繰り返す方法とでは
デバイスによってパフォーマンスが違うようだ。

CUDA の Stream 間の同期方法

CUDA の Stream 間の同期方法が API を見ても分からなかったが、下記サイトを読むと、cudaEventRecord と cudaStreamWaitEvent を使えばいいようだ。

Declaring dependencies with cudaStreamWaitEvent - Cedric Augonnet:
http://cedric-augonnet.com/declaring-dependencies-with-cudastreamwaitevent/

Java で例えるなら、cudaEventRecord が Object#notifyAll で、cudaStreamWaitEvent が Object#wait になるだろうか。

私の個人的な感覚では、cudaEventRecord はイベント API に属さずにストリームの API に属した方が直感的ではないかと思う。

Visual Studio Community 2013 に CUDA 6.5 をインストール

Visual Studio Express 2013 for Windows Desktop に CUDA 6.5 をインストールして使おうとしたのだが、設定などが上手くいかない。

Visual Studio Community 2013 Update 4 に CUDA 6.5 をインストールしたら、すんなりいった。

無理に Visual Studio Express を使わずに、素直に Visual Studio Community 2013 を使えばいいようだ。