[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

[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

[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

[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/