CUDA メモ: 2015

2015-11-25

OpenCL の C++ 実装と早見表

CUDA の話題ではないが、CUDA の C++ 簡易ラッパを作成するときに OpenCL の C++ バインディングを参考にしている。

また、OpenCL には早見表 (Reference Card) が公開されている。とても見やすくてよいと思う。CUDA の早見表も作ってくれないだろうか。

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

OpenCL 2.0 C++ Bindings Header File (this has not yet been updated for OpenCL 2.1)
https://www.khronos.org/registry/cl/api/2.1/cl2.hpp

OpenCL 2.0 Reference Card (November 14, 2013)
http://www.khronos.org/registry/cl/sdk/2.0/docs/OpenCL-2.0-refcard.pdf

2015-11-22

CPU で float(FP32) 型から half(FP16) 型へ変換する（逆変換もあり）方法

CPU で float(FP32) 型から half(FP16) 型へ変換する（逆変換もあり）ソースコードが、cuDNN v3 (September 8, 2015) のサンプルソースコード cudnn-sample-v3.tgz の中の fp16_emu.cpp にある。

half1 cpu_float2half_rn(float f); float cpu_half2float(half1 h);

2015-10-25

日本語の 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

2015-07-15

[GTX 970][CUDA 7.5 RC] half 型のデータで gemm

cublasSgemm 関数の half 型版である cublasHgemm 関数を試してみたかったが、Compute Capability が 5.3 以降のようだ。GTX 970 では使えないので、cublasSgemmEx 関数の引数にCUBLAS_DATA_HALF を指定して計算してみた。

サンプルソースコード matrixMulCUBLAS を簡易的に改造して実行した。cublasSgemm 関数では内部で maxwell_sgemm_128x64_nn 関数が呼び出されていたが、cublasSgemmEx 関数 (CUBLAS_DATA_HALF を指定) では maxwell_fp16_sgemm_fp16_128x64_nn 関数が呼び出されている。計算部分は float 型らしいので、性能はほとんど変わらないようだ（メモリ転送やキャッシュヒットの有無も考慮すれば、結果は違うのかも知れないが）。

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.5\bin\win64\Release>nvprof matrixMulCUBLAS_half.exe -sizemult=10
[Matrix Multiply CUBLAS] - Starting...
==6140== NVPROF is profiling process 6140, command: matrixMulCUBLAS_half.exe -sizemult=10
GPU Device 0: "GeForce GTX 970" with compute capability 5.2

MatrixA(640,1280), MatrixB(640,1280), MatrixC(640,1280)
Computing result using CUBLAS...done.
Performance= 2953.19 GFlop/s, Time= 0.355 msec, Size= 1048576000 Ops
Computing result using host CPU...done.
Listing first 100 Differences > 0.000010...

  Row 0:
    Loc(0,0)    CPU=160.44701   GPU=160.50000   Diff=0.052994
    Loc(1,0)    CPU=167.72478   GPU=167.75000   Diff=0.025223
    Loc(2,0)    CPU=171.19669   GPU=171.25000   Diff=0.053314
(省略)

  Total Errors = 819097
Comparing CUBLAS Matrix Multiply with CPU results: FAIL

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
==6140== Profiling application: matrixMulCUBLAS_half.exe -sizemult=10
==6140== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 77.29%  10.771ms        31  347.46us  344.64us  352.74us  maxwell_fp16_sgemm_fp16_128x64_nn
  9.45%  1.3174ms         3  439.14us  1.3120us  864.16us  [CUDA memcpy HtoD]
  6.19%  862.27us         1  862.27us  862.27us  862.27us  [CUDA memcpy DtoH]
  4.70%  655.58us         2  327.79us  327.17us  328.42us  copy_array_from_float_to_half(float const *, __half*, __int64)
  2.36%  329.02us         1  329.02us  329.02us  329.02us  copy_array_from_half_to_float(__half const *, float*, __int64)

