Fp8 gemm for gfx940 enabled. by wenchenvincent · Pull Request #8 · ROCm/TransformerEngine

wenchenvincent · 2023-04-20T15:13:13Z

No description provided.

…to NANOO mode.

HaiShaw

LGTM

May address later -

gfx940 to be extended
start to consider differentiate F8 type: kFloat8E4M3, kFloat8E5M2 between NANOO and OCP

HaiShaw · 2023-08-23T00:18:32Z

+#define HIP_HOST __host__ 
+
+#define E5M2_AMAX 57344.0
+#define E4M3_AMAX 240.0


We may consider 224 for better OCP interop

HaiShaw · 2023-08-23T00:18:55Z


  // constructor from float
-  explicit HIP_HOST_DEVICE hip_f8(float v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
+#if defined(__gfx940__)


May need to consider gfx941, etc.

This was addressed in a PR #9

HaiShaw · 2023-08-23T00:22:28Z

-  explicit HIP_HOST_DEVICE hip_f8(half v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) {
+#if defined(__gfx940__)
+  explicit HIP_DEVICE hip_f8(half v, hip_f8_rounding_mode rm=hip_f8_rounding_mode::standard, uint32_t rng=0) 
+	  : hip_f8((float)v, rm, rng)


have builtins for half ?
this seems to be slower.

No native instruction for fp8 upcasting to half.

HaiShaw · 2023-08-23T00:26:13Z

-  explicit inline HIP_HOST_DEVICE operator half() const {
+#if defined(__gfx940__)
+  explicit HIP_DEVICE inline operator half() const {
+    return __half(float(*this));


same concern of performance if there were half to fp8/bf8 builtins

HaiShaw · 2023-08-23T00:50:05Z

+	       //Currently does not support output of fp8 tensors
+		    "Only the following combinations of data types are enabled now!\n 1. input: fp16, output: fp16.\n \
+		    2. input: fp32, output: fp32.\n 3. input: fp8, output: fp32");
+


similar from previous review, this logic expression is too long, and may hurt performance

This only gets executed once per function call and it is on the CPU side. I wouldn't worry too much about the performance of it.

HaiShaw · 2023-08-23T00:53:46Z

    """

-    E4M3 = _FormatHelper(max_fwd=448, max_bwd=448)
+    E4M3 = _FormatHelper(max_fwd=240, max_bwd=240)


may add definition in one f8 header/place where E5M2_AMAX, E4M3_AMAX exist, and use them

This is at Python level and the f8 header is at C++ level. It is possible to get it done with some Python/C++ binding but it might require jumping through some hoops.

* Remove fp8_out from LN API Signed-off-by: Przemyslaw Tredak <ptredak@nvidia.com> * fix LN test Signed-off-by: Przemyslaw Tredak <ptredak@nvidia.com> * Fixes Signed-off-by: Przemyslaw Tredak <ptredak@nvidia.com> Signed-off-by: Przemyslaw Tredak <ptredak@nvidia.com> Co-authored-by: ksivamani <ksivamani@nvidia.com>

Applied review comments 1, 2, 3, 4, 6, 7, 8 from PR #667 in one pass. Comments 5 and 9 will be answered on-thread (see below). Applied ======= #7 cpp_extensions/gemm.py: gate NVTE_USE_GEMM_TRITON behind IS_HIP_EXTENSION. Our Triton kernels use gfx942/gfx950-specific MFMA instructions and autotune configs; refuse to enable on non-HIP builds. #1 + #4 + #8: deduplicate dtype utilities against triton_kernels/common.py, which is the authoritative source shared across all Triton kernel backends here. - gemm_common.py: drop our copies of _get_fp8_dtypes(), torch_to_te_dtype(), te_to_torch_dtype(). Keep is_fp8_dtype() (operates on tex.DType; common.py's is_fp8_torch_dtype operates on torch dtypes -- different signatures), reinterpret_as_fp8_tensor, getGemmOutputShape, product. - gemm_wrapper.py: import torch_dtype_to_te_dtype / te_dtype_to_torch_dtype from ..common; use those instead of the removed local names. - __init__.py: drop the redundant re-exports (torch_to_te_dtype, te_to_torch_dtype, _get_fp8_dtypes). Note in the file directing callers to triton_kernels.common. - test_gemm_kernel.py: import get_torch_e4m3_type, get_torch_e5m2_type, torch_dtype_to_te_dtype from triton_kernels.common; define a local _get_fp8_dtypes() shim over them to keep the test's parametrization tables unchanged. - test_gemm_mxfp8.py: previously used _get_e4m3_dtype() -- dropped along with the test that used it (see below). #2 + #3: add module docstrings to test_gemm.py, test_gemm_kernel.py, test_gemm_mxfp8.py making the scope of each file explicit and pointing at the sibling files. Prompted by wangye805's confusion between test_gemm.py and test_gemm_kernel.py, and by the question of whether MXFP8 is only tested in test_gemm_mxfp8.py. #6 test_numerics.py: rewrite the comment above and the reason= text for _skip_grouped_under_gemm_triton to make it explicit that the skip is a *backend mismatch* between the two sides of the comparison, not Triton non-determinism. - sequential side -> our Triton (via NVTE_USE_GEMM_TRITON=1) - grouped side -> hipBLASLt / CUTLASS / AITER-Triton grouped (controlled by the SEPARATE NVTE_USE_GROUPED_GEMM_TRITON env var) With just NVTE_USE_GEMM_TRITON=1 set, the two sides diverge in fp32 rounding because they run different kernels. Not our Triton being non-deterministic. Also removed test_mxfp8_kernel_with_simulated_data (an early kernel- bring-up smoke test that only asserted non-zero output; MXFP8 correctness is fully covered by test_gemm.py::test_triton_vs_*_mxfp8 with real MXFP8Tensor and both PyTorch/C++ references). To be answered on-thread ======================== #5 (conftest.py hook vs explicit skip lists) -- keep the hook: it fires only on three specific ValueError substrings from our own gate code, is self-retiring when the gates are relaxed, and explicit skips would require ~1300 marks. Detail in the PR reply. #9 (Float8TensorWrapper / MXFP8TensorWrapper vs the TE tensor types) -- defer; needs a side-by-side to see how much of the wrapper API is genuinely required for the Triton-kernel call surface (fields accessed, dimension reordering, missing-transpose fallbacks) vs. what could route through the TE tensor types directly. Verified on gfx950 / PyTorch 2.10 -- all three direct suites clean: test_gemm.py 212 pass / 72 skip test_gemm_kernel.py 450 pass / 606 skip test_gemm_mxfp8.py 2 pass (was 3; smoke test removed)

wenchenvincent added 3 commits April 5, 2023 04:31

Initial integration of rocBlas F8 GEMM without epilogues. Changed F8 …

32045db

…to NANOO mode.

Added support for bias data type different from matrix B or D.

8a4d985

Enabled fp8 native conversion instructions for gfx940.

ac69653

wenchenvincent requested a review from HaiShaw April 20, 2023 15:13

wenchenvincent changed the base branch from main to dev April 20, 2023 15:14

HaiShaw approved these changes Aug 23, 2023

View reviewed changes

wenchenvincent merged commit 7e5a89b into dev Aug 24, 2023

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Fp8 gemm for gfx940 enabled.#8

Fp8 gemm for gfx940 enabled.#8
wenchenvincent merged 3 commits into
devfrom
fp8_gemm_gfx940

wenchenvincent commented Apr 20, 2023

Uh oh!

HaiShaw left a comment

Uh oh!

HaiShaw Aug 23, 2023

Uh oh!

HaiShaw Aug 23, 2023

Uh oh!

wenchenvincent Aug 24, 2023

Uh oh!

HaiShaw Aug 23, 2023

Uh oh!

wenchenvincent Aug 24, 2023

Uh oh!

HaiShaw Aug 23, 2023

Uh oh!

HaiShaw Aug 23, 2023

Uh oh!

wenchenvincent Aug 24, 2023

Uh oh!

HaiShaw Aug 23, 2023

Uh oh!

wenchenvincent Aug 24, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

wenchenvincent commented Apr 20, 2023

Uh oh!

HaiShaw left a comment

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!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants