CUDA7.5

CUDA7.5が出ました．

 

16-bit floating point (FP16) data format

• Store up to 2x larger datasets in GPU memory
• Reduce memory bandwidth requirements by up to 2x
• New mixed precision cublasSgemmEX() routine supports 2x larger matrices

 

前から言われていたFP16．つまり単精度じゃなくて半精度演算．

倍精度よりも単精度の速いわけでしたが，半精度だと単精度の倍の速度が出るとか．

あと，メモリも半分で良い．行列も倍のサイズが持てる．

 

これ，たぶんDeep Nueral Netを超意識していて，，cuDNNの演算部分で sigmoid とかtanhとかReLUがたぶんそれほど高精度にいらないから，FP16導入してこれだけで実質x2の高速化とかになるんじゃないだろうか？

 

で，

とりあえずTool kit を落として来てSampleを見てみると，

　　6_Advanced/matrixMulDynlinkJIT/cuda_drvapi_dynlink_cuda.h

このサンプルだけに使われていて，

/**
* Array formats
*/
typedef enum CUarray_format_enum
{
CU_AD_FORMAT_UNSIGNED_INT8 = 0x01, /**< Unsigned 8-bit integers */
CU_AD_FORMAT_UNSIGNED_INT16 = 0x02, /**< Unsigned 16-bit integers */
CU_AD_FORMAT_UNSIGNED_INT32 = 0x03, /**< Unsigned 32-bit integers */
CU_AD_FORMAT_SIGNED_INT8 = 0x08, /**< Signed 8-bit integers */
CU_AD_FORMAT_SIGNED_INT16 = 0x09, /**< Signed 16-bit integers */
CU_AD_FORMAT_SIGNED_INT32 = 0x0a, /**< Signed 32-bit integers */
CU_AD_FORMAT_HALF = 0x10, /**< 16-bit floating point */
CU_AD_FORMAT_FLOAT = 0x20 /**< 32-bit floating point */
} CUarray_format;

・・・・・

CU_RES_VIEW_FORMAT_FLOAT_1X16 = 0x13, /**< 1 channel 16-bit floating point */
CU_RES_VIEW_FORMAT_FLOAT_2X16 = 0x14, /**< 2 channel 16-bit floating point */
CU_RES_VIEW_FORMAT_FLOAT_4X16 = 0x15, /**< 4 channel 16-bit floating point */
CU_RES_VIEW_FORMAT_UNSIGNED_BC6H = 0x20, /**< Block compressed 6 unsigned half-float */

 

けど色々と見たけど16bit演算しているような箇所は見当たらなかったので，float4型に代わる hfloat8とかできるのかと思ったけど，まだ不明．

---

追記

 New Features in CUDA 7.5 | Parallel Forall

もうちょっと詳しいページがあって，

「A new header, cuda_fp16.h defines the half and half2 datatypes and __half2float() and __float2half() functions for conversion to and from FP32 types, respectively.」

と書いてあった． half 型と half2はあるみたいなのと， half => floatの変換とその逆もあるらしい．

「The range of half-precision numbers is approximately 5.96 \times 10^{-8} \ldots 6.55 \times 10^4. half2 structures store two half values in the space of a single 32-bit word」

ということで値の範囲も書いてあった．