Fermi世代
GeForce GTX 580
単精度演算理論性能
2(1クロックあたりの演算量) × 32(SMあたりの単精度演算器数) × 1544MHz(動作周波数) × 16(SM数) = 1581Gflops
倍精度演算理論性能
2(1クロックあたりの演算量) × 16(SMあたりの倍精度演算器数) × 1544MHz(動作周波数) × 16(SM数) ÷ 4 (Teslaとの差別化のため係数) = 197Gflops
Kelper世代
GeForce GTX 680
単精度演算理論性能
2(1クロックあたりの演算量) × 192(SMXあたりの単精度演算器数) × 1006MHz(動作周波数) × 8(SMX数) = 3090Gflops
倍精度演算理論性能
2(1クロックあたりの演算量) × 64(SMXあたりの倍精度演算器数) × 1006MHz(動作周波数) × 8(SM数) ÷ 8 (Teslaとの差別化のため係数) = 128Gflops
Geforce GTX780
単精度演算理論性能
2(1クロックあたりの演算量) × 192(SMXあたりの単精度演算器数) × 863MHz(動作周波数) × 12(SMX数) = 3976Gflops
倍精度演算理論性能
2(1クロックあたりの演算量) × 64(SMXあたりの倍精度演算器数) × 863MHz(動作周波数) × 12(SM数) ÷ 8 (Teslaとの差別化のため係数) = 165Gflops
つまり、単精度ではFermi世代よりも2倍以上性能が向上しているにも拘わらず、倍精度に関してはKeplerの方が遅くなってしまっているのである。これはTeslaとの差別化のためのクロックダウンを1/4から1/8にしたためである。
これは当ブログのようにGPUを演算加速装置として使用するソフトウェアを一般向けに配布する立場からすると非常に困った事態である。もしMaxwell世代以降で更なる倍精度演算のクロックダウンが行われるならば、GeForceを新しいものにしてもソフトウェアの性能は向上しない。もしGPUの発展の恩恵に預かろうとすれば、20～30万円もする数値演算専用のTeslaを購入しなければならない。そのような選択は一般ユーザーには非現実的であろう。
そこで、今後のGeForceで倍精度演算が伸びず、単精度演算性能のみが向上した場合に備えて、擬似倍精度を用いた加算、乗算のコードを書いてみた。擬似倍精度とは、単精度浮動小数点数型を2つ用いて倍精度を表現する方法である。普通の倍精度は52ビットの仮数部を持つのに対し、単精度浮動小数点数型は仮数部が23ビットなので、擬似倍精度は46ビットの精度をもつ。よって擬似倍精度は通常の倍精度に比べて若干精度が落ちる。
以下にサンプルコードを示す。
作成にあたっては、
反復法ライブラリ向け4倍精度演算の実装とSSE2を用いた高速化
http://www.slis.tsukuba.ac.jp/~hasegawa.hidehiko.ga/GYOSEKI/IPSJ-TACS0101009A.pdf
と
Implementation of float-float operators on graphics hardware
http://hal.archives-ouvertes.fr/docs/00/06/33/56/PDF/float-float.pdf
を参考にした。
#include <stdio.h>
#include <time.h>
/* 倍精度から擬似倍精度に変換 */
__device__ void double_to_float_float(double a, float *a_hi, float *a_lo)
{
*a_hi = a;
*a_lo = a - (float)a;
}
/* 擬似倍精度から倍精度に変換 */
__device__ void float_float_to_double(float a_hi, float a_lo, double *a)
{
*a = a_hi + (double)a_lo;
}
/* 擬似倍精度加算 */
__device__ void float_float_add(float *a_hi, float *a_lo, float b_hi, float b_lo, float c_hi, float c_lo)
{ float sh, eh, v;
/* TWO_SUM */
sh = b_hi + c_hi;
v = sh - b_hi;
eh = (b_hi - (sh - v)) + (c_hi - v);
/* */
eh += (b_lo + c_lo);
/* FAST_TWO_SUM */
*a_hi = sh + eh;
v = *a_hi - sh;
*a_lo = (sh - (*a_hi - v)) + (eh - v);
}
/* 擬似倍精度乗算 */
__device__ void float_float_mul(float *a_hi, float *a_lo, float b_hi, float b_lo, float c_hi, float c_lo)
{
float p1, p2, d_hi, d_lo, e_hi, e_lo, t, v;
/* TWO_PROD */
p1 = b_hi * c_hi;
t = 4097.0 * b_hi;
d_hi = t - ( t - b_hi);
d_lo = b_hi - d_hi;
t = 4097.0 * c_hi;
e_hi = t - ( t - c_hi);
e_lo = c_hi - e_hi;
p2 = ( (d_hi * e_hi - p1) + d_hi * e_lo + d_lo * e_hi) + d_lo * e_lo;
/* */
p2 += (b_hi * c_lo) + (b_lo * c_hi);
/* FAST_TWO_SUM */
*a_hi = p1 + p2;
v = *a_hi - p1;
*a_lo = (p1 - (*a_hi - v)) + (p2 - v);
}
/* 倍精度、擬似倍精度比較 */
__global__ void kernel1()
{
size_t t_s, t_e;
double a, a_dash;
float a_hi, a_lo;
double theta_a = 2.21315648654123846246;
double b, b_dash;
float b_hi, b_lo;
double theta_b = 1.21315648654123846246;
double c, c_dash;
float c_hi, c_lo;
double theta_c = 0.21315648654123846246;
/* クロック関数呼び出しコスト */
t_s = clock();
t_e = clock();
printf("clock %ld clocks\n\n", t_e - t_s);
/* サイン関数呼び出しコスト */
t_s = clock();
a = sin(theta_a);
t_e = clock();
printf("sinf %ld clocks\n\n", t_e - t_s);
/* 倍精度 -> 擬似倍精度変換 */
t_s = clock();
double_to_float_float(a, &a_hi, &a_lo);
t_e = clock();
printf("double to float-float %ld clocks\n\n", t_e - t_s);
/* 擬似倍精度 -> 倍精度変換 */
t_s = clock();
float_float_to_double(a_hi, a_lo, &a_dash);
t_e = clock();
printf("float-float to double %ld clocks\n\n", t_e - t_s);
/* 変換結果 */
printf("a = %1.15e\n", a);
printf("a' = %1.15e\n\n", a_dash);
b = sin(theta_b);
c = sin(theta_c);
/* 加算比較 */
printf("Add\n");
t_s = clock();
a = b + c;
t_e = clock();
printf("double %ld clocks\n", t_e - t_s);
double_to_float_float(b, &b_hi, &b_lo);
double_to_float_float(c, &c_hi, &c_lo);
t_s = clock();
float_float_add(&a_hi, &a_lo, b_hi, b_lo, c_hi, c_lo);
t_e = clock();
printf("float-float %ld clocks\n", t_e - t_s);
float_float_to_double(a_hi, a_lo, &a_dash);
printf("a = %1.15e\n", a);
printf("a' = %1.15e\n\n", a_dash);N
/* 乗算比較 */
printf("Multiply\n");
b = -sin(theta_b);
c = sin(theta_c);
t_s = clock();
a = b * c;
t_e = clock();
printf("double %ld clocks\n", t_e - t_s);
double_to_float_float(b, &b_hi, &b_lo);
double_to_float_float(c, &c_hi, &c_lo);
t_s = clock();
float_float_mul(&a_hi, &a_lo, b_hi, b_lo, c_hi, c_lo);
t_e = clock();
printf("float-float %ld clocks\n", t_e - t_s);
float_float_to_double(a_hi, a_lo, &a_dash);
printf("a = %1.15e\n", a);
printf("a' = %1.15e\n\n", a_dash);
/* 絶対値計算コスト測定 */
long *lp_a_hi, *lp_a_lo;
lp_a_hi = (long*) &a_hi;
lp_a_lo = (long*) &a_lo;
printf("FABS\n");
t_s = clock();
*lp_a_hi &= 0x7fffffff;
*lp_a_lo &= 0x7fffffff;
t_e = clock();
printf("float-float %ld clocks\n", t_e - t_s);
float_float_to_double(a_hi, a_lo, &a_dash);
printf("a' = %1.15e\n", a_dash);
}
int main()
{
int device_id = 0; /* 複数GPUがある場合には0以外も設定可能 */
cudaSetDevice(device_id);
kernel<<<1,1>>>();
cudaThreadSynchronize();
return 0;
}
- ZOTAC GeForce GTX TITAN BLACK グラフィックスボード 日本正規代理.../ZOTAC
- Amazon.co.jp
- GeForceのラインナップとしては例外的にTeslaと同等の倍精度演算性能を持つボードです。CUDAでバリバリ開発するなら持っておいて損はありません。