Comments (8)
OK,
I come to answer my own question again. I tried --m=10240 --k=4096 --n=4096 on TITAN X with CUDA 9.2, cutlass performance is about more than 95% of cublas.
from cutlass.
Hmm, it is strange.
- I notice what wrong is for my first time. I should use sgemm_nn to compare with cublas.
- On TITAN X with CUDA 9.2, the performance is 95%.
- On TITAN Xp with CUDA 9.0, the performance is also 95%.
well, on jetson TX2 with CUDA 9.0, jetpack 3.2, the performance is 70%.
Could some one give some suggestions?
from cutlass.
I don't think cutlass is capable of doing the same automatic heuristic based tuning as cublas. it'd probably take some manual tweaking to get up to the 95% mark. even though both devices are pascal arch, the core count and probably other important factors are way different. But even then are you ever going to make up that 5% with a fused epilogue function? that's what I'm trying to figure out atm, but my guess is no.
from cutlass.
Hi Jeff,
I profiled the whole system with nvvp, it is worth to do that~
from cutlass.
NVVP is an awesome tool. I look forward to getting the same results then.
Let us know if you find a way to get good cutlass results on your TX2, perhaps with different block sizes.
from cutlass.
I just want to confirm what wtiandong stated (contrary to my original guess) that, in my case as well, it would be beneficial to implement cutlass for performance gains. My simple epilogue of merely a bias and relu is taking 9% of my forward pass time. Nearly all of that 9% is coming from mem reads that would be unnecessary with an epilogue built into the gemm kernel. There would be room for significant gains with cutlass going at 95% of cublass speed. Unfortunately I use strided batch gemms, so I'll need to wait for those to be implemented before I can adopt. Looking forward to it.
from cutlass.
Hmm, I'm little busy these days... I tried to change the block size/tile size a little, but had no good result - either build failed or slow speed. I think it needs more deep analysis.
If you use TITAN or Geforce, you can directly use cutlass. The slow speed only occurs on TX2.
from cutlass.
Sorry to let this issue languish for so long. Several comments.
> Unfortunately I use strided batch gemms, so I'll need to wait for those to be implemented before I can adopt. Looking forward to it.
Batched strided GEMMs are implemented in CUTLASS 2.x!
> I don't think cutlass is capable of doing the same automatic heuristic based tuning as cublas.
Correct. CUTLASS performs no heuristics or decision-making on its own, except in selecting the first functionally sufficient kernel in the recently added host-side API in the CUTLASS Library.
I'd like to close this issue as the code has undergone significant changes since it was opened. Feel free to re-open to resume this discussion.
Thanks!
from cutlass.
Related Issues (20)
- [BUG] print_layout seems to change the swizzle mode of its input HOT 5
- Warp Group MMA vs Warp MMA HOT 1
- Tiled copy misaligned, how to solve it?
- [QST] (BUG?)The stride of TensorNCxHWx seems to be confusing when C is smaller than Interleave HOT 4
- [QST] GEMM Epilogue Fusion: Element-wise Ops and Two-Tensor Element-wise Multiplication HOT 7
- [QST]Why fp8 convert only has float2fp8 function without ptx ? HOT 1
- [QST] GEMM Epilogue Fusion: Row-wise and Column-wise Multiplication HOT 1
- [FEA] Add cuTensorMapEncodeTiled to CudaHostAdapter HOT 1
- [QST] CUTLASS kernels appear to be significantly slower than CuBLAS for an fp16 gemm on `sm_75` HOT 4
- [QST] Equality of shapes HOT 1
- How to perform operations like crop, concat on tensors in CuTe? [QST] HOT 2
- [BUG] Failing to build on MSVC due to call to _div128 HOT 4
- [QST] Is there grouped_gemv
- Int8 multiplication with pytorch extension: namespace "torch" has no member "I8
- [QST]What is the difference between `WmmaTensorOp` and `TensorOp`? HOT 2
- [DOC] Incorrect link in main README file
- [BUG] Circular Dependency in Header Files HOT 2
- [QST] Tiling an MMA in the K dimension HOT 3
- [QST] How to run GEMM with CUDA Graph?
- [BUG] CUTE: zipped_divide function returns different results on cuda device and on host HOT 6
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.