Comments (5)
as to your q1. 884 is the only one available on a100 for fp64
from cutlass.
on ampere, tf32 is using 1688. fp16/bf16 is using 16816. s8/u8 is using 16832. s4/us is using 16864. b1 is using 168256.
from cutlass.
@hwu36 Thanks for answering. If I understand correctly, it means the more bits the element type has, the fewer elements the tensor core can process. To this, I have an extended question:
In example_14_ampere_tf32_tensorop_gemm
, I changed the datatype from float
to cutlass::half_t
, and I found m16n8k8
can still work, does this mean that there are some computing power wasted in each warp MMA? As it can do m16n8k16
, but I only give it m16n8k8
.
from cutlass.
1688 is turing era mma, it still can work on ampere. but the bigger, the better unless your problem size is small.
from cutlass.
@hwu36 Thanks!
from cutlass.
Related Issues (20)
- [QST] How to do Batched Gemm with python API
- [BUG] undefined symbol: cuTensorMapEncodeTiled on CUTLASS 3.5.1 HOT 3
- [QST] [CuTe] Why `right_inverse` does not work as expected when shape is not static? HOT 2
- [BUG] undefined reference to `cutlass::library::Handle::~Handle()' collect2: error: ld returned 1 exit status HOT 2
- [QST] How to implement non-interleaving partitioning in CuTe? HOT 5
- ldmatrix instruction HOT 1
- [QST] Why do we only need the result of the last k-loop in `cute::gemm` dispatch-5? HOT 5
- [BUG] ElementC=void kernel reads non-void in `GemmDescription` HOT 1
- undefined reference to `cutlass::library::Handle::~Handle()' collect2: error: ld returned 1 exit status HOT 6
- ldmatrix instruction HOT 1
- [QST] Error: Could not detect active GPU device ID [CUDA driver is a stub library] HOT 3
- [QST] TMA Reduce ADD HOT 3
- [FEA] Support .shared::cta to .shared::cluster TMA copy? HOT 1
- [BUG] example 57 Hopper Grouped GEMM Disposition failed
- [QST] Getting `CUDA error: misaligned address` when calculating GEMM with a small matrix, and multiple other questions on understanding Cutlass. HOT 4
- [QST] Understanding layouts in CuTe MMA atom traits HOT 1
- I also noticed that the option of FP8 operand and FP16 accumulator is excluded in `ss_op_selector` and `rs_op_selector`, even though this is accommodated by the PTX ISA and otherwise has `MMA_Traits` structs defined for it (F16E4M3E4M3 etc.). This is done by the following code: HOT 1
- [QST] About "Threadblock Swizzle", why need this? HOT 1
- [QST] Usage of CUTLASS in non-matmul kernels
- [QST] How to understand iter mode from ThrMMA::partition_fragment_{A,B,C}? HOT 3
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 cutlass.