[llvm-branch-commits] [llvm] 1508332 - Revert "[NVPTX] Fold symbol addresses into memory operands (#202379)"
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Jun 25 09:20:18 PDT 2026
Author: Krzysztof Parzyszek
Date: 2026-06-25T11:20:14-05:00
New Revision: 1508332f203092dda1753eafc206af0d2ebbd0c2
URL: https://github.com/llvm/llvm-project/commit/1508332f203092dda1753eafc206af0d2ebbd0c2
DIFF: https://github.com/llvm/llvm-project/commit/1508332f203092dda1753eafc206af0d2ebbd0c2.diff
LOG: Revert "[NVPTX] Fold symbol addresses into memory operands (#202379)"
This reverts commit 94c1a8bc3e6b790ccc5723803129b3b956f11d13.
Added:
Modified:
llvm/lib/Target/NVPTX/CMakeLists.txt
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
llvm/lib/Target/NVPTX/NVPTX.h
llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
Removed:
llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
llvm/test/CodeGen/NVPTX/address-folder.ll
llvm/test/CodeGen/NVPTX/address-folder.mir
################################################################################
diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt
index d2b5083a5f83f..e0e2b1a40bffc 100644
--- a/llvm/lib/Target/NVPTX/CMakeLists.txt
+++ b/llvm/lib/Target/NVPTX/CMakeLists.txt
@@ -12,7 +12,6 @@ tablegen(LLVM NVPTXGenSubtargetInfo.inc -gen-subtarget)
add_public_tablegen_target(NVPTXCommonTableGen)
set(NVPTXCodeGen_sources
- NVPTXAddressFolder.cpp
NVPTXAliasAnalysis.cpp
NVPTXAllocaHoisting.cpp
NVPTXAsmPrinter.cpp
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
index ef447e4ceed6b..5933f74fcb25f 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
@@ -24,7 +24,6 @@
using namespace llvm;
#define GET_INSTRINFO_MC_DESC
-#define GET_INSTRINFO_NAMED_OPS
#define ENABLE_INSTR_PREDICATE_VERIFIER
#include "NVPTXGenInstrInfo.inc"
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
index 58083e1248bb0..78f4e67455023 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
@@ -13,7 +13,6 @@
#ifndef LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXMCTARGETDESC_H
#define LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXMCTARGETDESC_H
-#include "llvm/Support/Compiler.h"
#include <stdint.h>
// Defines symbolic names for PTX registers.
@@ -23,7 +22,6 @@
// Defines symbolic names for the PTX instructions.
#define GET_INSTRINFO_ENUM
#define GET_INSTRINFO_MC_HELPER_DECLS
-#define GET_INSTRINFO_OPERAND_ENUM
#include "NVPTXGenInstrInfo.inc"
#define GET_SUBTARGETINFO_ENUM
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 57ed9bb43aeea..08ccc143dfb8a 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -59,7 +59,6 @@ FunctionPass *createNVPTXIRPeepholePass();
MachineFunctionPass *createNVPTXPeephole();
MachineFunctionPass *createNVPTXProxyRegErasurePass();
MachineFunctionPass *createNVPTXForwardParamsPass();
-MachineFunctionPass *createNVPTXAddressFolderPass();
void initializeNVVMReflectLegacyPassPass(PassRegistry &);
void initializeGenericToNVVMLegacyPassPass(PassRegistry &);
@@ -75,7 +74,6 @@ void initializeNVPTXLowerArgsLegacyPassPass(PassRegistry &);
void initializeNVPTXSetByValParamAlignLegacyPassPass(PassRegistry &);
void initializeNVPTXProxyRegErasurePass(PassRegistry &);
void initializeNVPTXForwardParamsPassPass(PassRegistry &);
-void initializeNVPTXAddressFolderPassPass(PassRegistry &);
void initializeNVVMIntrRangePass(PassRegistry &);
void initializeNVVMReflectPass(PassRegistry &);
void initializeNVPTXAAWrapperPassPass(PassRegistry &);
@@ -296,7 +294,6 @@ void initializeNVPTXDAGToDAGISelLegacyPass(PassRegistry &);
// Defines symbolic names for the NVPTX instructions.
#define GET_INSTRINFO_ENUM
#define GET_INSTRINFO_MC_HELPER_DECLS
-#define GET_INSTRINFO_OPERAND_ENUM
#include "NVPTXGenInstrInfo.inc"
#endif
diff --git a/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp b/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
deleted file mode 100644
index 5299dbe686611..0000000000000
--- a/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
+++ /dev/null
@@ -1,123 +0,0 @@
-//===- NVPTXAddressFolder.cpp - Fold symbol addresses into memory ops -----===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// SelectionDAG folds a symbol address (a kernel parameter, global variable, or
-// external symbol) directly into the address operand of a memory access, but
-// only within a single basic block. When the address is carried across a block
-// boundary in a register it is materialized with a generic `mov`
-// (MOV_B{32,64}_sym) and the accesses become register-relative:
-//
-// mov.b64 %rd1, kernel_param_0;
-// ld.param.b64 %rd2, [%rd1];
-// ld.param.b64 %rd3, [%rd1+8];
-//
-// This pass folds the symbol back into those address operands, eliminating the
-// redundant address arithmetic (and the `mov` itself once no use remains):
-//
-// ld.param.b64 %rd2, [kernel_param_0];
-// ld.param.b64 %rd3, [kernel_param_0+8];
-//
-// Shared-memory accesses are left alone: `mov`s of shared symbols are
-// deliberately kept CSE-able rather than duplicated into their uses, as
-// rematerializing them has caused performance regressions before (see
-// MovSymInst in NVPTXInstrInfo.td).
-//
-//===----------------------------------------------------------------------===//
-
-#include "NVPTX.h"
-#include "llvm/CodeGen/MachineFunctionPass.h"
-#include "llvm/CodeGen/MachineInstr.h"
-#include "llvm/CodeGen/MachineOperand.h"
-#include "llvm/CodeGen/MachineRegisterInfo.h"
-
-using namespace llvm;
-
-// Try to fold the definition of \p Addr into \p MI's address operand. The
-// defining instruction is erased once it has no uses left; it is kept for any
-// remaining use, e.g. because the address also feeds arithmetic or escapes.
-static bool foldAddress(MachineInstr &MI, MachineOperand &Addr,
- MachineRegisterInfo &MRI) {
- assert(Addr.isReg() && "Expected an address register");
- assert(Addr.getReg().isVirtual() && "Expected a virtual address register");
-
- MachineInstr *Mov = MRI.getVRegDef(Addr.getReg());
- if (!Mov || (Mov->getOpcode() != NVPTX::MOV_B32_sym &&
- Mov->getOpcode() != NVPTX::MOV_B64_sym))
- return false;
-
- const MachineOperand &Sym = Mov->getOperand(1);
- if (!Sym.isGlobal() && !Sym.isSymbol())
- return false;
-
- // The accessed address space must be known and must not be shared.
- const int AddrSpaceIdx =
- NVPTX::getNamedOperandIdx(MI.getOpcode(), NVPTX::OpName::addsp);
- if (AddrSpaceIdx < 0)
- return false;
- const auto AddrSpace = MI.getOperand(AddrSpaceIdx).getImm();
- if (AddrSpace == NVPTX::AddressSpace::Shared ||
- AddrSpace == NVPTX::AddressSpace::SharedCluster)
- return false;
-
- if (Sym.isGlobal()) {
- Addr.ChangeToGA(Sym.getGlobal(), Sym.getOffset(), Sym.getTargetFlags());
- } else {
- Addr.ChangeToES(Sym.getSymbolName(), Sym.getTargetFlags());
- }
-
- if (MRI.use_empty(Mov->getOperand(0).getReg()))
- Mov->eraseFromParent();
-
- return true;
-}
-
-static bool foldAddresses(MachineFunction &MF) {
- MachineRegisterInfo &MRI = MF.getRegInfo();
-
- bool Changed = false;
- for (MachineBasicBlock &MBB : MF)
- for (MachineInstr &MI : make_early_inc_range(MBB))
- if (MI.mayLoadOrStore()) {
- const int AddrIdx =
- NVPTX::getNamedOperandIdx(MI.getOpcode(), NVPTX::OpName::addr);
- if (AddrIdx >= 0 && MI.getOperand(AddrIdx).isReg())
- Changed |= foldAddress(MI, MI.getOperand(AddrIdx), MRI);
- }
-
- return Changed;
-}
-
-/// ----------------------------------------------------------------------------
-/// Pass (Manager) Boilerplate
-/// ----------------------------------------------------------------------------
-
-namespace {
-struct NVPTXAddressFolderPass : public MachineFunctionPass {
- static char ID;
- NVPTXAddressFolderPass() : MachineFunctionPass(ID) {}
-
- bool runOnMachineFunction(MachineFunction &MF) override;
-
- void getAnalysisUsage(AnalysisUsage &AU) const override {
- MachineFunctionPass::getAnalysisUsage(AU);
- }
-};
-} // namespace
-
-char NVPTXAddressFolderPass::ID = 0;
-
-INITIALIZE_PASS(NVPTXAddressFolderPass, "nvptx-address-folder",
- "NVPTX Address Folder", false, false)
-
-bool NVPTXAddressFolderPass::runOnMachineFunction(MachineFunction &MF) {
- return foldAddresses(MF);
-}
-
-MachineFunctionPass *llvm::createNVPTXAddressFolderPass() {
- return new NVPTXAddressFolderPass();
-}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td b/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
index 361f819a340ae..719be0300940e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
@@ -22,7 +22,6 @@ class NVPTXInst<dag outs, dag ins, string asmstr, list<dag> pattern = []>
dag InOperandList = ins;
let AsmString = asmstr;
let Pattern = pattern;
- let UseNamedOperandTable = true;
// TSFlagFields
bit IsLoad = false;
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index c9db5c5cec025..5f1c16dc7f0fb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2454,8 +2454,8 @@ let isTerminator=1 in {
let isBranch=1 in {
def CBranch : NVPTXInst<(outs),
- (ins B1:$a, brtarget:$target, BranchFlag:$negate),
- "@${negate}$a bra \t$target;",
+ (ins B1:$a, brtarget:$target, BranchFlag:$not),
+ "@${not}$a bra \t$target;",
[(brcond i1:$a, bb:$target)]>;
let isBarrier=1 in
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index 2bc2158fccae1..6054feeedf932 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -119,7 +119,6 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTarget() {
initializeNVPTXLowerAggrCopiesPass(PR);
initializeNVPTXProxyRegErasurePass(PR);
initializeNVPTXForwardParamsPassPass(PR);
- initializeNVPTXAddressFolderPassPass(PR);
initializeNVPTXDAGToDAGISelLegacyPass(PR);
initializeNVPTXAAWrapperPassPass(PR);
initializeNVPTXExternalAAWrapperPass(PR);
@@ -422,8 +421,6 @@ bool NVPTXPassConfig::addInstSelector() {
void NVPTXPassConfig::addPreRegAlloc() {
addPass(createNVPTXForwardParamsPass());
- if (getOptLevel() != CodeGenOptLevel::None)
- addPass(createNVPTXAddressFolderPass());
// Remove Proxy Register pseudo instructions used to keep `callseq_end` alive.
addPass(createNVPTXProxyRegErasurePass());
}
diff --git a/llvm/test/CodeGen/NVPTX/address-folder.ll b/llvm/test/CodeGen/NVPTX/address-folder.ll
deleted file mode 100644
index 239f21594265d..0000000000000
--- a/llvm/test/CodeGen/NVPTX/address-folder.ll
+++ /dev/null
@@ -1,113 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
-; RUN: llc < %s | FileCheck %s
-; RUN: llc -O0 < %s | FileCheck %s --check-prefix=O0
-
-target triple = "nvptx64-nvidia-cuda"
-
-; A byval kernel parameter loaded at a constant offset from a block other than
-; the one that defines its address must still be read directly from parameter
-; space (ld.param [param+off]), not via a materialized base register.
-define ptx_kernel void @test_kernel_cross_block(ptr byval([2 x i64]) align 8 %r, i64 %n, ptr addrspace(1) %out) {
-; CHECK-LABEL: test_kernel_cross_block(
-; CHECK: {
-; CHECK-NEXT: .reg .pred %p<2>;
-; CHECK-NEXT: .reg .b64 %rd<6>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
-; CHECK-NEXT: setp.lt.s64 %p1, %rd2, 1;
-; CHECK-NEXT: @%p1 bra $L__BB0_2;
-; CHECK-NEXT: // %bb.1: // %body
-; CHECK-NEXT: ld.param.b64 %rd1, [test_kernel_cross_block_param_2];
-; CHECK-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_0];
-; CHECK-NEXT: ld.param.b64 %rd4, [test_kernel_cross_block_param_0+8];
-; CHECK-NEXT: mul.lo.s64 %rd5, %rd3, %rd4;
-; CHECK-NEXT: st.global.b64 [%rd1], %rd5;
-; CHECK-NEXT: $L__BB0_2: // %done
-; CHECK-NEXT: ret;
-;
-; O0-LABEL: test_kernel_cross_block(
-; O0: {
-; O0-NEXT: .reg .pred %p<2>;
-; O0-NEXT: .reg .b64 %rd<8>;
-; O0-EMPTY:
-; O0-NEXT: // %bb.0:
-; O0-NEXT: ld.param.b64 %rd4, [test_kernel_cross_block_param_2];
-; O0-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_1];
-; O0-NEXT: mov.b64 %rd1, test_kernel_cross_block_param_0;
-; O0-NEXT: mov.b64 %rd2, %rd1;
-; O0-NEXT: setp.lt.s64 %p1, %rd3, 1;
-; O0-NEXT: @%p1 bra $L__BB0_2;
-; O0-NEXT: bra.uni $L__BB0_1;
-; O0-NEXT: $L__BB0_1: // %body
-; O0-NEXT: ld.param.b64 %rd5, [%rd1];
-; O0-NEXT: ld.param.b64 %rd6, [%rd1+8];
-; O0-NEXT: mul.lo.s64 %rd7, %rd5, %rd6;
-; O0-NEXT: st.global.b64 [%rd4], %rd7;
-; O0-NEXT: bra.uni $L__BB0_2;
-; O0-NEXT: $L__BB0_2: // %done
-; O0-NEXT: ret;
- %c = icmp sgt i64 %n, 0
- br i1 %c, label %body, label %done
-body:
- %v0 = load i64, ptr %r, align 8
- %p1 = getelementptr inbounds i8, ptr %r, i64 8
- %v1 = load i64, ptr %p1, align 8
- %m = mul i64 %v0, %v1
- store i64 %m, ptr addrspace(1) %out, align 8
- br label %done
-done:
- ret void
-}
-
-; When the address is also used for arithmetic (here a variable offset) it is
-; not just loaded, so the `mov` must be kept rather than folded into the load.
-define ptx_kernel void @test_kernel_var_offset(ptr byval([8 x i64]) align 8 %r, i64 %n, ptr addrspace(1) %out) {
-; CHECK-LABEL: test_kernel_var_offset(
-; CHECK: {
-; CHECK-NEXT: .reg .pred %p<2>;
-; CHECK-NEXT: .reg .b64 %rd<6>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd4, [test_kernel_var_offset_param_1];
-; CHECK-NEXT: setp.lt.s64 %p1, %rd4, 1;
-; CHECK-NEXT: @%p1 bra $L__BB1_2;
-; CHECK-NEXT: // %bb.1: // %body
-; CHECK-NEXT: ld.param.b64 %rd2, [test_kernel_var_offset_param_2];
-; CHECK-NEXT: mov.b64 %rd3, test_kernel_var_offset_param_0;
-; CHECK-NEXT: add.s64 %rd1, %rd3, %rd4;
-; CHECK-NEXT: ld.param.b64 %rd5, [%rd1];
-; CHECK-NEXT: st.global.b64 [%rd2], %rd5;
-; CHECK-NEXT: $L__BB1_2: // %done
-; CHECK-NEXT: ret;
-;
-; O0-LABEL: test_kernel_var_offset(
-; O0: {
-; O0-NEXT: .reg .pred %p<2>;
-; O0-NEXT: .reg .b64 %rd<7>;
-; O0-EMPTY:
-; O0-NEXT: // %bb.0:
-; O0-NEXT: ld.param.b64 %rd4, [test_kernel_var_offset_param_2];
-; O0-NEXT: ld.param.b64 %rd3, [test_kernel_var_offset_param_1];
-; O0-NEXT: mov.b64 %rd1, test_kernel_var_offset_param_0;
-; O0-NEXT: mov.b64 %rd2, %rd1;
-; O0-NEXT: setp.lt.s64 %p1, %rd3, 1;
-; O0-NEXT: @%p1 bra $L__BB1_2;
-; O0-NEXT: bra.uni $L__BB1_1;
-; O0-NEXT: $L__BB1_1: // %body
-; O0-NEXT: add.s64 %rd5, %rd1, %rd3;
-; O0-NEXT: ld.param.b64 %rd6, [%rd5];
-; O0-NEXT: st.global.b64 [%rd4], %rd6;
-; O0-NEXT: bra.uni $L__BB1_2;
-; O0-NEXT: $L__BB1_2: // %done
-; O0-NEXT: ret;
- %c = icmp sgt i64 %n, 0
- br i1 %c, label %body, label %done
-body:
- %p = getelementptr inbounds i8, ptr %r, i64 %n
- %v = load i64, ptr %p, align 8
- store i64 %v, ptr addrspace(1) %out, align 8
- br label %done
-done:
- ret void
-}
diff --git a/llvm/test/CodeGen/NVPTX/address-folder.mir b/llvm/test/CodeGen/NVPTX/address-folder.mir
deleted file mode 100644
index ac50925579506..0000000000000
--- a/llvm/test/CodeGen/NVPTX/address-folder.mir
+++ /dev/null
@@ -1,60 +0,0 @@
-# RUN: llc -mtriple=nvptx64-nvidia-cuda -run-pass=nvptx-address-folder -verify-machineinstrs -o - %s | FileCheck %s
-
---- |
- target triple = "nvptx64-nvidia-cuda"
-
- @g = addrspace(1) global [4 x i64] zeroinitializer
- @g_shared = addrspace(3) global [4 x i64] zeroinitializer
-
- define void @fold_global() { ret void }
- define void @skip_shared() { ret void }
- define void @fold_partial() { ret void }
-...
----
-# A global address materialized into a register is folded into the address
-# operand of loads and stores, and the dead `mov` is removed.
-# CHECK-LABEL: name: fold_global
-# CHECK-NOT: MOV_B64_sym
-# CHECK: %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, @g, 0
-# CHECK: %2:b64 = LD_i64 0, 0, 1, 3, 64, -1, @g, 8
-# CHECK: ST_i64 %1, 0, 0, 1, 64, @g, 16
-name: fold_global
-tracksRegLiveness: true
-body: |
- bb.0:
- %0:b64 = MOV_B64_sym @g
- %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, %0, 0 :: (load (s64), addrspace 1)
- %2:b64 = LD_i64 0, 0, 1, 3, 64, -1, %0, 8 :: (load (s64), addrspace 1)
- ST_i64 %1, 0, 0, 1, 64, %0, 16 :: (store (s64), addrspace 1)
- Return
-...
----
-# Shared-memory accesses are not folded: the `mov` of the shared symbol must
-# stay CSE-able rather than be duplicated into every access.
-# CHECK-LABEL: name: skip_shared
-# CHECK: %0:b64 = MOV_B64_sym @g_shared
-# CHECK: %1:b64 = LD_i64 0, 0, 3, 3, 64, -1, %0, 0
-name: skip_shared
-tracksRegLiveness: true
-body: |
- bb.0:
- %0:b64 = MOV_B64_sym @g_shared
- %1:b64 = LD_i64 0, 0, 3, 3, 64, -1, %0, 0 :: (load (s64), addrspace 3)
- Return
-...
----
-# When the address has non-memory uses too, memory uses are still folded but
-# the `mov` is kept for the remaining use.
-# CHECK-LABEL: name: fold_partial
-# CHECK: %0:b64 = MOV_B64_sym @g
-# CHECK: %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, @g, 0
-# CHECK: %2:b64 = MULT64rr %0, %0
-name: fold_partial
-tracksRegLiveness: true
-body: |
- bb.0:
- %0:b64 = MOV_B64_sym @g
- %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, %0, 0 :: (load (s64), addrspace 1)
- %2:b64 = MULT64rr %0, %0
- Return
-...
More information about the llvm-branch-commits
mailing list