caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
94e01d579c1954bed2dbd2d82a64ff72baf72223
clang-p2996/llvm/test/CodeGen/NVPTX
History
Steffen Larsen 3644726a78 [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions 
Adds NVPTX builtins and intrinsics for the CUDA PTX `wmma.load`, `wmma.store`, `wmma.mma`, and `mma` instructions added in PTX 6.5 and 7.0.

PTX ISA description of

  - `wmma.load`: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-ld
  - `wmma.store`: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-st
  - `wmma.mma`: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-wmma-mma
  - `mma`: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma

Overview of `wmma.mma` and `mma` matrix shape/type combinations added with specific PTX versions: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-shape

Authored-by: Steffen Larsen <steffen.larsen@codeplay.com>
Co-Authored-by: Stuart Adams <stuart.adams@codeplay.com>

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D104847
2021-06-29 15:44:07 -07:00
..
access-non-generic.ll
…
add-128bit.ll
…
addrspacecast-gvar.ll
…
addrspacecast.ll
…
aggr-param.ll
…
aggregate-return.ll
…
alias.ll
…
annotations.ll
…
arg-lowering.ll
…
arithmetic-fp-sm20.ll
…
arithmetic-int.ll
…
async-copy.ll
[Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async instructions
2021-05-17 09:46:59 -07:00
atomic-lower-local.ll
[NVPTX] Enable lowering of atomics on local memory
2021-04-26 20:12:12 -04:00
atomics-sm60.ll
…
atomics-with-scope.ll
…
atomics.ll
…
barrier.ll
…
bfe.ll
…
branch-fold.ll
…
bug17709.ll
…
bug21465.ll
Revert "[NFC] remove explicit default value for strboolattr attribute in tests"
2021-05-24 19:43:40 +02:00
bug22246.ll
…
bug22322.ll
Revert "[NFC] remove explicit default value for strboolattr attribute in tests"
2021-05-24 19:43:40 +02:00
bug26185-2.ll
…
bug26185.ll
…
bug41651.ll
…
bypass-div.ll
…
call-with-alloca-buffer.ll
…
callchain.ll
…
calling-conv.ll
…
calls-with-phi.ll
…
combine-min-max.ll
…
compare-int.ll
…
constant-vectors.ll
…
convergent-mir-call.ll
…
convert-fp.ll
…
convert-int-sm20.ll
…
ctlz.ll
…
ctpop.ll
…
cttz.ll
…
disable-opt.ll
…
div-ri.ll
…
divrem-combine.ll
…
envreg.ll
…
extloadv.ll
…
f16-instructions.ll
Update @llvm.powi to handle different int sizes for the exponent
2021-06-17 09:38:28 +02:00
f16x2-instructions.ll
Update @llvm.powi to handle different int sizes for the exponent
2021-06-17 09:38:28 +02:00
fast-math.ll
[DAGCombine] Remove the check for unsafe-fp-math when we are checking the AFN
2021-01-11 02:25:53 +00:00
fcos-no-fast-math.ll
…
fma-assoc.ll
…
fma-disable.ll
…
fma.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-ordering.ll
…
global-variable-big.ll
…
global-visibility.ll
…
globals_init.ll
…
globals_lowering.ll
…
gvar-init.ll
…
half.ll
…
i1-global.ll
…
i1-int-to-fp.ll
…
i1-param.ll
…
i8-param.ll
…
i128-global.ll
…
i128-param.ll
…
i128-retval.ll
…
i128-struct.ll
…
idioms.ll
…
imad.ll
…
inline-asm.ll
…
inlineasm-output-template.ll
…
intrin-nocapture.ll
…
intrinsic-old.ll
[NewPM][NVPTX] Port NVPTX opt passes
2021-01-07 15:12:35 -08:00
intrinsics.ll
…
isspacep.ll
…
ld-addrspace.ll
…
ld-generic.ll
…
ld-st-addrrspace.py
…
ldg-invariant.ll
…
ldparam-v4.ll
…
ldu-i8.ll
…
ldu-ldg.ll
…
ldu-reg-plus-offset.ll
…
libcall-fulfilled.ll
[NVPTX] CUDA does provide malloc/free since compute capability 2.X
2021-03-15 22:45:56 -05:00
libcall-instruction.ll
…
libcall-intrinsic.ll
Update @llvm.powi to handle different int sizes for the exponent
2021-06-17 09:38:28 +02:00
lit.local.cfg
[Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions
2021-06-29 15:44:07 -07:00
load-sext-i1.ll
…
load-store.ll
…
load-with-non-coherent-cache.ll
…
LoadStoreVectorizer.ll
…
local-stack-frame.ll
…
loop-vectorize.ll
…
lower-aggr-copies.ll
…
lower-alloca.ll
…
lower-args.ll
[NVPTX] Avoid temp copy of byval kernel parameters.
2021-03-15 14:27:22 -07:00
lower-byval-args.ll
[NVPTX] Handle bitcast and ASC(101) when trying to avoid argument copy.
2021-04-06 13:06:00 -07:00
lower-kernel-ptr-arg.ll
[NVPTX] Avoid temp copy of byval kernel parameters.
2021-03-15 14:27:22 -07:00
machine-sink.ll
…
MachineSink-call.ll
…
MachineSink-convergent.ll
…
managed.ll
…
match.ll
…
math-intrins.ll
…
mbarrier.ll
[Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async instructions
2021-05-17 09:46:59 -07:00
minmax-negative.ll
…
misaligned-vector-ldst.ll
…
module-inline-asm.ll
…
mulwide.ll
…
named-barriers.ll
…
noduplicate-syncthreads.ll
…
nofunc.ll
…
nounroll.ll
…
nvcl-param-align.ll
…
nvvm-reflect-arch.ll
[NVPTX][NewPM] Re-enable NVVMReflectPass
2021-02-08 13:58:17 -08:00
nvvm-reflect-module-flag.ll
[NewPM][NVPTX] Port NVPTX opt passes
2021-01-07 15:12:35 -08:00
nvvm-reflect.ll
[NewPM][NVPTX] Port NVPTX opt passes
2021-01-07 15:12:35 -08:00
param-align.ll
OpaquePtr: Bulk update tests to use typed byval
2020-11-20 14:00:46 -05:00
param-load-store.ll
…
pow2_mask_cmp.ll
[TargetLowering] Add i1 condition for bit comparison fold
2020-10-27 12:22:20 +00:00
pr13291-i1-store.ll
…
pr16278.ll
…
pr17529.ll
Revert "[NFC] remove explicit default value for strboolattr attribute in tests"
2021-05-24 19:43:40 +02:00
proxy-reg-erasure-mir.ll
…
proxy-reg-erasure-ptx.ll
…
read-global-variable-constant.ll
…
redux-sync.ll
[Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions
2021-05-17 09:46:59 -07:00
refl1.ll
…
reg-copy.ll
…
reg-types.ll
…
rotate.ll
…
sched1.ll
…
sched2.ll
…
sext-in-reg.ll
…
sext-params.ll
…
shfl-p.ll
…
shfl-sync-p.ll
…
shfl-sync.ll
…
shfl.ll
…
shift-parts.ll
…
simple-call.ll
…
sm-version-20.ll
…
sm-version-21.ll
…
sm-version-30.ll
…
sm-version-32.ll
…
sm-version-35.ll
…
sm-version-37.ll
…
sm-version-50.ll
…
sm-version-52.ll
…
sm-version-53.ll
…
sm-version-60.ll
…
sm-version-61.ll
…
sm-version-62.ll
…
sm-version-70.ll
…
speculative-execution-divergent-target.ll
…
sqrt-approx.ll
[DAGCombine] Remove the check for unsafe-fp-math when we are checking the AFN
2021-01-11 02:25:53 +00:00
st-addrspace.ll
…
st-generic.ll
…
surf-read-cuda.ll
…
surf-read.ll
…
surf-write-cuda.ll
…
surf-write.ll
…
symbol-naming.ll
…
TailDuplication-convergent.ll
…
tex-read-cuda.ll
…
tex-read.ll
…
texsurf-queries.ll
…
tid-range.ll
…
tuple-literal.ll
…
vec8.ll
…
vec-param-load.ll
…
vector-args.ll
…
vector-call.ll
…
vector-compare.ll
…
vector-global.ll
…
vector-loads.ll
…
vector-select.ll
…
vector-stores.ll
…
vectorize-misaligned.ll
…
vote.ll
…
weak-global.ll
…
weak-linkage.ll
…
wmma.py
[Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions
2021-06-29 15:44:07 -07:00
zeroext-32bit.ll
…