https://github.com/mariusz-sikora-at-amd created https://github.com/llvm/llvm-project/pull/210713
- builtins with documentation - instructions - new subtarget feature only for gfx13 - GlobalISel RegBankLegalize rules - tests >From fc6fa9a27419fb065b960f842f6d1f0c7a69e823 Mon Sep 17 00:00:00 2001 From: Mariusz Sikora <[email protected]> Date: Mon, 1 Jun 2026 07:49:47 -0400 Subject: [PATCH] [AMDGPU] Builtins and intrinsics for v_cvt_scalef32_pk32_(fp|bf)6_f32 - builtins with documentation - instructions - new subtarget feature only for gfx13 - GlobalISel RegBankLegalize rules - tests --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 12 + .../include/clang/Basic/BuiltinsAMDGPUDocs.td | 31 + .../CodeGen/amdgpu-builtin-is-invocable.c | 2 +- .../CodeGen/amdgpu-builtin-processor-is.c | 2 +- .../CodeGenCXX/dynamic-cast-address-space.cpp | 4 +- .../CodeGenOpenCL/builtins-amdgcn-gfx13.cl | 48 ++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 10 + llvm/lib/Target/AMDGPU/AMDGPU.td | 5 + .../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 6 + .../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 2 + llvm/lib/Target/AMDGPU/VOP3Instructions.td | 6 +- llvm/lib/TargetParser/AMDGPUTargetParser.cpp | 1 + .../llvm.amdgcn.cvt.scalef32.pk.gfx13.ll | 646 ++++++++++++++++++ 13 files changed, 768 insertions(+), 7 deletions(-) create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx13.ll diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 668646d8d0070..40fb90af62827 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -1274,6 +1274,18 @@ def __builtin_amdgcn_cooperative_atomic_store_16x8B : AMDGPUBuiltin<"void(_ExtVe def __builtin_amdgcn_cooperative_atomic_load_8x16B : AMDGPUBuiltin<"_ExtVector<4, int>(_ExtVector<4, int> *, _Constant int, char const *)", [Const], "gfx1250-insts,wavefrontsize32">; def __builtin_amdgcn_cooperative_atomic_store_8x16B : AMDGPUBuiltin<"void(_ExtVector<4, int> *, _ExtVector<4, int>, _Constant int, char const *)", [Const], "gfx1250-insts,wavefrontsize32">; +//===----------------------------------------------------------------------===// +// GFX13+ only builtins. +//===----------------------------------------------------------------------===// +def __builtin_amdgcn_cvt_scalef32_pk32_bf6_f32 : AMDGPUBuiltin<"_Vector<6, unsigned int>(_Vector<32, float>, float)", [Const], "f32-to-fp6bf6-cvt-scale-insts"> { + let Documentation = [DocCvtScaleF32Pk32BF6F32]; + let ArgNames = ["src", "scale"]; +} +def __builtin_amdgcn_cvt_scalef32_pk32_fp6_f32 : AMDGPUBuiltin<"_Vector<6, unsigned int>(_Vector<32, float>, float)", [Const], "f32-to-fp6bf6-cvt-scale-insts"> { + let Documentation = [DocCvtScaleF32Pk32FP6F32]; + let ArgNames = ["src", "scale"]; +} + //===----------------------------------------------------------------------===// // Image builtins //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td index c78c8ec291eef..88b709471bd44 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td @@ -71,6 +71,13 @@ on all AMDGCN targets unless otherwise noted. }]; } +def DocCatAMDGPUConversion : DocumentationCategory<"Conversion Builtins"> { + let Content = [{ +These builtins provide access to AMDGPU conversion instructions, including conversions +between standard floating-point types, integer types and packed low-precision formats. +}]; +} + //===----------------------------------------------------------------------===// // ABI / Special Register Builtins — Documentation records //===----------------------------------------------------------------------===// @@ -762,3 +769,27 @@ Loads data from a buffer using an SGPR buffer descriptor. The volatile bit is not permitted for this intrinsic. }]; } + +//===----------------------------------------------------------------------===// +// Conversion Builtins +//===----------------------------------------------------------------------===// + +def DocCvtScaleF32Pk32BF6F32 : Documentation { + let Category = DocCatAMDGPUConversion; + let Content = [{ +Scales a packed 32-component single-precision floating-point input ``src`` +using the scale factor ``scale``, then converts the values to a packed +32-component BF6 value, returned as a 6-element vector of 32-bit unsigned +integers. +}]; +} + +def DocCvtScaleF32Pk32FP6F32 : Documentation { + let Category = DocCatAMDGPUConversion; + let Content = [{ +Scales a packed 32-component single-precision floating-point input ``src`` +using the scale factor ``scale``, then converts the values to a packed +32-component FP6 value, returned as a 6-element vector of 32-bit unsigned +integers. +}]; +} diff --git a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c index c894dc5589bdf..f5e852be47721 100644 --- a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c +++ b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c @@ -59,7 +59,7 @@ void foo() { // AMDGCN-GFX1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } // AMDGCN-GFX1010: attributes #[[ATTR2]] = { noreturn nounwind } //. -// 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,+cvt-pknorm-vop3-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,+gfx13-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,+cvt-pknorm-vop3-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,+f32-to-fp6bf6-cvt-scale-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,+gfx13-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) } // AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind } diff --git a/clang/test/CodeGen/amdgpu-builtin-processor-is.c b/clang/test/CodeGen/amdgpu-builtin-processor-is.c index 01b347aac31f7..04c31e0f0da38 100644 --- a/clang/test/CodeGen/amdgpu-builtin-processor-is.c +++ b/clang/test/CodeGen/amdgpu-builtin-processor-is.c @@ -68,7 +68,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,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-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,+gfx13-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,+cvt-pknorm-vop3-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,+f32-to-fp6bf6-cvt-scale-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,+gfx13-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) } // AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind } diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp index c6543eae669e9..2c56cf697dfc0 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,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-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,+gfx13-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,+cvt-pknorm-vop3-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,+f32-to-fp6bf6-cvt-scale-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,+gfx13-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,+bvh-ray-tracing-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+cvt-pknorm-vop3-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,+gfx13-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,+cvt-pknorm-vop3-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,+f32-to-fp6bf6-cvt-scale-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,+gfx13-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-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl new file mode 100644 index 0000000000000..a6e0b15921860 --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl @@ -0,0 +1,48 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1310 -emit-llvm -o - %s | FileCheck %s + +// REQUIRES: amdgpu-registered-target + +typedef unsigned int __attribute__((ext_vector_type(6))) uint6; +typedef float __attribute__((ext_vector_type(32))) float32; + +// CHECK-LABEL: @test_cvt_scalef32_pk32_bf6_f32( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[SRCF32_ADDR:%.*]] = alloca <32 x float>, align 128, addrspace(5) +// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5) +// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store <32 x float> [[SRCF32:%.*]], ptr addrspace(5) [[SRCF32_ADDR]], align 128 +// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load <32 x float>, ptr addrspace(5) [[SRCF32_ADDR]], align 128 +// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]], align 4 +// CHECK-NEXT: [[TMP2:%.*]] = call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32(<32 x float> [[TMP0]], float [[TMP1]]) +// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store <6 x i32> [[TMP2]], ptr addrspace(1) [[TMP3]], align 32 +// CHECK-NEXT: ret void +// +void test_cvt_scalef32_pk32_bf6_f32(global uint6 *out, float32 srcf32, float src) +{ + *out = __builtin_amdgcn_cvt_scalef32_pk32_bf6_f32(srcf32, src); +} + +// CHECK-LABEL: @test_cvt_scalef32_pk32_fp6_f32( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[SRCF32_ADDR:%.*]] = alloca <32 x float>, align 128, addrspace(5) +// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5) +// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store <32 x float> [[SRCF32:%.*]], ptr addrspace(5) [[SRCF32_ADDR]], align 128 +// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(5) [[SRC_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load <32 x float>, ptr addrspace(5) [[SRCF32_ADDR]], align 128 +// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[SRC_ADDR]], align 4 +// CHECK-NEXT: [[TMP2:%.*]] = call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32(<32 x float> [[TMP0]], float [[TMP1]]) +// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store <6 x i32> [[TMP2]], ptr addrspace(1) [[TMP3]], align 32 +// CHECK-NEXT: ret void +// +void test_cvt_scalef32_pk32_fp6_f32(global uint6 *out, float32 srcf32, float src) +{ + *out = __builtin_amdgcn_cvt_scalef32_pk32_fp6_f32(srcf32, src); +} + diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 917730cea404d..565637b36131c 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -3272,6 +3272,16 @@ def int_amdgcn_global_prefetch : ClangBuiltin<"__builtin_amdgcn_global_prefetch" [SDNPMemOperand] >; +//===----------------------------------------------------------------------===// +// GFX13 Intrinsics +//===----------------------------------------------------------------------===// + +// v6i32 @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32 <src0> <scale> +def int_amdgcn_cvt_scalef32_pk32_bf6_f32 : AMDGPUCvtScaleF32Intrinsic<llvm_v6i32_ty, llvm_v32f32_ty, "cvt_scalef32_pk32_bf6_f32">; + +// v6i32 @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32 <src0> <scale> +def int_amdgcn_cvt_scalef32_pk32_fp6_f32 : AMDGPUCvtScaleF32Intrinsic<llvm_v6i32_ty, llvm_v32f32_ty, "cvt_scalef32_pk32_fp6_f32">; + //===----------------------------------------------------------------------===// // Deep learning intrinsics. //===----------------------------------------------------------------------===// diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index ed454806d7193..0a7115055291b 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -479,6 +479,10 @@ defm F16BF16ToFP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"f16bf16-to-fp "Has f16bf16 to fp6bf6 conversion scale instructions" >; +defm F32ToFP6BF6ConversionScaleInsts : AMDGPUSubtargetFeature<"f32-to-fp6bf6-cvt-scale-insts", + "Has f32 to fp6bf6 conversion scale instructions" +>; + defm F32ToF16BF16ConversionSRInsts : AMDGPUSubtargetFeature<"f32-to-f16bf16-cvt-sr-insts", "Has f32 to f16bf16 conversion scale instructions" >; @@ -2365,6 +2369,7 @@ def FeatureISAVersion13 : FeatureSet< FeatureGloballyAddressableScratch, FeatureCvtPkF16F32Inst, FeatureF16BF16ToFP6BF6ConversionScaleInsts, + FeatureF32ToFP6BF6ConversionScaleInsts, FeatureIEEEMinimumMaximumInsts, FeatureSWakeupBarrier, FeatureClusters, diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp index df5b5719afd05..b9bc59a79a642 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp @@ -2230,6 +2230,12 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST, .Any({{DivV6S32}, {{VgprV6S32}, {IntrId, VgprV32S16, Vgpr32}}}) .Any({{UniV6S32}, {{UniInVgprV6S32}, {IntrId, VgprV32S16, Vgpr32}}}); + addRulesForIOpcs( + {amdgcn_cvt_scalef32_pk32_bf6_f32, amdgcn_cvt_scalef32_pk32_fp6_f32}, + Standard) + .Any({{DivV6S32}, {{VgprV6S32}, {IntrId, VgprV32S32, Vgpr32}}}) + .Any({{UniV6S32}, {{UniInVgprV6S32}, {IntrId, VgprV32S32, Vgpr32}}}); + addRulesForIOpcs( {amdgcn_cvt_scalef32_pk_fp8_f32, amdgcn_cvt_scalef32_pk_bf8_f32}, Standard) diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp index 200234c23c886..21472bb886196 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -4761,6 +4761,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { case Intrinsic::amdgcn_cvt_scalef32_pk32_bf6_f16: case Intrinsic::amdgcn_cvt_scalef32_pk32_fp6_bf16: case Intrinsic::amdgcn_cvt_scalef32_pk32_bf6_bf16: + case Intrinsic::amdgcn_cvt_scalef32_pk32_fp6_f32: + case Intrinsic::amdgcn_cvt_scalef32_pk32_bf6_f32: case Intrinsic::amdgcn_cvt_scalef32_f16_fp8: case Intrinsic::amdgcn_cvt_scalef32_f16_bf8: case Intrinsic::amdgcn_cvt_scalef32_f32_fp8: diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td index df723dfb93c44..6611649207eea 100644 --- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td @@ -2072,9 +2072,9 @@ let SubtargetPredicate = HasAshrPkInsts, True16Predicate = UseRealTrue16Insts in defm : AshrPkU8Pat<V_ASHR_PK_U8_I32_t16_e64, 0, 255, 1>; } -let SubtargetPredicate = isGFX13Plus in { - defm V_CVT_SCALEF32_PK32_BF6_F32 : VOP3Inst<"v_cvt_scalef32_pk32_bf6_f32", VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V6I32_V32F32_F32>>; - defm V_CVT_SCALEF32_PK32_FP6_F32 : VOP3Inst<"v_cvt_scalef32_pk32_fp6_f32", VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V6I32_V32F32_F32>>; +let SubtargetPredicate = HasF32ToFP6BF6ConversionScaleInsts in { + defm V_CVT_SCALEF32_PK32_BF6_F32 : VOP3Inst<"v_cvt_scalef32_pk32_bf6_f32", VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V6I32_V32F32_F32>, int_amdgcn_cvt_scalef32_pk32_bf6_f32>; + defm V_CVT_SCALEF32_PK32_FP6_F32 : VOP3Inst<"v_cvt_scalef32_pk32_fp6_f32", VOP3_CVT_SCALEF32_PK_F864_Profile<VOP_V6I32_V32F32_F32>, int_amdgcn_cvt_scalef32_pk32_fp6_f32>; } //===----------------------------------------------------------------------===// diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp index 16acaaeb7b71c..ff752e5dbef5d 100644 --- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp +++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp @@ -463,6 +463,7 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["atomic-ds-pk-add-16-insts"] = true; Features["s-wakeup-barrier-inst"] = true; Features["f16bf16-to-fp6bf6-cvt-scale-insts"] = true; + Features["f32-to-fp6bf6-cvt-scale-insts"] = true; Features["clusters"] = true; Features["cube-insts"] = true; Features["lerp-inst"] = true; diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx13.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx13.ll new file mode 100644 index 0000000000000..4bbb893087686 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.scalef32.pk.gfx13.ll @@ -0,0 +1,646 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4 +; RUN: llc -global-isel=0 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefixes=GFX13-SDAG %s +; RUN: llc -global-isel=1 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefixes=GFX13-GISEL %s + +declare <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32(<32 x float> %src, float %scale) +declare <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32(<32 x float> %src, float %scale) + +define amdgpu_kernel void @test_cvt_scalef32_pk32_bf6_f32_v(ptr addrspace(1) %out, <32 x float> %src, float %scale) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_bf6_f32_v: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x3 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_load_b32 s2, s[4:5], 0x124 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], s2 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_bf6_f32_v: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x3 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x124 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], s0 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32(<32 x float> %src, float %scale) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_bf6_f32_s(ptr addrspace(1) %out, <32 x float> %src) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_bf6_f32_s: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x2 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_bf6_f32_s: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x2 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32(<32 x float> %src, float 100.0) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_bf6_f32_v_inreg_src(ptr addrspace(1) %out, <32 x float> inreg %src, float %scale) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_bf6_f32_v_inreg_src: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x3 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b32 s2, s[4:5], 0x124 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], s2 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_bf6_f32_v_inreg_src: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x3 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x124 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], s0 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32(<32 x float> %src, float %scale) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_bf6_f32_s_inreg_src(ptr addrspace(1) %out, <32 x float> inreg %src) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_bf6_f32_s_inreg_src: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x2 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_bf6_f32_s_inreg_src: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x2 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_bf6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.bf6.f32(<32 x float> %src, float 100.0) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_fp6_f32_v(ptr addrspace(1) %out, <32 x float> %src, float %scale) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_fp6_f32_v: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x3 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_load_b32 s2, s[4:5], 0x124 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], s2 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_fp6_f32_v: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x3 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x124 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], s0 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32(<32 x float> %src, float %scale) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_fp6_f32_s(ptr addrspace(1) %out, <32 x float> %src) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_fp6_f32_s: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x2 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_fp6_f32_s: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x2 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32(<32 x float> %src, float 100.0) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_fp6_f32_v_inreg_src(ptr addrspace(1) %out, <32 x float> inreg %src, float %scale) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_fp6_f32_v_inreg_src: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x3 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b32 s2, s[4:5], 0x124 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], s2 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_fp6_f32_v_inreg_src: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x3 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x124 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], s0 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32(<32 x float> %src, float %scale) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} + +define amdgpu_kernel void @test_cvt_scalef32_pk32_fp6_f32_s_inreg_src(ptr addrspace(1) %out, <32 x float> inreg %src) { +; GFX13-SDAG-LABEL: test_cvt_scalef32_pk32_fp6_f32_s_inreg_src: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_clause 0x2 +; GFX13-SDAG-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-SDAG-NEXT: s_load_b512 s[8:23], s[4:5], 0xe4 nv +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v0, s36 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v1, s37 :: v_dual_mov_b32 v2, s38 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v3, s39 :: v_dual_mov_b32 v4, s40 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v5, s41 :: v_dual_mov_b32 v6, s42 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v7, s43 :: v_dual_mov_b32 v8, s44 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v9, s45 :: v_dual_mov_b32 v10, s46 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v11, s47 :: v_dual_mov_b32 v12, s48 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v13, s49 :: v_dual_mov_b32 v14, s50 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v15, s51 :: v_dual_mov_b32 v16, s8 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v17, s9 :: v_dual_mov_b32 v18, s10 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v19, s11 :: v_dual_mov_b32 v20, s12 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v21, s13 :: v_dual_mov_b32 v22, s14 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v23, s15 :: v_dual_mov_b32 v24, s16 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v25, s17 :: v_dual_mov_b32 v26, s18 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v27, s19 :: v_dual_mov_b32 v28, s20 +; GFX13-SDAG-NEXT: v_dual_mov_b32 v29, s21 :: v_dual_mov_b32 v30, s22 +; GFX13-SDAG-NEXT: v_mov_b32_e32 v31, s23 +; GFX13-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-SDAG-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-SDAG-NEXT: s_clause 0x1 +; GFX13-SDAG-NEXT: global_store_b64 v32, v[4:5], s[0:1] offset:16 +; GFX13-SDAG-NEXT: global_store_b128 v32, v[0:3], s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: test_cvt_scalef32_pk32_fp6_f32_s_inreg_src: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_clause 0x2 +; GFX13-GISEL-NEXT: s_load_b512 s[36:51], s[4:5], 0xa4 nv +; GFX13-GISEL-NEXT: s_load_b512 s[52:67], s[4:5], 0xe4 nv +; GFX13-GISEL-NEXT: s_load_b64 s[4:5], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v0, s36 :: v_dual_mov_b32 v1, s37 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v2, s38 :: v_dual_mov_b32 v3, s39 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v4, s40 :: v_dual_mov_b32 v5, s41 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, s42 :: v_dual_mov_b32 v7, s43 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v8, s44 :: v_dual_mov_b32 v9, s45 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v10, s46 :: v_dual_mov_b32 v11, s47 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v12, s48 :: v_dual_mov_b32 v13, s49 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v14, s50 :: v_dual_mov_b32 v15, s51 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v16, s52 :: v_dual_mov_b32 v17, s53 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v18, s54 :: v_dual_mov_b32 v19, s55 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v20, s56 :: v_dual_mov_b32 v21, s57 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v22, s58 :: v_dual_mov_b32 v23, s59 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v24, s60 :: v_dual_mov_b32 v25, s61 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v26, s62 :: v_dual_mov_b32 v27, s63 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v28, s64 :: v_dual_mov_b32 v29, s65 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v30, s66 :: v_dual_mov_b32 v31, s67 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX13-GISEL-NEXT: v_cvt_scalef32_pk32_fp6_f32 v[0:5], v[0:31], 0x42c80000 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s0, v0 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s1, v1 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s2, v2 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s3, v3 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s6, v4 +; GFX13-GISEL-NEXT: v_readfirstlane_b32 s7, v5 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v6, 0 :: v_dual_mov_b32 v0, s0 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v1, s1 :: v_dual_mov_b32 v2, s2 +; GFX13-GISEL-NEXT: v_dual_mov_b32 v3, s3 :: v_dual_mov_b32 v4, s6 +; GFX13-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_4) +; GFX13-GISEL-NEXT: v_mov_b32_e32 v5, s7 +; GFX13-GISEL-NEXT: s_clause 0x1 +; GFX13-GISEL-NEXT: global_store_b128 v6, v[0:3], s[4:5] +; GFX13-GISEL-NEXT: global_store_b64 v6, v[4:5], s[4:5] offset:16 +; GFX13-GISEL-NEXT: s_endpgm + %cvt = tail call <6 x i32> @llvm.amdgcn.cvt.scalef32.pk32.fp6.f32(<32 x float> %src, float 100.0) + store <6 x i32> %cvt, ptr addrspace(1) %out, align 8 + ret void +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
