[llvm] [AMDGPU] Guard media instruction patterns with target feature (PR #217995)
Shilei Tian via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 21 11:08:00 PDT 2026
https://github.com/shiltian created https://github.com/llvm/llvm-project/pull/217995
As well as their corresponding intrinsics.
>From 6bc39cc1a0192c1f1f1f04516dee352ffb229220 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Fri, 21 Aug 2026 14:07:25 -0400
Subject: [PATCH] [AMDGPU] Guard media instruction patterns with target feature
As well as their corresponding intrinsics.
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 14 ++++--
llvm/lib/Target/AMDGPU/SIInstructions.td | 28 ++++++------
.../AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll | 44 +++++++++++++++++++
llvm/test/CodeGen/AMDGPU/sad-u32.ll | 40 +++++++++++++++++
4 files changed, 109 insertions(+), 17 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/sad-u32.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index b321b9155a199..191501b7862cb 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2471,14 +2471,11 @@ def int_amdgcn_lerp :
ClangBuiltin<"__builtin_amdgcn_lerp">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+let TargetFeatures = "sad-insts" in {
def int_amdgcn_sad_u8 :
ClangBuiltin<"__builtin_amdgcn_sad_u8">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
-def int_amdgcn_msad_u8 :
- ClangBuiltin<"__builtin_amdgcn_msad_u8">,
- PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
-
def int_amdgcn_sad_hi_u8 :
ClangBuiltin<"__builtin_amdgcn_sad_hi_u8">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
@@ -2486,15 +2483,24 @@ def int_amdgcn_sad_hi_u8 :
def int_amdgcn_sad_u16 :
ClangBuiltin<"__builtin_amdgcn_sad_u16">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+} // End TargetFeatures = "sad-insts"
+
+let TargetFeatures = "msad-insts" in
+def int_amdgcn_msad_u8 :
+ ClangBuiltin<"__builtin_amdgcn_msad_u8">,
+ PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+let TargetFeatures = "qsad-insts" in
def int_amdgcn_qsad_pk_u16_u8 :
ClangBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
PureIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty]>;
+let TargetFeatures = "mqsad-pk-insts" in
def int_amdgcn_mqsad_pk_u16_u8 :
ClangBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
PureIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty]>;
+let TargetFeatures = "mqsad-insts" in
def int_amdgcn_mqsad_u32_u8 :
ClangBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
PureIntrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty]>;
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index c2fe750371528..441098168f450 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -3100,20 +3100,22 @@ defm : SI_INDIRECT_Pattern <v32i32, i32, "V32">;
// SAD Patterns
//===----------------------------------------------------------------------===//
-def : GCNPat <
- (add (sub_oneuse (umax i32:$src0, i32:$src1),
- (umin i32:$src0, i32:$src1)),
- i32:$src2),
- (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
->;
+let SubtargetPredicate = HasSadInsts in {
+ def : GCNPat <
+ (add (sub_oneuse (umax i32:$src0, i32:$src1),
+ (umin i32:$src0, i32:$src1)),
+ i32:$src2),
+ (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
+ >;
-def : GCNPat <
- (add (select_oneuse (i1 (setugt i32:$src0, i32:$src1)),
- (sub i32:$src0, i32:$src1),
- (sub i32:$src1, i32:$src0)),
- i32:$src2),
- (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
->;
+ def : GCNPat <
+ (add (select_oneuse (i1 (setugt i32:$src0, i32:$src1)),
+ (sub i32:$src0, i32:$src1),
+ (sub i32:$src1, i32:$src0)),
+ i32:$src2),
+ (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
+ >;
+} // End SubtargetPredicate = HasSadInsts
//===----------------------------------------------------------------------===//
// Conversion Patterns
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll
new file mode 100644
index 0000000000000..3135b7953f60d
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll
@@ -0,0 +1,44 @@
+; RUN: not llc -global-isel=0 -mtriple=amdgpu12.5-amd-amdhsa -filetype=null < %s 2>&1 | FileCheck %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgpu12.5-amd-amdhsa -filetype=null < %s 2>&1 | FileCheck %s
+
+; CHECK: error: {{.*}} llvm.amdgcn.sad.u8 requires target feature 'sad-insts'
+define i32 @test_sad_u8(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.sad.u8(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.sad.hi.u8 requires target feature 'sad-insts'
+define i32 @test_sad_hi_u8(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.sad.hi.u8(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.sad.u16 requires target feature 'sad-insts'
+define i32 @test_sad_u16(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.sad.u16(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.msad.u8 requires target feature 'msad-insts'
+define i32 @test_msad_u8(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.msad.u8(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.qsad.pk.u16.u8 requires target feature 'qsad-insts'
+define i64 @test_qsad_pk_u16_u8(i64 %a, i32 %b, i64 %c) {
+ %result = call i64 @llvm.amdgcn.qsad.pk.u16.u8(i64 %a, i32 %b, i64 %c)
+ ret i64 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.mqsad.pk.u16.u8 requires target feature 'mqsad-pk-insts'
+define i64 @test_mqsad_pk_u16_u8(i64 %a, i32 %b, i64 %c) {
+ %result = call i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %a, i32 %b, i64 %c)
+ ret i64 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.mqsad.u32.u8 requires target feature 'mqsad-insts'
+define <4 x i32> @test_mqsad_u32_u8(i64 %a, i32 %b, <4 x i32> %c) {
+ %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %a, i32 %b, <4 x i32> %c)
+ ret <4 x i32> %result
+}
diff --git a/llvm/test/CodeGen/AMDGPU/sad-u32.ll b/llvm/test/CodeGen/AMDGPU/sad-u32.ll
new file mode 100644
index 0000000000000..4297a5963d1b0
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/sad-u32.ll
@@ -0,0 +1,40 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -global-isel=0 -mtriple=amdgpu9.00 < %s | FileCheck -check-prefix=GFX9 %s
+; RUN: llc -global-isel=1 -mtriple=amdgpu9.00 < %s | FileCheck -check-prefix=GFX9 %s
+; RUN: llc -global-isel=0 -mtriple=amdgpu12.5 < %s | FileCheck -check-prefix=GFX12-5 %s
+; RUN: llc -global-isel=1 -mtriple=amdgpu12.5 < %s | FileCheck -check-prefix=GFX12-5 %s
+
+define i32 @usad(i32 %a, i32 %b, i32 %c) {
+; GFX9-LABEL: usad:
+; GFX9: ; %bb.0:
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-NEXT: v_sad_u32 v0, v0, v1, v2
+; GFX9-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX12-5-SDAG-LABEL: usad:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: v_min_u32_e32 v3, v0, v1
+; GFX12-5-SDAG-NEXT: v_max_u32_e32 v0, v0, v1
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX12-5-SDAG-NEXT: v_sub_nc_u32_e32 v0, v0, v3
+; GFX12-5-SDAG-NEXT: v_add_nc_u32_e32 v0, v0, v2
+; GFX12-5-SDAG-NEXT: s_set_pc_i64 s[30:31]
+;
+; GFX12-5-GISEL-LABEL: usad:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: v_max_u32_e32 v3, v0, v1
+; GFX12-5-GISEL-NEXT: v_min_u32_e32 v0, v0, v1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX12-5-GISEL-NEXT: v_sub_nc_u32_e32 v0, v3, v0
+; GFX12-5-GISEL-NEXT: v_add_nc_u32_e32 v0, v0, v2
+; GFX12-5-GISEL-NEXT: s_set_pc_i64 s[30:31]
+ %max = call i32 @llvm.umax.i32(i32 %a, i32 %b)
+ %min = call i32 @llvm.umin.i32(i32 %a, i32 %b)
+ %diff = sub i32 %max, %min
+ %r = add i32 %diff, %c
+ ret i32 %r
+}
More information about the llvm-commits
mailing list