[clang] 2b4f6c4 - AMDGPU: Add missing mqsad-pk-insts to gfx13 frontend feature map (#213124)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 31 01:26:39 PDT 2026
Author: Matt Arsenault
Date: 2026-07-31T08:26:34Z
New Revision: 2b4f6c42a7dbc8b28c71a334a9fbecce99b68eaf
URL: https://github.com/llvm/llvm-project/commit/2b4f6c42a7dbc8b28c71a334a9fbecce99b68eaf
DIFF: https://github.com/llvm/llvm-project/commit/2b4f6c42a7dbc8b28c71a334a9fbecce99b68eaf.diff
LOG: AMDGPU: Add missing mqsad-pk-insts to gfx13 frontend feature map (#213124)
fillAMDGCNFeatureMap omitted mqsad-pk-insts for gfx1310/gfx13-generic,
so clang wrongly rejected __builtin_amdgcn_mqsad_pk_u16_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)
Added:
Modified:
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
llvm/lib/TargetParser/AMDGPUTargetParser.cpp
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.pk.u16.u8.ll
Removed:
################################################################################
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
index af40e044e2099..1c23e67eb3b94 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
@@ -100,3 +100,26 @@ void test_raw_ptr_buffer_atomic_fmin_fmax(global float *fout, global double *dou
*dout = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(d, rsrc, offset, soffset, 0);
}
+// CHECK-LABEL: @test_mqsad_pk_u16_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 i64, align 8, 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 i64 [[SRC2:%.*]], ptr addrspace(5) [[SRC2_ADDR]], align 8
+// 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 i64, ptr addrspace(5) [[SRC2_ADDR]], align 8
+// CHECK-NEXT: [[TMP3:%.*]] = call i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 [[TMP0]], i32 [[TMP1]], i64 [[TMP2]])
+// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store i64 [[TMP3]], ptr addrspace(1) [[TMP4]], align 8
+// CHECK-NEXT: ret void
+//
+void test_mqsad_pk_u16_u8(global unsigned long *out, unsigned long src0, unsigned int src1, unsigned long src2)
+{
+ *out = __builtin_amdgcn_mqsad_pk_u16_u8(src0, src1, src2);
+}
+
diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
index a244b1acecc53..3425d9b9e71e2 100644
--- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
@@ -527,6 +527,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T,
Features["lerp-inst"] = true;
Features["sad-insts"] = true;
Features["qsad-insts"] = true;
+ Features["mqsad-pk-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.pk.u16.u8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.pk.u16.u8.ll
index 97b3caa4e7790..374eb2d1043d6 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.pk.u16.u8.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.pk.u16.u8.ll
@@ -1,12 +1,17 @@
-; RUN: llc -mtriple=amdgpu6.00 < %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 i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64, i32, i64) #0
; GCN-LABEL: {{^}}v_mqsad_pk_u16_u8:
-; GCN: v_mqsad_pk_u16_u8 v[0:1], v[4:5], s{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
-; GCN-DAG: v_mov_b32_e32 v5, v1
-; GCN-DAG: v_mov_b32_e32 v4, v0
+; GFX600: v_mqsad_pk_u16_u8 v[0:1], v[4:5], s{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+; GFX600-DAG: v_mov_b32_e32 v5, v1
+; GFX600-DAG: v_mov_b32_e32 v4, v0
+; GFX803: v_mqsad_pk_u16_u8 v[0:1], v[4:5], s{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}]
+; GFX803-DAG: v_mov_b32_e32 v5, v1
+; GFX803-DAG: v_mov_b32_e32 v4, v0
+; GFX13: v_mqsad_pk_u16_u8 v[0:1], v[4:5], 0x64, 0x64
define amdgpu_kernel void @v_mqsad_pk_u16_u8(ptr addrspace(1) %out, i64 %src) {
%tmp = call i64 asm "v_lsrlrev_b64 $0, $1, 1", "={v[4:5]},v"(i64 %src) #0
%tmp1 = call i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %tmp, i32 100, i64 100) #0
@@ -17,8 +22,10 @@ define amdgpu_kernel void @v_mqsad_pk_u16_u8(ptr addrspace(1) %out, i64 %src) {
; GCN-LABEL: {{^}}v_mqsad_pk_u16_u8_non_immediate:
; GCN: v_mqsad_pk_u16_u8 v[0:1], v[2:3], v4, v[6:7]
-; GCN-DAG: v_mov_b32_e32 v3, v1
-; GCN-DAG: v_mov_b32_e32 v2, v0
+; GFX600-DAG: v_mov_b32_e32 v3, v1
+; GFX600-DAG: v_mov_b32_e32 v2, v0
+; GFX803-DAG: v_mov_b32_e32 v3, v1
+; GFX803-DAG: v_mov_b32_e32 v2, v0
define amdgpu_kernel void @v_mqsad_pk_u16_u8_non_immediate(ptr addrspace(1) %out, i64 %src, i32 %a, i64 %b) {
%tmp = call i64 asm "v_lsrlrev_b64 $0, $1, 1", "={v[2:3]},v"(i64 %src) #0
%tmp1 = call i32 asm "v_mov_b32 $0, $1", "={v4},v"(i32 %a) #0
More information about the cfe-commits
mailing list