Comments (5)
I think it should be $0, $1
instead of $0, $0
.
from triton.
Thank you for looking into this!
@ThomasRaoux , what instruction are you using for
ASM
? I triedcvt.rna.tf32.f32 $0, $0;
on my A100 and it failed (produced all 0s).Wrt truncation: do you know why
tf32
with truncation would be less accurate thanfp16
? They should ideally be using basically the same number of input bits of precision, so I wouldn't expect an almost order-of-magnitude accuracy loss.
Sorry I missed to copy it... Here is what I ran that matched exactly torch:
@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, N: tl.constexpr, allow_tf32: tl.constexpr):
ASM: tl.constexpr = "cvt.rna.tf32.f32 $0, $1;"
offs = tl.arange(0, N)[:, None] * N + tl.arange(0, N)
a = tl.load(a_ptr + offs)
b = tl.load(b_ptr + offs)
if allow_tf32:
a = tl.inline_asm_elementwise(ASM, "=r, r", [a], dtype=tl.float32, is_pure=True, pack=1)
b = tl.inline_asm_elementwise(ASM, "=r, r", [b], dtype=tl.float32, is_pure=True, pack=1)
c = tl.dot(a, b, allow_tf32=allow_tf32)
tl.store(c_ptr + offs, c.to(a.dtype))
from triton.
This has to do with the rounding applied to go from f32 to tf32. By default the tensorcore will truncate the f32. If you apply rounding before by doing something like that then you get the same precision as torch:
if allow_tf32:
a = tl.inline_asm_elementwise(ASM, "=r, r", [a], dtype=tl.float32, is_pure=True, pack=1)
b = tl.inline_asm_elementwise(ASM, "=r, r", [b], dtype=tl.float32, is_pure=True, pack=1)
It's not something we want to do implicitly do in triton as it can affect performance significantly and allowing user to control this sounds better.
On Hopper applying such rounding would break the pipeline as tensorcore inputs come from shared memory.
from triton.
This has to do with the rounding applied to go from f32 to tf32.
I'm looking into Triton TF32 precision issues from a related context, and just "discovered" the same: my experimenting involves CUTLASS too, that gives the same error as cuBLAS when tensor cores (i.e. TF32) used, so I was able to compare PTX generated from Triton and from CUTLASS, and it could be seen in the PTX that CUTLASS is applying rounding, while Triton is not.
from triton.
Thank you for looking into this!
@ThomasRaoux , what instruction are you using for ASM
? I tried cvt.rna.tf32.f32 $0, $0;
on my A100 and it failed (produced all 0s).
Wrt truncation: do you know why tf32
with truncation would be less accurate than fp16
? They should ideally be using basically the same number of input bits of precision, so I wouldn't expect an almost order-of-magnitude accuracy loss.
from triton.
Related Issues (20)
- Cannot find 2.0.0.dev20221202 version HOT 2
- make error
- FlexAttention Segmentation Fault HOT 2
- What is the use of divisibility in coalesce pass
- triton build failing to access https://tritonlang.blob.core.windows.net/llvm-builds/ with HTTP Error 409 HOT 17
- Confusion about memory of pointers HOT 1
- [Tutorial] Fused-attention `bwd` tutorial incorrect for `causal=True`
- The question of implementing the addmm_ operator based on triton HOT 1
- Nothing found at https://oaitriton.blob.core.windows.net/public/llvm-builds/llvm-ce80c80d-centos-x64.tar.gz
- Numeric errors with tf32 matmul on A100 GPU HOT 2
- DeepSpeed came to Windows natively - pip install deepspeed - where are you Triton?
- .so文件的作用? HOT 3
- The number of generated TTIR files HOT 2
- PTX version error with CUDA 12.6 HOT 1
- RuntimeError: Triton Error [CUDA]: an illegal memory access was encountered HOT 1
- Cannot build triton HOT 1
- The number of generated TTIR files
- fused_moe
- fused_moe SIGSEGV
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.
from triton.