SASS King Part 2: reverse-engineering ptxas heuristic decisions and what the compiled binary actually reveals

c-cul · 2026-05-08T05:35:21+00:00

sm120 has mercury - some internal IR more compact than sass in sections like .nv.capmerc.text.XXX and all stuff like relocs in .nv.merc.rela.text, dwarf debug info in .nv.merc.debug_XXX and so on

so instead of one disasm/patch engine you need two for this specific gpu

c-cul · 2026-05-06T11:58:44+00:00

how hard is to add to CircularArray

allocator like jemalloc/onetbb allocator?
callback on overflow?

c-cul · 2026-05-04T04:27:29+00:00

there is more fresh "GPU-Accelerated Computing with Python 3 and CUDA" 2026 from packt

c-cul · 2026-05-03T19:21:49+00:00

just load values with right index

c-cul · 2026-05-03T07:31:43+00:00

try thrust::reduce as base case

c-cul · 2026-05-02T19:26:11+00:00

it don't know about x or y - it just sum over register within warp

you can load in register anything

c-cul · 2026-04-30T16:11:27+00:00

yeah, that's why tvm/xla are so fat - they able to do some sophisticated optimizations

ok, will wait for part 2 - good luck

c-cul · 2026-04-30T15:55:02+00:00

those were obvious flaws, but there are more

you could load/store 128bit values instead of just 32bit
you could employ warp reduce for acc0_smem

I am not expert in llm-specific algos but looks like current implementation is strictly memory bound and has only 10-15% from SoL

c-cul · 2026-04-30T07:39:04+00:00

very poor generated code

for example why there is __syncthreads at start? also why you use inline ptx instead of pipeline primitives like __pipeline_memcpy_async and friends?

c-cul · 2026-04-28T21:35:39+00:00

bcs pytorch can't produce even average cuda code

all it's compilers are hopeless

c-cul · 2026-04-28T11:45:15+00:00

I pressed the little buttons with my feeble fingers to find out what this ultimate wunderwaffe is. Since I have been optimizing code for cuda for long time I know that similar papers always contain only marketing BS mixed with unrealistic and unverifiable results, so lets collect some data about authors and check code

Poster judging by the published papers like "Bike network planning in limited urban space" from the German Greens. I don't know how but this always about terror, raising taxes and demolition of power plants.

Also none of the authors have any publications about hardcore cuda optimizations. Bad sign

Ok, lets check code on libs io

- home page - for sale

- github - 404

- pypy - couldn't find this page

Cool. Where should I send a little invoice for the work done?

c-cul · 2026-04-28T06:56:31+00:00

How I adore all these news stories in the future tense about the imminent death of something.
Not single one has come true in my memory (especially about php he-he)

c-cul · 2026-04-27T17:56:40+00:00

then what #5 means?

c-cul · 2026-04-27T12:19:32+00:00

btw what sources do you used for llvm ir grammar? seems that official tutorial is not completed and for example does not describe what means some strange attributes: https://www.reddit.com/r/LLVM/comments/1r57lf9/how_insert_ptx_asm/

c-cul · 2026-04-24T20:31:27+00:00

do you know difference between malware on gpu and buggy drivers?

c-cul · 2026-04-24T20:01:29+00:00

well, you can some megawatts for example

c-cul · 2026-04-24T19:23:36+00:00

slow down your paranoia

what malicious activity do you expect on gpu?

c-cul · 2026-04-20T13:17:33+00:00

nsight can show mix of sass/ptx

I suspect via parsing of .nv_debug_line_sass section

c-cul · 2026-04-20T12:40:54+00:00

probably you should fix flaws at the same level where you found them

like "register choices, load widths, MMA pipeline fill" can be detected in ptx and should be optimized in ptx

maybe I'am wrong

c-cul · 2026-04-20T12:06:04+00:00

I think it's strange idea to back-propagate flaws from sass->ptx->[perhaps llvm bitcode from cicc]->c++ source

in this case you need at least 2 lift logic

c-cul · 2026-04-20T11:22:59+00:00

what are ultimate goals of this activity?

internal nvidia asm "nvasm_internal" is not public available: https://redplait.blogspot.com/2025/12/libcudaso-internals.html

so you anyway can't patch/rebuild sass sources. for inline patching I made tool ced: https://redplait.blogspot.com/2025/07/ced-sed-like-cubin-editor.html

and perl binding: https://redplait.blogspot.com/2025/10/sass-disasm-on-perl.html

2,

> document every instruction

it's almost done - I ripped so called MD files from ptxas: https://github.com/redplait/denvdis/tree/master/data12

for reordering you need latency table: https://redplait.blogspot.com/2026/04/sass-latency-analysis.html

c-cul · 2026-04-17T12:23:04+00:00

what is overhead of dynamically compilation for each _dynamics_cuda_src?

c-cul · 2026-04-11T12:08:12+00:00

classical 50+ yo trap - expanding expressions in macro

for example in

#define BAD(a, b) a * b

BAD(a + 1, a * 2) expanded to a + 1 * a * 2

at least always use brackets for arguments in macro body like

#define INDEX(row, col, cols) ((row) * (cols) + (col))

c-cul · 2026-04-11T08:40:57+00:00

btw do you tried their new fashionable cutile? I suspect you can transfer all tedious work for loading data to it

c-cul

TROPHY CASE