==6140== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 69.23%  331.75ms        10  33.175ms  22.937us  330.90ms  cudaFree
 27.33%  130.98ms         9  14.554ms  9.4200us  128.45ms  cudaMalloc
  2.22%  10.629ms         1  10.629ms  10.629ms  10.629ms  cudaEventSynchronize
  0.65%  3.1297ms         4  782.42us  22.937us  1.6253ms  cudaMemcpy
  0.21%  999.81us         2  499.91us  427.21us  572.61us  cudaGetDeviceProperties
  0.17%  809.76us       166  4.8780us       0ns  258.04us  cuDeviceGetAttribute
  0.09%  420.24us        34  12.360us  9.8300us  31.539us  cudaLaunch
  0.05%  235.52us         2  117.76us  88.882us  146.63us  cuDeviceGetName
  0.01%  71.272us       381     187ns       0ns  1.2290us  cudaSetupArgument
  0.01%  26.214us        16  1.6380us     409ns  9.8310us  cudaEventCreateWithFlags
  0.01%  25.804us         2  12.902us  4.9150us  20.889us  cudaEventRecord
  0.00%  21.708us         2  10.854us  9.8300us  11.878us  cudaThreadSynchronize
  0.00%  15.974us         1  15.974us  15.974us  15.974us  cudaEventElapsedTime
  0.00%  14.335us        16     895ns     409ns  2.8670us  cudaEventDestroy
  0.00%  10.649us        34     313ns       0ns  2.0480us  cudaConfigureCall
  0.00%  9.4200us         2  4.7100us  4.0960us  5.3240us  cuDeviceTotalMem
  0.00%  6.9640us         2  3.4820us  1.2290us  5.7350us  cudaGetDevice
  0.00%  5.7330us        31     184ns       0ns     410ns  cudaGetLastError
  0.00%  4.9150us        10     491ns     409ns  1.2280us  cudaDeviceGetAttribute
  0.00%  4.5060us         2  2.2530us     819ns  3.6870us  cudaEventCreate
  0.00%  2.8680us         3     956ns     410ns  2.0480us  cuDeviceGetCount
  0.00%     819ns         3     273ns       0ns     410ns  cuDeviceGet
  0.00%     819ns         1     819ns     819ns     819ns  cuInit
  0.00%     409ns         1     409ns     409ns     409ns  cuDriverGetVersion
======== Error: Application returned non-zero code 1

2015-07-14

[GTX 970][CUDA 7.5 RC] sgemm の性能

CUDA のサンプルソース matrixMulCUBLAS で計測した。
環境は、Windows 7 x64、GTX 970、CUDA 7.5 RC、ドライバ 353.30。
性能は 2.93 TFlop/s。

[関連ページ]
[GTX 970][CUDA 7] sgemm の性能
http://cuda-memo.blogspot.jp/2015/07/gtx-970cuda-7-sgemm.html

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.5\bin\win64\Release>nvprof matrixMulCUBLAS.exe -sizemult=10
[Matrix Multiply CUBLAS] - Starting...
==4876== NVPROF is profiling process 4876, command: matrixMulCUBLAS.exe -sizemult=10
GPU Device 0: "GeForce GTX 970" with compute capability 5.2

MatrixA(640,1280), MatrixB(640,1280), MatrixC(640,1280)
Computing result using CUBLAS...done.
Performance= 2931.65 GFlop/s, Time= 0.358 msec, Size= 1048576000 Ops
Computing result using host CPU...done.
Comparing CUBLAS Matrix Multiply with CPU results: PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
==4876== Profiling application: matrixMulCUBLAS.exe -sizemult=10
==4876== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 84.06%  10.850ms        31  349.99us  346.34us  355.75us  maxwell_sgemm_128x64_nn
  9.67%  1.2485ms         3  416.16us  1.3120us  806.45us  [CUDA memcpy HtoD]
  6.26%  808.59us         1  808.59us  808.59us  808.59us  [CUDA memcpy DtoH]

