[AMD] NFC: Tidy up FP8 variant support cases (#7267)

Author: antiagainst

python/test/unit/language/test_compile_errors.py

@@ -7,7 +7,7 @@
 import triton.language as tl
 from triton.compiler.errors import CompilationError, CompileTimeAssertionFailure
 import traceback
-from triton._internal_testing import is_cuda, is_hip, is_hip_cdna3, is_hip_cdna4
+from triton._internal_testing import is_cuda, is_hip, is_hip_cdna3
 
 
 def format_exception(type, value, tb):
@@ -364,10 +364,9 @@ def test_fp8_support(fresh_triton_cache, dtype):
         if cc >= (8, 9):
             supported_dtypes.append(tl.float8e4nv)
     elif is_hip():
+        supported_dtypes.append(tl.float8e4nv)
         if is_hip_cdna3():
-            supported_dtypes += [tl.float8e4nv, tl.float8e4b8, tl.float8e5b16]
-        if is_hip_cdna4():
-            supported_dtypes += [tl.float8e4nv]
+            supported_dtypes += [tl.float8e4b8, tl.float8e5b16]
 
     @triton.jit
     def dtype_kernel(dtype: tl.constexpr):

third_party/amd/backend/compiler.py

@@ -37,7 +37,9 @@ class HIPOptions:
     debug: bool = False
     sanitize_overflow: bool = True
     arch: str = None
-    supported_fp8_dtypes: Tuple[str] = ("fp8e5", )
+    # We have native support for OCP fp8 variants since CNDA4/RDNA4. For earlier generations,
+    # we software emulate the support for them.
+    supported_fp8_dtypes: Tuple[str] = ("fp8e4nv", "fp8e5")
     deprecated_fp8_dot_operand_dtypes: Tuple[str] = ()
     default_dot_input_precision: str = "ieee"
     allowed_dot_input_precisions: Tuple[str] = ("ieee", )
@@ -113,11 +115,8 @@ def parse_options(self, opts) -> Any:
         if "supported_fp8_dtypes" not in opts:
             supported_fp8_dtypes = set(HIPOptions.supported_fp8_dtypes)
             if self.target.arch == 'gfx942':
-                supported_fp8_dtypes.update({'fp8e4nv', 'fp8e4b8', 'fp8e5b16'})
-            elif self.target.arch == 'gfx950':
-                supported_fp8_dtypes.update({'fp8e4nv', 'fp8e5'})
-            elif 'gfx12' in self.target.arch:
-                supported_fp8_dtypes.update({'fp8e4nv', 'fp8e5'})
+                # CDNA3/gfx942 has native support for AMD specific FP8 types.
+                supported_fp8_dtypes.update({'fp8e4b8', 'fp8e5b16'})
             args["supported_fp8_dtypes"] = tuple(sorted(supported_fp8_dtypes))
 
         if "enable_fp_fusion" not in opts: