de en es fr nl pl pt sv zh

disasm

SASS latency table & instructions reordering

redp

In these difficult times, no one wants to report bad or simply weak results (and this will destroy this hypocritical civilization). Since this is my personal blog and I am not looking for grants, I don't care.Let's dissect one truly inspiring paper - they employed reinforcement learning and claim...

sass registers reusing

redp

Lets continue to compose some useful things based on perl driven Ced. This time I add couple of new options to test script dg.pl for registers reusingWhat is it at all? Nvidia as usually don't want you to know. It implemented in SASS as set of operand attributes "reuse_src_XX" and located usually...

barriers & registers tracking for sass disasm

redp

Finally I add registers tracking in my perl sass disasmNow I can do some full-featured analysis of sass - like find candidates pairs of instruction to swap/run them in so called "dual" mode - and all of this in barely 1200 LoC of perl codeLet's think what must mean for couple of instructions to be...

sass disasm on perl

redp

as an illustration of the use of the modules presented in my previous post I made yet another sass disasm - fully written on Perl. It is almost exact copy of my nvd - implemented just in 460 LoC, the only unsupported feature is registers tracking - bcs I still don't make perl binding for it. What...

perl modules for CUBINs patching

redp

After playing a bit with my ced I came to the conclusion that implemented DSL for editing is not enough - like it would be good to have subroutines to patch repeated/similar instructions, check that patched instruction is what I want, patch attributes/relocs etcIn other words, I need full-fledged...

addresses of cuda kernel functions

redp

Quote from official document:It is not allowed to take the address of a __device__ function in host code I haven't been surprised for a long time that entire CUDA is made up of ridiculous restrictions. What if I told you that paranoid nvidia lies as usually and actually you can get addresses of...