==4876== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 67.27%  311.92ms         7  44.559ms  18.432us  311.27ms  cudaFree
 29.27%  135.70ms         6  22.617ms  10.649us  134.12ms  cudaMalloc
  2.30%  10.664ms         1  10.664ms  10.664ms  10.664ms  cudaEventSynchronize
  0.57%  2.6268ms         4  656.69us  23.348us  1.1702ms  cudaMemcpy
  0.24%  1.1256ms         2  562.79us  478.41us  647.17us  cudaGetDeviceProperties
  0.17%  806.91us       166  4.8600us       0ns  212.58us  cuDeviceGetAttribute
  0.08%  370.28us        31  11.944us  9.8300us  35.635us  cudaLaunch
  0.05%  228.15us         2  114.07us  84.377us  143.77us  cuDeviceGetName
  0.01%  69.225us       372     186ns       0ns     819ns  cudaSetupArgument
  0.01%  27.851us        16  1.7400us     819ns  11.059us  cudaEventCreateWithFlags
  0.00%  19.251us         2  9.6250us  9.0110us  10.240us  cudaThreadSynchronize
  0.00%  17.203us         1  17.203us  17.203us  17.203us  cudaEventElapsedTime
  0.00%  14.745us        16     921ns     409ns  2.0480us  cudaEventDestroy
  0.00%  13.517us         2  6.7580us  5.3250us  8.1920us  cudaEventRecord
  0.00%  9.0120us         2  4.5060us  3.6870us  5.3250us  cuDeviceTotalMem
  0.00%  8.5970us        31     277ns       0ns  1.6380us  cudaConfigureCall
  0.00%  7.3730us         2  3.6860us  1.6390us  5.7340us  cudaGetDevice
  0.00%  6.9660us        31     224ns       0ns     410ns  cudaGetLastError
  0.00%  4.5060us        10     450ns     409ns     819ns  cudaDeviceGetAttribute
  0.00%  4.0960us         2  2.0480us  1.2290us  2.8670us  cudaEventCreate
  0.00%  2.0490us         3     683ns       0ns  1.6390us  cuDeviceGetCount
  0.00%     410ns         1     410ns     410ns     410ns  cuDriverGetVersion
  0.00%     409ns         1     409ns     409ns     409ns  cuInit
  0.00%     409ns         3     136ns       0ns     409ns  cuDeviceGet

2015-07-13

[GTX 970][CUDA 7] sgemm の性能

CUDA のサンプルソース matrixMulCUBLAS で計測した。
環境は、Windows 7 x64、GTX 970、CUDA 7、ドライバ 353.30。
性能は 2.92 TFlop/s。

下記の参考ページによると、GTX 970 は 2.95TFlop/s、GTX 980 は 3.87TFlop/s、GTX 780 Ti は 3.06TFlop/s らしい。

[参考ページ]
http://www.comphys.las.shibaura-it.ac.jp/matrixMulCUBLAS_2015

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.0\bin\win64\Release>nvprof matrixMulCUBLAS.exe -sizemult=10
[Matrix Multiply CUBLAS] - Starting...
==1864== NVPROF is profiling process 1864, command: matrixMulCUBLAS.exe -sizemult=10
GPU Device 0: "GeForce GTX 970" with compute capability 5.2

MatrixA(640,1280), MatrixB(640,1280), MatrixC(640,1280)
Computing result using CUBLAS...done.
Performance= 2919.77 GFlop/s, Time= 0.359 msec, Size= 1048576000 Ops
Computing result using host CPU...done.
Comparing CUBLAS Matrix Multiply with CPU results: PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
==1864== Profiling application: matrixMulCUBLAS.exe -sizemult=10
==1864== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 81.92%  10.896ms        31  351.50us  348.19us  355.84us  maxwell_sgemm_128x64_nn
 12.14%  1.6145ms         3  538.15us  1.3120us  1.0752ms  [CUDA memcpy HtoD]
  5.94%  789.79us         1  789.79us  789.79us  789.79us  [CUDA memcpy DtoH]

