https://github.com/mariusz-sikora-at-amd created https://github.com/llvm/llvm-project/pull/226971
None >From 933d454e4248531d8b55d10a7da9743f10c572fb Mon Sep 17 00:00:00 2001 From: Mariusz Sikora <[email protected]> Date: Fri, 11 Sep 2026 07:21:50 -0400 Subject: [PATCH] [AMDGPU] Add intrinsics and builtins for v_wave_match_b32 --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 4 + .../include/clang/Basic/BuiltinsAMDGPUDocs.td | 21 +++++ .../builtins-amdgcn-gfx12-err.cl | 2 + .../CodeGenOpenCL/builtins-amdgcn-gfx13.cl | 19 ++++ .../builtins-amdgcn-wave-match-err.cl | 9 ++ llvm/docs/AMDGPUUsage.rst | 5 + llvm/docs/ReleaseNotes.md | 4 + llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 ++ llvm/lib/Target/AMDGPU/AMDGPU.td | 4 +- .../AMDGPU/AMDGPUInstCombineIntrinsic.cpp | 12 +++ .../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 4 + .../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 1 + llvm/lib/Target/AMDGPU/VOP3Instructions.td | 2 +- .../AMDGPU/llvm.amdgcn.wave.match.b32.ll | 91 +++++++++++++++++++ .../InstCombine/AMDGPU/amdgcn-intrinsics.ll | 90 ++++++++++++++++-- 15 files changed, 262 insertions(+), 12 deletions(-) create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-wave-match-err.cl create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wave.match.b32.ll diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 4fd604390d3ce..13379d7d6a11e 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -928,6 +928,10 @@ def __builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64 : AMDGPUBuiltin<"_ExtVector def __builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64 : AMDGPUBuiltin<"_ExtVector<4, float>(int, _ExtVector<2, int>, _ExtVector<4, float>, int)", [Const], "swmmac-gfx1200-insts,wavefrontsize64">; def __builtin_amdgcn_prng_b32 : AMDGPUBuiltin<"unsigned int(unsigned int)", [Const], "prng-inst">; +def __builtin_amdgcn_wave_match_b32 : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const], "wave-match-insts"> { + let Documentation = [DocWaveMatchB32]; + let ArgNames = ["src0", "src1"]; +} def __builtin_amdgcn_cvt_scalef32_pk32_fp6_f16 : AMDGPUBuiltin<"_ExtVector<6, unsigned int>(_ExtVector<32, _Float16>, float)", [Const], "f16bf16-to-fp6bf6-cvt-scale-insts">; def __builtin_amdgcn_cvt_scalef32_pk32_bf6_f16 : AMDGPUBuiltin<"_ExtVector<6, unsigned int>(_ExtVector<32, _Float16>, float)", [Const], "f16bf16-to-fp6bf6-cvt-scale-insts">; def __builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16 : AMDGPUBuiltin<"_ExtVector<6, unsigned int>(_ExtVector<32, __bf16>, float)", [Const], "f16bf16-to-fp6bf6-cvt-scale-insts">; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td index bfccd7de6590a..6a73a979e6e79 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td @@ -74,6 +74,13 @@ between standard floating-point types, integer types and packed low-precision fo }]; } +def DocCatAMDGPUWaveIntrinsics : DocumentationCategory<"Wave Intrinsic Builtins"> { + let Content = [{ +These builtins expose wave-level intrinsics that operate across active +lanes in a wavefront. +}]; +} + def DocCatAMDGPUPrefetch : DocumentationCategory<"Prefetch Builtins"> { let Content = [{ These builtins provide access to AMDGPU prefetch instructions, including instructions @@ -877,6 +884,20 @@ integers. }]; } +//===----------------------------------------------------------------------===// +// Wave Intrinsic Builtins +//===----------------------------------------------------------------------===// + +def DocWaveMatchB32 : Documentation { + let Category = DocCatAMDGPUWaveIntrinsics; + let Content = [{ +Returns a 32-bit bitmask whose bit N is set iff lane N is active and its +``src0`` equals the current lane's ``src1``. Passing the same value as both +operands yields the mask of active lanes sharing that value. In wave64 mode +each 32-lane half is handled independently. +}]; +} + //===----------------------------------------------------------------------===// // Prefetch Builtins //===----------------------------------------------------------------------===// diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl index d8fb243865f76..90b426bf346a9 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-err.cl @@ -94,4 +94,6 @@ void builtin_test_unsupported(double a_double, float a_float, a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8' needs target feature fp8-insts}} a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8' needs target feature fp8-insts}} a_v16f = __builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8(a_v2i, a_v4i, a_v16f, a_int, 0, 0); // expected-error {{'__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8' needs target feature fp8-insts}} + + a = __builtin_amdgcn_wave_match_b32(a, b); // expected-error {{'__builtin_amdgcn_wave_match_b32' needs target feature wave-match-insts}} } diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl index 00e65563a1868..8b4fb588c28ae 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl @@ -211,3 +211,22 @@ void test_s_prefetch_data(global float *gp, unsigned int len) __builtin_amdgcn_s_prefetch_data(gp, len); } +// CHECK-LABEL: @test_wave_match_b32( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[SRC0_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[SRC1_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store i32 [[SRC0:%.*]], ptr addrspace(5) [[SRC0_ADDR]], align 4 +// CHECK-NEXT: store i32 [[SRC1:%.*]], ptr addrspace(5) [[SRC1_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[SRC0_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(5) [[SRC1_ADDR]], align 4 +// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.wave.match.b32(i32 [[TMP0]], i32 [[TMP1]]) +// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8 +// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4 +// CHECK-NEXT: ret void +// +void test_wave_match_b32(global unsigned int *out, unsigned int src0, unsigned int src1) +{ + *out = __builtin_amdgcn_wave_match_b32(src0, src1); +} diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-wave-match-err.cl b/clang/test/SemaOpenCL/builtins-amdgcn-wave-match-err.cl new file mode 100644 index 0000000000000..44a2ad0d67640 --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-wave-match-err.cl @@ -0,0 +1,9 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu13.10-unknown-unknown -verify -S -o - %s + +typedef unsigned int uint; + +void test_wave_match_b32(global uint* out, uint src0, uint src1) { + *out = __builtin_amdgcn_wave_match_b32(src0); // expected-error {{too few arguments to function call, expected 2, have 1}} + *out = __builtin_amdgcn_wave_match_b32(src0, src1, src0); // expected-error {{too many arguments to function call, expected 2, have 3}} +} diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst index 9254114e5044f..37dc0affd05b0 100644 --- a/llvm/docs/AMDGPUUsage.rst +++ b/llvm/docs/AMDGPUUsage.rst @@ -1926,6 +1926,11 @@ The AMDGPU backend implements the following LLVM IR intrinsics. bfloat, <2 x i16>, <2 x half>, <2 x bfloat>, i64, double, pointers, multiples of the 32-bit vectors. + llvm.amdgcn.wave.match.b32 Provides direct access to v_wave_match_b32. Returns a 32-bit mask whose bit N is + set when lane N is active and its first operand equals the current lane's second + operand. Passing the same value as both operands yields the mask of active lanes + sharing that value. In wave64 mode each 32-lane half is handled independently. + llvm.amdgcn.udot2 Provides direct access to v_dot2_u32_u16 across targets which support such instructions. This performs an unsigned dot product with two v2i16 operands, summed with the third i32 operand. The diff --git a/llvm/docs/ReleaseNotes.md b/llvm/docs/ReleaseNotes.md index 41b2fa83d380f..87848a64c13c9 100644 --- a/llvm/docs/ReleaseNotes.md +++ b/llvm/docs/ReleaseNotes.md @@ -247,6 +247,10 @@ Makes programs 10x faster by doing Special New Thing. * `llvm.amdgcn.icmp` * `llvm.amdgcn.fcmp` +* Added the `llvm.amdgcn.wave.match.b32` intrinsic (and matching + `__builtin_amdgcn_wave_match_b32` clang builtin) providing direct access to the + `v_wave_match_b32` instruction. + ### Changes to the ARM Backend * Using the hard-float procedure call standard without floating-point registers diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index e7829eb29ba34..8f65f96d68671 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -908,6 +908,12 @@ def int_amdgcn_prng_b32 : DefaultAttrsIntrinsic< [llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrNoCreateUndefOrPoison] >, ClangBuiltin<"__builtin_amdgcn_prng_b32">; +let TargetFeatures = "wave-match-insts" in +def int_amdgcn_wave_match_b32 : DefaultAttrsIntrinsic< + [llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], + [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCreateUndefOrPoison] +>, ClangBuiltin<"__builtin_amdgcn_wave_match_b32">; + def int_amdgcn_bitop3 : PureIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i32_ty], diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td index 5fbf30d55947c..bf3d8cfedd965 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.td +++ b/llvm/lib/Target/AMDGPU/AMDGPU.td @@ -3355,8 +3355,8 @@ def AMDGPUFrontendVisibleFeatures { FeatureSadInsts, FeatureSetPrioIncWgInst, FeatureSmemPrefetchInsts, FeatureTDMInsts, FeatureTanhInsts, FeatureTensorCvtLutInsts, FeatureTransposeLoadF4F6Insts, FeatureVMemToLDSLoad, FeatureVmemPrefInsts, - FeatureWMMA128bInsts, FeatureWMMA256bInsts, FeatureWMMAN16Insts, - FeatureWMMAF4Insts, FeatureXF32Insts, + FeatureWaveMatchInsts, FeatureWMMA128bInsts, FeatureWMMA256bInsts, + FeatureWMMAN16Insts, FeatureWMMAF4Insts, FeatureXF32Insts, FeatureWavefrontSize32, FeatureWavefrontSize64, FeatureSupportsWGP, FeatureSupportsWave32, FeatureFastFMAF32, FeatureFastDenormalF32, FeatureSupportsXNACK, FeatureSupportsSRAMECC, FeatureXNACKOnOffModes, diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp index 7721114fdd2f7..18a0505269002 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp @@ -1915,6 +1915,18 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const { return std::nullopt; } + case Intrinsic::amdgcn_wave_match_b32: { + const Use &Src0 = II.getArgOperandUse(0); + const Use &Src1 = II.getArgOperandUse(1); + if (Src0.get() == Src1.get() && isTriviallyUniform(Src0)) { + Function *NewF = Intrinsic::getOrInsertDeclaration( + II.getModule(), Intrinsic::amdgcn_ballot, II.getType()); + CallInst *NewCall = + IC.Builder.CreateCall(NewF, {IC.Builder.getInt1(true)}); + return IC.replaceInstUsesWith(II, NewCall); + } + break; + } case Intrinsic::amdgcn_writelane: { // TODO: Fold bitcast like readlane. if (simplifyDemandedLaneMaskArg(IC, II, 1)) diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp index 1d92bca3e12ab..ba1892ace6ad9 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp @@ -2094,6 +2094,10 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST, .Any({{UniS32}, {{UniInVgprS32}, {IntrId, Vgpr32}}}) .Any({{DivS32}, {{Vgpr32}, {IntrId, Vgpr32}}}); + addRulesForIOpcs({amdgcn_wave_match_b32}) + .Any({{UniS32}, {{UniInVgprS32}, {IntrId, Vgpr32, Vgpr32}}}) + .Any({{DivS32}, {{Vgpr32}, {IntrId, Vgpr32, Vgpr32}}}); + addRulesForIOpcs({amdgcn_sffbh}, Standard) .Uni(S32, {{Sgpr32}, {IntrId, Sgpr32}}) .Div(S32, {{Vgpr32}, {IntrId, Vgpr32}}); diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp index 91bd0d006cc64..6515a4fc974c2 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -4727,6 +4727,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { case Intrinsic::amdgcn_alignbyte: case Intrinsic::amdgcn_perm: case Intrinsic::amdgcn_prng_b32: + case Intrinsic::amdgcn_wave_match_b32: case Intrinsic::amdgcn_fdot2: case Intrinsic::amdgcn_sdot2: case Intrinsic::amdgcn_udot2: diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td b/llvm/lib/Target/AMDGPU/VOP3Instructions.td index 5e78ba730c20e..02f98cfc7ce3c 100644 --- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td @@ -2120,7 +2120,7 @@ def VOP_WAVE_MATCH_B32 : VOP3_Profile<VOP_I32_I32_I32> { } let isConvergent = 1, SubtargetPredicate = HasWaveMatchInsts in { - defm V_WAVE_MATCH_B32 : VOP3Inst<"v_wave_match_b32", VOP_WAVE_MATCH_B32>; + defm V_WAVE_MATCH_B32 : VOP3Inst<"v_wave_match_b32", VOP_WAVE_MATCH_B32, int_amdgcn_wave_match_b32>; } class VOP_EXCLUSIVE_SCAN<VOPProfile p, VOP3Features f = VOP3_REGULAR> : VOP3_Profile<p, f> { diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wave.match.b32.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wave.match.b32.ll new file mode 100644 index 0000000000000..490a21abcb6bb --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wave.match.b32.ll @@ -0,0 +1,91 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc -global-isel=0 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefixes=GFX13,GFX13-SDAG %s +; RUN: llc -global-isel=1 -mtriple=amdgpu13.10 < %s | FileCheck -check-prefixes=GFX13,GFX13-GISEL %s + +define amdgpu_kernel void @v_wave_match_sgpr(ptr addrspace(1) %out, i32 %src1, i32 %src2) { +; GFX13-SDAG-LABEL: v_wave_match_sgpr: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_wave_match_b32 v1, s2, s3 +; GFX13-SDAG-NEXT: global_store_b32 v0, v1, s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: v_wave_match_sgpr: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: v_mov_b32_e32 v1, 0 +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_wave_match_b32 v0, s2, s3 +; GFX13-GISEL-NEXT: global_store_b32 v1, v0, s[0:1] +; GFX13-GISEL-NEXT: s_endpgm + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 %src1, i32 %src2) + store i32 %v, ptr addrspace(1) %out + ret void +} + +define amdgpu_kernel void @v_wave_match_vgpr(ptr addrspace(1) %out) { +; GFX13-LABEL: v_wave_match_vgpr: +; GFX13: ; %bb.0: +; GFX13-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-NEXT: v_and_b32_e32 v1, 0x3ff, v0 +; GFX13-NEXT: v_bfe_u32 v0, v0, 10, 10 +; GFX13-NEXT: s_delay_alu instid0(VALU_DEP_1) +; GFX13-NEXT: v_wave_match_b32 v0, v1, v0 +; GFX13-NEXT: s_wait_kmcnt 0x0 +; GFX13-NEXT: global_store_b32 v1, v0, s[0:1] scale_offset +; GFX13-NEXT: s_endpgm + %tidx = call i32 @llvm.amdgcn.workitem.id.x() + %tidy = call i32 @llvm.amdgcn.workitem.id.y() + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 %tidx, i32 %tidy) + %out_ptr = getelementptr i32, ptr addrspace(1) %out, i32 %tidx + store i32 %v, ptr addrspace(1) %out_ptr + ret void +} + +define amdgpu_kernel void @v_wave_match_constant(ptr addrspace(1) %out) { +; GFX13-SDAG-LABEL: v_wave_match_constant: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX13-SDAG-NEXT: v_wave_match_b32 v1, 7, 5 +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: global_store_b32 v0, v1, s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: v_wave_match_constant: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: v_wave_match_b32 v0, 7, 5 +; GFX13-GISEL-NEXT: v_mov_b32_e32 v1, 0 +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: global_store_b32 v1, v0, s[0:1] +; GFX13-GISEL-NEXT: s_endpgm + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 7, i32 5) + store i32 %v, ptr addrspace(1) %out + ret void +} + +define amdgpu_kernel void @v_wave_match_poison(ptr addrspace(1) %out) { +; GFX13-SDAG-LABEL: v_wave_match_poison: +; GFX13-SDAG: ; %bb.0: +; GFX13-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-SDAG-NEXT: v_mov_b32_e32 v0, 0 +; GFX13-SDAG-NEXT: s_wait_kmcnt 0x0 +; GFX13-SDAG-NEXT: v_wave_match_b32 v1, s0, s0 +; GFX13-SDAG-NEXT: global_store_b32 v0, v1, s[0:1] +; GFX13-SDAG-NEXT: s_endpgm +; +; GFX13-GISEL-LABEL: v_wave_match_poison: +; GFX13-GISEL: ; %bb.0: +; GFX13-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv +; GFX13-GISEL-NEXT: v_mov_b32_e32 v1, 0 +; GFX13-GISEL-NEXT: s_wait_kmcnt 0x0 +; GFX13-GISEL-NEXT: v_wave_match_b32 v0, s0, s0 +; GFX13-GISEL-NEXT: global_store_b32 v1, v0, s[0:1] +; GFX13-GISEL-NEXT: s_endpgm + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 poison, i32 poison) + store i32 %v, ptr addrspace(1) %out + ret void +} diff --git a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll index 9cd0ee45bdfc1..8850eac213d6b 100644 --- a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll +++ b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll @@ -43,7 +43,7 @@ define double @test_constant_fold_rcp_f64_1() nounwind { define float @test_constant_fold_rcp_f32_half() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f32_half( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 5.000000e-01) #[[ATTR12:[0-9]+]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 5.000000e-01) #[[ATTR14:[0-9]+]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.rcp.f32(float 0.5) nounwind readnone @@ -52,7 +52,7 @@ define float @test_constant_fold_rcp_f32_half() nounwind { define double @test_constant_fold_rcp_f64_half() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f64_half( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 5.000000e-01) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 5.000000e-01) #[[ATTR14]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.rcp.f64(double 0.5) nounwind readnone @@ -61,7 +61,7 @@ define double @test_constant_fold_rcp_f64_half() nounwind { define float @test_constant_fold_rcp_f32_43() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f32_43( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR14]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) nounwind readnone @@ -70,7 +70,7 @@ define float @test_constant_fold_rcp_f32_43() nounwind { define double @test_constant_fold_rcp_f64_43() nounwind { ; CHECK-LABEL: @test_constant_fold_rcp_f64_43( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #[[ATTR14]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) nounwind readnone @@ -79,7 +79,7 @@ define double @test_constant_fold_rcp_f64_43() nounwind { define float @test_constant_fold_rcp_f32_43_strictfp() nounwind strictfp { ; CHECK-LABEL: @test_constant_fold_rcp_f32_43_strictfp( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR13:[0-9]+]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR15:[0-9]+]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) strictfp nounwind readnone @@ -190,7 +190,7 @@ define half @test_constant_fold_sqrt_f16_0() nounwind { define float @test_constant_fold_sqrt_f32_0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f32_0( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR14]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.sqrt.f32(float 0.0) nounwind readnone @@ -199,7 +199,7 @@ define float @test_constant_fold_sqrt_f32_0() nounwind { define double @test_constant_fold_sqrt_f64_0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f64_0( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR14]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.sqrt.f64(double 0.0) nounwind readnone @@ -216,7 +216,7 @@ define half @test_constant_fold_sqrt_f16_neg0() nounwind { define float @test_constant_fold_sqrt_f32_neg0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f32_neg0( -; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR14]] ; CHECK-NEXT: ret float [[VAL]] ; %val = call float @llvm.amdgcn.sqrt.f32(float -0.0) nounwind readnone @@ -225,7 +225,7 @@ define float @test_constant_fold_sqrt_f32_neg0() nounwind { define double @test_constant_fold_sqrt_f64_neg0() nounwind { ; CHECK-LABEL: @test_constant_fold_sqrt_f64_neg0( -; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -0.000000e+00) #[[ATTR12]] +; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -0.000000e+00) #[[ATTR14]] ; CHECK-NEXT: ret double [[VAL]] ; %val = call double @llvm.amdgcn.sqrt.f64(double -0.0) nounwind readnone @@ -5946,6 +5946,78 @@ define void @ds_bpermute_uniform_lane(ptr addrspace(1) %out, i32 %lanearg, i32 % ret void } +; -------------------------------------------------------------------- +; llvm.amdgcn.wave.match.b32 +; -------------------------------------------------------------------- + +; Two distinct operands do not fold: the match is a cross-comparison, not the +; mask of all active lanes. +define amdgpu_kernel void @v_wave_match_sgpr(ptr addrspace(1) %out, i32 %src1, i32 %src2) { +; CHECK-LABEL: @v_wave_match_sgpr( +; CHECK-NEXT: [[V:%.*]] = call i32 @llvm.amdgcn.wave.match.b32(i32 [[SRC1:%.*]], i32 [[SRC2:%.*]]) +; CHECK-NEXT: store i32 [[V]], ptr addrspace(1) [[OUT:%.*]], align 4 +; CHECK-NEXT: ret void +; + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 %src1, i32 %src2) + store i32 %v, ptr addrspace(1) %out + ret void +} + +; A divergent value does not fold even when passed as both operands. +define amdgpu_kernel void @v_wave_match_vgpr(ptr addrspace(1) %out) { +; CHECK-LABEL: @v_wave_match_vgpr( +; CHECK-NEXT: [[TIDX:%.*]] = call i32 @llvm.amdgcn.workitem.id.x() +; CHECK-NEXT: [[V:%.*]] = call i32 @llvm.amdgcn.wave.match.b32(i32 [[TIDX]], i32 [[TIDX]]) +; CHECK-NEXT: [[TMP1:%.*]] = zext nneg i32 [[TIDX]] to i64 +; CHECK-NEXT: [[OUT_PTR:%.*]] = getelementptr [4 x i8], ptr addrspace(1) [[OUT:%.*]], i64 [[TMP1]] +; CHECK-NEXT: store i32 [[V]], ptr addrspace(1) [[OUT_PTR]], align 4 +; CHECK-NEXT: ret void +; + %tidx = call i32 @llvm.amdgcn.workitem.id.x() + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 %tidx, i32 %tidx) + %out_ptr = getelementptr i32, ptr addrspace(1) %out, i32 %tidx + store i32 %v, ptr addrspace(1) %out_ptr + ret void +} + +; The same uniform value (an SGPR kernel argument, uniform by ABI) in both +; operands folds to ballot(true): every active lane holds that value, so it +; matches every active lane. +define amdgpu_kernel void @v_wave_match_same_sgpr(ptr addrspace(1) %out, i32 %src) { +; CHECK-LABEL: @v_wave_match_same_sgpr( +; CHECK-NEXT: [[V:%.*]] = call i32 @llvm.amdgcn.ballot.i32(i1 true) +; CHECK-NEXT: store i32 [[V]], ptr addrspace(1) [[OUT:%.*]], align 4 +; CHECK-NEXT: ret void +; + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 %src, i32 %src) + store i32 %v, ptr addrspace(1) %out + ret void +} + +; The same constant in both operands folds to ballot(true). +define amdgpu_kernel void @v_wave_match_constant(ptr addrspace(1) %out) { +; CHECK-LABEL: @v_wave_match_constant( +; CHECK-NEXT: [[V:%.*]] = call i32 @llvm.amdgcn.ballot.i32(i1 true) +; CHECK-NEXT: store i32 [[V]], ptr addrspace(1) [[OUT:%.*]], align 4 +; CHECK-NEXT: ret void +; + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 7, i32 7) + store i32 %v, ptr addrspace(1) %out + ret void +} + +; Two distinct constants do not fold. +define amdgpu_kernel void @v_wave_match_two_constants(ptr addrspace(1) %out) { +; CHECK-LABEL: @v_wave_match_two_constants( +; CHECK-NEXT: [[V:%.*]] = call i32 @llvm.amdgcn.wave.match.b32(i32 7, i32 8) +; CHECK-NEXT: store i32 [[V]], ptr addrspace(1) [[OUT:%.*]], align 4 +; CHECK-NEXT: ret void +; + %v = call i32 @llvm.amdgcn.wave.match.b32(i32 7, i32 8) + store i32 %v, ptr addrspace(1) %out + ret void +} + ; -------------------------------------------------------------------- ; llvm.amdgcn.make.buffer.rsrc.p8 ; -------------------------------------------------------------------- _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
