Update 2026-08-22 - #32
Merged
Merged
Conversation
Validate affine.max and affine.min operands before inlining their maps into affine.for bounds. This defers nested loops until enclosing induction variables become valid affine dimensions and keeps the exported pattern independent of greedy traversal order. Assisted-by: Codex
…7070) - `rintf128` - `lrintf128` - `llrintf128`
lit.util.memoize builds a (args, tuple(kwargs.items())) key and does two dict lookups per call. functools.lru_cache has a faster path that uses a single positional argument directly as the key, so it does one lookup instead of two and skips the tuple construction. Profiling lit with Scalene and comparing the two implementations on test suites showed the lru_cache version is faster for _caching_re_compile. runCommandCached takes multiple positional and keyword arguments, so it does not hit lru_cache's single-argument fast path, but the swap removes one more caller of the hand-rolled memoize implementation. memoize itself is now unused and will be removed in a follow-up.
Adds a newPM pass for CFIInstrInserter (cfi-instr-inserter). - Extracts the pass's working state (MBBVector, CSRLocMap) and logic into a CFIInstrInserterImpl class with a run method, called by both the legacy pass and the new pass manager pass. - Renames the old pass with the "Legacy" suffix. - Adds the new pass manager pass CFIInstrInserterPass, using RequiredPassInfoMixin: the legacy pass's runOnMachineFunction never calls skipFunction, so it always runs unconditionally and should not be skippable in the new PM either. run() unconditionally returns PreservedAnalyses::all(), matching the legacy pass's own AU.setPreservesAll() declaration -- the same shape CFIFixupPass (an already-ported sibling CFI pass) already uses. - Updates MachinePassRegistry.def, PassBuilder, and CodeGenPassBuilder. - Wires the pass into X86's and RISC-V's newPM pipelines, matching their existing legacy-PM gating conditions: - X86 replaces an existing TODO inside an already-correct conditional in addPreEmitPass2. - RISC-V's equivalent TODO had no gating at all. The legacy pipeline gates this pass behind -riscv-enable-cfi-instr-inserter, a flag private to RISCVTargetMachine.cpp and not reachable from RISCVCodeGenPassBuilder.cpp. Uses !TM.Options.EnableCFIFixup instead, which RISCVTargetMachine's constructor already sets to the exact inverse of that flag, avoiding any new cross-TU plumbing. - Adds -passes=cfi-instr-inserter RUN lines to all nine existing dedicated unit tests (eight X86 .mir, one RISC-V .mir), and updates llc-pipeline-npm.ll's O0/O2 X86 pipeline-dump expectations (verified empirically which triples actually enable the pass). Assisted-by: Claude Sonnet 5
…ount (llvm#217193) The SizeEmitter lambda in emitCommonOMPTargetDirective constructed OMPLoopScope as an unnamed temporary, so it was destroyed at the end of its own statement rather than at the end of the enclosing block. OMPLoopScope is a RunCleanupsScope, and destroying it emits the cleanups for the '.capture_expr.' variables that hold the loop bounds, including their llvm.lifetime.end calls. Those cleanups therefore ran before EmitScalarExpr(D.getNumIterations()) loaded the bounds, so Clang emitted a load from an alloca whose lifetime had already ended: store i32 %sub, ptr %.capture_expr.1 call void @llvm.lifetime.end.p0(ptr %.capture_expr.1) %0 = load i32, ptr %.capture_expr.1 ; read after lifetime ended This was harmless while PromoteMemToReg simply discarded lifetime intrinsics. Since llvm#191909, promotion inserts a 'store undef' in place of each marker, so the load now correctly resolves to undef and the trip count folds to zero. KernelArgsTy::Tripcount then reaches the runtime as 0, which can no longer size the launch: a 'target teams distribute parallel for' over 100000 iterations with OMP_NUM_TEAMS=50 launches one team instead of 50. Both the misplaced cleanup and llvm#191909 are present on main, so this currently breaks OpenMP target offloading at head: every target loop directive whose launch geometry or scheduling depends on the trip count is affected. It was caught by the AOMP smoke suite, where it took out all of the trip-count-sensitive tests at once. Name the scope so that it outlives the trip count computation, matching every other use of OMPLoopScope in this file. cc @isoard-amd @ronlieb
…path (llvm#216528) In `getTypeName()` the check for the template parameter is in the wrong place: ```cpp LLVM_GET_TYPE_NAME_CONSTEXPR std::string_view TemplateParamsStart = Name.substr(Name.find(Key)); static_assert(!TemplateParamsStart.empty(), "Unable to find the template parameter!"); ``` If `find` fails it returns `npos`, and `substr(npos)` throws `out_of_range` — so in the constexpr configuration the compile has already failed on the line above the assertion, and the message never appears. Checking the position before using it fixes all of that, and makes the check the one the message describes: ```cpp LLVM_GET_TYPE_NAME_CONSTEXPR std::string_view::size_type KeyPos = Name.find(Key); static_assert(KeyPos != std::string_view::npos, ...); LLVM_GET_TYPE_NAME_CONSTEXPR std::string_view TemplateParamsStart = Name.substr(KeyPos); ``` On Clang, GCC 9+, and MSVC 19.10+, the checks are all done at compile time. On versions of GCC < 9, the runtime asserts will be executed, but only if NDEBUG isn't defined. This is consistent with previous behavior. ## Test plan - GCC path, `static_assert` configuration: a TU instantiating `llvm::getTypeName<T>()` compiles and returns the expected name. - GCC path forced onto the `assert` configuration: compiles, same result. - The MSVC branch is unreachable on Linux, so I exercised it by preprocessor simulation, forcing that branch with `__FUNCSIG__` defined to a representative signature (`class llvm::StringRef __cdecl llvm::getTypeName<struct Foo>(void)`). Both configurations compile and return `Foo`. - With the key deliberately absent from the signature, the first diagnostic is now `error: static assertion failed: Unable to find the template parameter!` On the `assert` configuration it now aborts with that message instead of throwing `out_of_range`. Co-authored-by: Cursor <cursoragent@cursor.com>
…7783) Compare the generated properties structs instead of listing each inherent attribute individually. This preserves the original equality semantics when new properties are added. Follow-up to llvm#217233 Assisted-by: Codex
Rename the createPass function to have a legacy suffix, and also move the class definition for the NewPM version into RISCV.h to be consistent with other passes.
…lvm#217692) Add builtins and header wrappers for `__riscv_pssha`, `__riscv_psshar`, `__riscv_psshl` and `__riscv_psshlr`.
This operator is only used in DEBUG builds. Exposed by 732d841 which moved the Impl class into an anonymous namespace in the source file which allowed the compiler to reason about usage. Fixes https://lab.llvm.org/buildbot/#/builders/228/builds/8620
…ttr (llvm#217432) Adds support to the LLVM dialect `DICompileUnitAttr` for DWARF v6 source language name and version . `DICompileUnitAttr`'s `sourceLanugage` field now points to a `DISourceLanguageNameAttr`, which contains all source language information, including dialect, name, and version. This is done to map as closely as possible to `llvm::DISourceLanguageName`. `DISourceLanguageNameAttr` has parameters for: - An DW_LANG_* language. - A DWARF v6 DW_LNAME_* name, which can only be set if the language is not. - An optional language-dependent version. - An optional target-specific language dialect. Parsers and printers are added so if only a language is passed, it is parsed and printed in as before this change. Assistance from codex was used in this PR.
resolves llvm#216103 This change turns on VaryingLongVector for all HLSL intrinsics that map to an elementwise builtin. assisted by GPT 5.6-Sol via Copilot
…ver (llvm#217492) A direct `clang --target=spirv64-unknown-unknown --sycl-link a.bc` does not tell clang-sycl-linker which device it is finalizing for. Derive the triple from the target the driver was given, and the architecture from -march=, and pass them as -triple=/-arch=. An absent -march= means no specific device was requested. That is spelled as an absent -arch=, matching how clang-linker-wrapper renders an offload image that names no device. A caller may also name these via -Xlinker/-Wl, or (through clang-linker-wrapper) via -Xoffload-linker. clang-sycl-linker keeps the last value of each, so derived values are added before an explicitly given values. With the driver deriving both values, the SYCL-specific forwarding in clang-linker-wrapper becomes redundant - it spelled the triple and the architecture out through -Xlinker after having already passed the very same strings as --target= and -march= - so drop it. co-authored by claude
Keep DISubprogram linkage names when stripping non-line-table debug info if the compile unit has debugInfoForProfiling set. This mirrors -gline-tables-only, which keeps linkage names when -fdebug-info-for-profiling is enabled.
I was running into a Rust miscompilation using pretty standard code (see bugreport). I had Claude rootcause this to an LLVM bug, where SLL (on 64 bit) is marked as "move", even though it isn't - for full 64-bit values (with arbitrary upper 32 bits), SLL is not a move as the upper bits get just sign extended. The fix itself is tiny and makes sense to me (I know MIPS very well, but have little LLVM expertise myself). The test for this however is pretty convoluted as it is quite hard to trigger this bug reliably - it needs quite some register pressure to actually happen. On the llvm/test/CodeGen/Mips/madd-msub.ll change: Register indices changed here as the is the register allocator is now picking a new independent register instead of reusing an existing one (`sll $4, $4, 0` changes to `sll $1, $4, 0`). In this particular example that causes instruction counts to be the same, though in other examples a small increase is possible (cost of correctness). Fixes llvm#213419 Assisted-by: Claude Code (Opus 5)
When CIR's LoweringPrepare pass created the runtime variable for `__dso_handle` while emitting a guarded dtor region, it was creating the handle using the builtin `i8` type rather than the CIR equivalent. This caused the `__cxa_atexit` function to be created with a pointer to that type, which causes problems for the CIR calling convention lowering pass. This change updates the pass to use the CIR type.
…uctions (llvm#217316) A reductions final value can be simplified away to a constant, which llvm#201023 handled in replaceWithFinalIfReductionStore. However constant live ins can have other users, e.g. a different reductions start value, in which case the assertion will fail because we're comparing the backedge value of a different reduction entirely. Relax the assertion to allow any VPIRValue instead This an alternative to llvm#217306 that fixes llvm#215071
…17382) Example: ```fortran !$acc parallel loop worker reduction(+:s) num_workers(16) vector_length(8) ``` In this code, a ThreadY reduction forced blockDim.x to a warp, and blockDim.y was divided by the same factor to hold the thread count, so the launch ran as (32, 4, 1): 4 workers instead of 16, with the folded-away workers turned into ThreadX lanes redoing each other's work. Fix: skip the alignment when only the workers ask for it and the row width divides the subgroup size. Worker reductions combine their partials in the lowest blockDim.y threads of the block, so a row may sit anywhere inside a subgroup, and a per-row barrier is lane-masked to cover exactly that row. Thread-level array accumulates keep the padding: they combine through atomics on a shared array rather than worker shuffles.
…#190088) This first improves the structure of the compiler-rt BUILD.bazel, fixing bugs and exposing more carefully arranged source files. It also exposes compilation info for builtins and CRT files for use in compiling these source files.
…vm#217623) ## Summary The HowToCrossCompileLLVM option list spelled the CMake variable as `CMAKE_CROSSCOMPIILING` (extra I). The linked CMake docs use `CMAKE_CROSSCOMPILING`. Assisted-by: Grok (xAI)
This is to clarify which functions are creating legacy passes rather than NewPM passes.
## Motivation This PR is motivated by ongoing work of providing a fully hermetic clang-cl bazel toolchain built for windows with the MSVC runtime. The LLVM Bazel overlay does not select a Windows-compatible BLAKE3 assembly source set for an x86_64 Windows target compiled with the rules_cc `clang-cl` compiler dialect. The `windows_gnu.S` name can be misleading here. "GNU" describes the assembly syntax accepted by GNU-style assemblers and Clang's integrated assembler; it does **not** mean that the implementation uses the MinGW ABI or runtime. These files implement the Microsoft x64 calling convention and are intended to produce Windows COFF objects. A hermetic clang-cl toolchain can therefore assemble the GNU-syntax Windows sources directly. Selecting the MASM-syntax `windows_msvc.asm` files instead would require a separately declared MASM tool such as `ml64.exe`, which this compiler route does not provide. LLVM's CMake build makes the same distinction indirectly: its `MSVC` branch first enables the separate `ASM_MASM` language and then selects `windows_msvc.asm`; its other Windows assembler branch selects `windows_gnu.S`. The MASM choice is therefore contingent on an available MASM assembler, not inherent to every compiler targeting the MSVC ABI. ## Tests - Compiled each selected `windows_gnu.S` file with LLVM clang-cl 22.1.8 targeting `x86_64-pc-windows-msvc`; all four outputs were AMD64 COFF objects and defined the expected SSE2, SSE4.1, AVX2, and AVX-512 BLAKE3 symbols. - Inspected the selected assembly: it uses the Microsoft x64 argument registers and preserves the required Windows-nonvolatile general-purpose and XMM registers. - Confirmed the corresponding `windows_msvc.asm` inputs are not accepted by this clang-cl-only route. They require a separately declared MASM assembler. - A downstream full x86_64 Windows LLVM build linked successfully with this source selection. The identical baseline reached the final link with undefined BLAKE3 SIMD dispatch symbols. Assisted by: codex
…217792) Introduce Proxy<RetT(ArgTs...)>, a typed handle for invoking a controller-side operation from the executor. A Proxy abstracts over how a call reaches the controller: it holds an opaque callee tag and a dispatch function (supplied by a per-protocol spec) and forwards calls through the Session. This is a cut-down port of llvm/ExecutionEngine/Orc/Proxy.h, with two deliberate differences: - The callee is identified by an opaque tag (const void *, typically the address of a controller-side global) rather than an ExecutorAddr, since the executor->controller direction dispatches by tag. - Only the asynchronous (OnComplete) call operator is provided. The blocking convenience operator is omitted: the executor may be single-threaded or freestanding and cannot rely on std::promise/future or on blocking a dispatch thread. Add ProxyTest covering default construction, the void->Error and T->Expected<T> return mappings, and argument/tag forwarding.
…m#207146) With the +optimized-nfX-segment-load-store tuning flag, we cost a segmented store as a single wide load + some shuffle ops. However for e.g. a `<vscale x 5 x i64>` Factor=5 segmented load, a wide `<vscale x 5 x i64>` load gets costed as a full LMUL 8 load. From what I can see on https://camel-cdr.github.io/rvv-bench-results/spacemit_x100/index.html and on my own measurements on the spacemit-x60, uarchs likely don't do a full LMUL 8 load under the hood and instead dispatch the minimum number of DLEN sized ops needed for the full segment. This changes the wide load cost to be divideCeil(vector size, DLEN) ops so we don't overcost it. Whilst we're here, this also removes the LT.first legalization multiplier. We're computing the cost in terms of the unlegalized type so we shouldn't be scaling it by the legalization cost.
…gnment granule", which is applied to both the alignment and the allocated size. (llvm#203872) On CHERI targets, bounds are generally stored in a compressed format which imposes alignment requirements based on allocation size. When emitting globals, we need to align and tail-pad them as appropriate to ensure that the pointer to the global will have bounds that are precise, i.e. not covering any non-padding bytes that either precede or follow While this patch adds the plumbing necessary to support this feature on arbitrary targets, it only concretely implements it for RISCV Y-base and XCheriot. Support for other targets (such as AArch64/MTE and AArch64/Morello) is left as future work here. Based on llvm#121957 Co-authored-by: Florian Mayer <fmayer@google.com> Co-authored-by: Florian Mayer <fmayer@google.com>
It's not possible to compute an allocation granule for an unsized type.
Assisted-by: AI
…orms (llvm#217655) `wmemchr` on Apple platforms is incredibly slow. We can do significantly better by just implementing it on our own. Fixes llvm#205840
clang-repl cannot currently emit the LLVM IR it generates. The command-line emit actions are accepted but folded into `EmitLLVMOnlyAction` (producing no output), and `frontend::EmitLLVM` is rejected by the incremental frontend action when `-emit-llvm` is passed. This prevents users from inspecting IR produced per-input PTU. This patch teaches clang-repl to honor `-emit-llvm` by allowing the IncrementalAction to accept `frontend::EmitLLVM` , making the driver print each input's `llvm::Modul` to stdout instead of executing it, similar to `clang -emit-llvm`. This enables a lit-level test to FileCheck the IR produced by clang-repl and the effect of a statement on it. Exposing this functionality beyond the driver level to allow library consumers of `clang::Interpreter` to use it is a planned follow-up.
Pretty-print MSVC STL `std::valarray` from `_Myptr` / `_Mysize`. The size summary is emitted only when that layout is present so libstdc++ valarrays (same `std::` name, no inline namespace) are not stolen. References are dereferenced and the element type is taken from `_Myptr`. Tests: generic valarray suite's MSVC STL category (Windows), covering empty arrays, references, exact synthetic children, and the libstdc++ dispatch guard. Part of llvm#24834 Assisted-by: Grok 4.6 Assisted-by: codex-5.6-high --------- Co-authored-by: Bjorn Schobben <bjorn.schobben@aimsport.com>
Extends InferAlignment to infer and strengthen source/destination alignment on memcpy, memmove, memset, and related memory intrinsics. These changes were assisted by an LLM
…r" (llvm#218119) Reverts llvm#217727 due to buildbot failures.
Reviewers: Pull Request: llvm#218122
Summarize MSVC STL `std::source_location` from `_File` / `_Function` / `_Line` / `_Column`. LoadCommon last-match-wins, so this dispatches on the MSVC layout and otherwise keeps the existing libstdc++ formatter. A default-constructed location is left without a summary, and implementation children remain hidden. Tests: generic source_location suite's MSVC STL category (Windows), including child-hiding checks. Part of llvm#24834 Assisted-by: Grok 4.6 Assisted-by: codex-5.6-high --------- Co-authored-by: Bjorn Schobben <bjorn.schobben@aimsport.com>
This PR adds a LNT machine configuration that runs with Fast hardening on macOS. This is done by making it possible to pass a CMake cache to use for building the library to `run-benchmarks`. The obvious caveat is that we can't pass parameters to the test suite itself. In the future, I think the cleanest way to do that would be to make the test suite independent from the libc++ build and give it its own configuration options. We could then pass one CMake cache for the build, and one CMake cache for configuring the test suite. However, since this requires a lot more work and the hardening mode is baked into the library once configured, the current patch is sufficient to achieve our goals. Fixes llvm#218028
These cases were introduced before the poison value: llvm@59b6b7d.
…e-exception-escape (llvm#218067) `ExceptionAnalyzer` represents exceptions of unknown type with a null `Type`, but some consumers dereferenced it unconditionally, resulting in a crash when `TreatFunctionsWithoutSpecificationAsThrowing` is enabled. This commit fixes the problem by skipping type-dependent processing for unknown exceptions. Fixes llvm#217649
howManyGreaterThans computes the backedge-taken count as ((Start - End) + (Stride - 1)) /u Stride. The addition can overflow, causing incorrect results. Instead, use getUDivCeilSCEV if Start >= End, mirroring howManyLessThans. It also has additional handling for stride being a power of 2, which allows using the simpler formula in more cases. I'll check if we can unify the code (as implied by the FIXME I think), instead of duplicating more logic. Fixes llvm#217537. Fixes llvm#187472. PR: llvm#217744
The size in a dereferenceable bundle may have a different bitwidth than the access size. Both are interpreted as unsigned values. Update WidestTy to account for possibility of the type for the bundle being narrower. Fixes a crash similar to llvm#217770. PR: llvm#217939
… `#pragma omp declare variant` is followed by another OpenMP declarative directive containing a qualified identifier (llvm#217875) For code ```cpp void foo(); #pragma omp declare simd #pragma omp declare target to(foo) ``` Clang currently accepts this without rejection. The underlying cause is that OpenMP pragma parsing for directives like `declare target to(...)` performs name lookup without advancing the source location to create a new declaration. As a result, the parser fetches the existing Decl of foo and passes it up to declare simd, silently bypassing Sema diagnostics. https://godbolt.org/z/szPne5n45 --- However, when using a qualified name in a namespace: ```cpp namespace N { void foo(); } #pragma omp declare simd #pragma omp declare target to(N::foo) ``` This inherently invalid syntax causes `ActOnReenterFunctionContext` to an assertion failure because the `DeclContext` of the looked-up `N::foo` does not match the parser's current lexical context (`TranslationUnitDecl`). --- This patch adds a source location check in the parser for declare simd and declare variant to verify whether the returned Decl is a newly parsed declaration at the current position. If no new declaration was created, the parser intercepts it directly and emits `err_omp_decl_in_declare_simd_variant`. Newly parsed declarations (even non-function ones like int a;) are still passed through to preserve existing Sema diagnostics. Fixed llvm#217204 --------- Co-authored-by: Alexey Bataev <a.bataev@outlook.com>
Price the unfused fmul without a context instruction. Targets that model fma fusion price a contractable fmul as free, which discounted the scalar side of the comparison too and fmuladd never looked profitable.
…ile name. (llvm#205729) The `S_OBJNAME` field value gets by bypassing `CGDebugInfo` and its path prefix does not get remapped if requested `fdebug-prefix-map=` option as the other pathes in the debug info. This patch fixes it and does remapping for the object file path either.
…ption. NFC (llvm#207083) The `clang-cl -help` command show the help string for `/experimental:deterministic` options with the default metavar name value that should not be: /experimental:deterministic<value> This patch fixes this output and omits the `<value>` part for the option.
…signated initializers (llvm#215934) The problem is that we delete the necessary comma whenever we use implicit initializer lists. How we solve this is that whenever we see an implicit initializer list, we do not match ``InitListExpr`` nodes at all, so that we will not delete the necessary comma. Why we chose this path is detailed in Alternatives considered. This produces a fix that breaks valid code (llvm#214087) and one that never converges (llvm#214086). <details> <summary><b>Alternatives considered</b></summary> ### Why not repair the source ranges instead The synthesized nodes also carry misleading locations - their range is a snapshot of the designator that caused them to be created, so it need not cover their own children. The anonymous-struct node in llvm#214087 reports `9:5-9:10` while holding an initializer on line 10, which is why it measures as single-line. Repairing that would be a change to Clang rather than to the check, and it is not clear there is anything to repair: the semantic form exists to record which initializer belongs to which subobject, and nothing in Clang relies on its ranges bracketing their children. Only a source-rewriting tool needs that guarantee, and such a tool should not be inspecting nodes that were never written. Skipping them is both smaller and better scoped. ### Why not match on brace locations `getLBraceLoc()`/`getRBraceLoc()` are not a reliable discriminator. On the synthesized anonymous-struct node in llvm#214087 both are *valid*, pointing at the `.` and at `a`. A validity check would suppress only one of the two false positives in llvm#214087 and none of llvm#214086. This was in fact the previous implementation of `isExplicit()`, replaced in llvm#195175 for the same reason. </details> ## AI disclosure What Claude did - Explored the codebase to locate the check and the relevant Sema/AST machinery - Instrumented the check with temporary debug output and measured the actual InitListExpr properties on the reproducers (brace locations, isExplicit, computed policies) - Wrote the final two-line patch and all the added test cases - Found the isExplicit() history (PR llvm#195175) that the fix depends on What I did - Directed the approach and interrogated the reasoning at each step - Reviewed and verified the results locally Fixes: llvm#214086 Fixes: llvm#214087
…efined ops (llvm#216773) This PR adds standalone effect and speculatability specifiers for Python-defined operations. `NoMemoryEffect` and `AlwaysSpeculatable`, previously nested under `Pure`, are now public, and `RecursivelySpeculatable` is added. `Pure` remains a shorthand for attaching `NoMemoryEffect` and `AlwaysSpeculatable`. This also exposes `OpTrait::HasRecursiveMemoryEffects` through the C API and Python bindings as `ir.RecursiveMemoryEffectsTrait`, allowing region-bearing Python-defined operations to derive their memory effects from nested operations: ```python class LeafOp( TestDialect.Operation, name="leaf", traits=[NoMemoryEffect, AlwaysSpeculatable], ): pass class RegionOp( TestDialect.Operation, name="region", traits=[ ir.NoTerminatorTrait, ir.RecursiveMemoryEffectsTrait, RecursivelySpeculatable, ], ): body: Region ``` Tests cover direct interface queries and validate recursive memory effects with CSE and trivial DCE. Related to llvm#177735 and llvm#195505. Assisted-by: Codex / GPT-5.6 Sol
…u>= Y (llvm#216072) This teaches ValueTracking that `X` and `X urem Y` are non-equal when a dominating condition implies `X u>= Y`. It also handles cases where `X u>= Y` can be proven structurally, such as when `X` is an `add nuw` of `A` and `Y`. For a defined `urem`, let `R = X urem Y`. Then: ```text R == X <=> X u< Y R != X <=> X u>= Y ``` This allows InstSimplify's existing nonzero reasoning to simplify comparisons equivalent to `(X - R) != 0`, including: - `(X - R) u>= 1` and `(X - R) u< 1` - `(X - R) u> 0` and `(X - R) u<= 0` - `(X - R) != 0` and `(X - R) == 0` It also allows `umax(X - R, 1)` to simplify to `X - R` when `X u>= Y` is known. Alive2: https://alive2.llvm.org/ce/z/H83-ra Fixes llvm#216071
) Add lowering of `llvm.dx.resource.handlefromheap` intrinsic. It gets translated to DXIL ops `createHandleFromHeap` and `annotateHandle`. For example, the intrinsic call ```llvm %typed = call target("dx.TypedBuffer", <4 x float>, 1, 0, 0) @llvm.dx.resource.handlefromheap.tdx.TypedBuffer_v4f32_1_0_0(i32 3) ``` will lower to ```llvm %0 = call %dx.types.Handle @dx.op.createHandleFromHeap(i32 218, i32 3, i1 false, i1 false) %1 = call %dx.types.Handle @dx.op.annotateHandle(i32 216, %dx.types.Handle %0, %dx.types.ResourceProperties { i32 4106, i32 1033 }) ``` The value of non-uniform index flag is calculated using existing `hasNonUniformIndex` helper function. Fixes llvm#91406
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
20 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.
Another release, another update