https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/203956
>From 3b85daa82f12a69e38d1d814978bc5b91772dfdd Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Mon, 15 Jun 2026 12:09:29 -0400 Subject: [PATCH] [AMDGPU] Guard more intrinsics with target features --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 8 +-- .../CodeGen/amdgpu-builtin-is-invocable.c | 2 +- .../CodeGen/amdgpu-builtin-processor-is.c | 2 +- .../CodeGenCXX/dynamic-cast-address-space.cpp | 4 +- llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 17 +++++- llvm/lib/Target/AMDGPU/AMDGPU.td | 33 ++++++------ .../AMDGPU/AMDGPUInstructionSelector.cpp | 52 +------------------ .../lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp | 24 --------- .../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 7 ++- llvm/lib/Target/AMDGPU/GCNSubtarget.h | 6 --- llvm/lib/Target/AMDGPU/MIMGInstructions.td | 6 +-- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 42 --------------- .../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 4 -- llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 1 - llvm/lib/Target/AMDGPU/VOP3Instructions.td | 2 +- llvm/lib/TargetParser/AMDGPUTargetParser.cpp | 6 +++ .../inst-select-amdgcn.exp.compr.mir | 3 +- .../GlobalISel/llvm.amdgcn.intersect_ray.ll | 4 +- .../AMDGPU/llvm.amdgcn.dual_intersect_ray.ll | 4 +- .../CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll | 3 +- .../CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll | 2 +- .../AMDGPU/llvm.amdgcn.intersect_ray.ll | 7 ++- .../CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll | 2 +- .../CodeGen/AMDGPU/llvm.amdgcn.permlane.ll | 2 +- .../test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll | 2 +- .../test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll | 2 +- llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll | 5 ++ llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll | 5 ++ .../CodeGen/AMDGPU/unsupported-av-load.ll | 2 +- .../CodeGen/AMDGPU/unsupported-av-store.ll | 2 +- 30 files changed, 84 insertions(+), 177 deletions(-) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 8c64116066971..668646d8d0070 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -473,10 +473,10 @@ def __builtin_amdgcn_s_ttracedata_imm : AMDGPUBuiltin<"void(_Constant short)", [ // Postfix l indicates the 1st argument is i64. // Postfix h indicates the 4/5-th arguments are half4. //===----------------------------------------------------------------------===// -def __builtin_amdgcn_image_bvh_intersect_ray : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned int, float, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">; -def __builtin_amdgcn_image_bvh_intersect_ray_h : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned int, float, _ExtVector<4, float>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">; -def __builtin_amdgcn_image_bvh_intersect_ray_l : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, float, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">; -def __builtin_amdgcn_image_bvh_intersect_ray_lh : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, float, _ExtVector<4, float>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, unsigned int>)", [Const], "gfx10-insts">; +def __builtin_amdgcn_image_bvh_intersect_ray : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned int, float, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, unsigned int>)", [Const], "bvh-ray-tracing-insts">; +def __builtin_amdgcn_image_bvh_intersect_ray_h : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(unsigned int, float, _ExtVector<4, float>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, unsigned int>)", [Const], "bvh-ray-tracing-insts">; +def __builtin_amdgcn_image_bvh_intersect_ray_l : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, float, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, float>, _ExtVector<4, unsigned int>)", [Const], "bvh-ray-tracing-insts">; +def __builtin_amdgcn_image_bvh_intersect_ray_lh : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(uint64_t, float, _ExtVector<4, float>, _ExtVector<4, _Float16>, _ExtVector<4, _Float16>, _ExtVector<4, unsigned int>)", [Const], "bvh-ray-tracing-insts">; //===----------------------------------------------------------------------===// diff --git a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c index 9511dc84cc67a..8b1fb2c7f0c3a 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,+asynccnt,+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,+flat-global-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,+gfx1251-gemm-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 #[[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,+asynccnt,+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,+bvh-ray-tracing-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,+flat-global-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,+gfx1251-gemm-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 5c8017b918371..3d0b8ef65b098 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,+asynccnt,+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,+flat-global-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,+gfx1251-gemm-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 #[[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,+asynccnt,+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,+bvh-ray-tracing-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,+flat-global-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,+gfx1251-gemm-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/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp index 4050546954e13..7ec73d5825160 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,+asynccnt,+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,+flat-global-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,+gfx1251-gemm-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 #[[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,+asynccnt,+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,+bvh-ray-tracing-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,+flat-global-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,+gfx1251-gemm-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,+asynccnt,+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,+flat-global-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,+gfx1251-gemm-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 #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+asynccnt,+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,+bvh-ray-tracing-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,+flat-global-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,+gfx1251-gemm-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/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 21882e247c027..917730cea404d 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -318,6 +318,7 @@ def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">, // void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %barrier) // The %barrier argument must be uniform, otherwise behavior is undefined. +let TargetFeatures = "s-wakeup-barrier-inst" in def int_amdgcn_s_wakeup_barrier : ClangBuiltin<"__builtin_amdgcn_s_wakeup_barrier">, Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; @@ -475,6 +476,7 @@ def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">, // intended for use on subtargets that have the v_fma_legacy_f32 and/or // v_fmac_legacy_f32 instructions. (Note that v_fma_legacy_f16 is unrelated and // has a completely different kind of legacy behaviour.) +let TargetFeatures = "fma-legacy32-insts" in def int_amdgcn_fma_legacy : PureIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [Commutative] @@ -586,6 +588,7 @@ def int_amdgcn_fmad_ftz : [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>] >; +let TargetFeatures = "tanh-insts" in def int_amdgcn_tanh : PureIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>] >; @@ -915,8 +918,6 @@ class AMDGPUAVStore : Intrinsic < [SDNPMemOperand, SDNPMayStore] >; -def int_amdgcn_av_store_b128 : AMDGPUAVStore; - class AMDGPUAVLoad : Intrinsic < [llvm_v4i32_ty], [llvm_anyptr_ty, // Pointer to load from (flat or global) @@ -927,7 +928,10 @@ class AMDGPUAVLoad : Intrinsic < [SDNPMemOperand, SDNPMayLoad] >; +let TargetFeatures = "flat-global-insts" in { +def int_amdgcn_av_store_b128 : AMDGPUAVStore; def int_amdgcn_av_load_b128 : AMDGPUAVLoad; +} // End TargetFeatures = "flat-global-insts" } // TargetPrefix = "amdgcn" @@ -2159,6 +2163,7 @@ def int_amdgcn_exp_row : DefaultAttrsIntrinsic <[], [ >; // exp with compr bit set. Not supported on GFX11+. +let TargetFeatures = "-gfx11-insts" in def int_amdgcn_exp_compr : DefaultAttrsIntrinsic <[], [ llvm_i32_ty, // tgt, llvm_i32_ty, // en @@ -2882,6 +2887,7 @@ def int_amdgcn_wait_asyncmark : // GFX10 Intrinsics //===----------------------------------------------------------------------===// +let TargetFeatures = "permlane16-insts" in { // llvm.amdgcn.permlane16 <old> <src0> <src1> <src2> <fi> <bound_control> def int_amdgcn_permlane16 : Intrinsic<[llvm_any_ty], @@ -2895,10 +2901,12 @@ def int_amdgcn_permlanex16 : [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>, IntrNoCallback, IntrNoFree]>; +} // End TargetFeatures = "permlane16-insts" // llvm.amdgcn.mov.dpp8 <src> <sel> // <sel> is a 32-bit constant whose high 8 bits must be zero which selects // the lanes to read from. +let TargetFeatures = "dpp8" in def int_amdgcn_mov_dpp8 : Intrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty], @@ -2921,6 +2929,7 @@ class AMDGPUAtomicRtn<LLVMType vt, LLVMType pt = llvm_anyptr_ty> : Intrinsic < // <ray_dir>, <ray_inv_dir>, <texture_descr> // <node_ptr> is i32 or i64. // <ray_dir> and <ray_inv_dir> are both v3f16 or both v3f32. +let TargetFeatures = "bvh-ray-tracing-insts" in def int_amdgcn_image_bvh_intersect_ray : DefaultAttrsIntrinsic<[llvm_v4i32_ty], [llvm_anyint_ty, llvm_float_ty, llvm_v3f32_ty, llvm_anyvector_ty, @@ -3064,6 +3073,7 @@ def int_amdgcn_ds_bvh_stack_push8_pop1_rtn : IntDSBVHStackRtn<vdst = llvm_i32_t def int_amdgcn_ds_bvh_stack_push8_pop2_rtn : IntDSBVHStackRtn<vdst = llvm_i64_ty, data1 = llvm_v8i32_ty>; +let TargetFeatures = "bvh-dual-bvh-8-insts" in { // <vdata>, <ray_origin>, <ray_dir> // llvm.amdgcn.image.bvh.dual.intersect.ray <node_ptr>, <ray_extent>, // <instance_mask>, <ray_origin>, @@ -3085,6 +3095,7 @@ def int_amdgcn_image_bvh8_intersect_ray : [llvm_i64_ty, llvm_float_ty, llvm_i8_ty, llvm_v3f32_ty, llvm_v3f32_ty, llvm_i32_ty, llvm_v4i32_ty], [IntrReadMem, IntrWillReturn, IntrNoCallback, IntrNoFree]>; +} // End TargetFeatures = "bvh-dual-bvh-8-insts" // llvm.amdgcn.permlane16.var <old> <src0> <src1> <fi> <bound_control> def int_amdgcn_permlane16_var : ClangBuiltin<"__builtin_amdgcn_permlane16_var">, @@ -3404,6 +3415,7 @@ def int_amdgcn_udot4 : // a[i in 0. . . 3] = (%a_sign ? a.i8[i] : promoteToSigned(a.u8[i])); // b[i in 0. . . 3] = (%b_sign ? b.i8[i] : promoteToSigned(b.u8[i])); // %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c +let TargetFeatures = "dot8-insts" in def int_amdgcn_sudot4 : ClangBuiltin<"__builtin_amdgcn_sudot4">, PureIntrinsic< @@ -3457,6 +3469,7 @@ def int_amdgcn_udot8 : // b[i in 0. . . 7] = (%b_sign ? b.i4[i] : promoteToSigned(b.u4[i])); // %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + // %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c + let TargetFeatures = "dot8-insts" in def int_amdgcn_sudot8 : ClangBuiltin<"__builtin_amdgcn_sudot8">, PureIntrinsic< diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index a1298e5969ea0..2d0a367597f0b 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -179,6 +179,10 @@ defm FmaMixBF16Insts : AMDGPUSubtargetFeature<"fma-mix-bf16-insts", "Has v_fma_mix_f32_bf16, v_fma_mixlo_bf16, v_fma_mixhi_bf16 instructions" >; +defm FmaLegacy32Insts : AMDGPUSubtargetFeature<"fma-legacy32-insts", + "Has single-precision legacy FMA instructions" +>; + defm IEEEMinimumMaximumInsts : AMDGPUSubtargetFeature<"ieee-minimum-maximum-insts", "Has v_minimum/maximum_f16/f32/f64, v_minimummaximum/maximumminimum_f16/f32 and" "v_pk_minimum/maximum_f16 instructions" @@ -565,7 +569,9 @@ defm GFX13Insts : AMDGPUSubtargetFeature<"gfx13-insts", defm GFX10_3Insts : AMDGPUSubtargetFeature<"gfx10-3-insts", "Additional instructions for GFX10.3", - /*GenPredicate=*/0 + /*GenPredicate=*/0, + /*GenAssemblerPredicate=*/1, + [FeatureFmaLegacy32Insts] >; defm GFX11_7Insts : AMDGPUSubtargetFeature<"gfx11-7-insts", @@ -745,9 +751,8 @@ defm ExtendedImageInsts : AMDGPUSubtargetFeature<"extended-image-insts", "Support mips != 0, lod != 0, gather4, and get_lod" >; -defm GFX10_AEncoding : AMDGPUSubtargetFeature<"gfx10_a-encoding", - "Has BVH ray tracing instructions", - /*GenPredicate=*/0 +defm BVHRayTracingInsts : AMDGPUSubtargetFeature<"bvh-ray-tracing-insts", + "Has BVH ray tracing instructions" >; defm GFX10_BEncoding : AMDGPUSubtargetFeature<"gfx10_b-encoding", @@ -1572,7 +1577,7 @@ def FeatureGFX11 : GCNSubtargetFeatureGeneration<"GFX11", FeatureFlatAddressSpace, Feature16BitInsts, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts, - FeatureGFX10_AEncoding, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, + FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureTrue16BitInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, @@ -1600,7 +1605,7 @@ def FeatureGFX12 : GCNSubtargetFeatureGeneration<"GFX12", FeatureFlatAddressSpace, Feature16BitInsts, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts, - FeatureGFX10_AEncoding, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, + FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureGFX12Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, @@ -1627,7 +1632,7 @@ def FeatureGFX13 : GCNSubtargetFeatureGeneration<"GFX13", FeatureFlatAddressSpace, Feature16BitInsts, FeatureInv2PiInlineImm, FeatureApertureRegs, FeatureCIInsts, FeatureGFX8Insts, FeatureGFX9Insts, FeatureGFX10Insts, - FeatureGFX10_AEncoding, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, + FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureGFX11Insts, FeatureGFX12Insts, FeatureGFX13Insts, FeatureVOP3PInsts, FeatureVOPDInsts, FeatureMovrel, FeatureFastFMAF32, FeatureDPP, FeatureIntClamp, FeatureFlatInstOffsets, FeatureFlatGlobalInsts, FeatureFlatScratchInsts, @@ -1960,11 +1965,11 @@ def FeatureISAVersion10_1_2 : FeatureSet< def FeatureISAVersion10_1_3 : FeatureSet< !listconcat(FeatureISAVersion10_1_Common.Features, - [FeatureGFX10_AEncoding])>; + [FeatureBVHRayTracingInsts])>; def FeatureISAVersion10_3_0 : FeatureSet< !listconcat(FeatureISAVersion10_Common.Features, - [FeatureGFX10_AEncoding, + [FeatureBVHRayTracingInsts, FeatureGFX10_BEncoding, FeatureGFX10_3Insts, FeatureDot1Insts, @@ -2730,9 +2735,6 @@ def HasFlatScratchSTMode : Predicate<"Subtarget->hasFlatScratchSTMode()">, def HasFlatScratchSVSMode : Predicate<"Subtarget->hasFlatScratchSVSMode()">, AssemblerPredicate<(any_of FeatureGFX940Insts, FeatureGFX11Insts)>; -def HasGFX10_AEncoding : Predicate<"Subtarget->hasGFX10_AEncoding()">, - AssemblerPredicate<(all_of FeatureGFX10_AEncoding)>; - def HasGFX10_BEncoding : Predicate<"Subtarget->hasGFX10_BEncoding()">, AssemblerPredicate<(all_of FeatureGFX10_BEncoding)>; @@ -2845,11 +2847,8 @@ def NotHasMAIInsts : Predicate<"!Subtarget->hasMAIInsts()">, def NotHasFP8E5M3Insts : Predicate<"!Subtarget->hasFP8E5M3Insts()">, AssemblerPredicate<(all_of (not FeatureFP8E5M3Insts))>; -def HasFmaLegacy32 : Predicate<"Subtarget->hasGFX10_3Insts()">, - AssemblerPredicate<(any_of FeatureGFX10_3Insts)>; - -def HasFmacLegacy32 : Predicate<"Subtarget->hasGFX10_3Insts() && Subtarget->getGeneration() < AMDGPUSubtarget::GFX12">, - AssemblerPredicate<(all_of FeatureGFX10_3Insts, (not FeatureGFX12Insts))>; +def HasFmacLegacy32 : Predicate<"Subtarget->hasFmaLegacy32Insts() && Subtarget->getGeneration() < AMDGPUSubtarget::GFX12">, + AssemblerPredicate<(all_of FeatureFmaLegacy32Insts, (not FeatureGFX12Insts))>; def HasAtomicDsCondSubClampInsts : Predicate<"Subtarget->getGeneration() >= AMDGPUSubtarget::GFX12">, diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp index 9d44d4c5424a3..1101ab4be62d5 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp @@ -72,12 +72,6 @@ static Register getWaveAddress(const MachineInstr *Def) { : Register(); } -static void diagnoseUnsupportedIntrinsic(const MachineInstr &I) { - const Function &F = I.getMF()->getFunction(); - F.getContext().diagnose(DiagnosticInfoUnsupported( - F, "intrinsic not supported on subtarget", I.getDebugLoc(), DS_Error)); -} - bool AMDGPUInstructionSelector::isVCC(Register Reg, const MachineRegisterInfo &MRI) const { // The verifier is oblivious to s1 being a valid value for wavesize registers. @@ -1283,38 +1277,6 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC(MachineInstr &I) const { return selectPermlaneSwapIntrin(I, IntrinsicID); case Intrinsic::amdgcn_wave_shuffle: return selectWaveShuffleIntrin(I); - case Intrinsic::amdgcn_fma_legacy: - if (!STI.hasFmaLegacy32Insts()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - return selectImpl(I, *CoverageInfo); - case Intrinsic::amdgcn_sudot4: - case Intrinsic::amdgcn_sudot8: - if (!STI.hasDot8Insts()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - return selectImpl(I, *CoverageInfo); - case Intrinsic::amdgcn_permlane16: - case Intrinsic::amdgcn_permlanex16: - if (!STI.hasPermlane16Insts()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - return selectImpl(I, *CoverageInfo); - case Intrinsic::amdgcn_mov_dpp8: - if (!STI.hasDPP8()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - return selectImpl(I, *CoverageInfo); - case Intrinsic::amdgcn_tanh: - if (!STI.hasTanhInsts()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - return selectImpl(I, *CoverageInfo); default: return selectImpl(I, *CoverageInfo); } @@ -2497,12 +2459,6 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS( if (!Subtarget->hasAsyncMark()) return false; break; - case Intrinsic::amdgcn_exp_compr: - if (!STI.hasCompressedExport()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - break; case Intrinsic::amdgcn_ds_bvh_stack_rtn: case Intrinsic::amdgcn_ds_bvh_stack_push4_pop1_rtn: case Intrinsic::amdgcn_ds_bvh_stack_push8_pop1_rtn: @@ -2527,13 +2483,7 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS( case Intrinsic::amdgcn_s_barrier_init: case Intrinsic::amdgcn_s_barrier_signal_var: return selectNamedBarrierInit(I, IntrinsicID); - case Intrinsic::amdgcn_s_wakeup_barrier: { - if (!STI.hasSWakeupBarrier()) { - diagnoseUnsupportedIntrinsic(I); - return false; - } - return selectNamedBarrierInst(I, IntrinsicID); - } + case Intrinsic::amdgcn_s_wakeup_barrier: case Intrinsic::amdgcn_s_barrier_join: case Intrinsic::amdgcn_s_get_named_barrier_state: return selectNamedBarrierInst(I, IntrinsicID); diff --git a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp index 53b3cf9bb345d..a7421b5c4fe5b 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp @@ -7896,13 +7896,6 @@ bool AMDGPULegalizerInfo::legalizeBVHIntersectRayIntrinsic( Register RayInvDir = MI.getOperand(6).getReg(); Register TDescr = MI.getOperand(7).getReg(); - if (!ST.hasGFX10_AEncoding()) { - Function &Fn = B.getMF().getFunction(); - Fn.getContext().diagnose(DiagnosticInfoUnsupported( - Fn, "intrinsic not supported on subtarget", MI.getDebugLoc())); - return false; - } - RayExtent = B.buildBitcast(I32, RayExtent).getReg(0); const bool IsGFX11 = AMDGPU::isGFX11(ST); @@ -8056,13 +8049,6 @@ bool AMDGPULegalizerInfo::legalizeBVHDualOrBVH8IntersectRayIntrinsic( Register Offsets = MI.getOperand(9).getReg(); Register TDescr = MI.getOperand(10).getReg(); - if (!ST.hasBVHDualAndBVH8Insts()) { - Function &Fn = B.getMF().getFunction(); - Fn.getContext().diagnose(DiagnosticInfoUnsupported( - Fn, "intrinsic not supported on subtarget", MI.getDebugLoc())); - return false; - } - bool IsBVH8 = cast<GIntrinsic>(MI).getIntrinsicID() == Intrinsic::amdgcn_image_bvh8_intersect_ray; const unsigned NumVDataDwords = 10; @@ -8713,16 +8699,6 @@ bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper, return true; case Intrinsic::amdgcn_av_load_b128: case Intrinsic::amdgcn_av_store_b128: { - const GCNSubtarget &ST = B.getMF().getSubtarget<GCNSubtarget>(); - if (!ST.hasFlatGlobalInsts()) { - const char *Name = IntrID == Intrinsic::amdgcn_av_load_b128 - ? "llvm.amdgcn.av.load.b128" - : "llvm.amdgcn.av.store.b128"; - Function &Fn = B.getMF().getFunction(); - Fn.getContext().diagnose(DiagnosticInfoUnsupported( - Fn, Twine(Name) + " not supported on subtarget", MI.getDebugLoc())); - return false; - } assert(MI.hasOneMemOperand() && "Expected IRTranslator to set MemOp!"); if (IntrID == Intrinsic::amdgcn_av_load_b128) B.buildLoad(MI.getOperand(0), MI.getOperand(2), **MI.memoperands_begin()); diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp index 6747f60cb9c9d..bba9d3c45ba18 100644 --- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp +++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp @@ -1593,7 +1593,9 @@ class AMDGPUAsmParser : public MCTargetAsmParser { bool isGFX13Plus() const { return AMDGPU::isGFX13Plus(getSTI()); } - bool isGFX10_AEncoding() const { return AMDGPU::isGFX10_AEncoding(getSTI()); } + bool hasBVHRayTracingInsts() const { + return getFeatureBits()[AMDGPU::FeatureBVHRayTracingInsts]; + } bool isGFX10_BEncoding() const { return AMDGPU::isGFX10_BEncoding(getSTI()); @@ -4214,7 +4216,8 @@ bool AMDGPUAsmParser::validateMIMGDataSize(const MCInst &Inst, SMLoc IDLoc) { if (VDataIdx == -1 && isGFX10Plus()) // no return image_sample return true; - if ((DMaskIdx == -1 || TFEIdx == -1) && isGFX10_AEncoding()) // intersect_ray + if ((DMaskIdx == -1 || TFEIdx == -1) && + hasBVHRayTracingInsts()) // intersect_ray return true; unsigned VDataSize = getRegOperandSize(Desc, VDataIdx); diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h index 18d341c3b229b..170ec7f85aa23 100644 --- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h +++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h @@ -640,10 +640,6 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo, /// instructions. bool hasVCvtPkIU16F32() const { return HasGFX11Insts; } - /// Return true if the target's EXP instruction has the COMPR flag, which - /// affects the meaning of the EN (enable) bits. - bool hasCompressedExport() const { return !HasGFX11Insts; } - /// Return true if the target's EXP instruction supports the NULL export /// target. bool hasNullExportTarget() const { return !HasGFX11Insts; } @@ -753,8 +749,6 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo, bool hasSubClampInsts() const { return hasGFX10_3Insts(); } - bool hasFmaLegacy32Insts() const { return hasGFX10_3Insts(); } - /// \returns SGPR allocation granularity supported by the subtarget. unsigned getSGPRAllocGranule() const { return AMDGPU::IsaInfo::getSGPRAllocGranule(*this); diff --git a/llvm/lib/Target/AMDGPU/MIMGInstructions.td b/llvm/lib/Target/AMDGPU/MIMGInstructions.td index 90764c1a7850b..bfa6830869f4a 100644 --- a/llvm/lib/Target/AMDGPU/MIMGInstructions.td +++ b/llvm/lib/Target/AMDGPU/MIMGInstructions.td @@ -1909,17 +1909,17 @@ defm IMAGE_SAMPLE_C_CD_CL_O_G16 : MIMG_Sampler <mimgopc<MIMG.NOP, MIMG.NOP, MIMG //def IMAGE_RSRC256 : MIMG_NoPattern_RSRC256 <"image_rsrc256", mimgopc<0x7e>>; //def IMAGE_SAMPLER : MIMG_NoPattern_ <"image_sampler", mimgopc<0x7f>>; -let OtherPredicates = [HasImageInsts, HasGFX10_AEncoding, isGFX10Only] in +let OtherPredicates = [HasImageInsts, HasBVHRayTracingInsts, isGFX10Only] in defm IMAGE_MSAA_LOAD_X : MIMG_NoSampler <mimgopc<MIMG.NOP, MIMG.NOP, MIMG.NOP, 0x80>, "image_msaa_load", 1, 0, 0, 1>; -let OtherPredicates = [HasImageInsts, HasGFX10_AEncoding] in { +let OtherPredicates = [HasImageInsts, HasBVHRayTracingInsts] in { defm IMAGE_MSAA_LOAD : MIMG_MSAA_Load <mimgopc<0x80, 0x18, 0x18, MIMG.NOP>, "image_msaa_load">; defm IMAGE_BVH_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 0, 0>; defm IMAGE_BVH_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x19, 0x19, 0xe6>, "image_bvh_intersect_ray", 0, 1, 0>; defm IMAGE_BVH64_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 0, 0>; defm IMAGE_BVH64_INTERSECT_RAY_a16 : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x1a, 0x1a, 0xe7>, "image_bvh64_intersect_ray", 1, 1, 0>; -} // End OtherPredicates = [HasImageInsts, HasGFX10_AEncoding] +} // End OtherPredicates = [HasImageInsts, HasBVHRayTracingInsts] defm IMAGE_BVH_DUAL_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x80, MIMG.NOP, MIMG.NOP>, "image_bvh_dual_intersect_ray", 1, 0, 1>; defm IMAGE_BVH8_INTERSECT_RAY : MIMG_IntersectRay<mimgopc<MIMG.NOP, 0x81, MIMG.NOP, MIMG.NOP>, "image_bvh8_intersect_ray", 1, 0, 0, 1>; diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index 9cf33b4ccbb87..744b048d4874a 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -7970,10 +7970,6 @@ static SDValue lowerLaneOp(const SITargetLowering &TLI, SDNode *N, MVT IntVT = MVT::getIntegerVT(ValSize); const GCNSubtarget *ST = TLI.getSubtarget(); - if ((IsPermLane16 && !ST->hasPermlane16Insts()) || - (IID == Intrinsic::amdgcn_mov_dpp8 && !ST->hasDPP8())) - return emitRemovedIntrinsicError(DAG, SL, VT); - unsigned SplitSize = 32; if (IID == Intrinsic::amdgcn_update_dpp && (ValSize % 64 == 0) && ST->hasDPALU_DPP() && @@ -10956,17 +10952,9 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op, return emitRemovedIntrinsicError(DAG, DL, VT); return DAG.getNode(AMDGPUISD::RCP_LEGACY, DL, VT, Op.getOperand(1)); case Intrinsic::amdgcn_fma_legacy: - if (!Subtarget->hasFmaLegacy32Insts()) - return emitRemovedIntrinsicError(DAG, DL, VT); - return SDValue(); case Intrinsic::amdgcn_sudot4: case Intrinsic::amdgcn_sudot8: - if (!Subtarget->hasDot8Insts()) - return emitRemovedIntrinsicError(DAG, DL, VT); - return SDValue(); case Intrinsic::amdgcn_tanh: - if (!Subtarget->hasTanhInsts()) - return emitRemovedIntrinsicError(DAG, DL, VT); return SDValue(); case Intrinsic::amdgcn_rsq_clamp: { if (Subtarget->getGeneration() < AMDGPUSubtarget::VOLCANIC_ISLANDS) @@ -11836,11 +11824,6 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op, assert(NodePtr.getValueType() == MVT::i64); assert(RayDir.getValueType() == MVT::v3f32); - if (!Subtarget->hasBVHDualAndBVH8Insts()) { - emitRemovedIntrinsicError(DAG, DL, Op.getValueType()); - return SDValue(); - } - bool IsBVH8 = IntrID == Intrinsic::amdgcn_image_bvh8_intersect_ray; const unsigned NumVDataDwords = 10; const unsigned NumVAddrDwords = IsBVH8 ? 11 : 12; @@ -11881,11 +11864,6 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op, assert(RayDir.getValueType() == MVT::v3f16 || RayDir.getValueType() == MVT::v3f32); - if (!Subtarget->hasGFX10_AEncoding()) { - emitRemovedIntrinsicError(DAG, DL, Op.getValueType()); - return SDValue(); - } - const bool IsGFX11 = AMDGPU::isGFX11(*Subtarget); const bool IsGFX11Plus = AMDGPU::isGFX11Plus(*Subtarget); const bool IsGFX12Plus = AMDGPU::isGFX12Plus(*Subtarget); @@ -12077,14 +12055,6 @@ SDValue SITargetLowering::LowerINTRINSIC_W_CHAIN(SDValue Op, Chain, Ptr, MII->getMemOperand()); } case Intrinsic::amdgcn_av_load_b128: { - if (!Subtarget->hasFlatGlobalInsts()) { - DAG.getContext()->diagnose(DiagnosticInfoUnsupported( - DAG.getMachineFunction().getFunction(), - "llvm.amdgcn.av.load.b128 not supported on subtarget", - DL.getDebugLoc())); - return DAG.getMergeValues( - {DAG.getPOISON(Op->getValueType(0)), Op->getOperand(0)}, DL); - } MemIntrinsicSDNode *MII = cast<MemIntrinsicSDNode>(Op); SDValue Chain = Op->getOperand(0); SDValue Ptr = Op->getOperand(2); @@ -12274,11 +12244,6 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, switch (IntrinsicID) { case Intrinsic::amdgcn_exp_compr: { - if (!Subtarget->hasCompressedExport()) { - DAG.getContext()->diagnose(DiagnosticInfoUnsupported( - DAG.getMachineFunction().getFunction(), - "intrinsic not supported on subtarget", DL.getDebugLoc())); - } SDValue Src0 = Op.getOperand(4); SDValue Src1 = Op.getOperand(5); // Hack around illegal type on SI by directly selecting it. @@ -12793,13 +12758,6 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, Ptr, MII->getMemOperand()); } case Intrinsic::amdgcn_av_store_b128: { - if (!Subtarget->hasFlatGlobalInsts()) { - DAG.getContext()->diagnose(DiagnosticInfoUnsupported( - DAG.getMachineFunction().getFunction(), - "llvm.amdgcn.av.store.b128 not supported on subtarget", - DL.getDebugLoc())); - return Op->getOperand(0); // return the input chain - } MemIntrinsicSDNode *MII = cast<MemIntrinsicSDNode>(Op); SDValue Chain = Op->getOperand(0); SDValue Ptr = Op->getOperand(2); diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp index 71d95a23e30d3..e5d9c4237c2b4 100644 --- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp +++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp @@ -2640,10 +2640,6 @@ bool isGCN3Encoding(const MCSubtargetInfo &STI) { return STI.hasFeature(AMDGPU::FeatureGCN3Encoding); } -bool isGFX10_AEncoding(const MCSubtargetInfo &STI) { - return STI.hasFeature(AMDGPU::FeatureGFX10_AEncoding); -} - bool isGFX10_BEncoding(const MCSubtargetInfo &STI) { return STI.hasFeature(AMDGPU::FeatureGFX10_BEncoding); } diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h index 4b08c08c5a60d..430cff481ad74 100644 --- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h +++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h @@ -1519,7 +1519,6 @@ bool supportsWGP(const MCSubtargetInfo &STI); bool isNotGFX12Plus(const MCSubtargetInfo &STI); bool isNotGFX11Plus(const MCSubtargetInfo &STI); bool isGCN3Encoding(const MCSubtargetInfo &STI); -bool isGFX10_AEncoding(const MCSubtargetInfo &STI); bool isGFX10_BEncoding(const MCSubtargetInfo &STI); bool hasGFX10_3Insts(const MCSubtargetInfo &STI); bool isGFX10_3_GFX11(const MCSubtargetInfo &STI); diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td index ebbe89981e4ef..df723dfb93c44 100644 --- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td @@ -176,7 +176,7 @@ defm V_MAD_LEGACY_F32 : VOP3Inst <"v_mad_legacy_f32", VOP3_Profile<VOP_F32_F32_F defm V_MAD_F32 : VOP3Inst <"v_mad_f32", VOP3_Profile<VOP_F32_F32_F32_F32>, any_fmad>; } // End SubtargetPredicate = HasMadMacInsts -let SubtargetPredicate = HasFmaLegacy32 in +let SubtargetPredicate = HasFmaLegacy32Insts in defm V_FMA_LEGACY_F32 : VOP3Inst <"v_fma_legacy_f32", VOP3_Profile<VOP_F32_F32_F32_F32>, int_amdgcn_fma_legacy>; diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp index bf169c5b5ef42..2a698098c90a1 100644 --- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp +++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp @@ -520,6 +520,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["gfx12-insts"] = true; Features["atomic-fadd-rtn-insts"] = true; Features["image-insts"] = true; + Features["bvh-ray-tracing-insts"] = true; Features["cube-insts"] = true; Features["lerp-inst"] = true; Features["sad-insts"] = true; @@ -554,6 +555,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["gfx11-insts"] = true; Features["atomic-fadd-rtn-insts"] = true; Features["image-insts"] = true; + Features["bvh-ray-tracing-insts"] = true; Features["cube-insts"] = true; Features["lerp-inst"] = true; Features["sad-insts"] = true; @@ -597,6 +599,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["gfx11-insts"] = true; Features["atomic-fadd-rtn-insts"] = true; Features["image-insts"] = true; + Features["bvh-ray-tracing-insts"] = true; Features["cube-insts"] = true; Features["lerp-inst"] = true; Features["sad-insts"] = true; @@ -633,6 +636,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["gfx10-insts"] = true; Features["gfx10-3-insts"] = true; Features["image-insts"] = true; + Features["bvh-ray-tracing-insts"] = true; Features["s-memrealtime"] = true; Features["s-memtime-inst"] = true; Features["gws"] = true; @@ -660,6 +664,8 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, case GK_GFX1013: case GK_GFX1010: case GK_GFX10_1_GENERIC: + if (Kind == GK_GFX1013) + Features["bvh-ray-tracing-insts"] = true; Features["dl-insts"] = true; Features["ci-insts"] = true; Features["16-bit-insts"] = true; diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir index b74b68bf95b09..00297e8e948ff 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/inst-select-amdgcn.exp.compr.mir @@ -1,7 +1,6 @@ # RUN: llc -mtriple=amdgcn -mcpu=gfx600 -run-pass=instruction-select -verify-machineinstrs %s -o - | FileCheck %s -# RUN: not llc -mtriple=amdgcn -mcpu=gfx1100 -run-pass=instruction-select -verify-machineinstrs -global-isel-abort=0 %s -o - 2>&1 | FileCheck --check-prefix=ERR %s -# ERR: error: <unknown>:0:0: in function exp0 void (): intrinsic not supported on subtarget +# Note: gfx11+ rejection of llvm.amdgcn.exp.compr is now enforced at IR translation (see llvm.amdgcn.exp.compr.ll), not during instruction selection. --- name: exp0 diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll index 8cd03fd12ef2f..e8ef0336c2ff2 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll @@ -3,7 +3,7 @@ ; RUN: llc -global-isel -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1013 < %s | FileCheck -check-prefixes=GCN,GFX10,GFX1013 %s ; RUN: llc -global-isel -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-TRUE16 %s ; RUN: llc -global-isel -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1100 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GCN,GFX11,GFX11-FAKE16 %s -; RUN: not --crash llc -global-isel -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1012 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1012 < %s -filetype=null 2>&1 | FileCheck -check-prefix=ERR %s ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr) @@ -22,7 +22,7 @@ define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_ ; GCN-NEXT: image_bvh_intersect_ray v[0:3], v[0:10], s[0:3] ; GCN-NEXT: s_waitcnt vmcnt(0) ; GCN-NEXT: ; return to shader part epilog -; ERR: in function image_bvh_intersect_ray{{.*}}intrinsic not supported on subtarget +; ERR: in function @image_bvh_intersect_ray{{.*}}requires target feature 'bvh-ray-tracing-insts' %v = call <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(i32 %node_ptr, float %ray_extent, <3 x float> %ray_origin, <3 x float> %ray_dir, <3 x float> %ray_inv_dir, <4 x i32> %tdescr) %r = bitcast <4 x i32> %v to <4 x float> ret <4 x float> %r diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll index 49d1027ab11e3..6e31275e11032 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.dual_intersect_ray.ll @@ -1,12 +1,12 @@ ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py -; RUN: not --crash llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX12-SDAG %s ; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX12-GISEL %s declare {<10 x i32>, <3 x float>, <3 x float>} @llvm.amdgcn.image.bvh.dual.intersect.ray(i64, float, i8, <3 x float>, <3 x float>, <2 x i32>, <4 x i32>) -; ERR: in function image_bvh_dual_intersect_ray{{.*}}intrinsic not supported on subtarget +; ERR: in function @image_bvh_dual_intersect_ray{{.*}}requires target feature 'bvh-dual-bvh-8-insts' define amdgpu_ps <10 x float> @image_bvh_dual_intersect_ray(i64 %node_ptr, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, <2 x i32> %offsets, <4 x i32> inreg %tdescr, ptr addrspace(1) %origin, ptr addrspace(1) %dir) { ; GFX12-SDAG-LABEL: image_bvh_dual_intersect_ray: ; GFX12-SDAG: ; %bb.0: ; %main_body diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll index ef368bf06528a..e70f32c03a3e5 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll @@ -3,8 +3,9 @@ ; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -strict-whitespace -check-prefix=GCN %s ; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -strict-whitespace -check-prefix=GCN %s ; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -strict-whitespace -check-prefix=ERR %s +; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1100 < %s 2>&1 | FileCheck -strict-whitespace -check-prefix=ERR %s -; ERR: error: <unknown>:0:0: in function test_export_compr_zeroes_v2f16 void (): intrinsic not supported on subtarget +; ERR: error: <unknown>:0:0: in function @test_export_compr_zeroes_v2f16 void (): llvm.amdgcn.exp.compr requires target feature '-gfx11-insts' declare void @llvm.amdgcn.exp.compr.v2f16(i32, i32, <2 x half>, <2 x half>, i1, i1) #0 declare void @llvm.amdgcn.exp.compr.v2i16(i32, i32, <2 x i16>, <2 x i16>, i1, i1) #0 diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll index c1ad188671f64..c4cb98a25e866 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fma.legacy.ll @@ -13,7 +13,7 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s -; ERR: error: <unknown>:0:0: in function v_fma float (float, float, float): intrinsic not supported on subtarget +; ERR: error: <unknown>:0:0: in function @v_fma float (float, float, float): llvm.amdgcn.fma.legacy requires target feature 'fma-legacy32-insts' define float @v_fma(float %a, float %b, float %c) { ; GFX10-LABEL: v_fma: diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll index 77f0f00729433..0e2da14c84228 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll @@ -3,7 +3,7 @@ ; RUN: llc -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1013 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX10,GFX10-GISEL,GFX1013-GISEL %s ; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX10,GFX10-SDAG,GFX1030-SDAG %s ; RUN: llc -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX10,GFX10-GISEL,GFX1030-GISEL %s -; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX11-SDAG %s ; RUN: llc -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX11-GISEL %s ; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1100 -mattr=-real-true16 < %s | FileCheck -check-prefixes=PRE-GFX12,GFX11-SDAG %s @@ -13,6 +13,9 @@ ; RUN: llc -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1200 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s ; RUN: llc -global-isel=1 -global-isel-abort=2 -mtriple=amdgcn -mcpu=gfx1200 -mattr=-real-true16 < %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s +; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1012 < %s 2>&1 | FileCheck -check-prefix=ERR %s + ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f32(uint node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(uint node_ptr, float ray_extent, float3 ray_origin, half3 ray_dir, half3 ray_inv_dir, uint4 texture_descr) ; uint4 llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(ulong node_ptr, float ray_extent, float3 ray_origin, float3 ray_dir, float3 ray_inv_dir, uint4 texture_descr) @@ -23,7 +26,7 @@ declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i32.v4f16(i32, float, <3 declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f32(i64, float, <3 x float>, <3 x float>, <3 x float>, <4 x i32>) declare <4 x i32> @llvm.amdgcn.image.bvh.intersect.ray.i64.v4f16(i64, float, <3 x float>, <3 x half>, <3 x half>, <4 x i32>) -; ERR: in function image_bvh_intersect_ray{{.*}}intrinsic not supported on subtarget +; ERR: in function @image_bvh_intersect_ray{{.*}}requires target feature 'bvh-ray-tracing-insts' ; Arguments are flattened to represent the actual VGPR_A layout, so we have no ; extra moves in the generated kernel. define amdgpu_ps <4 x float> @image_bvh_intersect_ray(i32 %node_ptr, float %ray_extent, float %ray_origin_x, float %ray_origin_y, float %ray_origin_z, float %ray_dir_x, float %ray_dir_y, float %ray_dir_z, float %ray_inv_dir_x, float %ray_inv_dir_y, float %ray_inv_dir_z, <4 x i32> inreg %tdescr) { diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll index cf8c681df6d5a..3540e4dfb3c23 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mov.dpp8.ll @@ -8,7 +8,7 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s -; ERR: error: <unknown>:0:0: in function dpp8_test void (ptr addrspace(1), i32): intrinsic not supported on subtarget +; ERR: error: <unknown>:0:0: in function @dpp8_test void (ptr addrspace(1), i32): llvm.amdgcn.mov.dpp8 requires target feature 'dpp8' ; GFX10PLUS-LABEL: {{^}}dpp8_test: ; GFX10PLUS: v_mov_b32_e32 [[SRC:v[0-9]+]], s{{[0-9]+}} diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll index 6a799e5d71b7b..b3d2eeb94961d 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane.ll @@ -11,7 +11,7 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx900 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s -; ERR: error: <unknown>:0:0: in function v_permlane16_b32_vss_i32 void (ptr addrspace(1), i32, i32, i32): intrinsic not supported on subtarget +; ERR: error: <unknown>:0:0: in function @v_permlane16_b32_vss_i32 void (ptr addrspace(1), i32, i32, i32): llvm.amdgcn.permlane16 requires target feature 'permlane16-insts' declare i32 @llvm.amdgcn.permlane16(i32, i32, i32, i32, i1, i1) declare i32 @llvm.amdgcn.permlanex16(i32, i32, i32, i32, i1, i1) diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll index 1a9cc7615a518..30e6936245dac 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot4.ll @@ -8,7 +8,7 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s -; ERR: error: <unknown>:0:0: in function test_llvm_amdgcn_sudot4_uu i32 (i32, i32, i32): intrinsic not supported on subtarget +; ERR: error: <unknown>:0:0: in function @test_llvm_amdgcn_sudot4_uu i32 (i32, i32, i32): llvm.amdgcn.sudot4 requires target feature 'dot8-insts' declare i32 @llvm.amdgcn.sudot4(i1 %asign, i32 %a, i1 %bsign, i32 %b, i32 %c, i1 %clamp) diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll index 4c3dd27ae0300..eeef15d72ed0b 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sudot8.ll @@ -8,7 +8,7 @@ ; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s ; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgcn -mcpu=gfx950 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s -; ERR: error: <unknown>:0:0: in function test_llvm_amdgcn_sudot8_uu i32 (i32, i32, i32): intrinsic not supported on subtarget +; ERR: error: <unknown>:0:0: in function @test_llvm_amdgcn_sudot8_uu i32 (i32, i32, i32): llvm.amdgcn.sudot8 requires target feature 'dot8-insts' declare i32 @llvm.amdgcn.sudot8(i1 %asign, i32 %a, i1 %bsign, i32 %b, i32 %c, i1 %clamp) diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll index 184a89e8ce854..4bc6140ba5bb5 100644 --- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll @@ -8,6 +8,11 @@ ; xUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1310 -mattr=+real-true16 %s -o - | FileCheck -check-prefixes=GFX13,GFX13-GISEL-REAL16 %s ; xUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1310 -mattr=-real-true16 %s -o - | FileCheck -check-prefixes=GFX13,GFX13-GISEL-FAKE16 %s +; RUN: not llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR %s + +; ERR: error: <unknown>:0:0: in function @tanh_f32 float (float): llvm.amdgcn.tanh requires target feature 'tanh-insts' + define float @tanh_f32(float %src) { ; GFX1250-LABEL: tanh_f32: ; GFX1250: ; %bb.0: diff --git a/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll b/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll index 0b907ebfc1ed1..0988b72f9e101 100644 --- a/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll +++ b/llvm/test/CodeGen/AMDGPU/s-wakeup-barrier.ll @@ -2,6 +2,11 @@ ; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX1250-SDAG %s ; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX1250-GISEL %s +; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s +; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s + +; ERR: error: <unknown>:0:0: in function @kernel1 void (ptr addrspace(1), ptr addrspace(3)): llvm.amdgcn.s.wakeup.barrier requires target feature 's-wakeup-barrier-inst' + @bar = internal addrspace(3) global target("amdgcn.named.barrier", 0) poison define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) #0 { diff --git a/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll b/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll index 942a384184924..d6af59623d074 100644 --- a/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll +++ b/llvm/test/CodeGen/AMDGPU/unsupported-av-load.ll @@ -7,7 +7,7 @@ ; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s define <4 x i32> @av_load_b128(ptr addrspace(1) %addr) { -; ERR: error: {{.*}}: in function av_load_b128 {{.*}}: llvm.amdgcn.av.load.b128 not supported on subtarget +; ERR: error: {{.*}}: in function @av_load_b128 {{.*}}: llvm.amdgcn.av.load.b128 requires target feature 'flat-global-insts' entry: %data = call <4 x i32> @llvm.amdgcn.av.load.b128.p1(ptr addrspace(1) %addr, metadata !0) ret <4 x i32> %data diff --git a/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll b/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll index 4bd8945fea0b6..6015b489c8548 100644 --- a/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll +++ b/llvm/test/CodeGen/AMDGPU/unsupported-av-store.ll @@ -7,7 +7,7 @@ ; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx810 -filetype=null < %s 2>&1 | FileCheck -check-prefix=ERR %s define void @av_store_b128(ptr addrspace(1) %addr, <4 x i32> %data) { -; ERR: error: {{.*}}: in function av_store_b128 {{.*}}: llvm.amdgcn.av.store.b128 not supported on subtarget +; ERR: error: {{.*}}: in function @av_store_b128 {{.*}}: llvm.amdgcn.av.store.b128 requires target feature 'flat-global-insts' entry: call void @llvm.amdgcn.av.store.b128.p1(ptr addrspace(1) %addr, <4 x i32> %data, metadata !0) ret void _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
