[clang] 0da7c12 - [AMD-GPU] Fix smfmac builtin target (#193999)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 24 13:49:40 PDT 2026
Author: Kartik Ohlan
Date: 2026-04-24T16:49:35-04:00
New Revision: 0da7c12e43bf8b995c5fdd5e62180546fa007338
URL: https://github.com/llvm/llvm-project/commit/0da7c12e43bf8b995c5fdd5e62180546fa007338
DIFF: https://github.com/llvm/llvm-project/commit/0da7c12e43bf8b995c5fdd5e62180546fa007338.diff
LOG: [AMD-GPU] Fix smfmac builtin target (#193999)
Fixes #193882
For builins mentioned below, changed "mai-insts" to "gfx940-insts" so it
can match the backend isGFX940Plus predicate.
Added test coverage and kept the structure similar to
`clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl`
__builtin_amdgcn_smfmac_f32_16x16x32_f16
__builtin_amdgcn_smfmac_f32_32x32x16_f16
__builtin_amdgcn_smfmac_f32_16x16x32_bf16
__builtin_amdgcn_smfmac_f32_32x32x16_bf16
__builtin_amdgcn_smfmac_i32_16x16x64_i8
__builtin_amdgcn_smfmac_i32_32x32x32_i8
Added:
clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
Modified:
clang/include/clang/Basic/BuiltinsAMDGPU.td
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 8bd5cc5704b7f..fc910123560a9 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -627,12 +627,12 @@ def __builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<16, f
def __builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<16, float>(int64_t, int64_t, _ExtVector<16, float>, _Constant int, _Constant int, _Constant int)", [Const], "fp8-insts">;
def __builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<16, float>(int64_t, int64_t, _ExtVector<16, float>, _Constant int, _Constant int, _Constant int)", [Const], "fp8-insts">;
def __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8 : AMDGPUBuiltin<"_ExtVector<16, float>(int64_t, int64_t, _ExtVector<16, float>, _Constant int, _Constant int, _Constant int)", [Const], "fp8-insts">;
-def __builtin_amdgcn_smfmac_f32_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, _Float16>, _ExtVector<8, _Float16>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "mai-insts">;
-def __builtin_amdgcn_smfmac_f32_32x32x16_f16 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<4, _Float16>, _ExtVector<8, _Float16>, _ExtVector<16, float>, int, _Constant int, _Constant int)", [Const], "mai-insts">;
-def __builtin_amdgcn_smfmac_f32_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, short>, _ExtVector<8, short>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "mai-insts">;
-def __builtin_amdgcn_smfmac_f32_32x32x16_bf16 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<4, short>, _ExtVector<8, short>, _ExtVector<16, float>, int, _Constant int, _Constant int)", [Const], "mai-insts">;
-def __builtin_amdgcn_smfmac_i32_16x16x64_i8 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<4, int>, int, _Constant int, _Constant int)", [Const], "mai-insts">;
-def __builtin_amdgcn_smfmac_i32_32x32x32_i8 : AMDGPUBuiltin<"_ExtVector<16, int>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<16, int>, int, _Constant int, _Constant int)", [Const], "mai-insts">;
+def __builtin_amdgcn_smfmac_f32_16x16x32_f16 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, _Float16>, _ExtVector<8, _Float16>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "gfx940-insts">;
+def __builtin_amdgcn_smfmac_f32_32x32x16_f16 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<4, _Float16>, _ExtVector<8, _Float16>, _ExtVector<16, float>, int, _Constant int, _Constant int)", [Const], "gfx940-insts">;
+def __builtin_amdgcn_smfmac_f32_16x16x32_bf16 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, short>, _ExtVector<8, short>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "gfx940-insts">;
+def __builtin_amdgcn_smfmac_f32_32x32x16_bf16 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<4, short>, _ExtVector<8, short>, _ExtVector<16, float>, int, _Constant int, _Constant int)", [Const], "gfx940-insts">;
+def __builtin_amdgcn_smfmac_i32_16x16x64_i8 : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<4, int>, int, _Constant int, _Constant int)", [Const], "gfx940-insts">;
+def __builtin_amdgcn_smfmac_i32_32x32x32_i8 : AMDGPUBuiltin<"_ExtVector<16, int>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<16, int>, int, _Constant int, _Constant int)", [Const], "gfx940-insts">;
def __builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "fp8-insts">;
def __builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "fp8-insts">;
def __builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<2, int>, _ExtVector<4, int>, _ExtVector<4, float>, int, _Constant int, _Constant int)", [Const], "fp8-insts">;
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
index c972b1611b6f7..9f97265759ba0 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl
@@ -80,12 +80,12 @@ void builtin_test_unsupported(double a_double, float a_float,
a_v16f = __builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8(a_long, a_long, a_v16f, 0, 0, 0); // expected-error {{'__builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8' needs target feature fp8-insts}}
a_v16f = __builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8(a_long, a_long, a_v16f, 0, 0, 0); // expected-error {{'__builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8' needs target feature fp8-insts}}
a_v16f = __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8(a_long, a_long, a_v16f, 0, 0, 0); // expected-error {{'__builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8' needs target feature fp8-insts}}
- a_v4f = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a_v4h, a_v8h, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x32_f16' needs target feature mai-insts}}
- a_v16f = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a_v4h, a_v8h, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x16_f16' needs target feature mai-insts}}
- a_v4f = __builtin_amdgcn_smfmac_f32_16x16x32_bf16(a_v4s, a_v8s, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x32_bf16' needs target feature mai-insts}}
- a_v16f = __builtin_amdgcn_smfmac_f32_32x32x16_bf16(a_v4s, a_v8s, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x16_bf16' needs target feature mai-insts}}
- a_v4i = __builtin_amdgcn_smfmac_i32_16x16x64_i8(a_v2i, a_v4i, a_v4i, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_i32_16x16x64_i8' needs target feature mai-insts}}
- a_v16i = __builtin_amdgcn_smfmac_i32_32x32x32_i8(a_v2i, a_v4i, a_v16i, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_i32_32x32x32_i8' needs target feature mai-insts}}
+ a_v4f = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a_v4h, a_v8h, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x32_f16' needs target feature gfx940-insts}}
+ a_v16f = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a_v4h, a_v8h, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x16_f16' needs target feature gfx940-insts}}
+ a_v4f = __builtin_amdgcn_smfmac_f32_16x16x32_bf16(a_v4s, a_v8s, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x32_bf16' needs target feature gfx940-insts}}
+ a_v16f = __builtin_amdgcn_smfmac_f32_32x32x16_bf16(a_v4s, a_v8s, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x16_bf16' needs target feature gfx940-insts}}
+ a_v4i = __builtin_amdgcn_smfmac_i32_16x16x64_i8(a_v2i, a_v4i, a_v4i, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_i32_16x16x64_i8' needs target feature gfx940-insts}}
+ a_v16i = __builtin_amdgcn_smfmac_i32_32x32x32_i8(a_v2i, a_v4i, a_v16i, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_i32_32x32x32_i8' needs target feature gfx940-insts}}
a_v4f = __builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8(a_v2i, a_v4i, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8' needs target feature fp8-insts}}
a_v4f = __builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8(a_v2i, a_v4i, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8' needs target feature fp8-insts}}
a_v4f = __builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8(a_v2i, a_v4i, a_v4f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8' needs target feature fp8-insts}}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
new file mode 100644
index 0000000000000..487d31539cf4d
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
@@ -0,0 +1,34 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx908 -verify -S -o - %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a -verify -S -o - %s
+
+typedef int v2i __attribute__((ext_vector_type(2)));
+typedef float v4f __attribute__((ext_vector_type(4)));
+typedef half v4h __attribute__((ext_vector_type(4)));
+typedef short v4s __attribute__((ext_vector_type(4)));
+typedef int v4i __attribute__((ext_vector_type(4)));
+typedef half v8h __attribute__((ext_vector_type(8)));
+typedef short v8s __attribute__((ext_vector_type(8)));
+typedef float v16f __attribute__((ext_vector_type(16)));
+typedef int v16i __attribute__((ext_vector_type(16)));
+
+void test_amdgcn_smfmac(global v4f *out4f,
+ global v16f *out16f,
+ global v4i *out4i,
+ global v16i *out16i,
+ v4h a4h, v4s a4s,
+ v8h b8h, v8s b8s,
+ v2i b2i,
+ v4f c4f,
+ v16f c16f,
+ v4i c4i,
+ v16i c16i,
+ int index)
+{
+ *out4f = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a4h, b8h, c4f, index, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_16x16x32_f16' needs target feature gfx940-insts}}
+ *out4f = __builtin_amdgcn_smfmac_f32_16x16x32_bf16(a4s, b8s, c4f, index, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_16x16x32_bf16' needs target feature gfx940-insts}}
+ *out16f = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a4h, b8h, c16f, index, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x16_f16' needs target feature gfx940-insts}}
+ *out16f = __builtin_amdgcn_smfmac_f32_32x32x16_bf16(a4s, b8s, c16f, index, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_f32_32x32x16_bf16' needs target feature gfx940-insts}}
+ *out4i = __builtin_amdgcn_smfmac_i32_16x16x64_i8(b2i, c4i, c4i, index, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_i32_16x16x64_i8' needs target feature gfx940-insts}}
+ *out16i = __builtin_amdgcn_smfmac_i32_32x32x32_i8(b2i, c4i, c16i, index, 0, 0); // expected-error{{'__builtin_amdgcn_smfmac_i32_32x32x32_i8' needs target feature gfx940-insts}}
+}
More information about the cfe-commits
mailing list