==1864== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 66.45%  302.18ms         7  43.168ms  19.660us  301.50ms  cudaFree
 29.95%  136.21ms         6  22.702ms  11.469us  134.27ms  cudaMalloc
  2.36%  10.715ms         1  10.715ms  10.715ms  10.715ms  cudaEventSynchronize
  0.61%  2.7778ms         4  694.46us  23.347us  1.1751ms  cudaMemcpy
  0.26%  1.1714ms         2  585.71us  583.26us  588.17us  cudaGetDeviceProperties
  0.20%  925.67us       166  5.5760us       0ns  295.31us  cuDeviceGetAttribute
  0.08%  343.65us        31  11.085us  9.4200us  32.767us  cudaLaunch
  0.05%  219.54us         2  109.77us  78.232us  141.31us  cuDeviceGetName
  0.01%  58.161us       372     156ns       0ns     820ns  cudaSetupArgument
  0.01%  29.082us         2  14.541us  9.4210us  19.661us  cudaThreadSynchronize
  0.01%  26.623us        16  1.6630us     409ns  9.8300us  cudaEventCreateWithFlags
  0.00%  15.565us         1  15.565us  15.565us  15.565us  cudaEventElapsedTime
  0.00%  15.155us         2  7.5770us  6.1440us  9.0110us  cuDeviceTotalMem
  0.00%  14.336us         2  7.1680us  4.9150us  9.4210us  cudaEventRecord
  0.00%  13.517us        16     844ns     409ns  2.4570us  cudaEventDestroy
  0.00%  8.6020us         2  4.3010us  2.8670us  5.7350us  cudaGetDevice
  0.00%  6.9630us        31     224ns       0ns  1.2290us  cudaConfigureCall
  0.00%  4.5050us        10     450ns     409ns     819ns  cudaDeviceGetAttribute
  0.00%  4.0960us         2  2.0480us  1.2290us  2.8670us  cudaEventCreate
  0.00%  3.2760us        31     105ns       0ns     410ns  cudaGetLastError
  0.00%     820ns         3     273ns       0ns     410ns  cuDeviceGetCount
  0.00%     819ns         1     819ns     819ns     819ns  cuDriverGetVersion
  0.00%     819ns         1     819ns     819ns     819ns  cuInit
  0.00%     818ns         3     272ns       0ns     409ns  cuDeviceGet

2015-07-12

[GTX 970] cudaMallocPitch が返す pitch の値

GTX 970 の CUDA 7 において、cudaMallocPitch 関数が返す pitch の値は 512 （単位はバイト）である。
コンパイルオプションは compute_52, sm_52。cudaMallocPitch の引数 width と height はともに 1 とした。

なお、以下のページの情報によると、GPU や CUDA のバージョンによっては異なるようだ。

Returned pitch in cudaMallocPitch - NVIDIA Developer Forums:
https://devtalk.nvidia.com/default/topic/780542/returned-pitch-in-cudamallocpitch/

2015-06-14

cuDNN のダウンロード方法

cuDNN (NVIDIA CUDA Deep Neural Network library) をダウンロードするには、「CUDA Registered Developer Program」への登録が必要だ。

下記のページにアクセスして、[Register] ボタンを押して登録する。

https://developer.nvidia.com/cuDNN

しかし、登録してもすぐにはダウンロードできない。

登録完了後、しばらくして（私の場合は約6時間後に）「GPU Computing & Embedded Developer Program Application Appro」という表題のメールが送られてくる。その後にサイトからダウンロードできるようになる。

ライセンス情報は Windows の場合 cudnn-6.5-win-v2.zip ファイルの中に含まれている。

[追記] (2015-07-04) ライセンス情報について修正。

2015-03-13

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

2015-03-09

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 されるたびにメモリが確保されるのではなく、一度に最大ヒープサイズのメモリが確保されるようだ。

2015-03-07

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 となっている。

2015-03-06

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 が効かなくなるようだ。

2015-03-05

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

2015-03-04

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

2015-03-03

