Support float8 data type #64735

Wangzheee · 2024-05-30T03:56:30Z

PR Category

Performance Optimization

PR Types

New features

Description

pcard-71500

Add float8 data type
Support calculation(math) and cast, transpose of float8 in GPU/CPU
Support float8 matmul (cuBLASLt)

paddle-bot · 2024-05-30T03:56:35Z

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

CLAassistant · 2024-05-30T03:56:36Z

All committers have signed the CLA.

CLAassistant · 2024-05-30T03:56:37Z

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you all sign our Contributor License Agreement before we can accept your contribution.
2 out of 3 committers have signed the CLA.

✅ Wangzheee
✅ lizexu123
❌ Wanglongzhi2001

Wanglongzhi2001 seems not to be a GitHub user. You need a GitHub account to be able to sign the CLA. If you have already a GitHub account, please add the email address used for this commit to your account.
_{You have signed the CLA already but the status is still pending? Let us recheck it.}

phlrain · 2024-06-05T08:23:19Z

这里的文件是不是能拆分合入下

vivienfanghuagood · 2024-06-05T14:00:35Z

paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h

+      cudaGetDevice(&dev);
+      if (dev == 0) {
+        std::ofstream outfile;
+        outfile.open(config_filename_, std::ios::out | std::ios::trunc);


为啥要在析构函数里写文件？

vivienfanghuagood · 2024-06-05T14:01:10Z

paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h

+    infile.close();
+  }
+
+  std::string config_filename_{"/tmp/paddle_cublaslt_cache"};


这里采用全局唯一的命名，如果多实例运行怎么办？

vivienfanghuagood · 2024-06-05T14:01:34Z

paddle/phi/kernels/fusion/fp8_gemm/fp8_gemm_with_cublasLt/cublaslt_gemm.h

+  std::string config_filename_{"/tmp/paddle_cublaslt_cache"};
+  std::unordered_map<int64_t, cublasLtMatmulAlgo_t> map_;
+  int search_times_;
+  const int requested_algo_count_ = 100;


这些参数建议通过环境变量传入

Wangzheee · 2024-06-06T06:13:13Z

这里的文件是不是能拆分合入下

好的，已拆分

qingqing01 · 2024-06-13T04:58:50Z

paddle/fluid/pybind/tensor.cc

+  } else if (dst->place() == phi::CPUPlace() && place == phi::GPUPlace()) {
+    cudaMemcpy(
+        dst->Holder()->ptr(), src.data(), src.size(), cudaMemcpyDeviceToHost);
+  }


其他TensorCopyFrom接口不能用的原因是？

需要实现else分支，如果没有else if分支改成else，并判断place设置的正确性

这个api在开发阶段组网用的，最终发现不需要了，已删掉

qingqing01 · 2024-06-13T05:00:19Z

paddle/phi/backends/gpu/cuda/cudnn_desc.h

+#if CUDNN_VERSION_MIN(8, 6, 0) && CUDA_VERSION >= 11800
+    case phi::DataType::FLOAT8_E4M3FN:
+      type = CUDNN_DATA_FP8_E4M3;
+      break;


这里不支持 float8_e5m2 类型吗？

https://docs.nvidia.com/deeplearning/cudnn/latest/developer/graph-api.html#id33
ConvolutionFwd上建议E4M3
这里也加上E5M2，给后续的功能使用

qingqing01 · 2024-06-13T05:12:31Z

test/legacy_test/test_matmul_fp8_op.py

+    def config(self):
+        self.dtype = 'float8_e4m3fn'
+        self.rtol = 0.6
+        self.atol = 7.6


这里rtol、atol这么大？

yuanlehome · 2024-06-18T07:42:39Z

paddle/fluid/platform/e4m3.h

@@ -0,0 +1,24 @@
+/* Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.


这个文件名有问题吧，需要加上float_前缀？float8_e4m3fn.h

yuanlehome · 2024-06-18T07:43:00Z

paddle/fluid/platform/e5m2.h

@@ -0,0 +1,24 @@
+/* Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.


文件名改为float8_e5m2.h？

yuanlehome · 2024-06-18T07:44:02Z

paddle/fluid/pybind/tensor.cc

@@ -67,6 +67,7 @@ limitations under the License. */
 #include "paddle/fluid/imperative/amp_auto_cast.h"
 #include "paddle/fluid/imperative/layer.h"
 #include "paddle/fluid/memory/allocation/allocator_strategy.h"
+#include "paddle/fluid/memory/memcpy.h"


这个头文件是否是必须的？感觉可以删掉

可以去掉

yuanlehome · 2024-06-18T07:44:41Z

paddle/fluid/pybind/tensor_py.h

