[NVPTX] Integrate ptxas to LIT tests
authorAndrew Savonichev <andrew.savonichev@gmail.com>
Wed, 27 Apr 2022 19:43:55 +0000 (22:43 +0300)
committerAndrew Savonichev <andrew.savonichev@gmail.com>
Thu, 28 Apr 2022 11:59:45 +0000 (14:59 +0300)
commit0f1b5f115a7f6fd90989996ae514810773157b76
tree0847804a3b8c25a27bd76d5ed997142f4b953c41
parentb1f1688e90b22dedc829f5abc9a912f18c034fbc
[NVPTX] Integrate ptxas to LIT tests

ptxas is a proprietary compiler from Nvidia that can compile PTX to
machine code (SASS). It has a lot of diagnostics to catch errors
in PTX, which can be used to verify PTX output from llc.

Set -DPXTAS_EXECUTABLE=/path/to/ptxas CMake option to enable it.
If this option is not set, then ptxas is substituted to true which
effectively disables all ptxas RUN lines.

LLVM_PTXAS_EXECUTABLE environment variable takes precedence over
the CMake option, and allows to override ptxas executable that is used for LIT
without complete re-configuration.

Differential Revision: https://reviews.llvm.org/D121727
190 files changed:
llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll
llvm/test/CodeGen/NVPTX/MachineSink-call.ll
llvm/test/CodeGen/NVPTX/MachineSink-convergent.ll
llvm/test/CodeGen/NVPTX/TailDuplication-convergent.ll
llvm/test/CodeGen/NVPTX/access-non-generic.ll
llvm/test/CodeGen/NVPTX/add-128bit.ll
llvm/test/CodeGen/NVPTX/addrspacecast-gvar.ll
llvm/test/CodeGen/NVPTX/addrspacecast.ll
llvm/test/CodeGen/NVPTX/aggr-param.ll
llvm/test/CodeGen/NVPTX/aggregate-return.ll
llvm/test/CodeGen/NVPTX/annotations.ll
llvm/test/CodeGen/NVPTX/arg-lowering.ll
llvm/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll
llvm/test/CodeGen/NVPTX/arithmetic-int.ll
llvm/test/CodeGen/NVPTX/async-copy.ll
llvm/test/CodeGen/NVPTX/atomics-sm60.ll
llvm/test/CodeGen/NVPTX/atomics-with-scope.ll
llvm/test/CodeGen/NVPTX/atomics.ll
llvm/test/CodeGen/NVPTX/b52037.ll
llvm/test/CodeGen/NVPTX/barrier.ll
llvm/test/CodeGen/NVPTX/bfe.ll
llvm/test/CodeGen/NVPTX/branch-fold.ll
llvm/test/CodeGen/NVPTX/bug17709.ll
llvm/test/CodeGen/NVPTX/bug21465.ll
llvm/test/CodeGen/NVPTX/bug22246.ll
llvm/test/CodeGen/NVPTX/bug22322.ll
llvm/test/CodeGen/NVPTX/bug26185-2.ll
llvm/test/CodeGen/NVPTX/bug26185.ll
llvm/test/CodeGen/NVPTX/bug41651.ll
llvm/test/CodeGen/NVPTX/bypass-div.ll
llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
llvm/test/CodeGen/NVPTX/callchain.ll
llvm/test/CodeGen/NVPTX/calling-conv.ll
llvm/test/CodeGen/NVPTX/calls-with-phi.ll
llvm/test/CodeGen/NVPTX/combine-min-max.ll
llvm/test/CodeGen/NVPTX/compare-int.ll
llvm/test/CodeGen/NVPTX/constant-vectors.ll
llvm/test/CodeGen/NVPTX/convert-fp.ll
llvm/test/CodeGen/NVPTX/convert-int-sm20.ll
llvm/test/CodeGen/NVPTX/convert-sm80.ll
llvm/test/CodeGen/NVPTX/ctlz.ll
llvm/test/CodeGen/NVPTX/ctpop.ll
llvm/test/CodeGen/NVPTX/cttz.ll
llvm/test/CodeGen/NVPTX/disable-opt.ll
llvm/test/CodeGen/NVPTX/div-ri.ll
llvm/test/CodeGen/NVPTX/divrem-combine.ll
llvm/test/CodeGen/NVPTX/envreg.ll
llvm/test/CodeGen/NVPTX/extloadv.ll
llvm/test/CodeGen/NVPTX/f16-ex2.ll
llvm/test/CodeGen/NVPTX/f16-instructions.ll
llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
llvm/test/CodeGen/NVPTX/fast-math.ll
llvm/test/CodeGen/NVPTX/fma-assoc.ll
llvm/test/CodeGen/NVPTX/fma-disable.ll
llvm/test/CodeGen/NVPTX/fma.ll
llvm/test/CodeGen/NVPTX/fminimum-fmaximum.ll
llvm/test/CodeGen/NVPTX/fns.ll
llvm/test/CodeGen/NVPTX/fp-contract.ll
llvm/test/CodeGen/NVPTX/fp-literals.ll
llvm/test/CodeGen/NVPTX/fp16.ll
llvm/test/CodeGen/NVPTX/function-align.ll
llvm/test/CodeGen/NVPTX/generic-to-nvvm.ll
llvm/test/CodeGen/NVPTX/global-addrspace.ll
llvm/test/CodeGen/NVPTX/global-ordering.ll
llvm/test/CodeGen/NVPTX/global-variable-big.ll
llvm/test/CodeGen/NVPTX/global-visibility.ll
llvm/test/CodeGen/NVPTX/globals_init.ll
llvm/test/CodeGen/NVPTX/globals_lowering.ll
llvm/test/CodeGen/NVPTX/half.ll
llvm/test/CodeGen/NVPTX/i1-global.ll
llvm/test/CodeGen/NVPTX/i1-int-to-fp.ll
llvm/test/CodeGen/NVPTX/i1-param.ll
llvm/test/CodeGen/NVPTX/i128-global.ll
llvm/test/CodeGen/NVPTX/i128-param.ll
llvm/test/CodeGen/NVPTX/i128-retval.ll
llvm/test/CodeGen/NVPTX/i128-struct.ll
llvm/test/CodeGen/NVPTX/i8-param.ll
llvm/test/CodeGen/NVPTX/idioms.ll
llvm/test/CodeGen/NVPTX/imad.ll
llvm/test/CodeGen/NVPTX/inline-asm.ll
llvm/test/CodeGen/NVPTX/inlineasm-output-template.ll
llvm/test/CodeGen/NVPTX/intrinsic-old.ll
llvm/test/CodeGen/NVPTX/intrinsics.ll
llvm/test/CodeGen/NVPTX/isspacep.ll
llvm/test/CodeGen/NVPTX/ld-addrspace.ll
llvm/test/CodeGen/NVPTX/ld-generic.ll
llvm/test/CodeGen/NVPTX/ld-st-addrrspace.py
llvm/test/CodeGen/NVPTX/ldg-invariant.ll
llvm/test/CodeGen/NVPTX/ldparam-v4.ll
llvm/test/CodeGen/NVPTX/ldu-i8.ll
llvm/test/CodeGen/NVPTX/ldu-ldg.ll
llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll
llvm/test/CodeGen/NVPTX/libcall-fulfilled.ll
llvm/test/CodeGen/NVPTX/load-sext-i1.ll
llvm/test/CodeGen/NVPTX/load-store.ll
llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
llvm/test/CodeGen/NVPTX/local-stack-frame.ll
llvm/test/CodeGen/NVPTX/lower-aggr-copies.ll
llvm/test/CodeGen/NVPTX/lower-alloca.ll
llvm/test/CodeGen/NVPTX/lower-args.ll
llvm/test/CodeGen/NVPTX/lower-byval-args.ll
llvm/test/CodeGen/NVPTX/lower-kernel-ptr-arg.ll
llvm/test/CodeGen/NVPTX/machine-sink.ll
llvm/test/CodeGen/NVPTX/managed.ll
llvm/test/CodeGen/NVPTX/match.ll
llvm/test/CodeGen/NVPTX/math-intrins-sm53-ptx42.ll
llvm/test/CodeGen/NVPTX/math-intrins-sm80-ptx70.ll
llvm/test/CodeGen/NVPTX/math-intrins-sm86-ptx72.ll
llvm/test/CodeGen/NVPTX/math-intrins.ll
llvm/test/CodeGen/NVPTX/mbarrier.ll
llvm/test/CodeGen/NVPTX/minmax-negative.ll
llvm/test/CodeGen/NVPTX/misaligned-vector-ldst.ll
llvm/test/CodeGen/NVPTX/module-inline-asm.ll
llvm/test/CodeGen/NVPTX/mulwide.ll
llvm/test/CodeGen/NVPTX/named-barriers.ll
llvm/test/CodeGen/NVPTX/no-extra-parens.ll
llvm/test/CodeGen/NVPTX/nofunc.ll
llvm/test/CodeGen/NVPTX/nounroll.ll
llvm/test/CodeGen/NVPTX/nvcl-param-align.ll
llvm/test/CodeGen/NVPTX/nvvm-annotations-D120129.ll
llvm/test/CodeGen/NVPTX/param-align.ll
llvm/test/CodeGen/NVPTX/param-load-store.ll
llvm/test/CodeGen/NVPTX/param-vectorize-device.ll
llvm/test/CodeGen/NVPTX/param-vectorize-kernel.ll
llvm/test/CodeGen/NVPTX/pow2_mask_cmp.ll
llvm/test/CodeGen/NVPTX/pr13291-i1-store.ll
llvm/test/CodeGen/NVPTX/pr16278.ll
llvm/test/CodeGen/NVPTX/pr17529.ll
llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll
llvm/test/CodeGen/NVPTX/redux-sync.ll
llvm/test/CodeGen/NVPTX/refl1.ll
llvm/test/CodeGen/NVPTX/reg-copy.ll
llvm/test/CodeGen/NVPTX/reg-types.ll
llvm/test/CodeGen/NVPTX/rotate.ll
llvm/test/CodeGen/NVPTX/rotate_64.ll
llvm/test/CodeGen/NVPTX/sched1.ll
llvm/test/CodeGen/NVPTX/sched2.ll
llvm/test/CodeGen/NVPTX/sext-in-reg.ll
llvm/test/CodeGen/NVPTX/sext-params.ll
llvm/test/CodeGen/NVPTX/shfl-p.ll
llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
llvm/test/CodeGen/NVPTX/shfl-sync.ll
llvm/test/CodeGen/NVPTX/shfl.ll
llvm/test/CodeGen/NVPTX/shift-parts.ll
llvm/test/CodeGen/NVPTX/simple-call.ll
llvm/test/CodeGen/NVPTX/sqrt-approx.ll
llvm/test/CodeGen/NVPTX/st-addrspace.ll
llvm/test/CodeGen/NVPTX/st-generic.ll
llvm/test/CodeGen/NVPTX/store-retval.ll
llvm/test/CodeGen/NVPTX/surf-read-cuda.ll
llvm/test/CodeGen/NVPTX/surf-read.ll
llvm/test/CodeGen/NVPTX/surf-tex.py
llvm/test/CodeGen/NVPTX/surf-write-cuda.ll
llvm/test/CodeGen/NVPTX/surf-write.ll
llvm/test/CodeGen/NVPTX/symbol-naming.ll
llvm/test/CodeGen/NVPTX/tex-read-cuda.ll
llvm/test/CodeGen/NVPTX/tex-read.ll
llvm/test/CodeGen/NVPTX/texsurf-queries.ll
llvm/test/CodeGen/NVPTX/tid-range.ll
llvm/test/CodeGen/NVPTX/tuple-literal.ll
llvm/test/CodeGen/NVPTX/vec-param-load.ll
llvm/test/CodeGen/NVPTX/vec8.ll
llvm/test/CodeGen/NVPTX/vector-args.ll
llvm/test/CodeGen/NVPTX/vector-call.ll
llvm/test/CodeGen/NVPTX/vector-compare.ll
llvm/test/CodeGen/NVPTX/vector-global.ll
llvm/test/CodeGen/NVPTX/vector-loads.ll
llvm/test/CodeGen/NVPTX/vector-select.ll
llvm/test/CodeGen/NVPTX/vector-stores.ll
llvm/test/CodeGen/NVPTX/vectorize-misaligned.ll
llvm/test/CodeGen/NVPTX/vote.ll
llvm/test/CodeGen/NVPTX/weak-global.ll
llvm/test/CodeGen/NVPTX/weak-linkage.ll
llvm/test/CodeGen/NVPTX/wmma.py
llvm/test/CodeGen/NVPTX/zeroext-32bit.ll
llvm/test/DebugInfo/NVPTX/crash-missing-DISubprogram.ll
llvm/test/DebugInfo/NVPTX/cu-range-hole.ll
llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
llvm/test/DebugInfo/NVPTX/debug-addr-class.ll
llvm/test/DebugInfo/NVPTX/debug-empty.ll
llvm/test/DebugInfo/NVPTX/debug-file-loc-only.ll
llvm/test/DebugInfo/NVPTX/debug-file-loc.ll
llvm/test/DebugInfo/NVPTX/debug-info.ll
llvm/test/DebugInfo/NVPTX/debug-loc-offset.ll
llvm/test/DebugInfo/NVPTX/debug-name-table.ll
llvm/test/DebugInfo/NVPTX/dwarf-file-dir.ll
llvm/test/DebugInfo/NVPTX/packed_bitfields.ll
llvm/test/lit.cfg.py
llvm/test/lit.site.cfg.py.in