[Geforce GTX 970] cudaDeviceProp

        CUDA Device
            Name                                                                GeForce GTX 970
            Driver                                                              WDDM
            DeviceIndex                                                         0
            GPU Family                                                          GM204-A
            RmGpuId                                                             256
            Compute Major                                                       5
            Compute Minor                                                       2
            MAX_THREADS_PER_BLOCK                                               1024
            MAX_BLOCK_DIM_X                                                     1024
            MAX_BLOCK_DIM_Y                                                     1024
            MAX_BLOCK_DIM_Z                                                     64
            MAX_GRID_DIM_X                                                      2147483647
            MAX_GRID_DIM_Y                                                      65535
            MAX_GRID_DIM_Z                                                      65535
            MAX_SHARED_MEMORY_PER_BLOCK                                         49152
            TOTAL_CONSTANT_MEMORY                                               65536
            WARP_SIZE                                                           32
            MAX_PITCH                                                           2147483647
            MAX_REGISTERS_PER_BLOCK                                             65536
            CLOCK_RATE                                                          1253000
            TEXTURE_ALIGNMENT                                                   512
            GPU_OVERLAP                                                         1
            MULTIPROCESSOR_COUNT                                                13
            KERNEL_EXEC_TIMEOUT                                                 1
            INTEGRATED                                                          0
            CAN_MAP_HOST_MEMORY                                                 1
            COMPUTE_MODE                                                        0
            MAXIMUM_TEXTURE1D_WIDTH                                             65536
            MAXIMUM_TEXTURE2D_WIDTH                                             65536
            MAXIMUM_TEXTURE2D_HEIGHT                                            65536
            MAXIMUM_TEXTURE3D_WIDTH                                             4096
            MAXIMUM_TEXTURE3D_HEIGHT                                            4096
            MAXIMUM_TEXTURE3D_DEPTH                                             4096
            MAXIMUM_TEXTURE2D_LAYERED_WIDTH                                     16384
            MAXIMUM_TEXTURE2D_LAYERED_HEIGHT                                    16384
            MAXIMUM_TEXTURE2D_LAYERED_LAYERS                                    2048
            SURFACE_ALIGNMENT                                                   512
            CONCURRENT_KERNELS                                                  1
            ECC_ENABLED                                                         0
            PCI_BUS_ID                                                          1
            PCI_DEVICE_ID                                                       0
            TCC_DRIVER                                                          0
            MEMORY_CLOCK_RATE                                                   3505000
            GLOBAL_MEMORY_BUS_WIDTH                                             256
            L2_CACHE_SIZE                                                       1835008
            MAX_THREADS_PER_MULTIPROCESSOR                                      2048
            ASYNC_ENGINE_COUNT                                                  2
            UNIFIED_ADDRESSING                                                  1
            MAXIMUM_TEXTURE1D_LAYERED_WIDTH                                     16384
            MAXIMUM_TEXTURE1D_LAYERED_LAYERS                                    2048
            CAN_TEX2D_GATHER                                                    1
            MAXIMUM_TEXTURE2D_GATHER_WIDTH                                      16384
            MAXIMUM_TEXTURE2D_GATHER_HEIGHT                                     16384
            MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE                                   2048
            MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE                                  2048
            MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE                                   16384
            PCI_DOMAIN_ID                                                       0
            TEXTURE_PITCH_ALIGNMENT                                             32
            MAXIMUM_TEXTURECUBEMAP_WIDTH                                        16384
            MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH                                16384
            MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS                               2046
            MAXIMUM_SURFACE1D_WIDTH                                             65536
            MAXIMUM_SURFACE2D_WIDTH                                             65536
            MAXIMUM_SURFACE2D_HEIGHT                                            32768
            MAXIMUM_SURFACE3D_WIDTH                                             65536
            MAXIMUM_SURFACE3D_HEIGHT                                            32768
            MAXIMUM_SURFACE3D_DEPTH                                             2048
            MAXIMUM_SURFACE1D_LAYERED_WIDTH                                     65536
            MAXIMUM_SURFACE1D_LAYERED_LAYERS                                    2048
            MAXIMUM_SURFACE2D_LAYERED_WIDTH                                     65536
            MAXIMUM_SURFACE2D_LAYERED_HEIGHT                                    32768
            MAXIMUM_SURFACE2D_LAYERED_LAYERS                                    2048
            MAXIMUM_SURFACECUBEMAP_WIDTH                                        32768
            MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH                                32768
            MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS                               2046
            MAXIMUM_TEXTURE1D_LINEAR_WIDTH                                      134217728
            MAXIMUM_TEXTURE2D_LINEAR_WIDTH                                      65000
            MAXIMUM_TEXTURE2D_LINEAR_HEIGHT                                     65000
            MAXIMUM_TEXTURE2D_LINEAR_PITCH                                      1048544
            MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH                                   16384
            MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT                                  16384
            MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH                                   16384
            STREAM_PRIORITIES_SUPPORTED                                         0
            GLOBAL_L1_CACHE_SUPPORTED                                           0
            LOCAL_L1_CACHE_SUPPORTED                                            0
            MAX_SHARED_MEMORY_PER_MULTIPROCESSOR                                98304
            MAX_REGISTERS_PER_MULTIPROCESSOR                                    65536
            MANAGED_MEMORY                                                      1
            MULTI_GPU_BOARD                                                     0
            MULTI_GPU_BOARD_GROUP_ID                                            0
            DISPLAY_NAME                                                        GeForce GTX 970
            COMPUTE_CAPABILITY_MAJOR                                            5
            COMPUTE_CAPABILITY_MINOR                                            2
            TOTAL_MEMORY                                                        4294967296
            RAM_TYPE                                                            8
            RAM_LOCATION                                                        1
            GPU_PCI_DEVICE_ID                                                   331485406
            GPU_PCI_SUB_SYSTEM_ID                                               2231898179
            GPU_PCI_REVISION_ID                                                 161
            GPU_PCI_EXT_DEVICE_ID                                               5058
            GPU_PCI_EXT_GEN                                                     2
            GPU_PCI_EXT_GPU_GEN                                                 2
            GPU_PCI_EXT_GPU_LINK_RATE                                           8000
            GPU_PCI_EXT_GPU_LINK_WIDTH                                          16
            GPU_PCI_EXT_DOWNSTREAM_LINK_RATE                                    8000
            GPU_PCI_EXT_DOWNSTREAM_LINK_WIDTH                                   16

