[clang] [AMD-GPU] Fix smfmac builtin target (PR #193999)
Kartik Ohlan via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 24 10:58:05 PDT 2026
https://github.com/Ko496-glitch updated https://github.com/llvm/llvm-project/pull/193999
>From 87be41d478a526b54c0a89a2a716837288b14544 Mon Sep 17 00:00:00 2001
From: Kartik Ohlan <kartik7ohlan at gmail.com>
Date: Fri, 24 Apr 2026 12:01:27 -0400
Subject: [PATCH 1/4] Fixed main-insts -> gfx940 and test cases
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 12 ++++-----
.../builtins-amdgcn-smfmac-err.cl | 25 +++++++++++++++++++
2 files changed, 31 insertions(+), 6 deletions(-)
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index db0fbb4b048e2..e36b532ede829 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -623,12 +623,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-smfmac-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
new file mode 100644
index 0000000000000..16fe433417b0a
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
@@ -0,0 +1,25 @@
+// 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)));
+
+void test_amdgcn_smfmac(global v4f *out4f,
+ global v16f *out16f,
+ v4h a4h, v4s a4s,
+ v8h b8h, v8s b8s,
+ v4f c4f,
+ v16f c16f,
+ 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}}
+}
\ No newline at end of file
>From 5aa49c4ab6b609199095a821e79352876353b8b1 Mon Sep 17 00:00:00 2001
From: Kartik Ohlan <kartik7ohlan at gmail.com>
Date: Fri, 24 Apr 2026 12:33:36 -0400
Subject: [PATCH 2/4] replaced older expected-error
---
.../test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
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}}
>From e165a9465940758a8f3417285e45704685da39cd Mon Sep 17 00:00:00 2001
From: Kartik Ohlan <kartik7ohlan at gmail.com>
Date: Fri, 24 Apr 2026 13:42:50 -0400
Subject: [PATCH 3/4] Polised structure and added i8 variant
---
.../builtins-amdgcn-smfmac-err.cl | 22 +++++++++++++------
1 file changed, 15 insertions(+), 7 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
index 16fe433417b0a..86c7ad0036b40 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
@@ -9,17 +9,25 @@ 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 v16f *out16f,
+ global v4i *out4i,
+ global v16i *out16i,
v4h a4h, v4s a4s,
v8h b8h, v8s b8s,
+ v2i b2i,
v4f c4f,
- v16f c16f,
+ 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}}
-}
\ No newline at end of file
+*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}}
+}
>From d7f43a74f447d4af8739168e8a0ffc47d86eb5d2 Mon Sep 17 00:00:00 2001
From: Kartik Ohlan <kartik7ohlan at gmail.com>
Date: Fri, 24 Apr 2026 13:57:32 -0400
Subject: [PATCH 4/4] fixed the indentation and added //REQUIRES
---
.../CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl | 15 ++++++++-------
1 file changed, 8 insertions(+), 7 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
index 86c7ad0036b40..487d31539cf4d 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-smfmac-err.cl
@@ -1,3 +1,4 @@
+// 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
@@ -23,11 +24,11 @@ void test_amdgcn_smfmac(global v4f *out4f,
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}}
+{
+ *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