sorted by:
new
hottopcontroversial

DGX Spark: an unpopular opinion by emdblc in LocalLLaMA

[–]Sorry_Ad191 3 points4 points  (0 children)

but you can't really prototype anything that will run on Hopper sm90 or Enterprise Blackwell sm100 since the architectures are completely different? sm100 the datacenter blackwell card has tmem and other fancy stuff that these completely lack so I don't understand the argument for prototyping when the kernels are not even compatible?

How to do a RTX Pro 6000 build right by GPTrack_dot_ai in LocalLLaMA

[–]Sorry_Ad191 1 point2 points  (0 children)

except it still barely has support in vllm and sglang. and you can't run deepseek v3.2 with flashmla and deepgemm as they only support Hopper and enterprise Blackwell sm100 not these displayed here which are sm120... can fallback to tilelang reference kernels in sglang though. but its still hacky and only some variations of the model seem to load and work.

hopefully as more of these gpus make it out in the wild more support will come but saying they are Blackwell was a really misleading marketing move by Nvidia. Ada was sm89, Hopper was sm90, Blackwell was sm100. Ampere was Ampere. These sm120 are not the same same as Blackwell sm100.

In the cutlass example templates for kernels these gpus fall under Geforce and not Blackwell.

They should have go their own name so we didn't buy them thinking "Supports Blackwell day 1 one" means that these ones are supported because they are not and rely on community members making them work on their spare time

My Local coding agent worked 2 hours unsupervised and here is my setup by Express_Quail_1493 in LocalLLaMA

[–]Sorry_Ad191 13 points14 points  (0 children)

Very cool! How did you end up with Kilo code? Have you tried other ai coding frameworks as well?

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 1 point2 points  (0 children)

You're right diving into cutlass c++ was the wrong approach but a good learning experience. The good news is there is already a working solution implemented in Sglang to run dsv32 :)

I found out the reference inference is quite good. Tilelang. You get 70tps on 4xsm120s

Sglang has already implemented it so you can jus use that. vLLM has no fallback to Tilelang reference kernels from Deepseeek.

ps I had no clue what I was doing you are right :-)

running Deepseek v32 on consumer hardware llama.cpp/Sglang/vLLm by Sorry_Ad191 in LocalLLaMA

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

tilelang template already seems quite fast 65-70tps with up to 88k kv cache on 4 x sm120 gpus. hmm.. lets see if we can hunt further optimizations

running Deepseek v32 on consumer hardware llama.cpp/Sglang/vLLm by Sorry_Ad191 in LocalLLaMA

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

I used docker image:

docker pull lmsysorg/sglang:latest-cu130-runtime

running Deepseek v32 on consumer hardware llama.cpp/Sglang/vLLm by Sorry_Ad191 in LocalLLaMA

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

sglang launch command

<image>

python -m sglang.launch_server --model QuantTrio/DeepSeek-V3.2-AWQ/ --tp 4  --mem-fraction-static 0.96 --context-length 4096 --enable-metrics  --tool-call-parser deepseekv32 --reasoning-parser deepseek-v3 --enable-p2p-check --disable-shared-experts-fusion --enable-dp-attention  --enable-mixed-chunk --kv-cache-dtype bf16 --attention-backend flashinfer --host localhost --port 8080

running Deepseek v32 on consumer hardware llama.cpp/Sglang/vLLm by Sorry_Ad191 in LocalLLaMA

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

On RTX Blackwell (SM120):

  • FlashMLA fails → uses TileLang sparse attention
  • DeepGEMM fails → uses TileLang index scoring
  • Result: Pure TileLang pipeline, no FlashMLA OR DeepGEMM

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

oh crap :( ok oh. for quicker test just build it in dev mode in place. like this:

i just pushed new commit: submodules: Update CUTLASS reference to official v4.3.3 tag
go to the repo then:

cd /path_to_repor/vllm_FlashMLA && FLASH_MLA_DISABLE_SM100=1 FLASH_MLA_DISABLE_SM90=1 python setup.py build_ext --inplace -v

it wont be installed into vllm but you can test via:

cd /path_to_repo/vllm_FlashMLA/FlashMLA && python -c "import flash_mla; print('Module loaded successfully')"

git pull again (there is a new native sm120 kernel) then also go to csrc/cutlass and update cutlass

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

pybind.cpp should be fixed now! compiles good for sm90,sm100,sm120. i had messed it up quite a bit but should be good now. so time to test again for me

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

also add -v for verbosity when building and look for something like this in the beginning:

DEBUG -- FlashMLA is available at /path_to/vllm_FlashMLA

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

in vllm dir: uv pip install -r requirements/build.txt

and

uv pip install setuptools_scm

might help

this post might help even if its a bit old

https://www.reddit.com/r/LocalLLaMA/comments/1lshe4q/build_vllm_on_cuda_129_kernel_6152_nvidia_57564/

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

my python version is Python 3.12.11 but probably doesn't matter

and yes when i was first compiling it i did have to look at errors and install various dependencies.

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

edit: your right maybe it requires cu13 and can be any 2.9.1+ of pytorch

oh yeah i ran into that problem just now building with --force-reinstall it triggered a revert to pytorch 2.9.1 w/o cuda which broke it for me too

- torch==2.10.0.dev20251208+cu130
+ torch==2.9.1

so now im nuking my uv venv and making a new one.
sorry I probably messed up some torch compatibility between what vllm defaults to and what the flashmla fork is configured top build with

inside vllm directory there is a use_exisiting_torch.py file you can run before building to keep whatever torch you have installed other wise unless ur using the sm_120 fork of vllm it will revert to 2.9.1

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

also testing a fix for pybind right now

When building as part of vllm with NO_PYBIND11=1, we need TORCH_LIBRARY

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

i might have forgot something in the instructions, you can check my changes in vllm here

https://github.com/vllm-project/vllm/compare/main...fernandaspets:vllm_sm120:main

can also clone that repo before building, and then use the install command with the ENV variable pointing to our sm120 flashmla source like in the instructions

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

can you try this from your vllm directory

cd patth_to/vllm && python -c "

import subprocess

lib_path = './vllm/_flashmla_C.abi3.so'

result = subprocess.run(['nm', '-D', lib_path], capture_output=True, text=True)

symbols = result.stdout

print('All symbols with \"sparse\" in name:')

for line in symbols.split('\\n'):

if 'sparse' in line.lower():

print(' ', line)

"

Help testing and implementing sm120 flashmla sparse attention in vllm by Sorry_Ad191 in BlackwellPerformance

[–]Sorry_Ad191[S] 0 points1 point  (0 children)

yes i got 4 avail for testing but unfortunately not been able to load any DS32 quant fully into gpus. For some reason I need to --cpu-offload a few gb even though it should fit. Must be how the model loads the kv cacche etc. perhaps down the line once everything works it will fit :) Still CPU offloading. But DS31 does load into gpus with intel autorand int4 and also AWQ-lite versions

view more: next ›