Author: Matt Arsenault Date: 2026-07-31T14:30:44+02:00 New Revision: b4efecb0379353d51a172209193e500b6cb2a0f2
URL: https://github.com/llvm/llvm-project/commit/b4efecb0379353d51a172209193e500b6cb2a0f2 DIFF: https://github.com/llvm/llvm-project/commit/b4efecb0379353d51a172209193e500b6cb2a0f2.diff LOG: AMDGPU: Add missing msad-insts to gfx13 frontend feature map (#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) Added: Modified: clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl llvm/lib/TargetParser/AMDGPUTargetParser.cpp llvm/test/CodeGen/AMDGPU/llvm.amdgcn.msad.u8.ll Removed: ################################################################################ diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl index 20f415afe639b..f5d341b31991c 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl @@ -148,3 +148,26 @@ void test_mqsad_pk_u16_u8(global unsigned long *out, unsigned long src0, unsigne *out = __builtin_amdgcn_mqsad_pk_u16_u8(src0, src1, src2); } +// 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 920e92ec2679c..c5df20fb6b63a 100644 --- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp +++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp @@ -529,6 +529,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["sad-insts"] = true; Features["qsad-insts"] = true; Features["mqsad-pk-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..13828064df8bb 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,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 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]+}} +; GFX600: v_msad_u8 v{{[0-9]+}}, v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}} +; GFX803: v_msad_u8 v{{[0-9]+}}, v{{[0-9]+}}, s{{[0-9]+}}, s{{[0-9]+}} +; 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 @@ -12,7 +15,9 @@ 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]+}} +; GFX600: v_msad_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} +; GFX803: v_msad_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} +; 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
