<h3 dir="auto">System information</h3>
<ul dir="auto">
<li><strong>Have I written custom code (as opposed to using a stock example script provided in TensorFlow)</strong>:<br>
Partly</li>
<li><strong>OS Platform and Distribution (e.g., Linux Ubuntu 16.04)</strong>:<br>
RHEL 7</li>
<li><strong>TensorFlow installed from (source or binary)</strong>:<br>
unknown</li>
<li><strong>TensorFlow version (use command below)</strong>:<br>
('unknown', '1.4.0')</li>
<li><strong>Python version</strong>:<br>
python2.7</li>
<li><strong>Bazel version (if compiling from source)</strong>:<br>
unknown</li>
<li><strong>GCC/Compiler version (if compiling from source)</strong>:<br>
gcc 4.8</li>
<li><strong>CUDA/cuDNN version</strong>:<br>
9.0 / 7.0</li>
<li><strong>GPU model and memory</strong>:<br>
Pascal P-100</li>
<li><strong>Exact command to reproduce</strong>:<br>
N/A</li>
</ul>
<h3 dir="auto">Describe the problem</h3>
<p dir="auto">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.<br>
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.</p>
<p dir="auto">So the performance <strong>decreases</strong> when using FP16 in real-world examples.</p>
<p dir="auto">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.</p>
<p dir="auto">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 <strong>maxwell_fp16_segmemm_fp16_128x128_nn</strong> for FP16 and <strong>sgemm_128x128x8_NN_vec</strong> for FP32.</p>
<p dir="auto">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 <strong>maxwell_hgemm_256x128_nn</strong> for matrix multiplication of 16k x 16k square matrices in FP16. (again  according to the nvprof profiler)</p>
<p dir="auto">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.</p>
<h3 dir="auto">Source code</h3>
<p dir="auto">Tensorflow Code snippet:</p>
<div class="snippet-clipboard-content notranslate position-relative overflow-auto" data-snippet-clipboard-copy-content="    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(&quot;Initialized&quot;)
          for i in range(FLAGS.times):
              out = session.run([tf_output])#, feed_dict=feed_dict)
          print(&quot;Done&quot;)"><pre class="notranslate"><code class="notranslate">    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")
</code></pre></div>
<p dir="auto">cuBlas Code FP16 (Snippet):</p>
<div class="snippet-clipboard-content notranslate position-relative overflow-auto" data-snippet-clipboard-copy-content="        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 **) &amp;d_a, m * k * sizeof (*a));    // device memory alloc for a
        cudaStat = cudaMalloc ((void **) &amp;d_b, k * n * sizeof (*b));    // device memory alloc for b
        cudaStat = cudaMalloc ((void **) &amp;d_c, m * n * sizeof (*c));    // device memory alloc for c
        stat = cublasCreate (&amp;handle);  // initialize CUBLAS context
        // copy matrices from the host to the device
        stat = cublasSetMatrix (m, k, sizeof (*a), a, m, d_a, m);   //a -&gt; d_a
        stat = cublasSetMatrix (k, n, sizeof (*b), b, k, d_b, k);   //b -&gt; d_b
        stat = cublasSetMatrix (m, n, sizeof (*c), c, m, d_c, m);   //c -&gt; 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, &amp;al, d_a, CUDA_R_16F, m, d_b, CUDA_R_16F, k, &amp;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 - &gt;c"><pre class="notranslate"><code class="notranslate">        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 **) &amp;d_a, m * k * sizeof (*a));    // device memory alloc for a
        cudaStat = cudaMalloc ((void **) &amp;d_b, k * n * sizeof (*b));    // device memory alloc for b
        cudaStat = cudaMalloc ((void **) &amp;d_c, m * n * sizeof (*c));    // device memory alloc for c
        stat = cublasCreate (&amp;handle);  // initialize CUBLAS context
        // copy matrices from the host to the device
        stat = cublasSetMatrix (m, k, sizeof (*a), a, m, d_a, m);   //a -&gt; d_a
        stat = cublasSetMatrix (k, n, sizeof (*b), b, k, d_b, k);   //b -&gt; d_b
        stat = cublasSetMatrix (m, n, sizeof (*c), c, m, d_c, m);   //c -&gt; 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, &amp;al, d_a, CUDA_R_16F, m, d_b, CUDA_R_16F, k, &amp;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 - &gt;c
