[clang] [llvm] AMDGPU: Add missing mqsad-insts to gfx13 frontend feature map (PR #213292)

via cfe-commits cfe-commits at lists.llvm.org
Fri Jul 31 08:37:10 PDT 2026 

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

<details>
<summary>Changes</summary>

fillAMDGCNFeatureMap omitted mqsad-insts for gfx1310/gfx13-generic, so
clang wrongly rejected __builtin_amdgcn_mqsad_u32_u8 on those targets even
though the backend enables the feature. Add it to the gfx13 case.

Co-authored-by: Claude (Claude-Opus-4.8)

---
Full diff: https://github.com/llvm/llvm-project/pull/213292.diff


3 Files Affected:

- (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl (+24) 
- (modified) llvm/lib/TargetParser/AMDGPUTargetParser.cpp (+1) 
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll (+23-10) 


``````````diff
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
index f5d341b31991c..4e48f4b4d9f5c 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
@@ -3,6 +3,7 @@
 
 // REQUIRES: amdgpu-registered-target
 
+typedef unsigned int __attribute__((ext_vector_type(4))) uint4;
 typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
 typedef float __attribute__((ext_vector_type(32))) float32;
 typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
@@ -171,3 +172,26 @@ void test_msad_u8(global unsigned int *out, unsigned int src0, unsigned int src1
   *out = __builtin_amdgcn_msad_u8(src0, src1, src2);
 }
 
+// CHECK-LABEL: @test_mqsad_u32_u8(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT:    [[SRC0_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
+// CHECK-NEXT:    [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[SRC2_ADDR:%.*]] = alloca <4 x i32>, align 16, addrspace(5)
+// CHECK-NEXT:    store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT:    store i64 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 8
+// CHECK-NEXT:    store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT:    store <4 x i32> [[SRC2:%.*]], ptr addrspace(5) [[SRC2_ADDR]], align 16
+// CHECK-NEXT:    [[TMP0:%.*]] = load i64, ptr addrspace(5) [[SRC0_ADDR]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = load <4 x i32>, ptr addrspace(5) [[SRC2_ADDR]], align 16
+// CHECK-NEXT:    [[TMP3:%.*]] = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 [[TMP0]], i32 [[TMP1]], <4 x i32> [[TMP2]])
+// CHECK-NEXT:    [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT:    store <4 x i32> [[TMP3]], ptr addrspace(1) [[TMP4]], align 16
+// CHECK-NEXT:    ret void
+//
+void test_mqsad_u32_u8(global uint4 *out, unsigned long src0, unsigned int src1, uint4 src2)
+{
+  *out = __builtin_amdgcn_mqsad_u32_u8(src0, src1, src2);
+}
+
diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
index c5df20fb6b63a..7d4baa2c0776e 100644
--- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
@@ -530,6 +530,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T,
     Features["qsad-insts"] = true;
     Features["mqsad-pk-insts"] = true;
     Features["msad-insts"] = true;
+    Features["mqsad-insts"] = true;
     Features["cvt-pknorm-vop2-insts"] = true;
     Features["cvt-pknorm-vop3-insts"] = true;
     Features["image-insts"] = true;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll
index 5d93599f985cb..41ba33f089d24 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll
@@ -1,11 +1,15 @@
-; RUN: llc -mtriple=amdgpu7.04 < %s | FileCheck -check-prefix=GCN %s
-; RUN: llc -mtriple=amdgpu8.03 < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgpu6.00 < %s | FileCheck -check-prefixes=GCN,GFX600 %s
+; RUN: llc -mtriple=amdgpu8.03 < %s | FileCheck -check-prefixes=GCN,GFX803 %s
+; RUN: llc -mtriple=amdgpu13.10 < %s | FileCheck -check-prefixes=GCN,GFX13 %s
 
 declare <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64, i32, <4 x i32>) #0
 
 ; GCN-LABEL: {{^}}v_mqsad_u32_u8_inline_integer_immediate:
-; GCN-DAG: v_mov_b32_e32 v0, v2
-; GCN-DAG: v_mov_b32_e32 v1, v3
+; GFX600-DAG: v_mov_b32_e32 v0, v2
+; GFX600-DAG: v_mov_b32_e32 v1, v3
+; GFX803-DAG: v_mov_b32_e32 v0, v2
+; GFX803-DAG: v_mov_b32_e32 v1, v3
+; GFX13-NOT: v_mov_b32_e32
 ; GCN: v_mqsad_u32_u8 v[2:5], v[0:1], v6, v[{{[0-9]+:[0-9]+}}]
 define amdgpu_kernel void @v_mqsad_u32_u8_inline_integer_immediate(ptr addrspace(1) %out, i64 %src, i32 %a) {
   %tmp = call i64 asm "v_lsrlrev_b64 $0, $1, 1", "={v[2:3]},v"(i64 %src) #0
@@ -17,8 +21,11 @@ define amdgpu_kernel void @v_mqsad_u32_u8_inline_integer_immediate(ptr addrspace
 }
 
 ; GCN-LABEL: {{^}}v_mqsad_u32_u8_non_immediate:
-; GCN-DAG: v_mov_b32_e32 v0, v2
-; GCN-DAG: v_mov_b32_e32 v1, v3
+; GFX600-DAG: v_mov_b32_e32 v0, v2
+; GFX600-DAG: v_mov_b32_e32 v1, v3
+; GFX803-DAG: v_mov_b32_e32 v0, v2
+; GFX803-DAG: v_mov_b32_e32 v1, v3
+; GFX13-NOT: v_mov_b32_e32
 ; GCN: v_mqsad_u32_u8 v[2:5], v[0:1], v6, v[{{[0-9]+:[0-9]+}}]
 define amdgpu_kernel void @v_mqsad_u32_u8_non_immediate(ptr addrspace(1) %out, i64 %src, i32 %a, <4 x i32> %b) {
   %tmp = call i64 asm "v_lsrlrev_b64 $0, $1, 1", "={v[2:3]},v"(i64 %src) #0
@@ -30,8 +37,11 @@ define amdgpu_kernel void @v_mqsad_u32_u8_non_immediate(ptr addrspace(1) %out, i
 }
 
 ; GCN-LABEL: {{^}}v_mqsad_u32_u8_inline_fp_immediate:
-; GCN-DAG: v_mov_b32_e32 v0, v2
-; GCN-DAG: v_mov_b32_e32 v1, v3
+; GFX600-DAG: v_mov_b32_e32 v0, v2
+; GFX600-DAG: v_mov_b32_e32 v1, v3
+; GFX803-DAG: v_mov_b32_e32 v0, v2
+; GFX803-DAG: v_mov_b32_e32 v1, v3
+; GFX13-NOT: v_mov_b32_e32
 ; GCN: v_mqsad_u32_u8 v[2:5], v[0:1], v6, v[{{[0-9]+:[0-9]+}}]
 define amdgpu_kernel void @v_mqsad_u32_u8_inline_fp_immediate(ptr addrspace(1) %out, i64 %src, i32 %a) {
   %tmp = call i64 asm "v_lsrlrev_b64 $0, $1, 1", "={v[2:3]},v"(i64 %src) #0
@@ -43,8 +53,11 @@ define amdgpu_kernel void @v_mqsad_u32_u8_inline_fp_immediate(ptr addrspace(1) %
 }
 
 ; GCN-LABEL: {{^}}v_mqsad_u32_u8_use_sgpr_vgpr:
-; GCN-DAG: v_mov_b32_e32 v0, v2
-; GCN-DAG: v_mov_b32_e32 v1, v3
+; GFX600-DAG: v_mov_b32_e32 v0, v2
+; GFX600-DAG: v_mov_b32_e32 v1, v3
+; GFX803-DAG: v_mov_b32_e32 v0, v2
+; GFX803-DAG: v_mov_b32_e32 v1, v3
+; GFX13-NOT: v_mov_b32_e32
 ; GCN: v_mqsad_u32_u8 v[2:5], v[0:1], v6, v[{{[0-9]+:[0-9]+}}]
 define amdgpu_kernel void @v_mqsad_u32_u8_use_sgpr_vgpr(ptr addrspace(1) %out, i64 %src, i32 %a, ptr addrspace(1) %input) {
   %in = load <4 x i32>, ptr addrspace(1) %input

``````````

</details>


https://github.com/llvm/llvm-project/pull/213292

More information about the cfe-commits mailing list