[all-commits] [llvm/llvm-project] 369891: [NVPTX] use untyped loads and stores where ever po...

Sat May 10 08:26:47 PDT 2025

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 369891b6747e4ad4b5e4e6d06f3f7596f3ee3f02
      https://github.com/llvm/llvm-project/commit/369891b6747e4ad4b5e4e6d06f3f7596f3ee3f02
  Author: Alex MacLean <amaclean at nvidia.com>
  Date:   2025-05-10 (Sat, 10 May 2025)

  Changed paths:
    M clang/test/CodeGenCUDA/bf16.cu
    M clang/test/CodeGenCUDA/fp-contract.cu
    M clang/test/CodeGenCUDA/memcpy-libcall.cu
    M llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
    M llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
    M llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
    M llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
    M llvm/test/CodeGen/NVPTX/MachineSink-call.ll
    M llvm/test/CodeGen/NVPTX/MachineSink-convergent.ll
    M llvm/test/CodeGen/NVPTX/access-non-generic.ll
    M llvm/test/CodeGen/NVPTX/addr-mode.ll
    M llvm/test/CodeGen/NVPTX/addrspacecast-folding.ll
    M llvm/test/CodeGen/NVPTX/addrspacecast-ptx64.ll
    M llvm/test/CodeGen/NVPTX/addrspacecast.ll
    M llvm/test/CodeGen/NVPTX/aggregate-return.ll
    M llvm/test/CodeGen/NVPTX/and-or-setcc.ll
    M llvm/test/CodeGen/NVPTX/anonymous-fn-param.ll
    M llvm/test/CodeGen/NVPTX/applypriority.ll
    M llvm/test/CodeGen/NVPTX/atomics-sm70.ll
    M llvm/test/CodeGen/NVPTX/atomics-sm90.ll
    M llvm/test/CodeGen/NVPTX/atomics.ll
    M llvm/test/CodeGen/NVPTX/barrier.ll
    M llvm/test/CodeGen/NVPTX/bf16-instructions.ll
    M llvm/test/CodeGen/NVPTX/bf16.ll
    M llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll
    M llvm/test/CodeGen/NVPTX/bfe.ll
    M llvm/test/CodeGen/NVPTX/bmsk.ll
    M llvm/test/CodeGen/NVPTX/bswap.ll
    M llvm/test/CodeGen/NVPTX/bug21465.ll
    M llvm/test/CodeGen/NVPTX/bug22246.ll
    M llvm/test/CodeGen/NVPTX/bug26185-2.ll
    M llvm/test/CodeGen/NVPTX/bug26185.ll
    M llvm/test/CodeGen/NVPTX/byval-const-global.ll
    M llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
    M llvm/test/CodeGen/NVPTX/chain-different-as.ll
    M llvm/test/CodeGen/NVPTX/cmpxchg-sm60.ll
    M llvm/test/CodeGen/NVPTX/cmpxchg-sm70.ll
    M llvm/test/CodeGen/NVPTX/cmpxchg-sm90.ll
    M llvm/test/CodeGen/NVPTX/cmpxchg.ll
    M llvm/test/CodeGen/NVPTX/combine-mad.ll
    M llvm/test/CodeGen/NVPTX/convert-fp-i8.ll
    M llvm/test/CodeGen/NVPTX/convert-int-sm20.ll
    M llvm/test/CodeGen/NVPTX/convert-sm100.ll
    M llvm/test/CodeGen/NVPTX/convert-sm100a.ll
    M llvm/test/CodeGen/NVPTX/convert-sm80.ll
    M llvm/test/CodeGen/NVPTX/convert-sm90.ll
    M llvm/test/CodeGen/NVPTX/copysign.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-prefetch.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk.ll
    M llvm/test/CodeGen/NVPTX/ctlz.ll
    M llvm/test/CodeGen/NVPTX/dag-cse.ll
    M llvm/test/CodeGen/NVPTX/demote-vars.ll
    M llvm/test/CodeGen/NVPTX/discard.ll
    M llvm/test/CodeGen/NVPTX/disjoint-or-addr.ll
    M llvm/test/CodeGen/NVPTX/distributed-shared-cluster.ll
    M llvm/test/CodeGen/NVPTX/div.ll
    M llvm/test/CodeGen/NVPTX/dot-product.ll
    M llvm/test/CodeGen/NVPTX/dynamic-stackalloc-regression.ll
    M llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll
    M llvm/test/CodeGen/NVPTX/elect.ll
    M llvm/test/CodeGen/NVPTX/extloadv.ll
    M llvm/test/CodeGen/NVPTX/extractelement.ll
    M llvm/test/CodeGen/NVPTX/f16-instructions.ll
    M llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
    M llvm/test/CodeGen/NVPTX/f32-ex2.ll
    M llvm/test/CodeGen/NVPTX/f32-lg2.ll
    M llvm/test/CodeGen/NVPTX/fabs-intrinsics.ll
    M llvm/test/CodeGen/NVPTX/fexp2.ll
    M llvm/test/CodeGen/NVPTX/flo.ll
    M llvm/test/CodeGen/NVPTX/flog2.ll
    M llvm/test/CodeGen/NVPTX/fma-relu-contract.ll
    M llvm/test/CodeGen/NVPTX/fma-relu-fma-intrinsic.ll
    M llvm/test/CodeGen/NVPTX/fma-relu-instruction-flag.ll
    M llvm/test/CodeGen/NVPTX/fns.ll
    M llvm/test/CodeGen/NVPTX/forward-ld-param.ll
    M llvm/test/CodeGen/NVPTX/fp-contract.ll
    M llvm/test/CodeGen/NVPTX/fp128-storage-type.ll
    M llvm/test/CodeGen/NVPTX/frem.ll
    M llvm/test/CodeGen/NVPTX/funnel-shift-clamp.ll
    M llvm/test/CodeGen/NVPTX/generic-to-nvvm.ll
    M llvm/test/CodeGen/NVPTX/globals_lowering.ll
    M llvm/test/CodeGen/NVPTX/half.ll
    M llvm/test/CodeGen/NVPTX/i1-ext-load.ll
    M llvm/test/CodeGen/NVPTX/i1-icmp.ll
    M llvm/test/CodeGen/NVPTX/i1-load-lower.ll
    M llvm/test/CodeGen/NVPTX/i1-select.ll
    M llvm/test/CodeGen/NVPTX/i128-array.ll
    M llvm/test/CodeGen/NVPTX/i128-ld-st.ll
    M llvm/test/CodeGen/NVPTX/i128-param.ll
    M llvm/test/CodeGen/NVPTX/i128-retval.ll
    M llvm/test/CodeGen/NVPTX/i128.ll
    M llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
    M llvm/test/CodeGen/NVPTX/i8-param.ll
    M llvm/test/CodeGen/NVPTX/i8x2-instructions.ll
    M llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
    M llvm/test/CodeGen/NVPTX/idioms.ll
    M llvm/test/CodeGen/NVPTX/indirect_byval.ll
    M llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
    M llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
    M llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
    M llvm/test/CodeGen/NVPTX/intrinsics.ll
    M llvm/test/CodeGen/NVPTX/jump-table.ll
    M llvm/test/CodeGen/NVPTX/ld-addrspace.ll
    M llvm/test/CodeGen/NVPTX/ld-generic.ll
    M llvm/test/CodeGen/NVPTX/ld-st-addrrspace.py
    M llvm/test/CodeGen/NVPTX/ldg-invariant.ll
    M llvm/test/CodeGen/NVPTX/ldparam-v4.ll
    M llvm/test/CodeGen/NVPTX/ldu-i8.ll
    M llvm/test/CodeGen/NVPTX/ldu-ldg.ll
    M llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll
    M llvm/test/CodeGen/NVPTX/load-sext-i1.ll
    M llvm/test/CodeGen/NVPTX/load-store-scalars.ll
    M llvm/test/CodeGen/NVPTX/load-store-sm-70.ll
    M llvm/test/CodeGen/NVPTX/load-store-sm-90.ll
    M llvm/test/CodeGen/NVPTX/load-store-vectors.ll
    M llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
    M llvm/test/CodeGen/NVPTX/local-stack-frame.ll
    M llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll
    M llvm/test/CodeGen/NVPTX/lower-alloca.ll
    M llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
    M llvm/test/CodeGen/NVPTX/lower-args.ll
    M llvm/test/CodeGen/NVPTX/lower-byval-args.ll
    M llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll
    M llvm/test/CodeGen/NVPTX/machine-sink.ll
    M llvm/test/CodeGen/NVPTX/match.ll
    M llvm/test/CodeGen/NVPTX/math-intrins.ll
    M llvm/test/CodeGen/NVPTX/misaligned-vector-ldst.ll
    M llvm/test/CodeGen/NVPTX/misched_func_call.ll
    M llvm/test/CodeGen/NVPTX/mulhi-intrins.ll
    M llvm/test/CodeGen/NVPTX/nounroll.ll
    M llvm/test/CodeGen/NVPTX/nvvm-reflect-arch-O0.ll
    M llvm/test/CodeGen/NVPTX/param-add.ll
    M llvm/test/CodeGen/NVPTX/param-align.ll
    M llvm/test/CodeGen/NVPTX/param-load-store.ll
    M llvm/test/CodeGen/NVPTX/param-overalign.ll
    M llvm/test/CodeGen/NVPTX/param-vectorize-device.ll
    M llvm/test/CodeGen/NVPTX/param-vectorize-kernel.ll
    M llvm/test/CodeGen/NVPTX/pr13291-i1-store.ll
    M llvm/test/CodeGen/NVPTX/pr16278.ll
    M llvm/test/CodeGen/NVPTX/prefetch.ll
    M llvm/test/CodeGen/NVPTX/proxy-reg-erasure-ptx.ll
    M llvm/test/CodeGen/NVPTX/rcp-opt.ll
    M llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll
    M llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
    M llvm/test/CodeGen/NVPTX/redux-sync-f32.ll
    M llvm/test/CodeGen/NVPTX/reg-types.ll
    M llvm/test/CodeGen/NVPTX/rotate-add.ll
    M llvm/test/CodeGen/NVPTX/rotate.ll
    M llvm/test/CodeGen/NVPTX/rotate_64.ll
    M llvm/test/CodeGen/NVPTX/sad-intrins.ll
    M llvm/test/CodeGen/NVPTX/sched1.ll
    M llvm/test/CodeGen/NVPTX/sched2.ll
    M llvm/test/CodeGen/NVPTX/sext-params.ll
    M llvm/test/CodeGen/NVPTX/sext-setcc.ll
    M llvm/test/CodeGen/NVPTX/shfl-p.ll
    M llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
    M llvm/test/CodeGen/NVPTX/shfl-sync.ll
    M llvm/test/CodeGen/NVPTX/shfl.ll
    M llvm/test/CodeGen/NVPTX/short-ptr.ll
    M llvm/test/CodeGen/NVPTX/shuffle-vec-undef-init.ll
    M llvm/test/CodeGen/NVPTX/st-addrspace.ll
    M llvm/test/CodeGen/NVPTX/st-generic.ll
    M llvm/test/CodeGen/NVPTX/st-param-imm.ll
    M llvm/test/CodeGen/NVPTX/st_bulk.ll
    M llvm/test/CodeGen/NVPTX/stacksaverestore.ll
    M llvm/test/CodeGen/NVPTX/store-retval.ll
    M llvm/test/CodeGen/NVPTX/store-undef.ll
    M llvm/test/CodeGen/NVPTX/surf-read-cuda.ll
    M llvm/test/CodeGen/NVPTX/surf-read.ll
    M llvm/test/CodeGen/NVPTX/surf-write-cuda.ll
    M llvm/test/CodeGen/NVPTX/szext.ll
    M llvm/test/CodeGen/NVPTX/tag-invariant-loads.ll
    M llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll
    M llvm/test/CodeGen/NVPTX/tcgen05-commit.ll
    M llvm/test/CodeGen/NVPTX/tcgen05-cp.ll
    M llvm/test/CodeGen/NVPTX/tcgen05-ld.ll
    M llvm/test/CodeGen/NVPTX/tcgen05-shift.ll
    M llvm/test/CodeGen/NVPTX/tcgen05-st.ll
    M llvm/test/CodeGen/NVPTX/tex-read-cuda.ll
    M llvm/test/CodeGen/NVPTX/tex-read.ll
    M llvm/test/CodeGen/NVPTX/texsurf-queries.ll
    M llvm/test/CodeGen/NVPTX/unaligned-param-load-store.ll
    M llvm/test/CodeGen/NVPTX/unfold-masked-merge-vector-variablemask.ll
    M llvm/test/CodeGen/NVPTX/vaargs.ll
    M llvm/test/CodeGen/NVPTX/variadics-backend.ll
    M llvm/test/CodeGen/NVPTX/vec-param-load.ll
    M llvm/test/CodeGen/NVPTX/vec8.ll
    M llvm/test/CodeGen/NVPTX/vector-args.ll
    M llvm/test/CodeGen/NVPTX/vector-call.ll
    M llvm/test/CodeGen/NVPTX/vector-compare.ll
    M llvm/test/CodeGen/NVPTX/vector-loads.ll
    M llvm/test/CodeGen/NVPTX/vector-select.ll
    M llvm/test/CodeGen/NVPTX/vector-stores.ll
    M llvm/test/CodeGen/NVPTX/vectorize-misaligned.ll
    M llvm/test/DebugInfo/NVPTX/debug-info.ll
    M llvm/test/Transforms/NaryReassociate/NVPTX/nary-slsr.ll
    M llvm/test/Transforms/SeparateConstOffsetFromGEP/NVPTX/split-gep-and-gvn.ll
    M llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/reassociate-geps-and-slsr.ll
    M llvm/test/Transforms/StraightLineStrengthReduce/NVPTX/speculative-slsr.ll
    M llvm/test/tools/UpdateTestChecks/update_llc_test_checks/Inputs/nvptx-basic.ll.expected

  Log Message:
  -----------
  [NVPTX] use untyped loads and stores where ever possible (#137698)

In most cases, the type information attached to load and store
instructions is meaningless and inconsistently applied. We can usually
use ".b" loads and avoid the complexity of trying to assign the correct
type. The one expectation is sign-extending load, which will continue to
use ".s" to ensure the sign extension into a larger register is done
correctly.

To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications