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

[Matt Arsenault via cfe-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/213122

fillAMDGCNFeatureMap omitted msad-insts for gfx1310/gfx13-generic, so
clang wrongly rejected __builtin_amdgcn_msad_u8 on those targets even
though the backend enables the feature (FeatureGFX9 generation, inherited
through FeatureGFX13). Add it to the gfx13 case.

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

>From 334f18e315a6fd16b2bc7f5d4b471d72c8e10109 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 30 Jul 2026 21:10:02 +0200
Subject: [PATCH] AMDGPU: Add missing msad-insts to gfx13 frontend feature map

fillAMDGCNFeatureMap omitted msad-insts for gfx1310/gfx13-generic, so
clang wrongly rejected __builtin_amdgcn_msad_u8 on those targets even
though the backend enables the feature (FeatureGFX9 generation, inherited
through FeatureGFX13). Add it to the gfx13 case.

Co-authored-by: Claude (Claude-Opus-4.8)
---
 .../CodeGenOpenCL/builtins-amdgcn-gfx13.cl    | 23 +++++++++++++++++++
 llvm/lib/TargetParser/AMDGPUTargetParser.cpp  |  1 +
 .../CodeGen/AMDGPU/llvm.amdgcn.msad.u8.ll     |  9 ++++++--
 3 files changed, 31 insertions(+), 2 deletions(-)

diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
index a6e0b15921860..e8cea9baeda8f 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
@@ -46,3 +46,26 @@ void test_cvt_scalef32_pk32_fp6_f32(global uint6 *out, float32 srcf32, float src
   *out = __builtin_amdgcn_cvt_scalef32_pk32_fp6_f32(srcf32, src);
 }
 
+// CHECK-LABEL: @test_msad_u8(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT:    [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[SRC2_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT:    store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT:    store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT:    store i32 [[SRC2:%.*]], ptr addrspace(5) [[SRC2_ADDR]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4
+// CHECK-NEXT:    [[TMP2:%.*]] = load i32, ptr addrspace(5) [[SRC2_ADDR]], align 4
+// CHECK-NEXT:    [[TMP3:%.*]] = call i32 @llvm.amdgcn.msad.u8(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]])
+// CHECK-NEXT:    [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT:    store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
+// CHECK-NEXT:    ret void
+//
+void test_msad_u8(global unsigned int *out, unsigned int src0, unsigned int src1, unsigned int src2)
+{
+  *out = __builtin_amdgcn_msad_u8(src0, src1, src2);
+}
+
diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
index 6e16bc8dc9507..d401d8d4febba 100644
--- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
@@ -524,6 +524,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T,
     Features["lerp-inst"] = true;
     Features["sad-insts"] = true;
     Features["qsad-insts"] = true;
+    Features["msad-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.msad.u8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.msad.u8.ll
index b3951f34449d7..091ee08a461f1 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.msad.u8.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.msad.u8.ll
@@ -1,10 +1,13 @@
-; 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 %s
+; RUN: llc -mtriple=amdgpu8.03 < %s | FileCheck -check-prefixes=GCN %s
+; RUN: llc -mtriple=amdgpu13.10 < %s | FileCheck -check-prefixes=GFX13 %s
 
 declare i32 @llvm.amdgcn.msad.u8(i32, i32, i32) #0
 
 ; GCN-LABEL: {{^}}v_msad_u8:
 ; GCN: v_msad_u8 v{{[0-9]+}}, v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}
+; GFX13-LABEL: {{^}}v_msad_u8:
+; GFX13: v_msad_u8 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}
 define amdgpu_kernel void @v_msad_u8(ptr addrspace(1) %out, i32 %src) {
   %result= call i32 @llvm.amdgcn.msad.u8(i32 %src, i32 100, i32 100) #0
   store i32 %result, ptr addrspace(1) %out, align 4
@@ -13,6 +16,8 @@ define amdgpu_kernel void @v_msad_u8(ptr addrspace(1) %out, i32 %src) {
 
 ; GCN-LABEL: {{^}}v_msad_u8_non_immediate:
 ; GCN: v_msad_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+; GFX13-LABEL: {{^}}v_msad_u8_non_immediate:
+; GFX13: v_msad_u8 v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}
 define amdgpu_kernel void @v_msad_u8_non_immediate(ptr addrspace(1) %out, i32 %src, i32 %a, i32 %b) {
   %result= call i32 @llvm.amdgcn.msad.u8(i32 %src, i32 %a, i32 %b) #0
   store i32 %result, ptr addrspace(1) %out, align 4