Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

AMDGPU: Add missing mqsad-pk-insts to gfx13 frontend feature map (#213124)

Browse source 
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)
This commit is contained in:
Matt Arsenault 2026-07-31 10:26:34 +02:00 committed by GitHub
parent 5eb876fcaf
commit 2b4f6c42a7
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
3 changed files with 38 additions and 7 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

23
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
View file

@ -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);
}

1
llvm/lib/TargetParser/AMDGPUTargetParser.cpp
View file

@ -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;

21
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.pk.u16.u8.ll
View file

@ -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