[llvm] [SPIRV]Implementing PopCount for 16 and 64 bits (PR #191283)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Apr 9 19:06:59 PDT 2026
=?utf-8?q?João?= Saffran <joaosaffranllvm at gmail.com>,Joao Saffran
<joaosaffranllvm at gmail.com>,Joao Saffran <joaosaffranllvm at gmail.com>,Joao
Saffran <joaosaffranllvm at gmail.com>,Joao Saffran <joaosaffranllvm at gmail.com>,Joao
Saffran <joaosaffranllvm at gmail.com>,Joao Saffran <joaosaffranllvm at gmail.com>,Joao
Saffran <joaosaffranllvm at gmail.com>
Message-ID:
In-Reply-To: <llvm.org/llvm/llvm-project/pull/191283 at github.com>
https://github.com/joaosaffran updated https://github.com/llvm/llvm-project/pull/191283
>From 60c2c82d17416314c534e2872463a574d278e737 Mon Sep 17 00:00:00 2001
From: Joao Saffran <jderezende at microsoft.com>
Date: Tue, 7 Apr 2026 19:27:18 -0700
Subject: [PATCH 01/11] save work in progress
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 12 +++++++++++-
1 file changed, 11 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 9da8397efc22f..2f8f9277b2e87 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -235,6 +235,9 @@ class SPIRVInstructionSelector : public InstructionSelector {
bool selectOpIsNan(Register ResVReg, SPIRVTypeInst ResType,
MachineInstr &I) const;
+ bool selectPopCount(Register ResVReg, SPIRVTypeInst ResType,
+ MachineInstr &I, unsigned Opcode) const;
+
template <bool Signed>
bool selectDot4AddPacked(Register ResVReg, SPIRVTypeInst ResType,
MachineInstr &I) const;
@@ -1020,7 +1023,7 @@ bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
return selectIToF(ResVReg, ResType, I, false, SPIRV::OpConvertUToF);
case TargetOpcode::G_CTPOP:
- return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitCount);
+ return selectPopCount(ResVReg, ResType, I, SPIRV::OpBitCount);
case TargetOpcode::G_SMIN:
return selectExtInst(ResVReg, ResType, I, CL::s_min, GL::SMin);
case TargetOpcode::G_UMIN:
@@ -1559,6 +1562,13 @@ bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,
return true;
}
+bool SPIRVInstructionSelector::handlePopCount(Register ResVReg,
+ SPIRVTypeInst ResType,
+ MachineInstr &I,
+ unsigned Opcode) const {
+
+}
+
bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I,
>From dcbd23ff176f2b9c290da27d93be1f4778f9c952 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Wed, 8 Apr 2026 16:39:38 -0700
Subject: [PATCH 02/11] add popcount codegen
---
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 320 +++++++++++++++---
.../CodeGen/SPIRV/llvm-intrinsics/ctpop.ll | 12 +-
2 files changed, 277 insertions(+), 55 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 2f8f9277b2e87..4ff8935abdb11 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -238,6 +238,22 @@ class SPIRVInstructionSelector : public InstructionSelector {
bool selectPopCount(Register ResVReg, SPIRVTypeInst ResType,
MachineInstr &I, unsigned Opcode) const;
+ bool selectPopCount16(Register ResVReg, SPIRVTypeInst ResType,
+ MachineInstr &I, unsigned ExtOpcode,
+ unsigned Opcode) const;
+
+ bool selectPopCount32(Register ResVReg, SPIRVTypeInst ResType,
+ MachineInstr &I, Register SrcReg,
+ unsigned Opcode) const;
+
+ bool selectPopCount64(Register ResVReg, SPIRVTypeInst ResType,
+ MachineInstr &I, Register SrcReg,
+ unsigned Opcode) const;
+
+ bool selectPopCount64Overflow(Register ResVReg, SPIRVTypeInst ResType,
+ MachineInstr &I, Register SrcReg,
+ unsigned Opcode) const;
+
template <bool Signed>
bool selectDot4AddPacked(Register ResVReg, SPIRVTypeInst ResType,
MachineInstr &I) const;
@@ -498,6 +514,15 @@ bool sampledTypeIsSignedInteger(const llvm::Type *HandleType) {
#include "SPIRVGenGlobalISel.inc"
#undef GET_GLOBALISEL_IMPL
+static unsigned getVectorSizeOrOne(SPIRVTypeInst Type) {
+
+ if (Type->getOpcode() != SPIRV::OpTypeVector)
+ return 1;
+
+ // Operand(2) is the vector size
+ return Type->getOperand(2).getImm();
+}
+
SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
const SPIRVSubtarget &ST,
const RegisterBankInfo &RBI)
@@ -1562,11 +1587,204 @@ bool SPIRVInstructionSelector::selectOpWithSrcs(Register ResVReg,
return true;
}
-bool SPIRVInstructionSelector::handlePopCount(Register ResVReg,
- SPIRVTypeInst ResType,
- MachineInstr &I,
- unsigned Opcode) const {
+bool SPIRVInstructionSelector::selectPopCount16(Register ResVReg,
+ SPIRVTypeInst ResType,
+ MachineInstr &I,
+ unsigned ExtOpcode,
+ unsigned Opcode) const {
+ Register OpReg = I.getOperand(1).getReg();
+ SPIRVTypeInst OpRegType = GR.getSPIRVTypeForVReg(OpReg);
+
+ unsigned NumElems = getVectorSizeOrOne(OpRegType);
+
+ MachineIRBuilder MIRBuilder(I);
+ SPIRVTypeInst I32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
+ SPIRVTypeInst I32VectorType =
+ GR.getOrCreateSPIRVVectorType(I32Type, NumElems, MIRBuilder, false);
+
+ bool IsVector = NumElems > 1;
+ SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
+ Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
+ if (!selectOpWithSrcs(ExtReg, ExtType, I, {OpReg}, ExtOpcode))
+ return false;
+
+ Register PopCountReg = MRI->createVirtualRegister(GR.getRegClass(ExtType));
+ if (!selectPopCount32(PopCountReg, ExtType, I, ExtReg, Opcode))
+ return false;
+
+ return selectOpWithSrcs(ResVReg, ResType, I, {PopCountReg}, ExtOpcode);
+}
+
+bool SPIRVInstructionSelector::selectPopCount32(Register ResVReg,
+ SPIRVTypeInst ResType,
+ MachineInstr &I,
+ Register SrcReg,
+ unsigned Opcode) const {
+ return selectOpWithSrcs(ResVReg, ResType, I, {SrcReg}, Opcode);
+}
+
+bool SPIRVInstructionSelector::selectPopCount64Overflow(
+ Register ResVReg, SPIRVTypeInst ResType, MachineInstr &I, Register SrcReg,
+ unsigned int Opcode) const {
+
+ unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
+ assert(ComponentCount < 5 && "Vec 5+ will generate invalid SPIR-V ops");
+
+ MachineIRBuilder MIRBuilder(I);
+ SPIRVTypeInst BaseType = GR.retrieveScalarOrVectorIntType(ResType);
+ SPIRVTypeInst I32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
+ SPIRVTypeInst I32x2Type =
+ GR.getOrCreateSPIRVVectorType(I32Type, 2, MIRBuilder, false);
+ SPIRVTypeInst Vec2ResType =
+ GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);
+
+ std::vector<Register> PartialRegs;
+
+ // Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd
+ unsigned CurrentComponent = 0;
+ for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
+ // This register holds the firstbitX result for each of the i64x2 vectors
+ // extracted from SrcReg
+ Register PopCountResult =
+ MRI->createVirtualRegister(GR.getRegClass(I32x2Type));
+
+ auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
+ TII.get(SPIRV::OpVectorShuffle))
+ .addDef(PopCountResult)
+ .addUse(GR.getSPIRVTypeID(I32x2Type))
+ .addUse(SrcReg)
+ .addUse(SrcReg)
+ .addImm(CurrentComponent)
+ .addImm(CurrentComponent + 1);
+
+ MIB.constrainAllUses(TII, TRI, RBI);
+
+ Register SubVecReg =
+ MRI->createVirtualRegister(GR.getRegClass(Vec2ResType));
+
+ if (!selectPopCount64(SubVecReg, Vec2ResType, I, PopCountResult, Opcode))
+ return false;
+
+ PartialRegs.push_back(SubVecReg);
+ }
+
+ // On odd component counts we need to handle one more component
+ if (CurrentComponent != ComponentCount) {
+ bool ZeroAsNull = !STI.isShader();
+ Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I32Type));
+ Register ConstIntLastIdx = GR.getOrCreateConstInt(
+ ComponentCount - 1, I, BaseType, TII, ZeroAsNull);
+
+ if (!selectOpWithSrcs(FinalElemReg, I32Type, I, {SrcReg, ConstIntLastIdx},
+ SPIRV::OpVectorExtractDynamic))
+ return false;
+
+ Register FinalElemResReg =
+ MRI->createVirtualRegister(GR.getRegClass(BaseType));
+
+ if (!selectPopCount64(FinalElemResReg, BaseType, I, FinalElemReg, Opcode))
+ return false;
+
+ PartialRegs.push_back(FinalElemResReg);
+ }
+ // Join all the resulting registers back into the return type in order
+ // (ie i32x2, i32x2, i32x1 -> i32x5)
+ return selectOpWithSrcs(ResVReg, ResType, I, std::move(PartialRegs),
+ SPIRV::OpCompositeConstruct);
+}
+
+bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
+ SPIRVTypeInst ResType,
+ MachineInstr &I,
+ Register SrcReg,
+ unsigned Opcode) const {
+ unsigned ComponentCount = GR.getScalarOrVectorComponentCount(ResType);
+ if (ComponentCount > 2)
+ return selectPopCount64Overflow(ResVReg, ResType, I, SrcReg, Opcode);
+
+ bool ZeroAsNull = !STI.isShader();
+
+ MachineIRBuilder MIRBuilder(I);
+ SPIRVTypeInst I32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
+ SPIRVTypeInst VecI32Type = GR.getOrCreateSPIRVVectorType(
+ I32Type, 2 * ComponentCount, MIRBuilder, false);
+
+ Register BitcastReg = MRI->createVirtualRegister(GR.getRegClass(VecI32Type));
+ if (!selectOpWithSrcs(BitcastReg, VecI32Type, I, {SrcReg}, SPIRV::OpBitcast))
+ return false;
+
+ Register PopCountReg = MRI->createVirtualRegister(GR.getRegClass(VecI32Type));
+ if (!selectPopCount32(PopCountReg, VecI32Type, I, BitcastReg, Opcode))
+ return false;
+
+ bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;
+
+ Register HighReg, LowReg;
+ Register ConstIntZero =
+ GR.getOrCreateConstInt(0, I, I32Type, TII, ZeroAsNull);
+ Register ConstIntOne = GR.getOrCreateConstInt(1, I, I32Type, TII, ZeroAsNull);
+ if (IsScalarRes) {
+ HighReg = MRI->createVirtualRegister(GR.getRegClass(I32Type));
+ LowReg = MRI->createVirtualRegister(GR.getRegClass(I32Type));
+
+ if (!selectOpWithSrcs(HighReg, I32Type, I, {PopCountReg, ConstIntOne},
+ SPIRV::OpVectorExtractDynamic))
+ return false;
+ if (!selectOpWithSrcs(LowReg, I32Type, I, {PopCountReg, ConstIntZero},
+ SPIRV::OpVectorExtractDynamic))
+ return false;
+ } else {
+ SPIRVTypeInst HalfVecI32Type = GR.getOrCreateSPIRVVectorType(
+ I32Type, ComponentCount, MIRBuilder, false);
+
+ HighReg = MRI->createVirtualRegister(GR.getRegClass(HalfVecI32Type));
+ LowReg = MRI->createVirtualRegister(GR.getRegClass(HalfVecI32Type));
+
+ auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
+ TII.get(SPIRV::OpVectorShuffle))
+ .addDef(HighReg)
+ .addUse(GR.getSPIRVTypeID(HalfVecI32Type))
+ .addUse(PopCountReg)
+ .addUse(PopCountReg);
+ for (unsigned J = 1; J < ComponentCount * 2; J += 2)
+ MIB.addImm(J);
+ MIB.constrainAllUses(TII, TRI, RBI);
+
+ MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
+ TII.get(SPIRV::OpVectorShuffle))
+ .addDef(LowReg)
+ .addUse(GR.getSPIRVTypeID(HalfVecI32Type))
+ .addUse(PopCountReg)
+ .addUse(PopCountReg);
+ for (unsigned J = 0; J < ComponentCount * 2; J += 2)
+ MIB.addImm(J);
+ MIB.constrainAllUses(TII, TRI, RBI);
+ }
+
+ unsigned OpAdd = IsScalarRes ? SPIRV::OpIAddS : SPIRV::OpIAddV;
+ return selectOpWithSrcs(ResVReg, ResType, I, {HighReg, LowReg}, OpAdd);
+}
+
+bool SPIRVInstructionSelector::selectPopCount(Register ResVReg,
+ SPIRVTypeInst ResType,
+ MachineInstr &I,
+ unsigned Opcode) const {
+ Register OpReg = I.getOperand(1).getReg();
+ SPIRVTypeInst OpType = GR.getSPIRVTypeForVReg(OpReg);
+ unsigned ExtOpcode = GR.isScalarOrVectorSigned(ResType) ? SPIRV::OpSConvert
+ : SPIRV::OpUConvert;
+ switch (GR.getScalarOrVectorBitWidth(OpType)) {
+ case 8:
+ case 16:
+ return selectPopCount16(ResVReg, ResType, I, ExtOpcode, Opcode);
+ case 32:
+ return selectPopCount32(ResVReg, ResType, I, OpReg, Opcode);
+ case 64:
+ return selectPopCount64(ResVReg, ResType, I, OpReg, Opcode);
+ default:
+ report_fatal_error("unsupported operand bit width for popcount");
+ }
}
bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
@@ -1860,10 +2078,10 @@ bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
- report_fatal_error(
- "llvm.stackrestore intrinsic: this instruction requires the following "
- "SPIR-V extension: SPV_INTEL_variable_length_array",
- false);
+ report_fatal_error("llvm.stackrestore intrinsic: this instruction "
+ "requires the following "
+ "SPIR-V extension: SPV_INTEL_variable_length_array",
+ false);
if (!I.getOperand(0).isReg())
return false;
MachineBasicBlock &BB = *I.getParent();
@@ -1997,8 +2215,8 @@ bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
Register ScopeReg = buildI32Constant(Scope, I);
Register Ptr = I.getOperand(1).getReg();
- // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
- // auto ScSem =
+ // TODO: Changed as it's implemented in the translator. See
+ // test/atomicrmw.ll auto ScSem =
// getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
AtomicOrdering AO = MemOp->getSuccessOrdering();
uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
@@ -2301,8 +2519,9 @@ SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
// In SPIR-V address space casting can only happen to and from the Generic
// storage class. We can also only cast Workgroup, CrossWorkgroup, or Function
// pointers to and from Generic pointers. As such, we can convert e.g. from
-// Workgroup to Function by going via a Generic pointer as an intermediary. All
-// other combinations can only be done by a bitcast, and are probably not safe.
+// Workgroup to Function by going via a Generic pointer as an intermediary.
+// All other combinations can only be done by a bitcast, and are probably not
+// safe.
bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I) const {
@@ -2321,10 +2540,10 @@ bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);
if (isASCastInGVar(MRI, ResVReg)) {
- // AddrSpaceCast uses within OpVariable and OpConstantComposite instructions
- // are expressed by OpSpecConstantOp with an Opcode.
- // TODO: maybe insert a check whether the Kernel capability was declared and
- // so PtrCastToGeneric/GenericCastToPtr are available.
+ // AddrSpaceCast uses within OpVariable and OpConstantComposite
+ // instructions are expressed by OpSpecConstantOp with an Opcode.
+ // TODO: maybe insert a check whether the Kernel capability was declared
+ // and so PtrCastToGeneric/GenericCastToPtr are available.
unsigned SpecOpcode =
DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)
? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)
@@ -2385,8 +2604,8 @@ bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
return true;
}
- // Check if instructions from the SPV_INTEL_usm_storage_classes extension may
- // be applied
+ // Check if instructions from the SPV_INTEL_usm_storage_classes extension
+ // may be applied
if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)
return selectUnOp(ResVReg, ResType, I,
SPIRV::OpPtrCastToCrossWorkgroupINTEL);
@@ -2977,15 +3196,6 @@ bool SPIRVInstructionSelector::selectWaveActiveCountBits(
return true;
}
-unsigned getVectorSizeOrOne(SPIRVTypeInst Type) {
-
- if (Type->getOpcode() != SPIRV::OpTypeVector)
- return 1;
-
- // Operand(2) is the vector size
- return Type->getOperand(2).getImm();
-}
-
bool SPIRVInstructionSelector::selectWaveActiveAllEqual(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I) const {
@@ -3362,7 +3572,8 @@ bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
return selectBitreverseNative(ResVReg, ResType, I, OpReg);
// Expansion bitreverse using bit manipulation operations
- // Algo: https://graphics.stanford.edu/~seander/bithacks.html#ReverseParallel
+ // Algo:
+ // https://graphics.stanford.edu/~seander/bithacks.html#ReverseParallel
const unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
// TODO: add support for any bit width and bitwidth more than 64.
if (BitWidth > 64 || !isPowerOf2_32(BitWidth))
@@ -3429,8 +3640,8 @@ bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,
// There is no way to implement `freeze` correctly without support on SPIR-V
// standard side, but we may at least address a simple (static) case when
// undef/poison value presence is obvious. The main benefit of even
- // incomplete `freeze` support is preventing of translation from crashing due
- // to lack of support on legalization and instruction selection steps.
+ // incomplete `freeze` support is preventing of translation from crashing
+ // due to lack of support on legalization and instruction selection steps.
if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())
return false;
Register OpReg = I.getOperand(1).getReg();
@@ -3837,7 +4048,8 @@ bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
unsigned Opcode) const {
Register SrcReg = I.getOperand(1).getReg();
// We can convert bool value directly to float type without OpConvert*ToF,
- // however the translator generates OpSelect+OpConvert*ToF, so we do the same.
+ // however the translator generates OpSelect+OpConvert*ToF, so we do the
+ // same.
if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
SPIRVTypeInst TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
@@ -4078,9 +4290,9 @@ bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
MachineInstr &I) const {
const bool IsGEPInBounds = I.getOperand(2).getImm();
- // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only
- // relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,
- // we have to use Op[InBounds]AccessChain.
+ // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator
+ // only relies on PtrAccessChain, so we'll try not to deviate. For Vulkan
+ // however, we have to use Op[InBounds]AccessChain.
const unsigned Opcode = STI.isLogicalSPIRV()
? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
: SPIRV::OpAccessChain)
@@ -4121,8 +4333,9 @@ bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
OpDefine->getOpcode() == TargetOpcode::G_INTTOPTR ||
GR.isAggregateType(OpType)) {
- // The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
- // by selectAddrSpaceCast(), and G_INTTOPTR is processed by selectUnOp()
+ // The case of G_ADDRSPACE_CAST inside spv_const_composite() is
+ // processed by selectAddrSpaceCast(), and G_INTTOPTR is processed by
+ // selectUnOp()
CompositeArgs.push_back(OpReg);
continue;
}
@@ -4153,13 +4366,14 @@ bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
bool SPIRVInstructionSelector::selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &I,
const unsigned DPdOpCode) const {
- // TODO: This should check specifically for Fragment Execution Model, but STI
- // doesn't provide that information yet. See #167562
+ // TODO: This should check specifically for Fragment Execution Model, but
+ // STI doesn't provide that information yet. See #167562
errorIfInstrOutsideShader(I);
// If the arg/result types are half then we need to wrap the instr in
// conversions to float
- // This case occurs because a half arg/result is legal in HLSL but not spirv.
+ // This case occurs because a half arg/result is legal in HLSL but not
+ // spirv.
Register SrcReg = I.getOperand(2).getReg();
SPIRVTypeInst SrcType = GR.getSPIRVTypeForVReg(SrcReg);
unsigned BitWidth = std::min(GR.getScalarOrVectorBitWidth(SrcType),
@@ -4238,8 +4452,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
if (!selectGlobalValue(GVarVReg, *MI, Init))
return false;
// We violate SSA form by inserting OpVariable and still having a gMIR
- // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing
- // the duplicated definition.
+ // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by
+ // erasing the duplicated definition.
if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
GR.invalidateMachineInstr(MI);
MI->eraseFromParent();
@@ -4433,8 +4647,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
// The HLSL SV_GroupId semantic is lowered to
// llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.
//
- // In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`
- // builtin variable
+ // In SPIR-V backend, llvm.spv.group.id is now translated to a
+ // `WorkgroupId` builtin variable
return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
I);
case Intrinsic::spv_flattened_thread_id_in_group:
@@ -4442,8 +4656,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
// llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V
// backend.
//
- // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to
- // a `LocalInvocationIndex` builtin variable
+ // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated
+ // to a `LocalInvocationIndex` builtin variable
return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
ResType, I);
case Intrinsic::spv_workgroup_size:
@@ -4629,9 +4843,9 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);
case Intrinsic::spv_radians:
return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);
- // Discard intrinsics which we do not expect to actually represent code after
- // lowering or intrinsics which are not implemented but should not crash when
- // found in a customer's LLVM IR input.
+ // Discard intrinsics which we do not expect to actually represent code
+ // after lowering or intrinsics which are not implemented but should not
+ // crash when found in a customer's LLVM IR input.
case Intrinsic::instrprof_increment:
case Intrinsic::instrprof_increment_step:
case Intrinsic::instrprof_value_profile:
@@ -4747,8 +4961,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I) const {
- // The images need to be loaded in the same basic block as their use. We defer
- // loading the image to the intrinsic that uses it.
+ // The images need to be loaded in the same basic block as their use. We
+ // defer loading the image to the intrinsic that uses it.
if (ResType->getOpcode() == SPIRV::OpTypeImage)
return true;
@@ -4795,9 +5009,9 @@ bool SPIRVInstructionSelector::selectUpdateCounter(Register &ResVReg,
Register CounterHandleReg = Intr.getOperand(2).getReg();
Register IncrReg = Intr.getOperand(3).getReg();
- // The counter handle is a pointer to the counter variable (which is a struct
- // containing an i32). We need to get a pointer to that i32 member to do the
- // atomic operation.
+ // The counter handle is a pointer to the counter variable (which is a
+ // struct containing an i32). We need to get a pointer to that i32 member to
+ // do the atomic operation.
#ifndef NDEBUG
SPIRVTypeInst CounterVarType = GR.getSPIRVTypeForVReg(CounterHandleReg);
SPIRVTypeInst CounterVarPointeeType = GR.getPointeeType(CounterVarType);
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
index 98bc6e0f79e04..9eb400759cfd3 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
@@ -11,9 +11,12 @@
@g3 = addrspace(1) global i32 undef, align 4
@g4 = addrspace(1) global i64 undef, align 8
@g5 = addrspace(1) global <2 x i32> undef, align 4
+ at g6 = addrspace(1) global <2 x i64> undef, align 8
+ at g7 = addrspace(1) global <3 x i64> undef, align 8
+ at g8 = addrspace(1) global <4 x i64> undef, align 8
-define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32) local_unnamed_addr {
+define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64) local_unnamed_addr {
entry:
%0 = tail call i8 @llvm.ctpop.i8(i8 %x8)
store i8 %0, ptr addrspace(1) @g1, align 4
@@ -25,7 +28,12 @@ entry:
store i64 %3, ptr addrspace(1) @g4, align 8
%4 = tail call <2 x i32> @llvm.ctpop.v2i32(<2 x i32> %x2i32)
store <2 x i32> %4, ptr addrspace(1) @g5, align 4
-
+ %5 = tail call <2 x i64> @llvm.ctpop.v2i64(<2 x i64> %x2i64)
+ store <2 x i64> %5, ptr addrspace(1) @g6, align 4
+ %6 = tail call <3 x i64> @llvm.ctpop.v3i64(<3 x i64> %x3i64)
+ store <3 x i64> %6, ptr addrspace(1) @g7, align 4
+ %7 = tail call <4 x i64> @llvm.ctpop.v4i64(<4 x i64> %x4i64)
+ store <4 x i64> %7, ptr addrspace(1) @g8, align 4
ret void
}
>From c9b40afde20f1bc1380a712ab5dd25523a12a508 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 10:45:24 -0700
Subject: [PATCH 03/11] save ongoing progress
---
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 16 +++++++++-------
1 file changed, 9 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 4ff8935abdb11..552bed0ae1093 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1632,9 +1632,9 @@ bool SPIRVInstructionSelector::selectPopCount64Overflow(
MachineIRBuilder MIRBuilder(I);
SPIRVTypeInst BaseType = GR.retrieveScalarOrVectorIntType(ResType);
- SPIRVTypeInst I32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
- SPIRVTypeInst I32x2Type =
- GR.getOrCreateSPIRVVectorType(I32Type, 2, MIRBuilder, false);
+ SPIRVTypeInst I64Type = GR.getOrCreateSPIRVIntegerType(64, MIRBuilder);
+ SPIRVTypeInst I64x2Type =
+ GR.getOrCreateSPIRVVectorType(I64Type, 2, MIRBuilder, false);
SPIRVTypeInst Vec2ResType =
GR.getOrCreateSPIRVVectorType(BaseType, 2, MIRBuilder, false);
@@ -1646,12 +1646,12 @@ bool SPIRVInstructionSelector::selectPopCount64Overflow(
// This register holds the firstbitX result for each of the i64x2 vectors
// extracted from SrcReg
Register PopCountResult =
- MRI->createVirtualRegister(GR.getRegClass(I32x2Type));
+ MRI->createVirtualRegister(GR.getRegClass(I64x2Type));
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
.addDef(PopCountResult)
- .addUse(GR.getSPIRVTypeID(I32x2Type))
+ .addUse(GR.getSPIRVTypeID(I64x2Type))
.addUse(SrcReg)
.addUse(SrcReg)
.addImm(CurrentComponent)
@@ -1671,11 +1671,11 @@ bool SPIRVInstructionSelector::selectPopCount64Overflow(
// On odd component counts we need to handle one more component
if (CurrentComponent != ComponentCount) {
bool ZeroAsNull = !STI.isShader();
- Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I32Type));
+ Register FinalElemReg = MRI->createVirtualRegister(GR.getRegClass(I64Type));
Register ConstIntLastIdx = GR.getOrCreateConstInt(
ComponentCount - 1, I, BaseType, TII, ZeroAsNull);
- if (!selectOpWithSrcs(FinalElemReg, I32Type, I, {SrcReg, ConstIntLastIdx},
+ if (!selectOpWithSrcs(FinalElemReg, I64Type, I, {SrcReg, ConstIntLastIdx},
SPIRV::OpVectorExtractDynamic))
return false;
@@ -1706,6 +1706,8 @@ bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
bool ZeroAsNull = !STI.isShader();
MachineIRBuilder MIRBuilder(I);
+ bool IsSigned = GR.isScalarOrVectorSigned(ResType);
+
SPIRVTypeInst I32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
SPIRVTypeInst VecI32Type = GR.getOrCreateSPIRVVectorType(
I32Type, 2 * ComponentCount, MIRBuilder, false);
>From 9d0e4bfe9007c67005ecec7a9e38af95207e8706 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jo=C3=A3o=20Saffran?= <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 11:39:18 -0700
Subject: [PATCH 04/11] fix codegen and pass offload test
---
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 30 ++++++++++++-------
1 file changed, 19 insertions(+), 11 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 552bed0ae1093..57e3723afdd20 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1706,8 +1706,6 @@ bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
bool ZeroAsNull = !STI.isShader();
MachineIRBuilder MIRBuilder(I);
- bool IsSigned = GR.isScalarOrVectorSigned(ResType);
-
SPIRVTypeInst I32Type = GR.getOrCreateSPIRVIntegerType(32, MIRBuilder);
SPIRVTypeInst VecI32Type = GR.getOrCreateSPIRVVectorType(
I32Type, 2 * ComponentCount, MIRBuilder, false);
@@ -1723,12 +1721,16 @@ bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
bool IsScalarRes = ResType->getOpcode() != SPIRV::OpTypeVector;
Register HighReg, LowReg;
+ SPIRVTypeInst PartsType;
+
Register ConstIntZero =
GR.getOrCreateConstInt(0, I, I32Type, TII, ZeroAsNull);
Register ConstIntOne = GR.getOrCreateConstInt(1, I, I32Type, TII, ZeroAsNull);
+
if (IsScalarRes) {
- HighReg = MRI->createVirtualRegister(GR.getRegClass(I32Type));
- LowReg = MRI->createVirtualRegister(GR.getRegClass(I32Type));
+ PartsType = I32Type;
+ HighReg = MRI->createVirtualRegister(GR.getRegClass(PartsType));
+ LowReg = MRI->createVirtualRegister(GR.getRegClass(PartsType));
if (!selectOpWithSrcs(HighReg, I32Type, I, {PopCountReg, ConstIntOne},
SPIRV::OpVectorExtractDynamic))
@@ -1737,16 +1739,16 @@ bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
SPIRV::OpVectorExtractDynamic))
return false;
} else {
- SPIRVTypeInst HalfVecI32Type = GR.getOrCreateSPIRVVectorType(
- I32Type, ComponentCount, MIRBuilder, false);
+ PartsType = GR.getOrCreateSPIRVVectorType(I32Type, ComponentCount,
+ MIRBuilder, false);
- HighReg = MRI->createVirtualRegister(GR.getRegClass(HalfVecI32Type));
- LowReg = MRI->createVirtualRegister(GR.getRegClass(HalfVecI32Type));
+ HighReg = MRI->createVirtualRegister(GR.getRegClass(PartsType));
+ LowReg = MRI->createVirtualRegister(GR.getRegClass(PartsType));
auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
.addDef(HighReg)
- .addUse(GR.getSPIRVTypeID(HalfVecI32Type))
+ .addUse(GR.getSPIRVTypeID(PartsType))
.addUse(PopCountReg)
.addUse(PopCountReg);
for (unsigned J = 1; J < ComponentCount * 2; J += 2)
@@ -1756,7 +1758,7 @@ bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
.addDef(LowReg)
- .addUse(GR.getSPIRVTypeID(HalfVecI32Type))
+ .addUse(GR.getSPIRVTypeID(PartsType))
.addUse(PopCountReg)
.addUse(PopCountReg);
for (unsigned J = 0; J < ComponentCount * 2; J += 2)
@@ -1765,7 +1767,13 @@ bool SPIRVInstructionSelector::selectPopCount64(Register ResVReg,
}
unsigned OpAdd = IsScalarRes ? SPIRV::OpIAddS : SPIRV::OpIAddV;
- return selectOpWithSrcs(ResVReg, ResType, I, {HighReg, LowReg}, OpAdd);
+ Register AddReg = MRI->createVirtualRegister(GR.getRegClass(PartsType));
+ if (!selectOpWithSrcs(AddReg, PartsType, I, {HighReg, LowReg}, OpAdd))
+ return false;
+
+ bool IsSigned = GR.isScalarOrVectorSigned(PartsType);
+ return selectOpWithSrcs(ResVReg, ResType, I, {AddReg},
+ IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert);
}
bool SPIRVInstructionSelector::selectPopCount(Register ResVReg,
>From da496d1ea227507c1f079d7cdd0746f149768d2d Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 12:40:36 -0700
Subject: [PATCH 05/11] update test
---
.../CodeGen/SPIRV/llvm-intrinsics/ctpop.ll | 111 +++++++++++++++++-
1 file changed, 105 insertions(+), 6 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
index 9eb400759cfd3..4e69c353eb04d 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
@@ -1,10 +1,106 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s
-; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
-; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
-; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
-; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
-; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
+; CHECK-DAG: [[i8_t:%.+]] = OpTypeInt 8 0
+; CHECK-DAG: [[i16_t:%.+]] = OpTypeInt 16 0
+; CHECK-DAG: [[i32_t:%.+]] = OpTypeInt 32 0
+; CHECK-DAG: [[i64_t:%.+]] = OpTypeInt 64 0
+; CHECK-DAG: [[i32x2_t:%.+]] = OpTypeVector [[i32_t]] 2
+; CHECK-DAG: [[i32x3_t:%.+]] = OpTypeVector [[i32_t]] 3
+; CHECK-DAG: [[i32x4_t:%.+]] = OpTypeVector [[i32_t]] 4
+; CHECK-DAG: [[i64x2_t:%.+]] = OpTypeVector [[i64_t]] 2
+; CHECK-DAG: [[i64x3_t:%.+]] = OpTypeVector [[i64_t]] 3
+; CHECK-DAG: [[i64x4_t:%.+]] = OpTypeVector [[i64_t]] 4
+; CHECK-DAG: [[i16x3_t:%.+]] = OpTypeVector [[i16_t]] 3
+
+; CHECK-DAG: [[zero:%.*]] = OpConstantNull [[i32_t]]
+; CHECK-DAG: [[one:%.*]] = OpConstant [[i32_t]] 1
+; CHECK-DAG: [[two:%.*]] = OpConstant [[i64_t]] 2
+
+; CHECK-LABEL: ; -- Begin function test
+
+; CHECK: [[p8:%.+]] = OpFunctionParameter [[i8_t]]
+; CHECK: [[p16:%.+]] = OpFunctionParameter [[i16_t]]
+; CHECK: [[p32:%.+]] = OpFunctionParameter [[i32_t]]
+; CHECK: [[p64:%.+]] = OpFunctionParameter [[i64_t]]
+; CHECK: [[p32x2:%.+]] = OpFunctionParameter [[i32x2_t]]
+; CHECK: [[p64x2:%.+]] = OpFunctionParameter [[i64x2_t]]
+; CHECK: [[p64x3:%.+]] = OpFunctionParameter [[i64x3_t]]
+; CHECK: [[p64x4:%.+]] = OpFunctionParameter [[i64x4_t]]
+; CHECK: [[p16x3:%.+]] = OpFunctionParameter [[i16x3_t]]
+
+; p8
+; CHECK: [[p8_conversion_in:%.+]] = OpUConvert [[i32_t]] [[p8]]
+; CHECK: [[p8_bitcount:%.+]] = OpBitCount [[i32_t]] [[p8_conversion_in]]
+; CHECK: %[[#]] = OpUConvert [[i8_t]] [[p8_bitcount]]
+
+; p16
+; CHECK: [[p16_conversion_in:%.+]] = OpUConvert [[i32_t]] [[p16]]
+; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32_t]] [[p16_conversion_in]]
+; CHECK: %[[#]] = OpUConvert [[i16_t]] [[p16_bitcount]]
+
+; p32
+; CHECK: [[p32_bitcount:%.+]] = OpBitCount [[i32_t]] [[p32]]
+
+; p64
+; CHECK: [[p64_bitcast:%.+]] = OpBitcast [[i32x2_t]] [[p64]]
+; CHECK: [[p64_bitcount:%.+]] = OpBitCount [[i32x2_t]] [[p64_bitcast]]
+; CHECK: [[index_one:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[one]]
+; CHECK: [[index_zero:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[zero]]
+; CHECK: [[add:%.+]] = OpIAdd [[i32_t]] [[index_one]] [[index_zero]]
+; CHECK: [[#]] = OpUConvert [[i64_t]] [[add]]
+
+; p32x2
+; CHECK: [[#]] = OpBitCount [[i32x2_t]] [[p32x2]]
+
+; p64x2
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[p64x2]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[#]] = OpUConvert [[i64x2_t]] [[add]]
+
+; p64x3
+; CHECK: [[first_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x3]] [[p64x3]] 0 1
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[first_half]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[first_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
+
+; CHECK: [[second_half:%.+]] = OpVectorExtractDynamic [[i64_t]] [[p64x3]] [[two]]
+; CHECK: [[p64_bitcast:%.+]] = OpBitcast [[i32x2_t]] [[second_half]]
+; CHECK: [[p64_bitcount:%.+]] = OpBitCount [[i32x2_t]] [[p64_bitcast]]
+; CHECK: [[index_one:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[one]]
+; CHECK: [[index_zero:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[zero]]
+; CHECK: [[add:%.+]] = OpIAdd [[i32_t]] [[index_one]] [[index_zero]]
+; CHECK: [[second_half_result:%.+]] = OpUConvert [[i64_t]] [[add]]
+; CHECK: %[[#]] = OpCompositeConstruct [[i64x3_t]] [[first_half_result]] [[second_half_result]]
+
+; p64x4
+; CHECK: [[first_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x4]] [[p64x4]] 0 1
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[first_half]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[first_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
+
+; CHECK: [[second_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x4]] [[p64x4]] 2 3
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[second_half]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[second_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
+
+; CHECK: %[[#]] = OpCompositeConstruct [[i64x4_t]] [[first_half_result]] [[second_half_result]]
+
+; p16x3
+; CHECK: [[p16_conversion_in:%.+]] = OpUConvert [[i32x3_t]] [[p16x3]]
+; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32x3_t]] [[p16_conversion_in]]
+; CHECK: %[[#]] = OpUConvert [[i16x3_t]] [[p16_bitcount]]
@g1 = addrspace(1) global i8 undef, align 4
@g2 = addrspace(1) global i16 undef, align 4
@@ -14,9 +110,10 @@
@g6 = addrspace(1) global <2 x i64> undef, align 8
@g7 = addrspace(1) global <3 x i64> undef, align 8
@g8 = addrspace(1) global <4 x i64> undef, align 8
+ at g9 = addrspace(1) global <3 x i16> undef, align 4
-define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64) local_unnamed_addr {
+define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64, <3 x i16> %x3i16) local_unnamed_addr {
entry:
%0 = tail call i8 @llvm.ctpop.i8(i8 %x8)
store i8 %0, ptr addrspace(1) @g1, align 4
@@ -34,6 +131,8 @@ entry:
store <3 x i64> %6, ptr addrspace(1) @g7, align 4
%7 = tail call <4 x i64> @llvm.ctpop.v4i64(<4 x i64> %x4i64)
store <4 x i64> %7, ptr addrspace(1) @g8, align 4
+ %8 = tail call <3 x i16> @llvm.ctpop.v3i16(<3 x i16> %x3i16)
+ store <3 x i16> %8, ptr addrspace(1) @g9, align 4
ret void
}
>From 4d85bcc6d40885fd17f3b87b34a1050a0a5fdf65 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 12:48:31 -0700
Subject: [PATCH 06/11] clean up
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 3 ---
1 file changed, 3 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 57e3723afdd20..17fda503ccdd3 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1640,11 +1640,8 @@ bool SPIRVInstructionSelector::selectPopCount64Overflow(
std::vector<Register> PartialRegs;
- // Loops 0, 2, 4, ... but stops one loop early when ComponentCount is odd
unsigned CurrentComponent = 0;
for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
- // This register holds the firstbitX result for each of the i64x2 vectors
- // extracted from SrcReg
Register PopCountResult =
MRI->createVirtualRegister(GR.getRegClass(I64x2Type));
>From 08a6256931ddbb4f8658d22d8dc2459c99ee0b58 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 12:57:47 -0700
Subject: [PATCH 07/11] cleanup format issue
---
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 85 +++++++++----------
1 file changed, 40 insertions(+), 45 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 17fda503ccdd3..ff45c7232e896 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -2085,10 +2085,10 @@ bool SPIRVInstructionSelector::selectStackSave(Register ResVReg,
bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &I) const {
if (!STI.canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
- report_fatal_error("llvm.stackrestore intrinsic: this instruction "
- "requires the following "
- "SPIR-V extension: SPV_INTEL_variable_length_array",
- false);
+ report_fatal_error(
+ "llvm.stackrestore intrinsic: this instruction requires the following "
+ "SPIR-V extension: SPV_INTEL_variable_length_array",
+ false);
if (!I.getOperand(0).isReg())
return false;
MachineBasicBlock &BB = *I.getParent();
@@ -2222,8 +2222,8 @@ bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
Register ScopeReg = buildI32Constant(Scope, I);
Register Ptr = I.getOperand(1).getReg();
- // TODO: Changed as it's implemented in the translator. See
- // test/atomicrmw.ll auto ScSem =
+ // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
+ // auto ScSem =
// getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
AtomicOrdering AO = MemOp->getSuccessOrdering();
uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
@@ -2526,9 +2526,8 @@ SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &I, Register SrcPtr,
// In SPIR-V address space casting can only happen to and from the Generic
// storage class. We can also only cast Workgroup, CrossWorkgroup, or Function
// pointers to and from Generic pointers. As such, we can convert e.g. from
-// Workgroup to Function by going via a Generic pointer as an intermediary.
-// All other combinations can only be done by a bitcast, and are probably not
-// safe.
+// Workgroup to Function by going via a Generic pointer as an intermediary. All
+// other combinations can only be done by a bitcast, and are probably not safe.
bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I) const {
@@ -2547,10 +2546,10 @@ bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResType);
if (isASCastInGVar(MRI, ResVReg)) {
- // AddrSpaceCast uses within OpVariable and OpConstantComposite
- // instructions are expressed by OpSpecConstantOp with an Opcode.
- // TODO: maybe insert a check whether the Kernel capability was declared
- // and so PtrCastToGeneric/GenericCastToPtr are available.
+ // AddrSpaceCast uses within OpVariable and OpConstantComposite instructions
+ // are expressed by OpSpecConstantOp with an Opcode.
+ // TODO: maybe insert a check whether the Kernel capability was declared and
+ // so PtrCastToGeneric/GenericCastToPtr are available.
unsigned SpecOpcode =
DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC)
? static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric)
@@ -2611,8 +2610,8 @@ bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
return true;
}
- // Check if instructions from the SPV_INTEL_usm_storage_classes extension
- // may be applied
+ // Check if instructions from the SPV_INTEL_usm_storage_classes extension may
+ // be applied
if (isUSMStorageClass(SrcSC) && DstSC == SPIRV::StorageClass::CrossWorkgroup)
return selectUnOp(ResVReg, ResType, I,
SPIRV::OpPtrCastToCrossWorkgroupINTEL);
@@ -3579,8 +3578,7 @@ bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
return selectBitreverseNative(ResVReg, ResType, I, OpReg);
// Expansion bitreverse using bit manipulation operations
- // Algo:
- // https://graphics.stanford.edu/~seander/bithacks.html#ReverseParallel
+ // Algo: https://graphics.stanford.edu/~seander/bithacks.html#ReverseParallel
const unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
// TODO: add support for any bit width and bitwidth more than 64.
if (BitWidth > 64 || !isPowerOf2_32(BitWidth))
@@ -3647,8 +3645,8 @@ bool SPIRVInstructionSelector::selectFreeze(Register ResVReg,
// There is no way to implement `freeze` correctly without support on SPIR-V
// standard side, but we may at least address a simple (static) case when
// undef/poison value presence is obvious. The main benefit of even
- // incomplete `freeze` support is preventing of translation from crashing
- // due to lack of support on legalization and instruction selection steps.
+ // incomplete `freeze` support is preventing of translation from crashing due
+ // to lack of support on legalization and instruction selection steps.
if (!I.getOperand(0).isReg() || !I.getOperand(1).isReg())
return false;
Register OpReg = I.getOperand(1).getReg();
@@ -4055,8 +4053,7 @@ bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
unsigned Opcode) const {
Register SrcReg = I.getOperand(1).getReg();
// We can convert bool value directly to float type without OpConvert*ToF,
- // however the translator generates OpSelect+OpConvert*ToF, so we do the
- // same.
+ // however the translator generates OpSelect+OpConvert*ToF, so we do the same.
if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
SPIRVTypeInst TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
@@ -4297,9 +4294,9 @@ bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
MachineInstr &I) const {
const bool IsGEPInBounds = I.getOperand(2).getImm();
- // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator
- // only relies on PtrAccessChain, so we'll try not to deviate. For Vulkan
- // however, we have to use Op[InBounds]AccessChain.
+ // OpAccessChain could be used for OpenCL, but the SPIRV-LLVM Translator only
+ // relies on PtrAccessChain, so we'll try not to deviate. For Vulkan however,
+ // we have to use Op[InBounds]AccessChain.
const unsigned Opcode = STI.isLogicalSPIRV()
? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
: SPIRV::OpAccessChain)
@@ -4340,9 +4337,8 @@ bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
OpDefine->getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
OpDefine->getOpcode() == TargetOpcode::G_INTTOPTR ||
GR.isAggregateType(OpType)) {
- // The case of G_ADDRSPACE_CAST inside spv_const_composite() is
- // processed by selectAddrSpaceCast(), and G_INTTOPTR is processed by
- // selectUnOp()
+ // The case of G_ADDRSPACE_CAST inside spv_const_composite() is processed
+ // by selectAddrSpaceCast(), and G_INTTOPTR is processed by selectUnOp()
CompositeArgs.push_back(OpReg);
continue;
}
@@ -4373,14 +4369,13 @@ bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
bool SPIRVInstructionSelector::selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &I,
const unsigned DPdOpCode) const {
- // TODO: This should check specifically for Fragment Execution Model, but
- // STI doesn't provide that information yet. See #167562
+ // TODO: This should check specifically for Fragment Execution Model, but STI
+ // doesn't provide that information yet. See #167562
errorIfInstrOutsideShader(I);
// If the arg/result types are half then we need to wrap the instr in
// conversions to float
- // This case occurs because a half arg/result is legal in HLSL but not
- // spirv.
+ // This case occurs because a half arg/result is legal in HLSL but not spirv.
Register SrcReg = I.getOperand(2).getReg();
SPIRVTypeInst SrcType = GR.getSPIRVTypeForVReg(SrcReg);
unsigned BitWidth = std::min(GR.getScalarOrVectorBitWidth(SrcType),
@@ -4459,8 +4454,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
if (!selectGlobalValue(GVarVReg, *MI, Init))
return false;
// We violate SSA form by inserting OpVariable and still having a gMIR
- // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by
- // erasing the duplicated definition.
+ // instruction %vreg = G_GLOBAL_VALUE @gvar. We need to fix this by erasing
+ // the duplicated definition.
if (MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
GR.invalidateMachineInstr(MI);
MI->eraseFromParent();
@@ -4654,8 +4649,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
// The HLSL SV_GroupId semantic is lowered to
// llvm.spv.group.id intrinsic in LLVM IR for SPIR-V backend.
//
- // In SPIR-V backend, llvm.spv.group.id is now translated to a
- // `WorkgroupId` builtin variable
+ // In SPIR-V backend, llvm.spv.group.id is now translated to a `WorkgroupId`
+ // builtin variable
return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
I);
case Intrinsic::spv_flattened_thread_id_in_group:
@@ -4663,8 +4658,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
// llvm.spv.flattened.thread.id.in.group() intrinsic in LLVM IR for SPIR-V
// backend.
//
- // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated
- // to a `LocalInvocationIndex` builtin variable
+ // In SPIR-V backend, llvm.spv.flattened.thread.id.in.group is translated to
+ // a `LocalInvocationIndex` builtin variable
return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
ResType, I);
case Intrinsic::spv_workgroup_size:
@@ -4850,9 +4845,9 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
return selectExtInst(ResVReg, ResType, I, CL::step, GL::Step);
case Intrinsic::spv_radians:
return selectExtInst(ResVReg, ResType, I, CL::radians, GL::Radians);
- // Discard intrinsics which we do not expect to actually represent code
- // after lowering or intrinsics which are not implemented but should not
- // crash when found in a customer's LLVM IR input.
+ // Discard intrinsics which we do not expect to actually represent code after
+ // lowering or intrinsics which are not implemented but should not crash when
+ // found in a customer's LLVM IR input.
case Intrinsic::instrprof_increment:
case Intrinsic::instrprof_increment_step:
case Intrinsic::instrprof_value_profile:
@@ -4968,8 +4963,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
bool SPIRVInstructionSelector::selectHandleFromBinding(Register &ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I) const {
- // The images need to be loaded in the same basic block as their use. We
- // defer loading the image to the intrinsic that uses it.
+ // The images need to be loaded in the same basic block as their use. We defer
+ // loading the image to the intrinsic that uses it.
if (ResType->getOpcode() == SPIRV::OpTypeImage)
return true;
@@ -5016,9 +5011,9 @@ bool SPIRVInstructionSelector::selectUpdateCounter(Register &ResVReg,
Register CounterHandleReg = Intr.getOperand(2).getReg();
Register IncrReg = Intr.getOperand(3).getReg();
- // The counter handle is a pointer to the counter variable (which is a
- // struct containing an i32). We need to get a pointer to that i32 member to
- // do the atomic operation.
+ // The counter handle is a pointer to the counter variable (which is a struct
+ // containing an i32). We need to get a pointer to that i32 member to do the
+ // atomic operation.
#ifndef NDEBUG
SPIRVTypeInst CounterVarType = GR.getSPIRVTypeForVReg(CounterHandleReg);
SPIRVTypeInst CounterVarPointeeType = GR.getPointeeType(CounterVarType);
>From 60181602faff633b3b3a52d1630b175be4534353 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 13:55:02 -0700
Subject: [PATCH 08/11] remove undef
---
.../CodeGen/SPIRV/llvm-intrinsics/ctpop.ll | 18 +++++++++---------
1 file changed, 9 insertions(+), 9 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
index 4e69c353eb04d..95e844aecf82c 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
@@ -102,15 +102,15 @@
; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32x3_t]] [[p16_conversion_in]]
; CHECK: %[[#]] = OpUConvert [[i16x3_t]] [[p16_bitcount]]
- at g1 = addrspace(1) global i8 undef, align 4
- at g2 = addrspace(1) global i16 undef, align 4
- at g3 = addrspace(1) global i32 undef, align 4
- at g4 = addrspace(1) global i64 undef, align 8
- at g5 = addrspace(1) global <2 x i32> undef, align 4
- at g6 = addrspace(1) global <2 x i64> undef, align 8
- at g7 = addrspace(1) global <3 x i64> undef, align 8
- at g8 = addrspace(1) global <4 x i64> undef, align 8
- at g9 = addrspace(1) global <3 x i16> undef, align 4
+ at g1 = addrspace(1) global i8 0, align 4
+ at g2 = addrspace(1) global i16 0, align 4
+ at g3 = addrspace(1) global i32 0, align 4
+ at g4 = addrspace(1) global i64 0, align 8
+ at g5 = addrspace(1) global <2 x i32> zeroinitializer, align 4
+ at g6 = addrspace(1) global <2 x i64> zeroinitializer, align 8
+ at g7 = addrspace(1) global <3 x i64> zeroinitializer, align 8
+ at g8 = addrspace(1) global <4 x i64> zeroinitializer, align 8
+ at g9 = addrspace(1) global <3 x i16> zeroinitializer, align 4
define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64, <3 x i16> %x3i16) local_unnamed_addr {
>From 7adef7185c0357cc84df0233d60b791c1c584b19 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 15:38:29 -0700
Subject: [PATCH 09/11] update test
---
.../test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll | 17 ++++++-----------
1 file changed, 6 insertions(+), 11 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
index 95e844aecf82c..1b82f00039082 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
@@ -1,4 +1,7 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-linux %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv1.6-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv1.6-vulkan-compute %s -o - -filetype=obj | spirv-val %}
; CHECK-DAG: [[i8_t:%.+]] = OpTypeInt 8 0
; CHECK-DAG: [[i16_t:%.+]] = OpTypeInt 16 0
@@ -12,7 +15,7 @@
; CHECK-DAG: [[i64x4_t:%.+]] = OpTypeVector [[i64_t]] 4
; CHECK-DAG: [[i16x3_t:%.+]] = OpTypeVector [[i16_t]] 3
-; CHECK-DAG: [[zero:%.*]] = OpConstantNull [[i32_t]]
+; CHECK-DAG: [[zero:%.*]] = OpConstant [[i32_t]] 0
; CHECK-DAG: [[one:%.*]] = OpConstant [[i32_t]] 1
; CHECK-DAG: [[two:%.*]] = OpConstant [[i64_t]] 2
@@ -113,7 +116,7 @@
@g9 = addrspace(1) global <3 x i16> zeroinitializer, align 4
-define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64, <3 x i16> %x3i16) local_unnamed_addr {
+define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64, <3 x i16> %x3i16) local_unnamed_addr #1 {
entry:
%0 = tail call i8 @llvm.ctpop.i8(i8 %x8)
store i8 %0, ptr addrspace(1) @g1, align 4
@@ -136,12 +139,4 @@ entry:
ret void
}
-declare i8 @llvm.ctpop.i8(i8)
-
-declare i16 @llvm.ctpop.i16(i16)
-
-declare i32 @llvm.ctpop.i32(i32)
-
-declare i64 @llvm.ctpop.i64(i64)
-
-declare <2 x i32> @llvm.ctpop.v2i32(<2 x i32>)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
>From 3c37d23fee986f168abf0ad11b8127ecdf29513d Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 18:55:11 -0700
Subject: [PATCH 10/11] limit changes to shaders
---
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 3 +
.../CodeGen/SPIRV/llvm-intrinsics/ctpop-vk.ll | 141 +++++++++++++++++
.../CodeGen/SPIRV/llvm-intrinsics/ctpop.ll | 145 +++---------------
3 files changed, 166 insertions(+), 123 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop-vk.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index ff45c7232e896..365e9ef186f85 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1777,6 +1777,9 @@ bool SPIRVInstructionSelector::selectPopCount(Register ResVReg,
SPIRVTypeInst ResType,
MachineInstr &I,
unsigned Opcode) const {
+ if (!STI.isShader())
+ return selectUnOp(ResVReg, ResType, I, Opcode);
+
Register OpReg = I.getOperand(1).getReg();
SPIRVTypeInst OpType = GR.getSPIRVTypeForVReg(OpReg);
unsigned ExtOpcode = GR.isScalarOrVectorSigned(ResType) ? SPIRV::OpSConvert
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop-vk.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop-vk.ll
new file mode 100644
index 0000000000000..68a6f16e7f1de
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop-vk.ll
@@ -0,0 +1,141 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv1.6-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv1.6-vulkan-compute %s -o - -filetype=obj | spirv-val %}
+
+
+; CHECK-DAG: [[i8_t:%.+]] = OpTypeInt 8 0
+; CHECK-DAG: [[i16_t:%.+]] = OpTypeInt 16 0
+; CHECK-DAG: [[i32_t:%.+]] = OpTypeInt 32 0
+; CHECK-DAG: [[i64_t:%.+]] = OpTypeInt 64 0
+; CHECK-DAG: [[i32x2_t:%.+]] = OpTypeVector [[i32_t]] 2
+; CHECK-DAG: [[i32x3_t:%.+]] = OpTypeVector [[i32_t]] 3
+; CHECK-DAG: [[i32x4_t:%.+]] = OpTypeVector [[i32_t]] 4
+; CHECK-DAG: [[i64x2_t:%.+]] = OpTypeVector [[i64_t]] 2
+; CHECK-DAG: [[i64x3_t:%.+]] = OpTypeVector [[i64_t]] 3
+; CHECK-DAG: [[i64x4_t:%.+]] = OpTypeVector [[i64_t]] 4
+; CHECK-DAG: [[i16x3_t:%.+]] = OpTypeVector [[i16_t]] 3
+
+; CHECK-DAG: [[zero:%.*]] = OpConstant [[i32_t]] 0
+; CHECK-DAG: [[one:%.*]] = OpConstant [[i32_t]] 1
+; CHECK-DAG: [[two:%.*]] = OpConstant [[i64_t]] 2
+
+; CHECK-LABEL: ; -- Begin function test
+
+; CHECK: [[p8:%.+]] = OpFunctionParameter [[i8_t]]
+; CHECK: [[p16:%.+]] = OpFunctionParameter [[i16_t]]
+; CHECK: [[p32:%.+]] = OpFunctionParameter [[i32_t]]
+; CHECK: [[p64:%.+]] = OpFunctionParameter [[i64_t]]
+; CHECK: [[p32x2:%.+]] = OpFunctionParameter [[i32x2_t]]
+; CHECK: [[p64x2:%.+]] = OpFunctionParameter [[i64x2_t]]
+; CHECK: [[p64x3:%.+]] = OpFunctionParameter [[i64x3_t]]
+; CHECK: [[p64x4:%.+]] = OpFunctionParameter [[i64x4_t]]
+; CHECK: [[p16x3:%.+]] = OpFunctionParameter [[i16x3_t]]
+
+; p8
+; CHECK: [[p8_conversion_in:%.+]] = OpUConvert [[i32_t]] [[p8]]
+; CHECK: [[p8_bitcount:%.+]] = OpBitCount [[i32_t]] [[p8_conversion_in]]
+; CHECK: %[[#]] = OpUConvert [[i8_t]] [[p8_bitcount]]
+
+; p16
+; CHECK: [[p16_conversion_in:%.+]] = OpUConvert [[i32_t]] [[p16]]
+; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32_t]] [[p16_conversion_in]]
+; CHECK: %[[#]] = OpUConvert [[i16_t]] [[p16_bitcount]]
+
+; p32
+; CHECK: [[p32_bitcount:%.+]] = OpBitCount [[i32_t]] [[p32]]
+
+; p64
+; CHECK: [[p64_bitcast:%.+]] = OpBitcast [[i32x2_t]] [[p64]]
+; CHECK: [[p64_bitcount:%.+]] = OpBitCount [[i32x2_t]] [[p64_bitcast]]
+; CHECK: [[index_one:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[one]]
+; CHECK: [[index_zero:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[zero]]
+; CHECK: [[add:%.+]] = OpIAdd [[i32_t]] [[index_one]] [[index_zero]]
+; CHECK: [[#]] = OpUConvert [[i64_t]] [[add]]
+
+; p32x2
+; CHECK: [[#]] = OpBitCount [[i32x2_t]] [[p32x2]]
+
+; p64x2
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[p64x2]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[#]] = OpUConvert [[i64x2_t]] [[add]]
+
+; p64x3
+; CHECK: [[first_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x3]] [[p64x3]] 0 1
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[first_half]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[first_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
+
+; CHECK: [[second_half:%.+]] = OpVectorExtractDynamic [[i64_t]] [[p64x3]] [[two]]
+; CHECK: [[p64_bitcast:%.+]] = OpBitcast [[i32x2_t]] [[second_half]]
+; CHECK: [[p64_bitcount:%.+]] = OpBitCount [[i32x2_t]] [[p64_bitcast]]
+; CHECK: [[index_one:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[one]]
+; CHECK: [[index_zero:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[zero]]
+; CHECK: [[add:%.+]] = OpIAdd [[i32_t]] [[index_one]] [[index_zero]]
+; CHECK: [[second_half_result:%.+]] = OpUConvert [[i64_t]] [[add]]
+; CHECK: %[[#]] = OpCompositeConstruct [[i64x3_t]] [[first_half_result]] [[second_half_result]]
+
+; p64x4
+; CHECK: [[first_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x4]] [[p64x4]] 0 1
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[first_half]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[first_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
+
+; CHECK: [[second_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x4]] [[p64x4]] 2 3
+; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[second_half]]
+; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
+; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
+; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
+; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
+; CHECK: [[second_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
+
+; CHECK: %[[#]] = OpCompositeConstruct [[i64x4_t]] [[first_half_result]] [[second_half_result]]
+
+; p16x3
+; CHECK: [[p16_conversion_in:%.+]] = OpUConvert [[i32x3_t]] [[p16x3]]
+; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32x3_t]] [[p16_conversion_in]]
+; CHECK: %[[#]] = OpUConvert [[i16x3_t]] [[p16_bitcount]]
+
+ at g1 = addrspace(1) global i8 0, align 4
+ at g2 = addrspace(1) global i16 0, align 4
+ at g3 = addrspace(1) global i32 0, align 4
+ at g4 = addrspace(1) global i64 0, align 8
+ at g5 = addrspace(1) global <2 x i32> zeroinitializer, align 4
+ at g6 = addrspace(1) global <2 x i64> zeroinitializer, align 8
+ at g7 = addrspace(1) global <3 x i64> zeroinitializer, align 8
+ at g8 = addrspace(1) global <4 x i64> zeroinitializer, align 8
+ at g9 = addrspace(1) global <3 x i16> zeroinitializer, align 4
+
+
+define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64, <3 x i16> %x3i16) local_unnamed_addr #1 {
+entry:
+ %0 = tail call i8 @llvm.ctpop.i8(i8 %x8)
+ store i8 %0, ptr addrspace(1) @g1, align 4
+ %1 = tail call i16 @llvm.ctpop.i16(i16 %x16)
+ store i16 %1, ptr addrspace(1) @g2, align 4
+ %2 = tail call i32 @llvm.ctpop.i32(i32 %x32)
+ store i32 %2, ptr addrspace(1) @g3, align 4
+ %3 = tail call i64 @llvm.ctpop.i64(i64 %x64)
+ store i64 %3, ptr addrspace(1) @g4, align 8
+ %4 = tail call <2 x i32> @llvm.ctpop.v2i32(<2 x i32> %x2i32)
+ store <2 x i32> %4, ptr addrspace(1) @g5, align 4
+ %5 = tail call <2 x i64> @llvm.ctpop.v2i64(<2 x i64> %x2i64)
+ store <2 x i64> %5, ptr addrspace(1) @g6, align 4
+ %6 = tail call <3 x i64> @llvm.ctpop.v3i64(<3 x i64> %x3i64)
+ store <3 x i64> %6, ptr addrspace(1) @g7, align 4
+ %7 = tail call <4 x i64> @llvm.ctpop.v4i64(<4 x i64> %x4i64)
+ store <4 x i64> %7, ptr addrspace(1) @g8, align 4
+ %8 = tail call <3 x i16> @llvm.ctpop.v3i16(<3 x i16> %x3i16)
+ store <3 x i16> %8, ptr addrspace(1) @g9, align 4
+ ret void
+}
+
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
index 1b82f00039082..4bc2ae4c089e8 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctpop.ll
@@ -1,122 +1,20 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-linux %s -o - -filetype=obj | spirv-val %}
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv1.6-vulkan-compute %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv1.6-vulkan-compute %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-linux %s -o - -filetype=obj | spirv-val %}
-; CHECK-DAG: [[i8_t:%.+]] = OpTypeInt 8 0
-; CHECK-DAG: [[i16_t:%.+]] = OpTypeInt 16 0
-; CHECK-DAG: [[i32_t:%.+]] = OpTypeInt 32 0
-; CHECK-DAG: [[i64_t:%.+]] = OpTypeInt 64 0
-; CHECK-DAG: [[i32x2_t:%.+]] = OpTypeVector [[i32_t]] 2
-; CHECK-DAG: [[i32x3_t:%.+]] = OpTypeVector [[i32_t]] 3
-; CHECK-DAG: [[i32x4_t:%.+]] = OpTypeVector [[i32_t]] 4
-; CHECK-DAG: [[i64x2_t:%.+]] = OpTypeVector [[i64_t]] 2
-; CHECK-DAG: [[i64x3_t:%.+]] = OpTypeVector [[i64_t]] 3
-; CHECK-DAG: [[i64x4_t:%.+]] = OpTypeVector [[i64_t]] 4
-; CHECK-DAG: [[i16x3_t:%.+]] = OpTypeVector [[i16_t]] 3
+; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
+; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
+; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
+; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
+; CHECK: %[[#]] = OpBitCount %[[#]] %[[#]]
-; CHECK-DAG: [[zero:%.*]] = OpConstant [[i32_t]] 0
-; CHECK-DAG: [[one:%.*]] = OpConstant [[i32_t]] 1
-; CHECK-DAG: [[two:%.*]] = OpConstant [[i64_t]] 2
+ at g1 = addrspace(1) global i8 undef, align 4
+ at g2 = addrspace(1) global i16 undef, align 4
+ at g3 = addrspace(1) global i32 undef, align 4
+ at g4 = addrspace(1) global i64 undef, align 8
+ at g5 = addrspace(1) global <2 x i32> undef, align 4
-; CHECK-LABEL: ; -- Begin function test
-; CHECK: [[p8:%.+]] = OpFunctionParameter [[i8_t]]
-; CHECK: [[p16:%.+]] = OpFunctionParameter [[i16_t]]
-; CHECK: [[p32:%.+]] = OpFunctionParameter [[i32_t]]
-; CHECK: [[p64:%.+]] = OpFunctionParameter [[i64_t]]
-; CHECK: [[p32x2:%.+]] = OpFunctionParameter [[i32x2_t]]
-; CHECK: [[p64x2:%.+]] = OpFunctionParameter [[i64x2_t]]
-; CHECK: [[p64x3:%.+]] = OpFunctionParameter [[i64x3_t]]
-; CHECK: [[p64x4:%.+]] = OpFunctionParameter [[i64x4_t]]
-; CHECK: [[p16x3:%.+]] = OpFunctionParameter [[i16x3_t]]
-
-; p8
-; CHECK: [[p8_conversion_in:%.+]] = OpUConvert [[i32_t]] [[p8]]
-; CHECK: [[p8_bitcount:%.+]] = OpBitCount [[i32_t]] [[p8_conversion_in]]
-; CHECK: %[[#]] = OpUConvert [[i8_t]] [[p8_bitcount]]
-
-; p16
-; CHECK: [[p16_conversion_in:%.+]] = OpUConvert [[i32_t]] [[p16]]
-; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32_t]] [[p16_conversion_in]]
-; CHECK: %[[#]] = OpUConvert [[i16_t]] [[p16_bitcount]]
-
-; p32
-; CHECK: [[p32_bitcount:%.+]] = OpBitCount [[i32_t]] [[p32]]
-
-; p64
-; CHECK: [[p64_bitcast:%.+]] = OpBitcast [[i32x2_t]] [[p64]]
-; CHECK: [[p64_bitcount:%.+]] = OpBitCount [[i32x2_t]] [[p64_bitcast]]
-; CHECK: [[index_one:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[one]]
-; CHECK: [[index_zero:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[zero]]
-; CHECK: [[add:%.+]] = OpIAdd [[i32_t]] [[index_one]] [[index_zero]]
-; CHECK: [[#]] = OpUConvert [[i64_t]] [[add]]
-
-; p32x2
-; CHECK: [[#]] = OpBitCount [[i32x2_t]] [[p32x2]]
-
-; p64x2
-; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[p64x2]]
-; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
-; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
-; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
-; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
-; CHECK: [[#]] = OpUConvert [[i64x2_t]] [[add]]
-
-; p64x3
-; CHECK: [[first_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x3]] [[p64x3]] 0 1
-; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[first_half]]
-; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
-; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
-; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
-; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
-; CHECK: [[first_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
-
-; CHECK: [[second_half:%.+]] = OpVectorExtractDynamic [[i64_t]] [[p64x3]] [[two]]
-; CHECK: [[p64_bitcast:%.+]] = OpBitcast [[i32x2_t]] [[second_half]]
-; CHECK: [[p64_bitcount:%.+]] = OpBitCount [[i32x2_t]] [[p64_bitcast]]
-; CHECK: [[index_one:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[one]]
-; CHECK: [[index_zero:%.+]] = OpVectorExtractDynamic [[i32_t]] [[p64_bitcount]] [[zero]]
-; CHECK: [[add:%.+]] = OpIAdd [[i32_t]] [[index_one]] [[index_zero]]
-; CHECK: [[second_half_result:%.+]] = OpUConvert [[i64_t]] [[add]]
-; CHECK: %[[#]] = OpCompositeConstruct [[i64x3_t]] [[first_half_result]] [[second_half_result]]
-
-; p64x4
-; CHECK: [[first_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x4]] [[p64x4]] 0 1
-; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[first_half]]
-; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
-; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
-; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
-; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
-; CHECK: [[first_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
-
-; CHECK: [[second_half:%.+]] = OpVectorShuffle [[i64x2_t]] [[p64x4]] [[p64x4]] 2 3
-; CHECK: [[p64x2_bitcast:%.+]] = OpBitcast [[i32x4_t]] [[second_half]]
-; CHECK: [[p64x2_bitcount:%.+]] = OpBitCount [[i32x4_t]] [[p64x2_bitcast]]
-; CHECK: [[odd_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 1 3
-; CHECK: [[even_indexes:%.+]] = OpVectorShuffle [[i32x2_t]] [[p64x2_bitcount]] [[p64x2_bitcount]] 0 2
-; CHECK: [[add:%.+]] = OpIAdd [[i32x2_t]] [[odd_indexes]] [[even_indexes]]
-; CHECK: [[second_half_result:%.+]] = OpUConvert [[i64x2_t]] [[add]]
-
-; CHECK: %[[#]] = OpCompositeConstruct [[i64x4_t]] [[first_half_result]] [[second_half_result]]
-
-; p16x3
-; CHECK: [[p16_conversion_in:%.+]] = OpUConvert [[i32x3_t]] [[p16x3]]
-; CHECK: [[p16_bitcount:%.+]] = OpBitCount [[i32x3_t]] [[p16_conversion_in]]
-; CHECK: %[[#]] = OpUConvert [[i16x3_t]] [[p16_bitcount]]
-
- at g1 = addrspace(1) global i8 0, align 4
- at g2 = addrspace(1) global i16 0, align 4
- at g3 = addrspace(1) global i32 0, align 4
- at g4 = addrspace(1) global i64 0, align 8
- at g5 = addrspace(1) global <2 x i32> zeroinitializer, align 4
- at g6 = addrspace(1) global <2 x i64> zeroinitializer, align 8
- at g7 = addrspace(1) global <3 x i64> zeroinitializer, align 8
- at g8 = addrspace(1) global <4 x i64> zeroinitializer, align 8
- at g9 = addrspace(1) global <3 x i16> zeroinitializer, align 4
-
-
-define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32, <2 x i64> %x2i64, <3 x i64> %x3i64, <4 x i64> %x4i64, <3 x i16> %x3i16) local_unnamed_addr #1 {
+define dso_local spir_kernel void @test(i8 %x8, i16 %x16, i32 %x32, i64 %x64, <2 x i32> %x2i32) local_unnamed_addr {
entry:
%0 = tail call i8 @llvm.ctpop.i8(i8 %x8)
store i8 %0, ptr addrspace(1) @g1, align 4
@@ -128,15 +26,16 @@ entry:
store i64 %3, ptr addrspace(1) @g4, align 8
%4 = tail call <2 x i32> @llvm.ctpop.v2i32(<2 x i32> %x2i32)
store <2 x i32> %4, ptr addrspace(1) @g5, align 4
- %5 = tail call <2 x i64> @llvm.ctpop.v2i64(<2 x i64> %x2i64)
- store <2 x i64> %5, ptr addrspace(1) @g6, align 4
- %6 = tail call <3 x i64> @llvm.ctpop.v3i64(<3 x i64> %x3i64)
- store <3 x i64> %6, ptr addrspace(1) @g7, align 4
- %7 = tail call <4 x i64> @llvm.ctpop.v4i64(<4 x i64> %x4i64)
- store <4 x i64> %7, ptr addrspace(1) @g8, align 4
- %8 = tail call <3 x i16> @llvm.ctpop.v3i16(<3 x i16> %x3i16)
- store <3 x i16> %8, ptr addrspace(1) @g9, align 4
+
ret void
}
-attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
+declare i8 @llvm.ctpop.i8(i8)
+
+declare i16 @llvm.ctpop.i16(i16)
+
+declare i32 @llvm.ctpop.i32(i32)
+
+declare i64 @llvm.ctpop.i64(i64)
+
+declare <2 x i32> @llvm.ctpop.v2i32(<2 x i32>)
>From b4569bf615b520537b2f38358eb5246139bf33b3 Mon Sep 17 00:00:00 2001
From: Joao Saffran <joaosaffranllvm at gmail.com>
Date: Thu, 9 Apr 2026 19:06:40 -0700
Subject: [PATCH 11/11] fix extent bug
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 365e9ef186f85..60d21b805b2a7 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -1605,7 +1605,8 @@ bool SPIRVInstructionSelector::selectPopCount16(Register ResVReg,
bool IsVector = NumElems > 1;
SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
Register ExtReg = MRI->createVirtualRegister(GR.getRegClass(ResType));
- if (!selectOpWithSrcs(ExtReg, ExtType, I, {OpReg}, ExtOpcode))
+ // Always use OpUConvert to always use a 0 extend
+ if (!selectOpWithSrcs(ExtReg, ExtType, I, {OpReg}, SPIRV::OpUConvert))
return false;
Register PopCountReg = MRI->createVirtualRegister(GR.getRegClass(ExtType));
More information about the llvm-commits
mailing list