Issue 1:
System information

Have I written custom code (as opposed to using a stock example script provided in TensorFlow):

Partly

OS Platform and Distribution (e.g., Linux Ubuntu 16.04):

RHEL 7

TensorFlow installed from (source or binary):

unknown

TensorFlow version (use command below):

('unknown', '1.4.0')

Python version:

python2.7

Bazel version (if compiling from source):

unknown

GCC/Compiler version (if compiling from source):

gcc 4.8

CUDA/cuDNN version:

9.0 / 7.0

GPU model and memory:

Pascal P-100

Exact command to reproduce:

N/A

Describe the problem

I tried running the tutorial/rnn/ptb and tutorial/images/mnist examples with the --use_fp16 option, to see if speedup was achievable by utilizing the new half-precision features.

It turned out that a single training step for MNIST with FP32 took 3.3ms, with FP16 it was 4ms. For PTB small (I had to use lstm_cell=basic, because other types are not yet supported in FP16), the WPS dropped from 24000 to 22000 when switching to FP16.

So the performance decreases when using FP16 in real-world examples.

To check this, I've created a small benchmark for myself (since I can't get the nightly tf build on my machine, I can't run the official benchmarks), which is basically one big matrix multiplication in Tensorflow.

I've run it with square matrices with a width of 8k, 16k and 32k. In each case, FP16 and FP32 yielded nearly the same runtime. Profiling with nvprof I found out that the used CUDA function is maxwell_fp16_segmemm_fp16_128x128_nn for FP16 and sgemm_128x128x8_NN_vec for FP32.

Since I wanted to double check if matrix multiplication in FP16 is really slower than in FP32 on my GPU, I tried to directly benchmark the GPU using cuBlas with a similar operation. It turns out that here, FP16 is nearly twice as fast as FP32. CuBlas internally uses maxwell_hgemm_256x128_nn for matrix multiplication of 16k x 16k square matrices in FP16. (again according to the nvprof profiler)

So I'm wondering why Tensorflow is unable to achieve similar results in terms of speed, or if I'm doing something wrong in my tests.

Source code

Tensorflow Code snippet:

graph = tf.Graph()
with graph.as_default():
tf_input1 = tf.Variable(tf.truncated_normal([FLAGS.size, FLAGS.size], dtype=get_dtype()))
tf_input2 = tf.Variable(tf.truncated_normal([FLAGS.size, FLAGS.size], dtype=get_dtype()))
tf_output = tf.matmul(tf_input1, tf_input2)

with tf.Session(graph=graph) as session:
tf.global_variables_initializer().run()
print("Initialized")
for i in range(FLAGS.times):
out = session.run([tf_output])#, feed_dict=feed_dict)
print("Done")

cuBlas Code FP16 (Snippet):

uint16_t *d_a; // d_a - a on the device
uint16_t *d_b; // d_b - b on the device
uint16_t *d_c; // d_c - c on the device
cudaStat = cudaMalloc ((void **) &d_a, m * k * sizeof (*a)); // device memory alloc for a
cudaStat = cudaMalloc ((void **) &d_b, k * n * sizeof (*b)); // device memory alloc for b
cudaStat = cudaMalloc ((void **) &d_c, m * n * sizeof (*c)); // device memory alloc for c
stat = cublasCreate (&handle); // initialize CUBLAS context
// copy matrices from the host to the device
stat = cublasSetMatrix (m, k, sizeof (*a), a, m, d_a, m); //a -> d_a
stat = cublasSetMatrix (k, n, sizeof (*b), b, k, d_b, k); //b -> d_b
stat = cublasSetMatrix (m, n, sizeof (*c), c, m, d_c, m); //c -> d_c
uint16_t al = FP_16_ONE; // al = 1
uint16_t bet = FP_16_ONE; // bet =1
// matrix - matrix multiplication : d_c = al*d_a *d_b + bet *d_c
// d_a -mxk matrix , d_b -kxn matrix , d_c -mxn matrix ;
// al ,bet -scalars