+
+  static std::string format() {
+    // Note: "E4M3" represents float8_e4m3fn.
+    return "c";


return "c" 还是 "E4M3FN"?

yuanlehome · 2024-06-18T07:45:28Z

paddle/fluid/pybind/tensor_py.h

+
+  static std::string format() {
+    // Note: "E5M2" represents float8_e5m2.
+    return "E4M3";


return "E4M3" 还是 "E5M2" ？

Tom-Zheng · 2024-06-07T01:15:16Z

paddle/fluid/pybind/tensor_py.h

+
+  static std::string format() {
+    // Note: "E5M2" represents float8_e5m2.
+    return "E4M3";


Why E5M2's format is "E4M3"?

已修改～

Tom-Zheng · 2024-06-07T01:16:57Z

paddle/fluid/framework/data_type.cc

@@ -61,6 +61,8 @@ static DataTypeMap* InitDataTypeMap() {
  _ForEachDataType_(RegType);
  // Register pstring individually
  RegType(pstring, proto::VarType::PSTRING);
+  RegType(::paddle::platform::float8_e5m2, proto::VarType::FP8_E5M2);
+  RegType(::paddle::platform::float8_e4m3fn, proto::VarType::FP8_E4M3FN);


Sorry, what does "FN" stand for?

表示是否支持NaN、0、Inf, 与ONNX格式相同：https://onnx.ai/onnx/technical/float8.html#e4m3fn-and-e5m2

Aurelius84

LGTM for full/cast kernel add more dtype but full_grad/cast_grad have not registered these dtypes

sunzhongkai588

LGTM，后续在 API 的 doctoring 里，和 dcos 仓库里对应的 API 中文文档中，增加一下新增数据类型支持的描述，以便用户可以感知

* Support float8 data type

Wangzheee requested review from XiaoguangHu01, zhiqiu, Xreki, qili93 and Aurelius84 as code owners May 30, 2024 03:56

Wangzheee requested review from SigureMo, gouzil, cxxly, xiaoguoguo626807, changeyoung98, risemeup1, zhangbo9674, XieYunshen, zhwesky2010, wanghuancoder, LiYuRio, ForFishes, phlrain, zyfncg, YuanRisheng, Charles-hit, cyber-pioneer and JiabinYang as code owners June 5, 2024 07:51

Wangzheee changed the title ~~FP8 Gemm Fusion(api, op, kernel, type define)~~ Support FP8 Jun 5, 2024

Wangzheee changed the title ~~Support FP8~~ Support float8 Jun 5, 2024

vivienfanghuagood reviewed Jun 5, 2024

View reviewed changes

Wangzheee changed the title ~~Support float8~~ Support float8 data type Jun 6, 2024

Wangzheee force-pushed the fp8_cutlass branch from 3754355 to 8354c7c Compare June 11, 2024 11:16

qingqing01 reviewed Jun 13, 2024

View reviewed changes

qingqing01 previously approved these changes Jun 17, 2024

View reviewed changes

Wangzheee dismissed qingqing01’s stale review via 8fa6be3 June 18, 2024 03:06

Wangzheee added 9 commits June 18, 2024 03:33

Support float8 data type

c83e846

fix

b2e9e90

fix

3ecffb3

fix

d02c137

fix

d6a92af

fix

e10a6ce

fix

84d1c95

fix

9fb5fa8

fix

8f3d2bc

Wangzheee force-pushed the fp8_cutlass branch from 8fa6be3 to 8f3d2bc Compare June 18, 2024 03:35

yuanlehome reviewed Jun 18, 2024

View reviewed changes

fix

81380a4

wanghuancoder previously approved these changes Jun 19, 2024

View reviewed changes

fix

8992f36

Wangzheee dismissed wanghuancoder’s stale review via 8992f36 June 19, 2024 02:06

Merge branch 'develop' into fp8_cutlass

fc8ceff

Tom-Zheng reviewed Jun 19, 2024

View reviewed changes

Aurelius84 approved these changes Jun 19, 2024

View reviewed changes

qingqing01 approved these changes Jun 19, 2024

View reviewed changes

sunzhongkai588 approved these changes Jun 19, 2024

View reviewed changes

raindrops2sea approved these changes Jun 19, 2024

View reviewed changes

Wangzheee merged commit 65944e9 into PaddlePaddle:develop Jun 19, 2024
32 of 33 checks passed

co63oc pushed a commit to co63oc/Paddle that referenced this pull request Jun 25, 2024

Support float8 data type (PaddlePaddle#64735)

b0638ab

* Support float8 data type

		@@ -0,0 +1,24 @@
		/* Copyright (c) 2024 PaddlePaddle Authors. All Rights Reserved.

Support float8 data type #64735

Support float8 data type #64735

Uh oh!

Conversation

Wangzheee commented May 30, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

PR Category

PR Types

Description

Uh oh!

paddle-bot bot commented May 30, 2024

Uh oh!

CLAassistant commented May 30, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

CLAassistant commented May 30, 2024

Uh oh!

phlrain commented Jun 5, 2024

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Wangzheee commented Jun 6, 2024

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Aurelius84 left a comment

Choose a reason for hiding this comment

Uh oh!

sunzhongkai588 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Wangzheee commented May 30, 2024 •

edited

Loading

CLAassistant commented May 30, 2024 •

edited

Loading