redp
Optimal instructions scheduling is NP-hard task. For this reason almost all compilers implement metaheuristic methods like list scheduling/Gibbons–Muchnick algorithm etc. ptxas is no exception - it also generates non-optimal scheduling, and this opens some opportunity for automatic optimization....
redp
Official documentation doesn't disclose ConstBank0 (c[0x0]) memory layout used at the SASS level for kernel arguments and special registers (like %gridid, %nctaid) So I've spent last week trying to solve this deceptively simple problem. Names of params are documented in official doc - seems that...
redp
Parts 1, 2 & 3First of all, it should be noted that the mask of instruction attributes has size 20 bytes, so I updated dump for them.structure for this attributes descriptor has size 0xd8 bytes and some fields:mask at offset 0name of instruction at 0xC8index at 0xD0 Instructions selecting...
redp
Parts 1 & 2Pseudo instructionsSurprise-surprise - some PTX instructions not mapped directly to underlying SASS 1:1. Instead they generate lots of another PTX code. I already extracted their decrypted bodies, so it's time to describe how they connected to specific PTX pseudo instructions There...
redp
Part 1PTX instructions that cicc cannot generateWhile reverse-engineering Nvidia's compilation pipeline, I extracted the set of PTX instructions that cicc (the CUDA C++ frontend) is capable of emitting. The next logical step is to intersect them with full set of instructions accepted by ptxas - so...
redp
DisclaimerHighly likely that author is an illiterate, inattentive, and incompetent lazy person with a poor imagination - therefore his hypotheses may be questionable, ideas delusional and his analysis simply incorrect. Also maybe I still haven't mastered ida pro in 28 years so extracted data can...
redp
After extracting latency table I became curious how good the code produced by ptxas. Projects like CuAsmRL never estimated limits of profit after rescheduling - it's strange and looks even worse than famous "proof left as an exercise to the reader" - what if ptxas generates perfect code and there...
redp
method 1requires building of hijacked .so with appropriate LLVM version. I am too lazy for this method 2cool, but does not work - cicc claims on bad arguments. I've tried many combinations with no luck But hey - we are under linux and can make many hacks, for example check what arguments genuine...
redp
I've add some support of DWARF debug info from nvidia nvcc to my dwarfdump. As everyone knows dwarf is over-complicated, fat and just disgusting - however, nvidia was able to take his nausea to a new levelrelocstheir cuda-gdb does not contains reloc_howto_type for CUDA relocs - it's special kind...
redp
In my first attempt I used latency tables extracted from MD file (located inside nvdisasm) and nothing good came out of itObvious reason is that real latency table should be located not in disassembler - it must be inside ptxas. But the problem with that file is that it is really huge - in SDK 13...
redp
I've done some research of libcudadebugger.so internals - seems that it has exactly the same patterns:functions table returned by GetCUDADebuggerAPI located in .data section so you can patch any callback addressand each API function has logger This last fact is strange - while loggers from...
redp
For some reason cuda-gdb from cuda sdk gives on my machine list of errors likeTraceback (most recent call last): File "/usr/share/gdb/python/gdb/__init__.py", line 169, in _auto_load_packages __import__(modname) File "/usr/share/gdb/python/gdb/command/explore.py", line 746, in...
redp
inconvenient cuda-gdb can't automatically processing them - you need explicitly say something liketarget cudacore /full/path/to/coredumpand then type lots of info cuda XXX So last weekend I wrote tool to parse/dump CUDA coredumps and it even works on machine without CUDA SDK (what might be useful...
redp
As illustration of ideas from my previous blogpost I made PoC for logging all libcuda.so calls - as the cuda-gdb debugger sees themIt just installs own debug handler and receives all messages. Note:only x86_64 linux supported, but logic can be easily extended for x86 32bit and highly likely for...
redp
Previous partI've noticed that almost all real API functions has the same prologues like: mov eax, cs:dword_5E14C00 ; unique for each API function mov [rbp+var_D0], 3E7h mov [rbp+var_C0], 0 mov [rbp+var_C8], 0 test eax, eax jz short loc_39603B lea ...
redp
The first question that comes to mind when looking at them is "why they are so huge?". For example libcuda.so from cuda 10.1 has size 28Mb and from 13.1 already 96Mb. So I rejected the idea that they are just yet another victims of vibe-coding and made some preliminary RE. The answer is - because...
redp
Spent couple of days in debugging rare bug in my sass disasm. I tested it on thousands of .cubin files and got bad instruction decoding for one. Btw I never saw papers about testing of disassemblers - compilers like gcc/clang has huge set of tests to detect regressions, so probably I should do the...
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...
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...
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...
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...
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...
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...