Comments (2)
An early study shows that this is due to vectorized loading, odd dimensions may cause unaligned address accessing hence the error. Perhaps the best solution is to set ScalarsPerLdg
to 1 to disable vectorized loading.
from cutlass.
@herberteuler 's comment is correct.
Note, we have substantially refactored CUTLASS with version 2.0 released in November 2019. The HGEMM kernel should not be performing vectorized loads and is not expected to have this issue.
There also exist kernels targeting Tensor Cores for mixed-precision matrix multiply with varying alignment restrictions (making performance tradeoffs). You could consider trying those to see if they address your use-case.
Closing this issue. Feel free to re-open if you have any issues with the current codebaase.
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.