</code></pre></div>
<p dir="auto"><a class="issue-link js-issue-link" data-error-text="Failed to load title" data-id="136722135" data-permission-text="Title is private" data-url="https://github.com/tensorflow/tensorflow/issues/1300" data-hovercard-type="issue" data-hovercard-url="/tensorflow/tensorflow/issues/1300/hovercard" href="https://github.com/tensorflow/tensorflow/issues/1300">#1300</a></p>
<p dir="auto">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?</p>
<p dir="auto">userid@ubuntu-WK-4xP100:<del>/weike/tensorflow-r0.11/tensorflow/models/image/cifar10$ vi cifar10_train.py<br>
userid@ubuntu-WK-4xP100:</del>/weike/tensorflow-r0.11/tensorflow/models/image/cifar10$ python cifar10_train.py<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcublas.so.8.0 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcudnn.so.5 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcufft.so.8.0 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcuda.so.1 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcurand.so.8.0 locally</p>
<blockquote>
<blockquote>
<p dir="auto">Downloading cifar-10-binary.tar.gz 100.0%<br>
...<br>
2016-11-14 00:07:57.143739: step 0, loss = 4.67 (12.1 examples/sec; 10.549 sec/batch)<br>
2016-11-14 00:07:58.395209: step 10, loss = 4.57 (1693.6 examples/sec; 0.076 sec/batch)<br>
2016-11-14 00:07:59.177525: step 20, loss = 4.97 (1668.9 examples/sec; 0.077 sec/batch)<br>
2016-11-14 00:07:59.957789: step 30, loss = 4.43 (1588.3 examples/sec; 0.081 sec/batch)<br>
2016-11-14 00:08:00.738431: step 40, loss = 4.52 (1690.5 examples/sec; 0.076 sec/batch)<br>
2016-11-14 00:08:01.501940: step 50, loss = 4.33 (1680.4 examples/sec; 0.076 sec/batch)<br>
2016-11-14 00:08:02.241604: step 60, loss = 4.20 (1733.4 examples/sec; 0.074 sec/batch)<br>
2016-11-14 00:08:03.001845: step 70, loss = 4.27 (1706.0 examples/sec; 0.075 sec/batch)<br>
2016-11-14 00:08:03.765522: step 80, loss = 4.18 (1601.3 examples/sec; 0.080 sec/batch)<br>
2016-11-14 00:08:04.516780: step 90, loss = 4.25 (1646.3 examples/sec; 0.078 sec/batch)</p>
</blockquote>
</blockquote>
<p dir="auto">userid@ubuntu-WK-4xP100:~/weike/tensorflow-r0.11/tensorflow/models/image/cifar10$ python cifar10_train.py --use_fp16<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcublas.so.8.0 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcudnn.so.5 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcufft.so.8.0 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcuda.so.1 locally<br>
I tensorflow/stream_executor/dso_loader.cc:111] successfully opened CUDA library libcurand.so.8.0 locally<br>
Filling queue with 20000 CIFAR images before starting to train. This will take a few minutes.<br>
..<br>
2016-11-14 00:09:09.382854: step 0, loss = 4.67 (12.3 examples/sec; 10.411 sec/batch)<br>
2016-11-14 00:09:11.923842: step 10, loss = 4.58 (659.6 examples/sec; 0.194 sec/batch)<br>
2016-11-14 00:09:13.918448: step 20, loss = 6.62 (460.6 examples/sec; 0.278 sec/batch)<br>
2016-11-14 00:09:16.211809: step 30, loss = 4.37 (583.7 examples/sec; 0.219 sec/batch)<br>
2016-11-14 00:09:18.327690: step 40, loss = 4.30 (618.3 examples/sec; 0.207 sec/batch)<br>
2016-11-14 00:09:20.395409: step 50, loss = 4.37 (643.8 examples/sec; 0.199 sec/batch)<br>
2016-11-14 00:09:22.466230: step 60, loss = 4.32 (574.0 examples/sec; 0.223 sec/batch)<br>
2016-11-14 00:09:24.533225: step 70, loss = 4.17 (646.2 examples/sec; 0.198 sec/batch)<br>
2016-11-14 00:09:26.609277: step 80, loss = 2.59 (601.5 examples/sec; 0.213 sec/batch)<br>
2016-11-14 00:09:28.703648: step 90, loss = 4.15 (625.5 examples/sec; 0.205 sec/batch)</p>