[Ventus] Add SFU/MMA lowering, rrsource metadata updates and others#213
Open
Humber-186 wants to merge 41 commits into
Open
[Ventus] Add SFU/MMA lowering, rrsource metadata updates and others#213Humber-186 wants to merge 41 commits into
Humber-186 wants to merge 41 commits into
Conversation
Add a pre-RA Ventus keepalive pass that extends SGPR-class vregs to the join of divergent vbranch regions so sibling paths cannot clobber scalar values that must survive until reconvergence. The pass now handles three important Ventus-specific edge cases: shared joins across nested divergence, join-only PHI edge uses, and explicit fatal errors for divergent branches that do not have an immediate post-dominator. It also filters keepalive candidates with MachineDominatorTree so we do not create non-dominating SSA uses on loop exits. Tests cover the basic keepalive insertion, join-PHI users, no-divergence and loop-header negative cases, the verifier regression from vbranch-join.ll, and the no-join crash path with not --crash.
Background: Ventus emits per-kernel .ventus.resource.<kernel> metadata, but the old VGPR/SGPR accounting was based on deduplicated physreg count. That does not match the hardware requirement of allocating a contiguous register window [0, usage), so sparse high-numbered register use under-reports the required window size. Implementation: switch VGPR/SGPR usage semantics to highest physical register index plus one, remove the old dedup-tracking state, and resolve physical subregs so tuple/window-style registers contribute through their leaf registers. To avoid missing PEI/eliminateFrameIndex-introduced base registers such as X4/X8/V32/V8, recompute final register usage from the post-PEI MachineFunction in RISCVAsmPrinter before emitting .ventus.resource.* while preserving LDS/PDS accounting from frame lowering. Validation: update resource-usage expectations and add a local_addressed_variables resource check covering post-PEI frame-base usage; llvm-lit passes for both Ventus resource tests.
Emit one .ventus.resobj payload per regular object from RISCVAsmPrinter and record function-level resource facts, direct-call edges, tail-call information, and static LDS/PDS references for later aggregation. Teach lld to consume .ventus.resobj after symbol resolution, GC, and COMDAT selection, then synthesize final .ventus.resource.<kernel> sections for non-relocatable Ventus links. The parser also accepts multiple payloads in a merged .ventus.resobj section so ld -r outputs can be finalized later. Distinguish unresolved callees from resolved-but-unsummarized callees: only the former keep HasUnknownExternalCallee, while both conservatively preserve RegisterUsageIncomplete and StackPeakUnavailable. Reject mixed legacy/new input formats, drop transient summaries from final outputs, and add LLVM/lld tests for resobj emission, tail edges, finalized aggregation, linker scripts, ld -r merges, mixed formats, and unsummarized helpers.
Humber-186
requested review from
Jules-Kong,
dodohack and
ivan-lei
as code owners
April 28, 2026 02:31
Ventus lowers OpenCL barrier, work_group_barrier and sub_group_barrier builtins to RISCV target intrinsics. The OpenCL builtins are convergent operations, but the Ventus target intrinsic definitions only carried IntrNoMem and IntrHasSideEffects. That let normal IR optimization treat the barrier calls as non-convergent side-effecting calls, so a barrier between divergent regions could be duplicated into both branch arms before reconvergence. Add IntrConvergent to all Ventus barrier-related RISCV intrinsics: llvm.riscv.ventus.barrier, barrier.with.scope, subgroup.barrier and subgroup.barrier.with.scope. This preserves the source-level convergence contract at the IR layer where mid-end optimization can see it. Update the RISCV OpenCL builtin CodeGen test to check the generated intrinsic declarations. Add an intrinsic-table smoke test for reconstructed declaration attributes. Add an -O1 OpenCL regression that verifies a barrier between lane-dependent regions is not cloned into extra call sites. Validation performed: - ninja -C llvm/build clang llvm-as llvm-dis llc llvm-objdump - timeout 60s llvm/build/bin/llvm-lit llvm/clang/test/CodeGenOpenCL/builtins-riscv.cl llvm/clang/test/CodeGenOpenCL/ventus-barrier-convergent-opt.cl llvm/llvm/test/CodeGen/RISCV/VentusGPGPU/barrier-intrinsic-attrs.ll - ninja -C llvm/build install - timeout 60s /work/issues/ventus-llvm-opencl-barrier-convergent-2026-04-28/reproduce.sh The issue reproduction no longer reports the missing convergent attribute, extra lud_diagonal barrier call site, or split barrier pattern.
Ventus divergent execution can interleave sibling paths before reconvergence. A scalar GPR value that is live from before a vector branch may therefore be read by lanes on one sibling while another sibling has already reused the same physical SGPR. Register allocation can expose this when a logical SGPR is split or reloaded into different vregs and the final physreg assignments collide across sibling paths. Add a post-register-allocation VentusSGPRSIMTChecker pass before VentusRemoveSGPRKeepAlive. The checker walks each divergent branch region up to its immediate post-dominator, records physical GPR live-in uses in each sibling, records physical GPR defs in other siblings, and reports a fatal backend error when the two overlap. Regmask operands, such as call and inline-asm clobbers, are also treated as physical GPR defs so caller-saved clobbers are not missed. The checker also models keepalive-only uses that do not appear as ordinary sibling instructions. Join block PseudoSGPRKeepAlive instructions are checked for both direct and non-direct joins, and PseudoSGPRKeepAlive instructions in edge keepalive blocks are treated as live-in uses while the marker PseudoSGPRKeepAliveBlock remains ignored. Add focused MIR tests for explicit sibling defs, regmask clobbers, direct-join keepalives, non-direct join keepalives, and edge keepalive blocks. These tests turn the previously silent SIMT SGPR clobber cases into deterministic checker failures.
Ventus SIMT divergent branches can execute sibling paths in a dynamic order while the scalar register file is not saved/restored by the SIMT stack. Ordinary LLVM liveness lets an SGPR child die after an early sibling use, so another sibling can later reuse the same physical SGPR before the first sibling resumes. This caused cfd_i1 compute_flux to use a clobbered a1/x11 value for nelr and derive invalid vector load addresses. Add a target hook for register-allocation extra interference and implement VentusSIMTInterference. For SGPR values whose reaching def is outside a divergent scope and whose use is inside a sibling component, compute a shadow live range over sibling components that may execute before the protected use. Insert that shadow range into LiveRegMatrix on assignment so greedy RA sees the real physreg pressure and can still choose preserve, evict, split, spill/reload, or rematerialize instead of being forced into a backend workaround. Teach LiveRegMatrix to query target extra interference, combine it with ordinary interference severity, and maintain extra ranges across assign/unassign. Add LiveIntervalUnion extraction support for removing all segments owned by a vreg so shadow segments do not leave stale interference. Shadow virtual-interference queries avoid the cached LiveRegMatrix::query path because they operate on temporary LiveRanges. Wire the Ventus implementation through RISCVRegisterInfo for Ventus kernels only. Keep non-Ventus targets on the default null hook. Strengthen VentusSGPRSIMTChecker as a safety checker: treat PseudoSGPRKeepAlive as a synthetic live-in use, keep keepalive block markers out of def collection, and use CFG may-dataflow so SGPR spill reloads only clear clobbers on paths they actually dominate. This preserves diagnostics for sibling clobbers, regmask clobbers, keepalive-only hazards, join reloads, and diamond-shaped sibling control flow. Add MIR tests for RA shadow interference, sibling clobber/regmask cases, join reload, reload-after-clobber, diamond reload path sensitivity, and keepalive hazard coverage. Validated with: ninja -C llvm/build libLLVMRISCVCodeGen.so bin/llc bin/clang; timeout 60s llvm/build/bin/llvm-lit over the Ventus SGPR SIMT tests; timeout 60s clang rebuild of rodinia/opencl/cfd compute_flux; and a bounded objdump check confirming the old vadd.vx ... a1 early-use hazard pattern is gone.
Refine the register-allocation SIMT shadow interference support added for Ventus SGPR hazards. LiveRegMatrix now lets the target unassign hook report whether it actually inserted shadow segments for the virtual register. This keeps the normal fast extract path for ordinary assignments and only falls back to removing all matrix entries owned by a vreg when extra shadow segments were present. VentusSIMTInterference tracks assigned shadow owners explicitly so invalidateVirtRegs() does not lose the information needed by later unassign calls. Enable the Ventus extra-interference hook based on the ventus-gpgpu CPU instead of only SPIR/VENTUS kernel calling conventions. Device helper functions compiled for Ventus can still contain divergent SIMT control flow and SGPR uses, so they need the same RA interference modeling as entry kernels. Turn legacy SGPR keepalive insertion into an explicit hidden option. The pass still converges returns before RA, but keepalive marker insertion is disabled by default now that RA models the shadow interference directly. Existing keepalive tests are updated to exercise both the default disabled behavior and the legacy opt-in path. Update Ventus codegen checks for the new allocation decisions and add MIR coverage for spir_func and ordinary helper functions. The updated tests also cover the revised keepalive option spelling and the RA shadow-interference behavior after virtual-interference conflicts are treated as non-evictable. Validation: ninja -C build libLLVMRISCVCodeGen.so bin/llc; timeout 60s build/bin/llvm-lit over the modified VentusGPGPU tests and sgpr-simt checker tests.
Tighten the Ventus SIMT SGPR handling in both the post-RA checker and the register allocator extra interference hook. The checker now uses the same Ventus subtarget predicate as the RA hook and exits early for non-Ventus functions. This keeps the fatal checker aligned with the mechanism that prevents these hazards during allocation. Join-block local def tracking now records only physical SGPR definitions, and it accounts for regmask clobbers explicitly. This avoids treating virtual, non-SGPR, or noreg operands as physical SGPR defs when suppressing later join live-in uses. The RA shadow model now indexes scopes by join block in addition to component block. A use in a block that is both an outer component and an inner divergent join is processed in both roles, so nested divergent joins no longer skip the inner join shadow range. Undef uses are ignored because they do not require preserving the original SGPR value. Tests cover join-local def filtering and add a nested divergent join case to the SIMT shadow interference MIR coverage. Verification: - ninja -C build bin/llc - timeout 60 ./build/bin/llvm-lit -q llvm/test/CodeGen/RISCV/VentusGPGPU/ventus-simt-ra-shadow-interference.mir llvm/test/CodeGen/RISCV/VentusGPGPU/sgpr-simt-join-local-def-filter-checker.mir llvm/test/CodeGen/RISCV/VentusGPGPU/sgpr-simt-join-reload-checker.mir llvm/test/CodeGen/RISCV/VentusGPGPU/sgpr-simt-direct-join-clobber-checker.mir llvm/test/CodeGen/RISCV/VentusGPGPU/sgpr-simt-sibling-clobber-checker.mir llvm/test/CodeGen/RISCV/VentusGPGPU/sgpr-simt-sibling-regmask-clobber-checker.mir
The recently added post-RA SGPR SIMT checker is currently reporting false positives during normal Ventus compilation. Stop adding it to the default post-register-allocation pipeline for now, while keeping the pass registered so it can still be invoked explicitly with -run-pass=ventus-sgpr-simt-checker for debugging. The SGPR keepalive removal pass remains enabled so the surrounding Ventus pipeline behavior is otherwise unchanged.
Ventus private pointers are raw private stack offsets, but OpenCL 2.0 generic helper signatures erase that provenance to flat addrspace(0). Once those private-derived generic pointers flow through helper calls or generic memory ops, codegen silently selects ordinary flat/global/local addressing and miscompiles private stack accesses. Add a Ventus-only generic address space specialization module pass before VentusPromoteAlloca. The pass tracks private-to-flat addrspace casts, rebuilds equivalent addrspace(5) values through GEP/phi/select chains, rewrites loads, stores, memintrinsics, helper calls, cast-back patterns and pointer compares to use private pointers again, and drops metadata-only intrinsic users that would otherwise keep dead generic casts alive. When a private-derived generic pointer crosses a direct helper call, clone the callee with specialized private pointer parameters, preserve operand bundles and call properties, and recursively narrow the clone body. Reject unresolved or ABI-sensitive escapes explicitly instead of silently falling back: external or indirect calls, variadic forwarding, pointer stores, ptrtoint/inttoptr style escapes, mixed-address-space phi/select merges, kernel cloning, recursion, and ABI-sensitive pointer attributes now fail with a fatal diagnostic. Add focused Ventus CodeGen coverage for helper cloning, operand bundles, memset/memcpy style narrowing, loop phi/gep rebuilding, cast-back elimination, pointer compare rewriting, lifetime intrinsic cleanup, available_externally helpers, and the expected hard failures for unsupported escape patterns.
Purpose: The Ventus generic address-space specialization pass rejects private-derived generic pointers that reach external calls. The existing libclc build rule also emitted per-source .bc.o files before helper definitions were linked into the module, so math helpers such as __clc_argReductionS and __clc_ep_log appeared as external calls and stopped the riscv32 libclc build. Changes: - Stop the CLC language rule from producing per-source target objects; it now emits LLVM bitcode only. - Generate riscv32clc.o from the linked builtins.link.riscv32--.bc module, where libclc helper definitions are visible to the Ventus address-space specialization pass. - Add explicit failure when the linked libclc bitcode is missing. - Add Zfinx selection patterns for plain riscv_fcvt_x/riscv_fcvt_xu nodes and a regression test covering the float-to-integer conversion path used by generated convert builtins. Impact: This removes the intermediate *.bc.o artifacts from the normal libclc CLC compile rule, so scripts that consumed those files must use the linked bitcode/object flow instead. Building the final riscv32clc.o now runs code generation over the whole linked libclc module, which may take longer and can expose real backend issues that were previously hidden by per-file codegen. The new FCVT patterns broaden Zfinx instruction selection for non-wrapper FCVT nodes while preserving the existing Ventus divergent vector path.
Originally authored by Yuyi and rebased onto the current branch by WangYuhan. Co-authored-by: whywhy <yiyu33129@gmail.com>
Jules-Kong
force-pushed
the
dev-thu-sfu-mma
branch
from
May 13, 2026 12:19
4abacd0 to
518173e
Compare
Ventus private memory load/store instructions encode their frame offset as a
signed 11-bit immediate. The previous frame-index elimination path used a
signed 12-bit check before accepting the folded offset, so offsets such as
-1028 were left directly on VSW/VLW. That immediate is outside the real
[-1024, 1023] range and can either trip the machine verifier or be encoded as a
wrong address for a real Ventus binary.
Fix the private-memory X4 path to check simm11 instead of simm12. When the
folded offset is outside simm11 but still within the signed 12-bit fragment
produced by the existing frame lowering code, materialize a one-step VGPR base
adjustment with vadd12.vi/vsub12.vi by 1024 bytes and leave only the remaining
legal simm11 offset on the memory instruction. This turns e.g. a direct
vsw.v v0, -1028(v32)
into
vsub12.vi v1, v32, 1024
vsw.v v0, -4(v1)
Also return immediately after handling the private X4 case. Without this, the
same instruction can continue into the later uniform/local memory branches after
its base register and offset have already been rewritten.
The X2 private-memory conversion is left to the existing uniform/local handling
because the converted opcode has a signed 12-bit immediate and is no longer the
private-memory simm11 encoding that needs this legalization.
Add a focused llc regression that requires a -1028 private-frame store. The old
backend fails with -verify-machineinstrs:
*** Bad machine code: Invalid immediate ***
VSW ... -1028
With this change, the regression emits vsub12.vi followed by a VSW with -4 and
passes the verifier.
Additional spike reproducer used during review:
__attribute__((noinline))
void keep_private_frame(volatile private int *buf) {
if (get_global_id(0) == (size_t)-1)
buf[0] = buf[0];
}
__kernel void private_frame_simm11(__global int *out) {
volatile private int buf[257];
int gid = get_global_id(0);
buf[0] = gid + 1234;
keep_private_frame(buf);
out[gid] = buf[0] - gid;
}
The helper keeps the 257-int private frame alive so the access to buf[0] reaches
the problematic large negative private-frame offset. Running this through the
Ventus POCL spike backend with two temporary installs produced:
old LLVM: /work/ventus-env-gh/install-review-old-simm11/bin/clang
private_frame_simm11 result=0 expected=1234
new LLVM: /work/ventus-env-gh/install-review-new-simm11/bin/clang
private_frame_simm11 result=1234 expected=1234
Verification:
ninja -C /work/ventus-env-gh/llvm/build
/work/ventus-env-gh/llvm/build/bin/llvm-lit -v \
llvm/test/CodeGen/RISCV/VentusGPGPU/private-frame-simm11-offset.ll
The linked libclc final-object path exposed Ventus values that had been broadcast from scalar registers into VGPRs and later consumed by scalar-only users. When those reverse-domain COPYs survived to physical register copying, the backend printed Impossible RISCV physreg copy for VGPR -> GPR/GPRF32. The previous build script also hid this final object failure with || true. Keep the architectural invariant intact: generic VGPR -> GPR/GPRF32 COPY remains illegal. Do not add read-lane behavior, do not restore VMV_X_S/VFMV_F_S as an implicit fallback, and do not return to the old per-.bc.o libclc flow. Teach the Ventus domain helper that divergent f16 values live in VGPRs, matching the existing divergent i32/f32 treatment. Extend VentusBroadcastCopyPropagation with a narrow f16 scalar round-trip rewrite for the pattern VGPR -> scalar GPR -> FMV_H_X -> VGPR. The rewrite only fires when each intermediate has a single non-debug user and the final consumer is VGPR, so it removes a provable broadcast round trip instead of legalizing arbitrary VGPR -> scalar copies. Make build-ventus.sh fail fast for libclc final object generation by removing the stale || true guard. If linked bitcode codegen regresses again, the top-level libclc build now reports the real backend failure. Add a Ventus GPGPU f16 regression that checks a divergent private half load returns directly as a vector half load and does not lower through fmv.h.x. Validation performed: timeout 60 ninja -C /work/ventus-llvm/build llc clang; git diff --check; timeout 60 build/bin/llvm-lit -sv on the Ventus broadcast-copy-propagation, mixed-PHI, register-domain-verifier, and divergent-f16-return tests; timeout 60 build/bin/clang -target riscv32 -mcpu=ventus-gpgpu -cl-std=CL2.0 -Wno-override-module -ffunction-sections -fdata-sections -c build-libclc/builtins.link.riscv32--.bc; timeout 60 build/bin/llc -mtriple=riscv32 -mcpu=ventus-gpgpu build-libclc/builtins.link.riscv32--.bc. The linked bitcode clang and llc logs were empty and both exited 0. Follow-up planning is recorded outside this repository in /work/issues/ventus-libclc-linked-bitcode-vgpr-gpr-copy/next-phase-targets-2026-05-18.md. The next phase should move more domain decisions into SelectionDAG/TableGen selection so MachineSSA repair becomes rare while the verifier remains an invariant check.
Ventus values have two execution domains: uniform values belong in scalar GPR/GPRF registers, while divergent values belong in VGPR. The previous patterns often only checked whether the result node was divergent. When a divergent f32 operation had a uniform operand, instruction selection could satisfy scalar/vector register constraints by inserting copies or by broadcasting the scalar operand and selecting a VV instruction. That left some domain mistakes for VentusBroadcastCopyPropagation or the domain verifier to catch later. This change makes the f32 selection rules look at the operands, not just the result. New PatFrags distinguish both-divergent, divergent-with-uniform-RHS, and uniform-LHS-with-divergent-RHS forms. The f32 arithmetic and compare patterns now select VV only for two divergent operands, VF for divergent plus uniform scalar, and reverse VF forms for uniform scalar on the left. The scalar Zfinx fcvt_x/fcvt_xu patterns are also restricted to uniform nodes, and divergent f32-to-i32 conversion is selected through the vector conversion path. CopyFromReg divergence classification now queries the original IR value first when available, then falls back to live-in physical registers and finally register class. What this fixes: for the covered f32 arithmetic, compare, copysign, and conversion cases, selection produces the domain-specific instruction before ventus-broadcast-copy-propagation runs. Examples include vfadd.vf for VGPR+uniform, vfrsub.vf/vfrdiv.vf for uniform-minus/divide-VGPR, vmf*.vf for mixed-domain comparisons, and VFCVT_RTZ_X*_F_V for divergent saturating f32-to-i32 lowering. The new selection-domain test stops before the broadcast-copy propagation pass to check that these cases no longer create VGPR-to-GPR/GPRF copies. What this does not fix: this is not a global proof that selection can never create a reverse-domain copy. It only covers the pattern families updated here. Other scalar-only patterns, call/return lowering, PHI edges, load/store addressing, f16/f64 paths, or custom lowering can still need separate domain-aware selection work. It also does not remove VentusBroadcastCopyPropagation: linked libclc can still contain pure broadcast round trips such as COPY $x0 -> VGPR followed by scalar use, which are safely rewritten by that pass. The verifier remains the hard guard for any real VGPR-to-scalar value flow.
PEI finalizes the frame layout before post-PEI frame-index scavenging can insert emergency spill/reload pairs. For Ventus private memory accesses, legalizing a large private-stack offset may itself need a VGPR scratch. If the scavenger spill is placed after normal private frame objects, addressing that spill can recursively require another VGPR scratch and PEI can fail while lowering libcalls or wide conversions with live VGPR values. Reserve dedicated private-stack VGPR scavenging slots before frame layout is finalized whenever the function has private or VGPR stack objects. Track those slots in RISCVMachineFunctionInfo, lay them out at the top of the VGPR stack, and exclude them from the normal frame-object traversal so ordinary private objects keep stable offsets after the reserved words. Teach RISCVRegisterInfo::saveScavengerRegister to spill and reload VGPRs through the dedicated private slots, while leaving non-VGPR scavenging on the generic emergency-slot path. The slot tracker records spill/reload insertion intervals per basic block so non-overlapping scavenger lifetimes can reuse a slot, and overlapping lifetimes fail explicitly instead of silently clobbering a live spill. Also tag VGPR stack slots consistently from both storeRegToStackSlot and loadRegFromStackSlot, update Ventus CodeGen expectations for the reserved private words, and add libcall stress coverage for VGPR scavenging with aligned and large private frames. Tested with: - timeout 60s build/bin/llvm-lit -sv llvm/test/CodeGen/RISCV/VentusGPGPU/private-frame-simm11-offset.ll llvm/test/CodeGen/RISCV/VentusGPGPU/frameindex-large-offset-no-clobber.mir - timeout 60s build/bin/llvm-lit -sv llvm/test/CodeGen/RISCV/VentusGPGPU/addr-space.ll llvm/test/CodeGen/RISCV/VentusGPGPU/addr-space2.ll llvm/test/CodeGen/RISCV/VentusGPGPU/builtin-noverify.ll llvm/test/CodeGen/RISCV/VentusGPGPU/builtin.ll llvm/test/CodeGen/RISCV/VentusGPGPU/eliminate_call_frame.ll llvm/test/CodeGen/RISCV/VentusGPGPU/inline-asm-branch.ll llvm/test/CodeGen/RISCV/VentusGPGPU/local-static-lds.ll llvm/test/CodeGen/RISCV/VentusGPGPU/local_addressed_variables.ll llvm/test/CodeGen/RISCV/VentusGPGPU/local_addressed_variables_multifunc.ll llvm/test/CodeGen/RISCV/VentusGPGPU/localmem-large-offset.ll llvm/test/CodeGen/RISCV/VentusGPGPU/parameter-vector-struct-types.ll llvm/test/CodeGen/RISCV/VentusGPGPU/regexti.ll llvm/test/CodeGen/RISCV/VentusGPGPU/var-arg.ll llvm/test/CodeGen/RISCV/VentusGPGPU/vbranch-join.ll llvm/test/CodeGen/RISCV/VentusGPGPU/vgpr-scavenging-libcall.ll - rm -r llvm2/build-libclc && LLVM_DIR=$PWD/llvm2 bash build-ventus.sh --build "llvm;libclc"
Backport upstream fix for dropping color vectors when StackSlotColoring sees a lower non-zero StackID after a higher StackID. Upstream commit: 40edb37bb370541b71657fc308ba26b1ec2deb16 Upstream URL: llvm/llvm-project@40edb37
Lower the exact OpenCL wait_group_events(int, event_t *) external call to the Ventus local+global barrier intrinsic while preserving fatal diagnostics for other private-derived generic pointer escapes. Fixes the basic async copy CTS aborts for async_strided_copy_global_to_local, async_strided_copy_local_to_global, async_copy_global_to_local, and async_copy_local_to_global. Verified with llvm-lit generic-as-specialization.ll and generic-as-specialization-errors.ll, including the near-miss negative matcher test, plus all four CTS targets through the temporary install-async-copy-worker prefix.
Ventus encodes non-kernel stack argument locations as negative offsets from the incoming private stack pointer. CCValAssign stores memory locations in an unsigned field, so these downward-growing offsets are represented as wrapped 32-bit values. unpackFromMemLoc previously created fixed stack objects with offset 0, which made split stack arguments alias each other. Changing that directly to VA.getLocMemOffset() preserved distinct slots but recorded values such as -4 as 4294967292 in MachineFrameInfo because CreateFixedObject takes a signed int64_t offset. Recover the signed Ventus ABI offset before creating fixed stack objects, and use the same conversion for the RV32 soft-f64 stack path. Add a regression test where a leading double16 consumes V0-V31 and the following i64 must be loaded from two distinct stack slots at -4 and 0.
RegAllocFast reloads live-through virtual registers when a physical register is clobbered by an instruction. For calls, the previous insertion point was immediately after the call instruction. That can still be before the target's call-frame destroy pseudo, so the reload observes the temporary outgoing-call frame rather than the caller frame restored after the call sequence. This matters for Ventus because the private stack is based on tp and the ADJCALLSTACKUP expansion adjusts tp back after the call before resynchronizing the private-frame VGPR base. Reloading a VGPR value from a private-stack spill before that pseudo uses the callee-call-frame tp value and can make the restored VGPR base stale for the following private-memory accesses. Teach RegAllocFast to choose a call-aware reload insertion point: after a call, skip debug instructions and insert after a following call-frame destroy pseudo. Non-call displacement keeps the original next-instruction insertion point. Add a Ventus regression test with repeated calls to get_el where the live vector base must be reloaded only after each tp restore, checking that the generated sequence is call, tp restore, regext, then vmv.v.x.
.bss / .data 的初始化仅在主机端驱动程序加载 ELF 时完成一次，
且仅限主机端能访问到的 DRAM / global memory 不能初始化 LDS / smem
设备端 _start 不做任何内存初始化
警告：这并非完整修复
现在 ventus llvm 能正确识别区分 static LDS 与 stack-style LDS 内存块
编译器会将 static LDS 当作全局变量处理并放到 .bss 段
这就暴露了设备端 _start 中错误地初始化 bss 的问题，造成运行缓慢
本修复能移除 bss 初始化修复运行缓慢，但 static LDS 仍在 .bss 中
这能造成的潜在不良影响：
1. 误导：占用普通内存 DRAM 而非 LDS
2. 误导：符号地址问题，使用 CSR_LDS + offset 而非 bss 中记录的地址
3. 误导：只对专用重定位正确
.ventus.resobj 能修正的是 Ventus 专用 LDS 重定位：
R_RISCV_VENTUS_LDS_HI20
R_RISCV_VENTUS_LDS_LO12_*
但 ELF 里其他机制仍会按普通 .bss symbol 理解它
比如普通符号查询、debug/map、链接脚本符号、可能的非专用重定位
把它留在 .bss 会让“看起来正确”和“实际访问方式”不一致。
4. 误导：资源统计会重复或混乱
实际 LDS 占用由 .ventus.resobj / lld finalizer 计算。
如果同一个对象还放在 .bss，ELF 普通数据区也会把它算进去。这不是实际硬件资源模型。
可能的正确解决方案：将 LDS 符号移入新的 ELF section 中，或合并到 .ventus.resobj
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Title
[Ventus] Add SFU/MMA lowering and resource metadata updates
Summary
This PR merges the
dev-thu-sfu-mmabranch intomain.It brings in the latest Ventus compiler backend updates, including SFU/FPU builtin support, MMA intrinsic/lowering support, resource metadata improvements, and several correctness fixes around SGPR lifetime, resource accounting, dynamic local memory arguments, and large frame-index offsets.
Main changes
Ventus SFU/FPU and shuffle builtins
Ventus MMA support
Resource metadata and linking
.ventus.resobjemission and lld-side finalization into.ventus.resource.<kernel>sections.Correctness fixes
CSR_LDSoffsets.Build improvements
LLVM_ENABLE_ASSERTIONSLLVM_ENABLE_EXPENSIVE_CHECKSCLANG_TOOLING_BUILD_AST_INTROSPECTIONTesting
The branch adds and updates tests under:
clang/test/CodeGenOpenCLclang/test/CodeGen/RISCVllvm/test/CodeGen/RISCVllvm/test/MC/RISCVlld/test/ELFReview focus
I would especially appreciate review on: