[clang] [llvm] [AMDGPU] Add intrinsics and builtins for v_wave_match_b32 (PR #226971)
Mariusz Sikora via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 28 05:27:15 PDT 2026
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 <mariusz.sikora at amd.com>
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
; --------------------------------------------------------------------
More information about the cfe-commits
mailing list