From c48bc2f096bcfd65dcb7c61543d6d9859e094d77 Mon Sep 17 00:00:00 2001 From: Shilei Tian Date: Mon, 18 May 2026 19:59:44 -0400 Subject: [PATCH] [AMDGPU] Add three target features msad-insts, mqsad-pk-insts, and mqsad-insts --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 6 +-- .../CodeGen/amdgpu-builtin-is-invocable.c | 2 +- .../CodeGen/amdgpu-builtin-processor-is.c | 2 +- clang/test/CodeGen/link-builtin-bitcode.c | 6 +-- .../CodeGenCXX/dynamic-cast-address-space.cpp | 4 +- .../builtins-amdgcn-gfx12-5-generic-err.cl | 6 +++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 18 +++---- llvm/lib/Target/AMDGPU/AMDGPU.td | 50 ++++++++++++++++--- llvm/lib/Target/AMDGPU/VOP3Instructions.td | 12 ++--- llvm/lib/TargetParser/TargetParser.cpp | 29 +++++++++++ .../AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll | 2 +- .../MC/AMDGPU/gfx12_5_generic_asm_vop3_err.s | 11 ++++ 12 files changed, 113 insertions(+), 35 deletions(-) create mode 100644 llvm/test/MC/AMDGPU/gfx12_5_generic_asm_vop3_err.s diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index fc910123560a9..b15a36df6c08f 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -234,13 +234,13 @@ def __builtin_amdgcn_cvt_pk_i16 : AMDGPUBuiltin<"_ExtVector<2, short>(int, int)" def __builtin_amdgcn_cvt_pk_u16 : AMDGPUBuiltin<"_ExtVector<2, unsigned short>(unsigned int, unsigned int)", [Const]>; def __builtin_amdgcn_cvt_pk_u8_f32 : AMDGPUBuiltin<"unsigned int(float, unsigned int, unsigned int)", [Const]>; def __builtin_amdgcn_cvt_off_f32_i4 : AMDGPUBuiltin<"float(int)", [Const]>; -def __builtin_amdgcn_msad_u8 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int)", [Const]>; +def __builtin_amdgcn_msad_u8 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int)", [Const], "msad-insts">; def __builtin_amdgcn_sad_u8 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int)", [Const], "sad-insts">; def __builtin_amdgcn_sad_hi_u8 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int)", [Const], "sad-insts">; def __builtin_amdgcn_sad_u16 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int, unsigned int)", [Const], "sad-insts">; def __builtin_amdgcn_qsad_pk_u16_u8 : AMDGPUBuiltin<"uint64_t(uint64_t, unsigned int, uint64_t)", [Const], "qsad-insts">; -def __builtin_amdgcn_mqsad_pk_u16_u8 : AMDGPUBuiltin<"uint64_t(uint64_t, unsigned int, uint64_t)", [Const]>; -def __builtin_amdgcn_mqsad_u32_u8 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, unsigned int, _ExtVector<4, unsigned int>)", [Const]>; +def __builtin_amdgcn_mqsad_pk_u16_u8 : AMDGPUBuiltin<"uint64_t(uint64_t, unsigned int, uint64_t)", [Const], "mqsad-pk-insts">; +def __builtin_amdgcn_mqsad_u32_u8 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, unsigned int, _ExtVector<4, unsigned int>)", [Const], "mqsad-insts">; def __builtin_amdgcn_make_buffer_rsrc : AMDGPUBuiltin<"__amdgpu_buffer_rsrc_t(void *, short, int64_t, int)", [Const]>; def __builtin_amdgcn_raw_buffer_store_b8 : AMDGPUBuiltin<"void(unsigned char, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">; diff --git a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c index 5c6d7b25f3bf1..3f988cbdf7cee 100644 --- a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c +++ b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c @@ -54,7 +54,7 @@ void foo() { // AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1010" } // AMDGCN-GFX1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } //. -// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } // AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind } // AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } //. diff --git a/clang/test/CodeGen/amdgpu-builtin-processor-is.c b/clang/test/CodeGen/amdgpu-builtin-processor-is.c index b2f04f11ecb54..308eec1d212a5 100644 --- a/clang/test/CodeGen/amdgpu-builtin-processor-is.c +++ b/clang/test/CodeGen/amdgpu-builtin-processor-is.c @@ -63,7 +63,7 @@ void foo() { //. // AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1010" } //. -// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } // AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind } // AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } //. diff --git a/clang/test/CodeGen/link-builtin-bitcode.c b/clang/test/CodeGen/link-builtin-bitcode.c index d03fd6fc66d03..ed5abca7f34b0 100644 --- a/clang/test/CodeGen/link-builtin-bitcode.c +++ b/clang/test/CodeGen/link-builtin-bitcode.c @@ -44,6 +44,6 @@ int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_in // CHECK-SAME: () #[[ATTR_INCOMPATIBLE:[0-9]+]] { // CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "target-cpu"="gfx90a" } -// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } -// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+extended-image-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } -// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" } +// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } +// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+extended-image-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" } +// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" } diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp index 355dde6f0cfda..695c290caf7d7 100644 --- a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp +++ b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp @@ -107,9 +107,9 @@ const B& f(A *a) { // CHECK: attributes #[[ATTR3]] = { nounwind } // CHECK: attributes #[[ATTR4]] = { noreturn } //. -// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } // WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) } -// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } // WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind } // WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn } //. diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl index f3f7e1f33dd04..3bb9ab8b74910 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-5-generic-err.cl @@ -2,6 +2,7 @@ // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic -verify -S -o - %s typedef short v2s __attribute__((ext_vector_type(2))); +typedef unsigned int v4u __attribute__((ext_vector_type(4))); typedef float v4f __attribute__((ext_vector_type(4))); typedef half v4h __attribute__((ext_vector_type(4))); typedef float v8f __attribute__((ext_vector_type(8))); @@ -14,10 +15,12 @@ void test_amdgcn_swmmac_gfx12_5_generic(global v8f *out8f, global v4f *out4f, global float *outf, global unsigned int *outu, global unsigned long *outul, + global v4u *outv4u, global v2s *outv2s, v8h a8h, v16h b16h, v4h a4h, v8h b8h, v8f c8f, v4f c4f, v16bf16 a16bf16, v32bf16 b32bf16, + v4u v4u0, float x, float y, float z, unsigned int u0, unsigned int u1, unsigned int u2, unsigned long l0, @@ -27,6 +30,9 @@ void test_amdgcn_swmmac_gfx12_5_generic(global v8f *out8f, global v4f *out4f, *outu = __builtin_amdgcn_lerp(u0, u1, u2); // expected-error{{'__builtin_amdgcn_lerp' needs target feature lerp-inst}} *outul = __builtin_amdgcn_qsad_pk_u16_u8(l0, u0, l1); // expected-error{{'__builtin_amdgcn_qsad_pk_u16_u8' needs target feature qsad-insts}} *outu = __builtin_amdgcn_sad_u8(u0, u1, u2); // expected-error{{'__builtin_amdgcn_sad_u8' needs target feature sad-insts}} + *outu = __builtin_amdgcn_msad_u8(u0, u1, u2); // expected-error{{'__builtin_amdgcn_msad_u8' needs target feature msad-insts}} + *outul = __builtin_amdgcn_mqsad_pk_u16_u8(l0, u0, l1); // expected-error{{'__builtin_amdgcn_mqsad_pk_u16_u8' needs target feature mqsad-pk-insts}} + *outv4u = __builtin_amdgcn_mqsad_u32_u8(l0, u0, v4u0); // expected-error{{'__builtin_amdgcn_mqsad_u32_u8' needs target feature mqsad-insts}} *out8f = __builtin_amdgcn_swmmac_f32_16x16x32_f16_w32(a8h, b16h, c8f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32' needs target feature swmmac-gfx1200-insts,wavefrontsize32}} *out4f = __builtin_amdgcn_swmmac_f32_16x16x32_f16_w64(a4h, b8h, c4f, index); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64' needs target feature swmmac-gfx1200-insts,wavefrontsize64}} *out8f = __builtin_amdgcn_swmmac_f32_16x16x64_bf16(false, a16bf16, false, b32bf16, c8f, index, false, true); // expected-error{{'__builtin_amdgcn_swmmac_f32_16x16x64_bf16' needs target feature swmmac-gfx1250-insts,wavefrontsize32}} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 2d645a968f2fd..11fbfdde92fa1 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -1288,23 +1288,21 @@ kernel void test_cvt_pk_u8_f32(global uint* out, float src0, uint src1, uint src *out = __builtin_amdgcn_cvt_pk_u8_f32(src0, src1, src2); } -// CHECK-LABEL: test_msad_u8( -// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.msad.u8(i32 %src0, i32 %src1, i32 %src2) +// CHECK-AMDGCN-LABEL: test_msad_u8( +// CHECK-AMDGCN: {{.*}}call{{.*}} i32 @llvm.amdgcn.msad.u8(i32 %src0, i32 %src1, i32 %src2) +#if !defined(__SPIRV__) kernel void test_msad_u8(global uint* out, uint src0, uint src1, uint src2) { *out = __builtin_amdgcn_msad_u8(src0, src1, src2); } +#endif -// CHECK-LABEL: @test_mqsad_pk_u16_u8( -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %src0, i32 %src1, i64 %src2) +// CHECK-AMDGCN-LABEL: @test_mqsad_pk_u16_u8( +// CHECK-AMDGCN: {{.*}}call{{.*}} i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %src0, i32 %src1, i64 %src2) +#if !defined(__SPIRV__) kernel void test_mqsad_pk_u16_u8(global ulong* out, ulong src0, uint src1, ulong src2) { *out = __builtin_amdgcn_mqsad_pk_u16_u8(src0, src1, src2); } - -// CHECK-LABEL: test_mqsad_u32_u8( -// CHECK: {{.*}}call{{.*}} <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %src0, i32 %src1, <4 x i32> %src2) -kernel void test_mqsad_u32_u8(global uint4* out, ulong src0, uint src1, uint4 src2) { - *out = __builtin_amdgcn_mqsad_u32_u8(src0, src1, src2); -} +#endif // CHECK-LABEL: test_s_setreg( // CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setreg(i32 65535, i32 %val) diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index 7f15cc6ceeae9..961763f0e1a5f 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -858,6 +858,18 @@ defm QsadInsts : AMDGPUSubtargetFeature<"qsad-insts", "Has v_qsad* instructions" >; +defm MsadInsts : AMDGPUSubtargetFeature<"msad-insts", + "Has v_msad_u8 instruction" +>; + +defm MqsadPkInsts : AMDGPUSubtargetFeature<"mqsad-pk-insts", + "Has v_mqsad_pk_u16_u8 instruction" +>; + +defm MqsadInsts : AMDGPUSubtargetFeature<"mqsad-insts", + "Has v_mqsad_u32_u8 instruction" +>; + defm CvtNormInsts : AMDGPUSubtargetFeature<"cvt-norm-insts", "Has v_cvt_norm* instructions" >; @@ -1401,8 +1413,9 @@ def FeatureSouthernIslands : GCNSubtargetFeatureGeneration<"SOUTHERN_ISLANDS", FeatureGDS, FeatureGWS, FeatureDefaultComponentZero, FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts, FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst, - FeatureSadInsts, FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode, - FeatureMTBUFInsts, FeatureFormattedMUBUFInsts + FeatureSadInsts, FeatureMsadInsts, FeatureMqsadPkInsts, + FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode, FeatureMTBUFInsts, + FeatureFormattedMUBUFInsts ] >; @@ -1417,7 +1430,8 @@ def FeatureSeaIslands : GCNSubtargetFeatureGeneration<"SEA_ISLANDS", FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts, FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst, - FeatureSadInsts, FeatureQsadInsts, FeatureCvtPkNormVOP2Insts, + FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, FeatureMqsadPkInsts, + FeatureMqsadInsts, FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode, FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts ] @@ -1437,7 +1451,8 @@ def FeatureVolcanicIslands : GCNSubtargetFeatureGeneration<"VOLCANIC_ISLANDS", FeatureDsSrc2Insts, FeatureExtendedImageInsts, FeatureFastDenormalF32, FeatureUnalignedBufferAccess, FeatureImageInsts, FeatureGDS, FeatureGWS, FeatureDefaultComponentZero, FeatureVmemWriteVgprInOrder, FeatureCubeInsts, - FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, + FeatureMqsadPkInsts, FeatureMqsadInsts, FeatureCvtPkNormVOP2Insts, FeatureDX10ClampAndIEEEMode, FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts ] @@ -1460,6 +1475,7 @@ def FeatureGFX9 : GCNSubtargetFeatureGeneration<"GFX9", FeatureUnalignedDSAccess, FeatureNegativeScratchOffsetBug, FeatureGWS, FeatureDefaultComponentZero,FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad, FeatureCubeInsts, FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureMsadInsts, FeatureMqsadPkInsts, FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode, FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts @@ -1487,7 +1503,8 @@ def FeatureGFX10 : GCNSubtargetFeatureGeneration<"GFX10", FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF64GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts, FeatureAtomicFMinFMaxF64FlatInsts, FeatureVmemWriteVgprInOrder, FeatureVMemToLDSLoad, FeatureCubeInsts, - FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, + FeatureMqsadPkInsts, FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, FeatureDX10ClampAndIEEEMode, FeatureFlatOffsetBits12, FeatureInstCacheLineSize64, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts @@ -1514,7 +1531,8 @@ def FeatureGFX11 : GCNSubtargetFeatureGeneration<"GFX11", FeatureDefaultComponentZero, FeatureMaxHardClauseLength32, FeatureAtomicFMinFMaxF32GlobalInsts, FeatureAtomicFMinFMaxF32FlatInsts, FeatureVmemWriteVgprInOrder, FeatureCubeInsts, FeatureLerpInst, - FeatureSadInsts, FeatureQsadInsts, FeatureCvtNormInsts, + FeatureSadInsts, FeatureQsadInsts, FeatureMsadInsts, FeatureMqsadPkInsts, + FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, FeatureInstCacheLineSize128, FeatureMTBUFInsts, FeatureFormattedMUBUFInsts ] @@ -1826,7 +1844,10 @@ def FeatureISAVersion10_Common : FeatureSet< FeatureLDSBankCount32, FeatureDLInsts, FeatureNSAEncoding, - FeatureBackOffBarrier]>; + FeatureBackOffBarrier, + FeatureMsadInsts, + FeatureMqsadPkInsts, + FeatureMqsadInsts]>; def FeatureISAVersion10_1_Common : FeatureSet< !listconcat(FeatureISAVersion10_Common.Features, @@ -1921,6 +1942,9 @@ def FeatureISAVersion11_Common : FeatureSet< FeatureMemoryAtomicFAddF32DenormalSupport, FeatureRealTrue16Insts, FeatureD16Writes32BitVgpr, + FeatureMsadInsts, + FeatureMqsadPkInsts, + FeatureMqsadInsts, ]>; // There are few workarounds that need to be added to all targets. This @@ -2054,6 +2078,9 @@ def FeatureISAVersion12 : FeatureSet< FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureMsadInsts, + FeatureMqsadPkInsts, + FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, @@ -2155,6 +2182,9 @@ def FeatureISAVersion12_50 : FeatureSet< FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureMsadInsts, + FeatureMqsadPkInsts, + FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, @@ -2169,6 +2199,9 @@ def FeatureISAVersion12_51 : FeatureSet< FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureMsadInsts, + FeatureMqsadPkInsts, + FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, @@ -2241,6 +2274,9 @@ def FeatureISAVersion13 : FeatureSet< FeatureLerpInst, FeatureSadInsts, FeatureQsadInsts, + FeatureMsadInsts, + FeatureMqsadPkInsts, + FeatureMqsadInsts, FeatureCvtNormInsts, FeatureCvtPkNormVOP2Insts, FeatureCvtPkNormVOP3Insts, diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td index 1c6447de407f3..489ab184385de 100644 --- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td @@ -364,13 +364,12 @@ let mayRaiseFPException = 0 in { // Seems suspicious but manual doesn't say it d defm V_DIV_SCALE_F64 : VOP3Inst <"v_div_scale_f64", VOP3b_F64_I1_F64_F64_F64>; } // End mayRaiseFPException = 0 -let isReMaterializable = 1 in +let isReMaterializable = 1, SubtargetPredicate = HasMsadInsts in defm V_MSAD_U8 : VOP3Inst <"v_msad_u8", VOP3_Profile>; -let Constraints = "@earlyclobber $vdst" in { +let SubtargetPredicate = HasMqsadPkInsts, Constraints = "@earlyclobber $vdst" in { defm V_MQSAD_PK_U16_U8 : VOP3Inst <"v_mqsad_pk_u16_u8", VOP3_Profile>; -} // End Constraints = "@earlyclobber $vdst" - +} // End SubtargetPredicate = HasMqsadPkInsts, Constraints = "@earlyclobber $vdst" let isReMaterializable = 1 in { let SchedRW = [WriteDouble], IsDPMACCInstruction = 1 in { @@ -425,13 +424,12 @@ def VOPProfileMQSAD : VOP3_Profile { let HasModifiers = 0; } -let SubtargetPredicate = isGFX7Plus in { let Constraints = "@earlyclobber $vdst", SchedRW = [WriteQuarterRate32] in { let SubtargetPredicate = HasQsadInsts in defm V_QSAD_PK_U16_U8 : VOP3Inst <"v_qsad_pk_u16_u8", VOP3_Profile>; -defm V_MQSAD_U32_U8 : VOP3Inst <"v_mqsad_u32_u8", VOPProfileMQSAD>; +let SubtargetPredicate = HasMqsadInsts in + defm V_MQSAD_U32_U8 : VOP3Inst <"v_mqsad_u32_u8", VOPProfileMQSAD>; } // End Constraints = "@earlyclobber $vdst", SchedRW = [WriteQuarterRate32] -} // End SubtargetPredicate = isGFX7Plus let isCommutable = 1, SchedRW = [WriteIntMul, WriteSALUDummy] in { let SubtargetPredicate = isGFX7Plus in { diff --git a/llvm/lib/TargetParser/TargetParser.cpp b/llvm/lib/TargetParser/TargetParser.cpp index e84e73572a4c0..3c2958d46521a 100644 --- a/llvm/lib/TargetParser/TargetParser.cpp +++ b/llvm/lib/TargetParser/TargetParser.cpp @@ -260,6 +260,9 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["lerp-inst"] = true; Features["qsad-insts"] = true; Features["sad-insts"] = true; + Features["msad-insts"] = true; + Features["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; [[fallthrough]]; case GK_GFX12_5_GENERIC: Features["ci-insts"] = true; @@ -334,6 +337,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; Features["fp8-conversion-insts"] = true; Features["wmma-128b-insts"] = true; @@ -363,6 +369,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; Features["gws"] = true; Features["dot11-insts"] = true; @@ -401,6 +410,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; Features["gws"] = true; Features["wmma-256b-insts"] = true; @@ -439,6 +451,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; break; case GK_GFX1012: @@ -471,6 +486,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; break; case GK_GFX950: @@ -529,6 +547,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; break; case GK_GFX90A: @@ -578,6 +599,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; break; case GK_GFX705: @@ -591,6 +615,9 @@ 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["mqsad-pk-insts"] = true; + Features["mqsad-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; Features["image-insts"] = true; Features["s-memtime-inst"] = true; @@ -611,6 +638,8 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["cube-insts"] = true; Features["lerp-inst"] = true; Features["sad-insts"] = true; + Features["msad-insts"] = true; + Features["mqsad-pk-insts"] = true; Features["cvt-pknorm-vop2-insts"] = true; break; case GK_NONE: 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 a271bcd67e018..c95cc7cd9a4b9 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mqsad.u32.u8.ll @@ -1,4 +1,4 @@ -; RUN: llc -mtriple=amdgcn < %s | FileCheck -check-prefix=GCN %s +; RUN: llc -mtriple=amdgcn -mcpu=bonaire < %s | FileCheck -check-prefix=GCN %s ; RUN: llc -mtriple=amdgcn -mcpu=fiji < %s | FileCheck -check-prefix=GCN %s declare <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64, i32, <4 x i32>) #0 diff --git a/llvm/test/MC/AMDGPU/gfx12_5_generic_asm_vop3_err.s b/llvm/test/MC/AMDGPU/gfx12_5_generic_asm_vop3_err.s new file mode 100644 index 0000000000000..6e1ed343f36d1 --- /dev/null +++ b/llvm/test/MC/AMDGPU/gfx12_5_generic_asm_vop3_err.s @@ -0,0 +1,11 @@ +// NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --unique --sort --version 6 +// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx12-5-generic %s -filetype=null 2>&1 | FileCheck --implicit-check-not=error: %s + +v_mqsad_pk_u16_u8 v[0:1], v[2:3], v4, v[6:7] +// CHECK: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx12-5-generic): v_mqsad_pk_u16_u8 + +v_mqsad_u32_u8 v[0:3], s[2:3], v4, v[0:3] +// CHECK: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx12-5-generic): v_mqsad_u32_u8 + +v_msad_u8 v0, v1, v2, v3 +// CHECK: :[[@LINE-1]]:1: error: instruction not supported on this GPU (gfx12-5-generic): v_msad_u8