如何在张量流上使用 fp16(Eigen::half) 进行卷积答案

【问题标题】：how to do convolution with fp16(Eigen::half) on tensorflow如何在张量流上使用 fp16(Eigen::half) 进行卷积
【发布时间】：2019-12-26 16:54:53
【问题描述】：

如何使用 tensorflow 在 GPU 上使用 fp16 进行卷积？（使用 __half 或 Eigen::half 的 python api）。

我想在 tensorflow 上用 fp16 测试一个模型，但我卡住了。实际上，我发现 tensorflow 中的 fp16 卷积似乎将 fp32 卷积的结果转换为 fp16，这不是我需要的。

我尝试为 tf.nn.conv2d 提供 fp16 格式的 fp16 输入，并为 tf.nn.conv2d 提供 fp32 格式的 fp16 输入（tf.cast 将其转换为 fp32）然后 tf.cast 将结果转换为 fp16 ，他们给出了完全相同的结果。但正如我所想，在 fp16 中进行卷积与在 fp32 中进行卷积然后将其转换为 fp16 是不同的，我错了吗？请帮助我，谢谢。

environment:
ubuntu 16.04
tensorflow 1.9.0
cuda 9.0
Tesla V100

import tensorflow as tf
import numpy as np
import os

def conv16_32(input, kernel): # fake fp16 convolution
    input = tf.cast(input, tf.float16)
    kernel = tf.cast(kernel, tf.float16)
    input = tf.cast(input, tf.float32)
    kernel = tf.cast(kernel, tf.float32)
    out = tf.nn.conv2d(input, kernel, [1,1,1,1], padding='VALID')
    out = tf.cast(out, tf.float16)
    out = tf.cast(out, tf.float64)
    return out

def conv16(input, kernel): # real fp16 convolution
    input = tf.cast(input, tf.float16)
    kernel = tf.cast(kernel, tf.float16)
    out = tf.nn.conv2d(input, kernel, [1,1,1,1], padding='VALID')
    out = tf.cast(out, tf.float64)
    return out

x = np.random.rand(16, 32, 32, 16).astype('float64')
w = np.random.rand(3, 3, 16, 16).astype('float64')
x = tf.get_variable('input', dtype=tf.float64, initializer=x)
w = tf.get_variable('weight', dtype=tf.float64, initializer=w)

out_16 = conv16(x, w)
out_16_32 = conv16_32(x, w)

os.environ['CUDA_VISIBLE_DEVICES'] = '1'
config = tf.ConfigProto()
config.gpu_options.allow_growth = True
sess = tf.Session(config = config)
sess.run(tf.global_variables_initializer())
sess.run(tf.local_variables_initializer())
print(sess.run(tf.reduce_max(out_16_32 - out_16)))

上述两个函数给出相同的结果，假设最终的“打印”结果为零。

fp16卷积和fp32卷积的结果应该不一样（在我看来）。如何使用 tensorflow 在 GPU 上使用真正的 fp16 进行卷积？（使用 __half 或 Eigen::half 的 python api）

【问题讨论】：

还有，当我将tensorflow设置为在CPU而不是GPU上工作时，结果还是一样，所以我认为我使用tensorflow fp16的方式错误。
感谢您的澄清，现在帖子已经完成，您有一个合法的关注/问题。
可能是小卷积的意外结果。使用 pytorch 再次测试。
就像评论一样，CUDA_VISIBLE_DEVICES 只有在导入之前更改它才会对 TensorFlow 产生任何影响。

标签： python c++ tensorflow half-precision-float

【解决方案1】：

我认为您正确地使用了这些操作。在您的示例中，您可以检查卷积操作是否确实具有正确的类型。

conv2d_op_16 = out_16.op.inputs[0].op
print(conv2d_op_16.name, conv2d_op_16.type, conv2d_op_16.get_attr('T'))
# Conv2D Conv2D <dtype: 'float16'>
conv2d_op_16_32 = out_16_32.op.inputs[0].op.inputs[0].op
print(conv2d_op_16_32.name, conv2d_op_16_32.type, conv2d_op_16_32.get_attr('T'))
# Conv2D_1 Conv2D <dtype: 'float32'>

而且 TensorFlow 确实为 fp16 for CPU 和 for GPU 注册了内核，因此没有理由认为正在做其他事情。我对 fp16 没有太多经验，所以我不确定零差是否“正常”，但似乎没有任何方式 conv16 使用 fp16 卷积以外的任何方法。

【讨论】：

在 pytorch 上也是如此。可能与 cuda fp16 实现中的 FMA 有关。无论如何，我认为 cuda 使用 fp16 优化了计算，并使其中间结果几乎完全精确。感谢您的回答。
如果你构造一个简单的卷积例子，应该在 f16 中溢出并作为 f16 执行你会发现它没有溢出，这意味着底层的算术确实是在 f32 中执行的，只有结果被转换到 f16。
@Szabolcs 我明白你的意思，谢谢你的例子。我可能错了，但是我认为，从 TensorFlow 的角度来看，一切都是 fp16。我不知道是否存在较低级别的堆栈，其中 Eigen 或 CUDA 由于某种原因在计算中的某个时刻切换到 fp32...例如，某些 GPU 不支持 fp16，或者速度较慢（见here)。
@jdehesa 我相信是这样，即使两个卷积输入的类型都是 f16，出于性能或兼容性原因，从事 Tensorflow 工作的人决定静默转换为 f32。我认为这不是一个好主意，我想测试 f16 卷积的准确性，但现在我做不到。

【解决方案2】：

我正在尝试解决相同的问题。下面是一些可以用来测试卷积的简单代码：

import tensorflow as tf
tf.enable_eager_execution()
input = tf.cast([[[[65519], [65519], [65519], [65519]]]], tf.float16) #BHWC
filter = tf.cast([[[[65519]], [[-65519]]]], tf.float16) #HWIO
tf.print(tf.nn.conv2d(input, filter, [1,1,1,1], "VALID"))

如果卷积在 fp16 中完成，这应该会溢出，但实际上不会在 Tensorflow 中溢出。我得到的结果是[[[[0][0][0]]]]，这表明卷积是在fp32中进行的。

编辑：解决方法是设置环境变量：

TF_FP16_CONV_USE_FP32_COMPUTE=0

这给出了结果[[[[inf][inf][inf]]]]，表明这次卷积是在fp16中进行的。看来您至少需要一个 10x0 的 GPU。

【讨论】：