[GeForce GTX 970] deviceQuery

deviceQuery.exe Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 970"
  CUDA Driver Version / Runtime Version          6.5 / 6.5
  CUDA Capability Major/Minor version number:    5.2
  Total amount of global memory:                 4096 MBytes (4294967296 bytes)
  (13) Multiprocessors, (128) CUDA Cores/MP:     1664 CUDA Cores
  GPU Clock rate:                                1253 MHz (1.25 GHz)
  Memory Clock rate:                             3505 Mhz
  Memory Bus Width:                              256-bit
  L2 Cache Size:                                 1835008 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  CUDA Device Driver Mode (TCC or WDDM):         WDDM (Windows Display Driver Model)
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Bus ID / PCI location ID:           1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 6.5, CUDA Runtime Version = 6.5, NumDevs = 1, Device0 = GeForce GTX 970
Result = PASS

2015-03-01

OpenCL のソースコード

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

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

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

2015-02-28

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 に属した方が直感的ではないかと思う。

2015-02-27

CUDA セッションのタイムラインを見たい

Visual Studio 上で、実行される関数がどれくらい
時間がかかっているかを視覚化したタイムラインを見るには、

メニューの [Nsight]-[Start Performance Analysis...] を選択して、

[Trace Settings] の CUDA にチェックを入れて、

[Application Control] の [Launch] ボタンを押す。

プログラムの実行後にレポートが表示されるので、

[Session Overview] の [Timeline] を選択する。

タイムラインの時間軸をズームするには、
[Ctrl] を押しながらマウスのホイールを回す。

2015-02-26

CUDA 6.5 でプロジェクトを作って実行してみた

Windows 7 に Visual Studio Community 2013 と CUDA 6.5 をインストールする。
言語パックがエラーで導入できなかったので、英語のIDEのまま使用することにした。

Visual Studio を起動する。

メニューの [File]-[Project ...] を選択する。

[Templates] で [NVIDIA]-[CUDA 6.5] を選択し、[CUDA 6.5 Runtime] を選び、
[Name] にプロジェクト名、[Location] に配置フォルダを指定して [OK] を押す。

プロジェクトが作成され、kernel.cu ファイルが作成される。
kernel.cu ファイルには main 関数といくつかのサンプルコードがある。

メニューの [Debug]-[Start Debugging] からプログラムを実行する。

上手く動けば、設定を変更して、ソースコードを入力していく。

CUDA 6.5 のビルドで出る大量の warning C4819 を消す

Visual Studio 2013 の CUDA 6.5 でビルドすると、次のような警告が大量に出る。


1>c:\program files\nvidia gpu computing toolkit\cuda\v6.5\include\math_functions.h : warning C4819: The file contains a character that cannot be represented in the current code page (932). Save the file in Unicode format to prevent data loss

これを表示しないようにするには、プロジェクトのプロパティを開き、
[CUDA C/C++]-[Command Line] の [Additional Options] に


-Xcompiler "/wd 4819"

を追加する。

Geforce GTX 9xx 用の CUDA

CUDA 6.5 をインストールして、設定の [Code Generation] に「compute_52,sm_52」を指定してビルドすると、次のエラーが出る。


nvcc fatal : Unsupported gpu architecture 'compute_52'.

GTX9xx 用に別の CUDA があるようなので、下記をダウンロードしてインストールするとビルドできるようになった。

CUDA Downloads For GeForce GTX9xx
https://developer.nvidia.com/cuda-downloads-geforce-gtx9xx

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 を使えばいいようだ。