cc @EikanWang . I'm disabling this for now since it broke with the H100
merge, but please feel free to fix the compilation errors and submit a
PR.

Also fixes a bug exposed in convertLayout lowering for float16. We
shouldn't be using cvt.pack.sat.u16.s32 to pack 16bits values as this
needs to take a 32bits register. Also this prevented optimization at
llvm ir level.

Issue triton-lang#1973

Co-authored-by: Philippe Tillet <phil@openai.com>

…2051)

…on-lang#2040)

Make sure that other threads within CTA do not operate on mbarrier until
it is initialized by thread 0.

Co-authored-by: Philippe Tillet <phil@openai.com>

Use camel case accessors ("getStaticOffsets" etc.) for `ExtractSliceOp`.
This change works with and without the changes from D156857. After
D156857 has landed, only camel case accessors work for ops that
implement the `OffsetSizeAndStrideOpInterface`.

https://reviews.llvm.org/D156857

Co-authored-by: Philippe Tillet <phil@openai.com>

We are interested in having python wheels for triton built for Linux
arm64 platforms, such as NVIDIA's Grace CPU.

This change is fairly simple, however:
- It requires a linux arm64 build of LLVM to be available (see MR here:
ptillet/triton-llvm-releases#15)
- For now my changes use the LLVM build hosted here:
https://github.com/acollins3/triton-llvm-releases/releases/tag/llvm-17.0.0-c5dede880d17
- The Triton release process will need to be updated to include arm64
wheels. Is this something you have time to work on @ptillet? It would be
difficult for me to update this part without more access permissions.

With these changes, I managed to build a set of python wheels and have
hosted them here for us to use in the meantime:
https://github.com/acollins3/triton/releases/tag/triton-2.1.0-arm64

Co-authored-by: Philippe Tillet <phil@openai.com>

…r than Q's (triton-lang#2033)

Implemented this situation with and without causal mask.
My implementation with causal mask looks like:
111000
111100
111110
Where only the right upper triangle part will be masked.
I added `P_SEQ` for the notation of extra sequence length for KV.

Co-authored-by: Philippe Tillet <phil@openai.com>

This allows the AOT client to tune the number of stages for the
generated kernel. set the default number to 3 to match the triton
compiler.

…in hopper tests (triton-lang#2041)

Co-authored-by: goostavz <gzhu@nvidia.com>
Co-authored-by: Philippe Tillet <phil@openai.com>
Co-authored-by: ben-zhang-609 <110140741+ben-zhang-609@users.noreply.github.com>

Co-authored-by: Allen Zhao <allzhao@nvidia.com>

Improve error messaging for block shape and value shape mismatch.

…#2050)

Co-authored-by: Philippe Tillet <phil@openai.com>

Rename "rocm" -> "hip", to comply with other uses in compiler.py.

…2063)

…riton-lang#2057)

Co-authored-by: Biao Wang <biaow@nvidia.com>

…lang#2067)

…m. (triton-lang#2068)

No functional changes intended, and it might slightly speed up the
build.

This allows a downstream Bazel build of Triton to avoid building a
number of dialects and passes that Triton doesn't need.

`getScratchSizeInBytes` was assuming that the size of all types in bits
is
a multiple of 8. If it is not, it would return 0. This caused a bug for
boolean
(i1) type, where the reduction lowering would attempt to use shared
memory,
which was not assigned to the op.

Fix this issue by setting the number of bytes per element to `ceil(bits
/ 8)`.

libtriton.so is pretty large these days and hashing it is slow.
Switching the hash from md5 to sha1 shaves close to 300ms off the time
for me (as well as being a better hash, for whatever that's worth).

As far as I could tell, sha1 is the fastest stable hash in the Python
standard library, including things like zlib.crc32

Realised I could do this right after my first PR got merged. This saves
another 100ms

…ng#2075)

remove unnecessary skips. decompose UTs in
persistent-warp-specialized-gemm into vintage and stylish

…g#2117)

…ernel names in cache to compare artifacts (triton-lang#2111)

`offset + ptr` and `ptr + offset` both work now

…n-lang#2020)

…g#2125)

Also stop promoting integer types as it doesn't give better perf this
will allow more vectorization oportuinity in the future.

…#2019)

When I'm using Kaggle GPU (https://www.kaggle.com/), I find that
`ldconfig -p` does not show libcuda.so, but requires `ldconfig` (run
with sudo) to refresh the cache to find the libcuda.so.

Therefore, I added this informative message to help users find
libcuda.so.

Forgot to remove this line

Before this PR, the determination of `TritonGPUToLLVMIRPass` to generate
NVVM-compatible LLVM or ROCDL-compatible LLVM is controlled by a boolean
`isROCM`. This method is hard to scale.
This PR changes it to use an enum instead, where new target can be added
easily when needed.

---------

Signed-off-by: Tsang, Whitney <whitney.tsang@intel.com>
Co-authored-by: Philippe Tillet <phil@openai.com>

…ymore (triton-lang#2116)

…on-lang#2136)

Add a new operation to be able to implement packed inline assembly for
elementwise operations. This way inline assembly can be used to control
elementwise operations. It also allows to pack elements to be able to
manually vectorize operations.

…rn (triton-lang#2137)

Use getEffect instead to tell passes whether the op has side effects or
not. This doesn't change functionality otherwise.

Co-authored-by: Philippe Tillet <phil@openai.com>

Replace the Turing version for the dot operation from following Volta
version to following Ampere version.

Update code generator to produce two m16.n8.k8 MMAs for Turing instead
of one m16.n8.k16 MMA we have for Ampere.

A little code cleanup

Disable tf32 if run on sm75 and below
Fix the pattern match to compare the generated ptx against if run on
sm75

`if _unwrap_if_constexpr(cond)` then enters `node.body` is wrong when
cond is a tensor since we cannot statically evaluate a dynamic tensor's
value.

The right way to solve the problem is probably:

1. visit the ast of IfExp (do not build IRs)
2. get the type of the last statement
3. initialize the return value and assign it to livein
4. call visit_If

…n-lang#2145)

…-lang#2143)

Simplify the code by using inline asm to implement globaltimer and smid
instead of relying on bc file.

)

For warp specialized persistent kernel, the instruction sequence for
Warp Groups are
```
// warp group 0
for wave in 0..num_waves:
    idx = wave * num_inner_loop_steps;
    for k_tile_idx in 0..num_k_tiles:
        mbarrier.wait EB[idx];
        W0;
        mbarrier.arrive FB[idx];
        idx++;
```
```
// warp group 1
for wave in 0..num_waves:
    idx = wave * num_inner_loop_steps;
    for k_tile_idx in 0..num_k_tiles:
        mbarrier.wait FB[idx];
        R0;
        mbarrier.arrive EB[idx];
        idx++;
```
then this would form a sequence of morally-strong relations W0 -> R0 ->
W1 -> R1 in causality order.
But if GEMM K is small than K-TileShape, then the num_inner_loop_steps
of persistent kernel is 0. The buffer id and mbarrier id will always be
0 in this case. And it may form W0 -> W1 -> R0 -> R1 order, which is
contradicts with the atomicity --
"If a read R precedes an overlapping write W in causality order, then R
cannot read from W."

…n-lang#2135)

1. Optimize the conversion and packing for 2xf32 -> 2xf16.
2. Split TMA store block into multiple slices of size 64x64.
3. Distribute the TMA store to all the warps.
4. Fix some naming issue.