[Do Not Merge] Float16 benchmark

paddle/fluid/inference/tests/book/CMakeLists.txt

@@ -26,6 +26,7 @@ endfunction(inference_test)
 
 inference_test(fit_a_line)
 inference_test(image_classification ARGS vgg resnet)
+inference_test(image_classification_float16 ARGS vgg resnet)
 inference_test(label_semantic_roles)
 inference_test(recognize_digits ARGS mlp conv)
 inference_test(recommender_system)

paddle/fluid/inference/tests/book/test_inference_image_classification.cc

@@ -19,6 +19,7 @@ limitations under the License. */
 DEFINE_string(dirname, "", "Directory of the inference model.");
 DEFINE_int32(batch_size, 1, "Batch size of input data");
 DEFINE_int32(repeat, 1, "Running the inference program repeat times");
+DEFINE_string(data_set, "cifar10", "Data set to use");
 
 TEST(inference, image_classification) {
   if (FLAGS_dirname.empty() || FLAGS_batch_size < 1 || FLAGS_repeat < 1) {
@@ -35,22 +36,32 @@ TEST(inference, image_classification) {
   paddle::framework::LoDTensor input;
   // Use normilized image pixels as input data,
   // which should be in the range [0.0, 1.0].
-  SetupTensor<float>(input,
-                     {FLAGS_batch_size, 3, 32, 32},
-                     static_cast<float>(0),
-                     static_cast<float>(1));
+  if (FLAGS_data_set == "cifar10") {
+    SetupTensor<float>(input,
+                       {FLAGS_batch_size, 3, 32, 32},
+                       static_cast<float>(0),
+                       static_cast<float>(1));
+  } else {
+    SetupTensor<float>(input,
+                       {FLAGS_batch_size, 3, 224, 224},
+                       static_cast<float>(0),
+                       static_cast<float>(1));
+  }
+
   std::vector<paddle::framework::LoDTensor*> cpu_feeds;
   cpu_feeds.push_back(&input);
 
   paddle::framework::LoDTensor output1;
   std::vector<paddle::framework::LoDTensor*> cpu_fetchs1;
   cpu_fetchs1.push_back(&output1);
 
-  // Run inference on CPU
-  LOG(INFO) << "--- CPU Runs: ---";
-  TestInference<paddle::platform::CPUPlace>(
-      dirname, cpu_feeds, cpu_fetchs1, FLAGS_repeat);
-  LOG(INFO) << output1.dims();
+// Run inference on CPU
+/*
+LOG(INFO) << "--- CPU Runs: ---";
+TestInference<paddle::platform::CPUPlace>(
+    dirname, cpu_feeds, cpu_fetchs1, FLAGS_repeat);
+LOG(INFO) << output1.dims();
+*/
 
 #ifdef PADDLE_WITH_CUDA
   paddle::framework::LoDTensor output2;
@@ -63,6 +74,6 @@ TEST(inference, image_classification) {
       dirname, cpu_feeds, cpu_fetchs2, FLAGS_repeat);
   LOG(INFO) << output2.dims();
 
-  CheckError<float>(output1, output2);
+// CheckError<float>(output1, output2);
 #endif
 }

paddle/fluid/inference/tests/book/test_inference_image_classification_float16.cc

@@ -0,0 +1,82 @@
+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include <gtest/gtest.h>
+#include "gflags/gflags.h"
+#include "paddle/fluid/inference/tests/test_helper.h"
+#include "paddle/fluid/platform/float16.h"
+
+DEFINE_string(dirname, "", "Directory of the inference model.");
+DEFINE_int32(batch_size, 1, "Batch size of input data");
+DEFINE_int32(repeat, 1, "Running the inference program repeat times");
+DEFINE_string(data_set, "cifar10", "Data set to use");
+
+TEST(inference, image_classification) {
+  using float16 = paddle::platform::float16;
+
+  if (FLAGS_dirname.empty() || FLAGS_batch_size < 1 || FLAGS_repeat < 1) {
+    LOG(FATAL) << "Usage: ./example --dirname=path/to/your/model "
+                  "--batch_size=1 --repeat=1";
+  }
+
+  LOG(INFO) << "FLAGS_dirname: " << FLAGS_dirname << std::endl;
+  std::string dirname = FLAGS_dirname;
+
+  // 0. Call `paddle::framework::InitDevices()` initialize all the devices
+  // In unittests, this is done in paddle/testing/paddle_gtest_main.cc
+
+  paddle::framework::LoDTensor input;
+  // Use normilized image pixels as input data,
+  // which should be in the range [0.0, 1.0].
+  if (FLAGS_data_set == "cifar10") {
+    SetupTensor<float16>(input,
+                         {FLAGS_batch_size, 3, 32, 32},
+                         static_cast<float16>(0),
+                         static_cast<float16>(1));
+  } else {
+    SetupTensor<float16>(input,
+                         {FLAGS_batch_size, 3, 224, 224},
+                         static_cast<float16>(0),
+                         static_cast<float16>(1));
+  }
+
+  std::vector<paddle::framework::LoDTensor*> cpu_feeds;
+  cpu_feeds.push_back(&input);
+
+  paddle::framework::LoDTensor output1;
+  std::vector<paddle::framework::LoDTensor*> cpu_fetchs1;
+  cpu_fetchs1.push_back(&output1);
+
+// Run inference on CPU
+/*
+LOG(INFO) << "--- CPU Runs: ---";
+TestInference<paddle::platform::CPUPlace>(
+    dirname, cpu_feeds, cpu_fetchs1, FLAGS_repeat);
+LOG(INFO) << output1.dims();
+*/
+
+#ifdef PADDLE_WITH_CUDA
+  paddle::framework::LoDTensor output2;
+  std::vector<paddle::framework::LoDTensor*> cpu_fetchs2;
+  cpu_fetchs2.push_back(&output2);
+
+  // Run inference on CUDA GPU
+  LOG(INFO) << "--- GPU Runs: ---";
+  TestInference<paddle::platform::CUDAPlace>(
+      dirname, cpu_feeds, cpu_fetchs2, FLAGS_repeat);
+  LOG(INFO) << output2.dims();
+
+// CheckError<float>(output1, output2);
+#endif
+}

paddle/fluid/operators/conv_cudnn_op.cu.cc

@@ -12,6 +12,8 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 See the License for the specific language governing permissions and
 limitations under the License. */
 
+#include <iostream>
+#include <typeindex>
 #include "paddle/fluid/framework/eigen.h"
 #include "paddle/fluid/framework/op_registry.h"
 #include "paddle/fluid/memory/memory.h"
@@ -112,6 +114,7 @@ class CUDNNConvOpKernel : public framework::OpKernel<T> {
     int group_offset_out =
         output_channels / groups * output_height * output_width * output_depth;
     int group_offset_filter = filter->numel() / groups;
+
     // ------------------- cudnn conv workspace ---------------------
     void* cudnn_workspace = nullptr;
     size_t workspace_size_in_bytes;  // final workspace to allocate.
@@ -128,13 +131,45 @@ class CUDNNConvOpKernel : public framework::OpKernel<T> {
         handle, cudnn_input_desc, cudnn_filter_desc, cudnn_conv_desc,
         cudnn_output_desc, CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
         workspace_size_limit, &algo));
+
+#if CUDA_VERSION >= 9000 && CUDNN_VERSION_MIN(7, 0, 1)
+    // Tensor core is supported since the volta GPU and
+    // is only enabled when input and filter data are float16
+    if (dev_ctx.GetComputeCapability() >= 70 &&
+        std::type_index(typeid(T)) ==
+            std::type_index(typeid(platform::float16))) {
+      PADDLE_ENFORCE(platform::dynload::cudnnSetConvolutionMathType(
+          cudnn_conv_desc, CUDNN_TENSOR_OP_MATH));
+      // Currently tensor core is only enabled using this algo
+      algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM;
+    } else {
+      PADDLE_ENFORCE(platform::dynload::cudnnSetConvolutionMathType(
+          cudnn_conv_desc, CUDNN_DEFAULT_MATH));
+    }
+#endif
+
+    //    std::cout << "The chosen algorithm is " << static_cast<int>(algo)
+    //              << std::endl;
+
     // get workspace size able to allocate
     PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardWorkspaceSize(
         handle, cudnn_input_desc, cudnn_filter_desc, cudnn_conv_desc,
         cudnn_output_desc, algo, &workspace_size_in_bytes));
+
+    // It is possible for float16 on Volta GPU to allocate more memory than
+    // the limit because the algo is overrided to use tensor core.
+    PADDLE_ENFORCE_LE(workspace_size_in_bytes, workspace_size_limit,
+                      "workspace_size to be allocated exceeds the limit");
+
+    // if (workspace_size_in_bytes > workspace_size_limit) {
+    // std::cout << "Workspace size is " << workspace_size_in_bytes
+    // << std::endl;
+    // }
+
     // Allocate on GPU memory
     platform::CUDAPlace gpu = boost::get<platform::CUDAPlace>(ctx.GetPlace());
     cudnn_workspace = paddle::memory::Alloc(gpu, workspace_size_in_bytes);
+
     // ------------------- cudnn conv forward ---------------------
     ScalingParamType<T> alpha = 1.0f, beta = 0.0f;
     for (int i = 0; i < groups; i++) {

paddle/fluid/operators/math/math_function.cu

@@ -39,18 +39,33 @@ void gemm<platform::CUDADeviceContext, float16>(
   cublasOperation_t cuTransB =
       (transB == CblasNoTrans) ? CUBLAS_OP_N : CUBLAS_OP_T;
 
-  const half h_alpha = static_cast<const half>(alpha);
-  const half h_beta = static_cast<const half>(beta);
-  const half* h_A = reinterpret_cast<const half*>(A);
-  const half* h_B = reinterpret_cast<const half*>(B);
-  half* h_C = reinterpret_cast<half*>(C);
+  float h_alpha = float(alpha);
+  float h_beta = float(beta);
 
   // TODO(kexinzhao): add processing code for compute capability < 53 case
   PADDLE_ENFORCE_GE(context.GetComputeCapability(), 53,
-                    "cublas Hgemm requires GPU compute capability >= 53");
-  PADDLE_ENFORCE(platform::dynload::cublasHgemm(
-      context.cublas_handle(), cuTransB, cuTransA, N, M, K, &h_alpha, h_B, ldb,
-      h_A, lda, &h_beta, h_C, N));
+                    "cublas fp16 gemm requires GPU compute capability >= 53");
+
+  cublasGemmAlgo_t algo = CUBLAS_GEMM_DFALT;
+#if CUDA_VERSION >= 9000
+  if (context.GetComputeCapability() >= 70) {
+    PADDLE_ENFORCE(platform::dynload::cublasSetMathMode(context.cublas_handle(),
+                                                        CUBLAS_TENSOR_OP_MATH));
+    algo = CUBLAS_GEMM_DFALT_TENSOR_OP;
+  } else {
+    PADDLE_ENFORCE(platform::dynload::cublasSetMathMode(context.cublas_handle(),
+                                                        CUBLAS_DEFAULT_MATH));
+  }
+#endif
+
+  // cublasHgemm does true FP16 computation which is slow for non-Volta
+  // GPUs. So use cublasGemmEx instead which does pesudo FP16 computation:
+  // input/output in fp16, computation in fp32, which can also be accelerated
+  // using tensor cores in volta GPUs.
+  PADDLE_ENFORCE(platform::dynload::cublasGemmEx(
+      context.cublas_handle(), cuTransB, cuTransA, N, M, K, &h_alpha, B,
+      CUDA_R_16F, ldb, A, CUDA_R_16F, lda, &h_beta, C, CUDA_R_16F, N,
+      CUDA_R_32F, algo));
 }
 
 template <>

paddle/fluid/operators/save_load_op_test.cc

@@ -33,9 +33,9 @@ TEST(SaveLoadOp, CPU) {
   expect_lod[0].push_back(3);
 
   tensor->set_lod(expect_lod);
-  int* expect = tensor->mutable_data<int>(place);
+  float* expect = tensor->mutable_data<float>(place);
   for (int64_t i = 0; i < tensor->numel(); ++i) {
-    expect[i] = static_cast<int>(i);
+    expect[i] = static_cast<float>(i);
   }
   paddle::framework::AttributeMap attrs;
   attrs.insert({"file_path", std::string("tensor.save")});
@@ -49,7 +49,7 @@ TEST(SaveLoadOp, CPU) {
   auto load_op = paddle::framework::OpRegistry::CreateOp(
       "load", {}, {{"Out", {"out_var"}}}, attrs);
   load_op->Run(scope, place);
-  int* actual = target->data<int>();
+  float* actual = target->data<float>();
   for (int64_t i = 0; i < tensor->numel(); ++i) {
     EXPECT_EQ(expect[i], actual[i]);
   }

paddle/fluid/operators/save_op.cc

@@ -18,6 +18,7 @@ limitations under the License. */
 #include <numeric>
 
 #include "paddle/fluid/framework/data_type.h"
+#include "paddle/fluid/framework/data_type_transform.h"
 #include "paddle/fluid/framework/framework.pb.h"
 #include "paddle/fluid/framework/lod_tensor.h"
 #include "paddle/fluid/framework/op_registry.h"
@@ -96,7 +97,45 @@ class SaveOp : public framework::OperatorBase {
     platform::DeviceContextPool &pool = platform::DeviceContextPool::Instance();
     auto &dev_ctx = *pool.Get(place);
 
-    framework::SerializeToStream(fout, tensor, dev_ctx);
+    auto in_dtype =
+        static_cast<framework::proto::VarType::Type>(Attr<int>("in_dtype"));
+    auto out_dtype =
+        static_cast<framework::proto::VarType::Type>(Attr<int>("out_dtype"));
+
+    PADDLE_ENFORCE_EQ(
+        static_cast<int>(framework::ToDataType(tensor.type())),
+        static_cast<int>(in_dtype),
+        "the tensor dtype does not match the attr of the save op");
+    /*
+        std::cout << std::endl
+                  << "filename is " << filename << ", var name is " << iname
+                  << std::endl
+                  << "in_dtype is " << static_cast<int>(in_dtype)
+                  << ", out_dtype is " << static_cast<int>(out_dtype) <<
+       std::endl;
+
+        std::cout << "before the conversion or not, the dtype is "
+                  << static_cast<int>(framework::ToDataType(tensor.type()))
+                  << std::endl;
+    */
+
+    if (in_dtype != out_dtype) {
+      // std::cout << "in_dtype and out_dtype not equal, start converting..."
+      //          << std::endl;
+      auto in_kernel_type = framework::OpKernelType(in_dtype, place);
+      auto out_kernel_type = framework::OpKernelType(out_dtype, place);
+      framework::LoDTensor out;
+      framework::TransDataType(in_kernel_type, out_kernel_type, tensor, &out);
+      std::cout << "after the conversion, the dtype is "
+                << static_cast<int>(framework::ToDataType(out.type()))
+                << std::endl;
+      framework::SerializeToStream(fout, out, dev_ctx);
+    } else {
+      // std::cout << "no conversion performed, the dtype is "
+      //          << static_cast<int>(framework::ToDataType(tensor.type()))
+      //          << std::endl;
+      framework::SerializeToStream(fout, tensor, dev_ctx);
+    }
   }
 };
 
@@ -114,6 +153,14 @@ This operator will serialize and write a tensor variable to file on disk.
                   "(boolean, default true)"
                   "Overwrite the output file if exist")
         .SetDefault(true);
+    AddAttr<int>("in_dtype",
+                 "(int, default 5)"
+                 "The data type of the input tensor")
+        .SetDefault(static_cast<int>(framework::proto::VarType::FP32));
+    AddAttr<int>("out_dtype",
+                 "(int, default 5)"
+                 "The data type of the converted tensor to be saved")
+        .SetDefault(static_cast<int>(framework::proto::VarType::FP32));
     AddAttr<std::string>("file_path",
                          "(string)"
                          "The \"file_path\" where the variable will be saved.")

paddle/fluid/platform/cudnn_helper.h

@@ -257,9 +257,12 @@ class ScopedConvolutionDescriptor {
     }
 #endif
 
+    cudnnDataType_t compute_type =
+        (type == CUDNN_DATA_DOUBLE) ? CUDNN_DATA_DOUBLE : CUDNN_DATA_FLOAT;
     PADDLE_ENFORCE(dynload::cudnnSetConvolutionNdDescriptor(
         desc_, pads.size(), pads.data(), strides.data(), dilations.data(),
-        CUDNN_CROSS_CORRELATION, type));
+        CUDNN_CROSS_CORRELATION, compute_type));
+
     return desc_;
   }
 

paddle/fluid/platform/dynload/cublas.cc

@@ -24,6 +24,10 @@ void *cublas_dso_handle = nullptr;
 
 CUBLAS_BLAS_ROUTINE_EACH(DEFINE_WRAP);
 
+#ifdef CUBLAS_BLAS_ROUTINE_EACH_R2
+CUBLAS_BLAS_ROUTINE_EACH_R2(DEFINE_WRAP);
+#endif
+
 }  // namespace dynload
 }  // namespace platform
 }  // namespace paddle