stat = cublasGemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &al, d_a, CUDA_R_16F, m, d_b, CUDA_R_16F, k, &bet, d_c, CUDA_R_16F, m, CUDA_R_16F, CUBLAS_GEMM_DEFAULT);

stat = cublasGetMatrix (m, n, sizeof (*c), d_c, m, c, m); // cp d_c - >c

Issue 2:
#1300

I have got a nvidia p100 GPU which is support fp16, and I run the TF case 'cifar10_train.py'. Without option '--use_fp16', the performance is also 1600 examples/sec, and with the option '--use_fp16', the performance down to 500 examples/sec. Any ideas about this issue?

userid@ubuntu-WK-4xP100:/weike/tensorflow-r0.11/tensorflow/models/image/cifar10$ vi cifar10_train.py

userid@ubuntu-WK-4xP100:/weike/tensorflow-r0.11/tensorflow/models/image/cifar10$ python cifar10_train.py

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcublas.so.8.0 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcudnn.so.5 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcufft.so.8.0 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcuda.so.1 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcurand.so.8.0 locally

Downloading cifar-10-binary.tar.gz 100.0%

...

2016-11-14 00:07:57.143739: step 0, loss = 4.67 (12.1 examples/sec; 10.549 sec/batch)

2016-11-14 00:07:58.395209: step 10, loss = 4.57 (1693.6 examples/sec; 0.076 sec/batch)

2016-11-14 00:07:59.177525: step 20, loss = 4.97 (1668.9 examples/sec; 0.077 sec/batch)

2016-11-14 00:07:59.957789: step 30, loss = 4.43 (1588.3 examples/sec; 0.081 sec/batch)

2016-11-14 00:08:00.738431: step 40, loss = 4.52 (1690.5 examples/sec; 0.076 sec/batch)

2016-11-14 00:08:01.501940: step 50, loss = 4.33 (1680.4 examples/sec; 0.076 sec/batch)

2016-11-14 00:08:02.241604: step 60, loss = 4.20 (1733.4 examples/sec; 0.074 sec/batch)

2016-11-14 00:08:03.001845: step 70, loss = 4.27 (1706.0 examples/sec; 0.075 sec/batch)

2016-11-14 00:08:03.765522: step 80, loss = 4.18 (1601.3 examples/sec; 0.080 sec/batch)

2016-11-14 00:08:04.516780: step 90, loss = 4.25 (1646.3 examples/sec; 0.078 sec/batch)

userid@ubuntu-WK-4xP100:~/weike/tensorflow-r0.11/tensorflow/models/image/cifar10$ python cifar10_train.py --use_fp16

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcublas.so.8.0 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcudnn.so.5 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcufft.so.8.0 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcuda.so.1 locally

I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcurand.so.8.0 locally

Filling queue with 20000 CIFAR images before starting to train. This will take a few minutes.

..

2016-11-14 00:09:09.382854: step 0, loss = 4.67 (12.3 examples/sec; 10.411 sec/batch)

2016-11-14 00:09:11.923842: step 10, loss = 4.58 (659.6 examples/sec; 0.194 sec/batch)

2016-11-14 00:09:13.918448: step 20, loss = 6.62 (460.6 examples/sec; 0.278 sec/batch)

2016-11-14 00:09:16.211809: step 30, loss = 4.37 (583.7 examples/sec; 0.219 sec/batch)

2016-11-14 00:09:18.327690: step 40, loss = 4.30 (618.3 examples/sec; 0.207 sec/batch)

2016-11-14 00:09:20.395409: step 50, loss = 4.37 (643.8 examples/sec; 0.199 sec/batch)

2016-11-14 00:09:22.466230: step 60, loss = 4.32 (574.0 examples/sec; 0.223 sec/batch)

2016-11-14 00:09:24.533225: step 70, loss = 4.17 (646.2 examples/sec; 0.198 sec/batch)

2016-11-14 00:09:26.609277: step 80, loss = 2.59 (601.5 examples/sec; 0.213 sec/batch)

2016-11-14 00:09:28.703648: step 90, loss = 4.15 (625.5 examples/sec; 0.205 sec/batch)