Durgadoss R d5e9691509 [NVPTX] Add elect.sync Intrinsic (#104780) ...

This patch adds an NVVM intrinsic and NVPTX codegen for the elect.sync
PTX instruction. Lit tests are
added in elect.ll and verified through ptxas.

PTX ISA reference:

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-elect-sync

Signed-off-by: Durgadoss R <durgadossr@nvidia.com>

2024-08-20 20:46:05 +05:30

..

access-non-generic.ll

…

activemask.ll

[NVPTX] Add 'activemask' builtin and intrinsic support (#79768)

2024-01-29 14:07:30 -06:00

add-sub-128bit.ll

…

addr-mode.ll

[NVPTX] enforce signed 32 bit type for immediate offset (#99682)

2024-07-20 08:54:57 -07:00

addrspacecast-gvar.ll

…

addrspacecast.ll

[SDISel] Teach the type legalizer about ADDRSPACECAST (#90969)

2024-05-07 11:08:33 +02:00

aggr-param.ll

…

aggregate-return.ll

…

alias-errors.ll

…

alias.ll

[NVPTX] Add support for calling aliases (#81170)

2024-02-08 17:14:13 -06:00

annotations.ll

…

anonymous-fn-param.ll

…

APIntLoadStore.ll

…

APIntParam.ll

…

APIntSextParam.ll

…

APIntZextParam.ll

…

arg-lowering.ll

…

arithmetic-fp-sm20.ll

…

arithmetic-int.ll

…

async-copy.ll

…

atomic-lower-local.ll

…

atomicrmw-expand.ll

Set MaxAtomicSizeInBitsSupported for remaining targets. (#75703)

2024-01-08 22:34:28 -05:00

atomics-sm60.ll

…

atomics-sm70.ll

Reland "[NVPTX] Add support for atomic add for f16 type" (#85197)

2024-03-15 08:13:59 +01:00

atomics-sm90.ll

[NVPTX] Add support for atomic add for bf16 type (#89586)

2024-05-08 10:22:02 +02:00

atomics-with-scope.ll

…

atomics.ll

Reland "[NVPTX] Add support for atomic add for f16 type" (#85197)

2024-03-15 08:13:59 +01:00

b52037.ll

[Tests] Drop inrange attribute from some tests (NFC)

2024-03-13 11:49:16 +01:00

barrier.ll

…

bf16-instructions.ll

[NVPTX] Fix bugs involving maximum/minimum and bf16

2024-08-20 05:09:49 +00:00

bf16.ll

…

bf16x2-instructions-approx.ll

[test] Convert text files from CRLF to LF

2024-05-03 10:09:52 -07:00

bf16x2-instructions.ll

[test] Convert text files from CRLF to LF

2024-05-03 10:09:52 -07:00

bfe.ll

…

boolean-patterns.ll

[NVPTX] improve Boolean ISel (#80166)

2024-01-31 14:37:27 -08:00

branch-fold.ll

…

branch-fold.mir

…

brkpt.ll

[LLVM][NVPTX] Add support for brkpt instruction (#104470)

2024-08-17 10:00:24 +05:30

bswap.ll

[NVPTX] Use PTX prmt for llvm.bswap (#85545)

2024-03-19 15:18:53 -07:00

bug17709.ll

…

bug21465.ll

…

bug22246.ll

…

bug22322.ll

…

bug26185-2.ll

…

bug26185.ll

…

bug41651.ll

…

bug52623.ll

…

bypass-div.ll

…

call_bitcast_byval.ll

…

call-with-alloca-buffer.ll

…

callchain.ll

…

calling-conv.ll

…

calls-with-phi.ll

…

chain-different-as.ll

…

combine-mad.ll

[NVPTX] Fold (add (select 0, (mul a, b)), c) -> (select c, (mad a, b, c)) (#96352)

2024-06-26 09:01:13 -07:00

combine-min-max.ll

…

common-linkage.ll

Reland "[NVPTX] Use .common linkage for common globals" (#86824)

2024-03-29 11:58:41 -07:00

compare-int.ll

…

constant-vectors.ll

…

convergent-mir-call.ll

…

convert-fp.ll

…

convert-int-sm20.ll

…

convert-sm80.ll

…

convert-sm89.ll

[NVPTX] Add conversion intrinsics from/to fp8 types (e4m3, e5m2) (#102969)

2024-08-19 10:34:22 -07:00

ctlz.ll

…

ctpop.ll

…

cttz.ll

…

dag-cse.ll

…

demote-vars.ll

…

disable-opt.ll

…

div-ri.ll

…

divrem-combine.ll

…

dot-product.ll

[NVPTX] Add idp2a, idp4a intrinsics (#102763)

2024-08-14 13:31:16 -07:00

dynamic_stackalloc.ll

[llvm][NVPTX] Don't emit unused var 'temp_param_reg' (NFC) (#89004)

2024-04-17 14:45:33 -04:00

dynamic-stackalloc-regression.ll

[NVPTX] Emit NVPTXISD::DYNAMIC_STACKALLOC's chain (#101714)

2024-08-05 10:53:33 -07:00

elect.ll

[NVPTX] Add elect.sync Intrinsic (#104780)

2024-08-20 20:46:05 +05:30

empty-type.ll

…

envreg.ll

…

extloadv.ll

…

extractelement.ll

[NVPTX] unbreak extract_elt lowering (#102688)

2024-08-12 10:28:56 -07:00

f16-ex2.ll

…

f16-instructions.ll

Enforce parameter order in f16 call, flipped call unit tests (#96258)

2024-06-25 14:07:40 -07:00

f16x2-instructions.ll

…

fast-math.ll

…

fcos-no-fast-math.ll

…

fence-proxy-tensormap.ll

[LLVM][NVPTX] Add NVPTX codegen support for fence.proxy.tensormap (#100748)

2024-08-07 20:25:45 +05:30

fence.ll

[NVPTX] Support fence instruction (#99649)

2024-07-24 10:59:03 -07:00

filetype-null.ll

…

fma-assoc.ll

…

fma-disable.ll

…

fma.ll

…

fminimum-fmaximum.ll

…

fns.ll

…

fp16.ll

…

fp-contract.ll

…

fp-literals.ll

…

fsin-no-fast-math.ll

…

function-align.ll

…

generic-to-nvvm-ir.ll

…

generic-to-nvvm.ll

…

global-addrspace.ll

…

global-ctor-empty.ll

…

global-ctor.ll

…

global-dtor.ll

…

global-incomplete-init.ll

[NVPTX] use incomplete aggregate initializers (#79062)

2024-01-24 09:24:28 -08:00

global-ordering.ll

…

global-variable-big.ll

[test] Fix check prefixes

2024-05-13 14:01:00 -07:00

global-visibility.ll

…

globals_init.ll

[NVPTX] use incomplete aggregate initializers (#79062)

2024-01-24 09:24:28 -08:00

globals_lowering.ll

…

gvar-init.ll

…

half.ll

…

i1-array-global.ll

Fix i1 array global crash in NVPTXAsmPrinter. (#92506)

2024-05-17 12:06:35 +02:00

i1-ext-load.ll

[NVPTX] fix emission for i1 load and extload (#99392)

2024-07-18 10:29:24 -07:00

i1-global.ll

…

i1-int-to-fp.ll

…

i1-load-lower.ll

[NVPTX] fix emission for i1 load and extload (#99392)

2024-07-18 10:29:24 -07:00

i1-param.ll

…

i8-param.ll

…

i8x4-instructions.ll

[llvm][NVPTX] Don't emit unused var 'temp_param_reg' (NFC) (#89004)

2024-04-17 14:45:33 -04:00

i16x2-instructions.ll

[NVPTX] Remove sub.s16x2 instruction

2024-03-05 12:43:02 +01:00

i128-global.ll

[NVPTX] use incomplete aggregate initializers (#79062)

2024-01-24 09:24:28 -08:00

i128-param.ll

…

i128-retval.ll

…

i128-struct.ll

…

i128.ll

[DAGCombiner] Freeze maybe poison operands when folding select to logic (#84924)

2024-07-22 17:19:46 +02:00

idioms.ll

…

imad.ll

…

indirect_byval.ll

[NVPTX] Fix internal indirect call prototypes not obeying the ABI (#100131)

2024-07-23 12:54:00 -05:00

inline-asm-b128-test1.ll

[NVPTX] Update 32-bit NVPTX tests to use 64-bit

2024-07-10 16:01:12 -05:00

inline-asm-b128-test2.ll

[NVPTX] Update 32-bit NVPTX tests to use 64-bit

2024-07-10 16:01:12 -05:00

inline-asm-b128-test3.ll

[NVPTX] Update 32-bit NVPTX tests to use 64-bit

2024-07-10 16:01:12 -05:00

inline-asm.ll

…

inlineasm-output-template.ll

…

intr-range.ll

[NVPTX] Revamp NVVMIntrRange pass (#94422)

2024-06-06 06:42:46 -07:00

intrin-nocapture.ll

…

intrinsic-old.ll

[NVPTX] Revamp NVVMIntrRange pass (#94422)

2024-06-06 06:42:46 -07:00

intrinsics-sm90.ll

[LLVM][NVPTX] Add cp.async.bulk.commit/wait intrinsics (#78698)

2024-01-19 10:42:33 -08:00

intrinsics.ll

[LLVM] Add __builtin_readsteadycounter intrinsic and builtin for realtime clocks (#81331)

2024-02-13 10:06:25 -06:00

isspacep.ll

…

jump-table.ll

[NVPTX] support switch statement with brx.idx (reland) (#102550)

2024-08-09 11:46:37 -07:00

ld-addrspace.ll

[CodeGen] Convert tests to opaque pointers (NFC)

2024-02-05 14:07:09 +01:00

ld-generic.ll

…

ld-st-addrrspace.py

…

ldg-invariant.ll

…

ldparam-v4.ll

…

ldu-i8.ll

…

ldu-ldg.ll

[CodeGen] Convert tests to opaque pointers (NFC)

2024-02-05 14:07:09 +01:00

ldu-reg-plus-offset.ll

…

lit.local.cfg

…

load-sext-i1.ll

…

load-store-sm-70.ll

[NVPTX] Add Volta Atomic SequentiallyConsistent Load and Store Operations (#98551)

2024-08-06 10:32:51 -07:00

load-store.ll

[NVPTX] Add Volta Atomic SequentiallyConsistent Load and Store Operations (#98551)

2024-08-06 10:32:51 -07:00

load-with-non-coherent-cache.ll

…

LoadStoreVectorizer.ll

…

local-stack-frame.ll

…

loop-vectorize.ll

…

lower-aggr-copies.ll

[LowerMemIntrinsics] Lower llvm.memmove to wide memory accesses (#100122)

2024-07-26 08:43:30 +02:00

lower-alloca.ll

…

lower-args-gridconstant.ll

[NVPTX] Improved support for grid_constant (#97112)

2024-06-30 00:50:37 -07:00

lower-args.ll

…

lower-byval-args.ll

…

lower-ctor-dtor.ll

…

lower-kernel-ptr-arg.ll

…

machine-sink.ll

…

MachineSink-call.ll

…

MachineSink-convergent.ll

…

managed.ll

…

match.ll

…

math-intrins-sm53-ptx42.ll

…

math-intrins-sm80-ptx70-autoupgrade.ll

[NVPTX][NFC] Update tests to use bfloat type (#101493)

2024-08-09 13:59:56 -07:00

math-intrins-sm80-ptx70-instcombine.ll

…

math-intrins-sm80-ptx70.ll

…

math-intrins-sm86-ptx72-autoupgrade.ll

[NVPTX][NFC] Update tests to use bfloat type (#101493)

2024-08-09 13:59:56 -07:00

math-intrins-sm86-ptx72.ll

…

math-intrins.ll

[NVPTX] Make minimum/maximum work on older GPUs

2024-07-31 22:38:19 +00:00

max-align.ll

[NVPTX] cap param alignment at 128 (max supported by ptx) (#96117)

2024-06-25 10:04:45 -07:00

maxclusterrank.ll

…

mbarrier.ll

…

minmax-negative.ll

…

misaligned-vector-ldst.ll

…

mma-no-sink-after-laneid-check.ll

[NVPTX] Make nvptx mma instructions convergent. (#96521)

2024-06-24 22:15:58 -04:00

module-inline-asm.ll

…

mulhi-intrins.ll

[NVPTX] extend type support for nvvm.{min,max,mulhi,sad} (#78385)

2024-01-17 16:18:39 -08:00

mulwide.ll

…

named-barriers.ll

…

nanosleep.ll

[NVPTX] pass correct GPU arch to ptxas test (#81535)

2024-02-12 13:18:08 -08:00

no-extra-parens.ll

…

noduplicate-syncthreads.ll

…

nofunc.ll

…

noreturn.ll

[CodeGen] Convert tests to opaque pointers (NFC)

2024-02-05 14:07:09 +01:00

nounroll.ll

…

nvcl-param-align.ll

…

nvptx-aa.ll

…

nvvm-annotations-D120129.ll

…

nvvm-reflect-arch-O0.ll

[NVVMReflect] Improve folding inside of the NVVMReflect pass (#81253)

2024-02-09 13:39:03 -06:00

nvvm-reflect-arch.ll

[NVVMReflect][Reland] Force dead branch elimination in NVVMReflect (#81189)

2024-02-08 20:09:44 -06:00

nvvm-reflect-module-flag.ll

…

nvvm-reflect-ocl.ll

[CodeGen] Convert tests to opaque pointers (NFC)

2024-02-05 14:07:09 +01:00

nvvm-reflect-opaque.ll

[NVPTX] Add clang builtin for __nvvm_reflect intrinsic (#81277)

2024-02-09 14:11:01 -06:00

nvvm-reflect.ll

[NVPTX] Add clang builtin for __nvvm_reflect intrinsic (#81277)

2024-02-09 14:11:01 -06:00

packed-aggr.ll

…

param-align.ll

[NVPTX] Don't use underlying alignment to align param (#96793)

2024-06-27 10:19:05 +01:00

param-load-store.ll

[NVPTX] fixup support for unaligned parameters and returns (#82562)

2024-02-22 17:27:28 -08:00

param-overalign.ll

NVPTX: fix test failures under ptxas-12.0.

2024-05-20 23:09:26 -04:00

param-vectorize-device.ll

…

param-vectorize-kernel.ll

…

pass-name.ll

…

pow2_mask_cmp.ll

…

pr13291-i1-store.ll

…

pr16278.ll

…

pr17529.ll

…

proxy-reg-erasure-mir.ll

…

proxy-reg-erasure-ptx.ll

…

rcp-opt.ll

[NVPTX] Lower -1/x to neg.f64(rcp.rn.f64) instead of fdiv (#98343)

2024-07-16 12:45:21 +05:30

read-global-variable-constant.ll

…

redux-sync.ll

…

refl1.ll

…

reg-copy.ll

…

reg-types.ll

…

rotate_64.ll

…

rotate.ll

[NVPTX] Fix 64 bits rotations with large shift values (#89399)

2024-05-01 10:15:52 -07:00

rsqrt-opt.ll

[NVPTX] Improve support for rsqrt.approx (#89417)

2024-04-23 08:56:39 -07:00

rsqrt.ll

[NVPTX] Improve support for rsqrt.approx (#89417)

2024-04-23 08:56:39 -07:00

sad-intrins.ll

[NVPTX] extend type support for nvvm.{min,max,mulhi,sad} (#78385)

2024-01-17 16:18:39 -08:00

sched1.ll

…

sched2.ll

…

setmaxnreg.ll

[LLVM][NVPTX]: Add intrinsic for setmaxnreg (#77289)

2024-01-09 12:04:13 -08:00

sext-in-reg.ll

…

sext-params.ll

…

shfl-p.ll

…

shfl-sync-p.ll

…

shfl-sync.ll

…

shfl.ll

…

shift-parts.ll

…

short-ptr.ll

[CodeGen] Convert tests to opaque pointers (NFC)

2024-02-05 14:07:09 +01:00

shuffle-vec-undef-init.ll

[NVPTX] Fix generating permute bytes from register pair when the initial values are undefined (#74437)

2024-01-16 11:05:41 -08:00

simple-call.ll

…

sm-version.ll

…

speculative-execution-divergent-target.ll

…

sqrt-approx.ll

…

st-addrspace.ll

[CodeGen] Convert tests to opaque pointers (NFC)

2024-02-05 14:07:09 +01:00

st-generic.ll

…

st-param-imm.ll

NVPTX: fix test failures under ptxas-12.0.

2024-05-20 23:09:26 -04:00

store-retval.ll

…

store-undef.ll

[SelectionDAG] Preserve volatile undef stores. (#99918)

2024-07-24 08:41:56 -04:00

surf-read-cuda.ll

…

surf-read.ll

…

surf-tex.py

[NVPTX] Add tex.grad.cube{array} intrinsics (#77693)

2024-01-17 10:41:11 -08:00

surf-write-cuda.ll

…

surf-write.ll

…

symbol-naming.ll

[NVPTX] improve identifier renaming for PTX (#79459)

2024-01-26 13:49:00 -08:00

TailDuplication-convergent.ll

…

tex-read-cuda.ll

…

tex-read.ll

…

texsurf-queries.ll

…

tid-range.ll

…

unaligned-param-load-store.ll

[NVPTX] fixup support for unaligned parameters and returns (#82562)

2024-02-22 17:27:28 -08:00

unfold-masked-merge-vector-variablemask.ll

…

unreachable.ll

…

vaargs.ll

…

variadics-backend.ll

[NVPTX] Implement variadic functions using IR lowering (#96015)

2024-07-12 17:09:48 -05:00

variadics-lowering.ll

[IRBuilder] Generate nuw GEPs for struct member accesses (#99538)

2024-08-09 13:25:04 +01:00

vec8.ll

…

vec-param-load.ll

…

vector-args.ll

…

vector-call.ll

…

vector-compare.ll

…

vector-global.ll

…

vector-loads.ll

[SeperateConstOffsetFromGEP] Handle or disjoint flags (#76997)

2024-01-26 09:56:06 -06:00

vector-select.ll

…

vector-stores.ll

Revert "[NVPTX] Lower 16xi8 and 8xi8 stores efficiently (#73646)" (#74518)

2023-12-06 10:48:43 -08:00

vectorize-misaligned.ll

…

vote.ll

…

weak-global.ll

Reland "[NVPTX] Use .common linkage for common globals" (#86824)

2024-03-29 11:58:41 -07:00

weak-linkage.ll

…

wmma-ptx60-sm70.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma-ptx61-sm70.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma-ptx63-sm72.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma-ptx63-sm75.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma-ptx64-sm70.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma-ptx65-sm75.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma-ptx71-sm80.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

wmma.py

[NVPTX] split long-running wmma.py test into smaller chunks. (#72331)

2023-11-15 09:55:47 -08:00

zeroext-32bit.ll

[test] Convert text files from CRLF to LF

2024-05-03 10:09:52 -07:00