cuDNN v5 (May 27, 2016), for CUDA 8.0 RC

cuDNN v5 (May 27, 2016), for CUDA 8.0 RC が出ているようだ。

CUDA 8 RC を入れてみた

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

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

cuDNN v5 Release Candidate (RC) (April, 2016), for CUDA 7.5 and later

cuDNN v5 Release Candidate (RC) (April, 2016) が出ているようだ。

https://developer.nvidia.com/cudnn

word2vec_cbow を Windows で実行する

word2vec_cbow を Visual Studio 2013 でビルドして、Windows 上で実行してみた。

プロジェクトを作成する手順は次のとおり。

1. Visual Studio 2013 で CUDA 7.5 Runtime プロジェクトを作成する。
2. https://github.com/ChenglongChen/word2vec_cbow から word2vec.cu と cbow.cu の2つのファイルをダウンロードして、プロジェクトのフォルダに入れる。
3. https://github.com/zhangyafeikimi/word2vec-win32 から win32-port.h ファイルをダウンロードして、プロジェクトのフォルダに入れる。
4. word2vec.cu を開いて、#include <pthread .h> を #include "win32-port.h" に変更する。
5. プロジェクトに word2vec.cu を追加する。

私の実行環境は Geforce GTX 970 なので、[Configuration Properties]-[CUDA C/C++]-[Device]-[Code Generation] に「compute_52,sm_52」を指定した。

[Solution Platforms] は「x64」でビルドした。

実行結果は次のとおり。

>word2vec_cbow.exe -train text8 -output vectors.bin -cbow 1 -size 200 -window 7 -negative 1 -hs 1 -sample 1e-3 -threads 1 -binary 1 -save-vocab voc
Starting training using file text8
Vocab size: 71290
Words in train file: 16718843
vocab size = 71290
Alpha: 0.000009  Progress: 99.99%  Words/thread/sec: 1528.12k

実行中の GPU の負荷は 50％程度だった。

なお、処理結果が正しいかどうかは確認していないので、注意されたい。

cuDNN v4 (Feb 10, 2016)

cuDNN v4 の正式版が出ているようだ。

https://developer.nvidia.com/cudnn

以下、リリースノートから引用。

UPDATES SINCE RELEASE CANDIDATE

• The API of cudnnBatchNormalizationBackward has been changed to include an additional set of scaling parameters (alphaParamsDiff and betaParamsDiff) applied to the dBnScaleResult and dBnBiasResult outputs of the function.

• The prior restriction of batch size 512 in all Batch Normalization routines has been removed.

• Numerical stability and performance of cudnnBatchNormalizationBackward in some cases has been improved.

• Performance of cudnnConvolutionBackwardFilter when using Algo 1 has been improved for some cases. This code path now also requires a workspace.

cuDNN v4 Release Candidate (December 10, 2015)

cuDNN v4 Release Candidate (December 10, 2015) が出ていたようだ。

日本語のリリース情報は次のリンク先を参照。
https://ja-jp.facebook.com/NVIDIAGPUComputing/posts/458731480981793

以下、リリースノートから引用。

NEW FEATURES

• Batch Normalization routines have been added.
• Convolution forward and backward now supports NHWC tensor format.
• FFT Tiling algorithm has been added for cudnnConvolutionForward and cudnnConvolutionBackwardData routines
• cudnnConvolutionForward now supports computation in FP16 when run on GPU with a compute capability >= 5.3
• cudnnConvolutionForward has been optimized for batch size = 1
• Pooling and activation routines have a descriptor option to propagate NaN numbers.

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

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);

関連

• CPU の FP16C を使って float (FP32) と half (FP16) を変換する方法

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

[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