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
While doing some reverse engineering of ptxas I discovered that their lexer was generated by lex in fast mode (lex -f). Knowing that nvidia trying to hide from us as much as possible it would be good to extract what tokens their lexer able to consume. Surprisingly I was unable to find in google...
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
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
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
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...
0x434b
Last time we left off with a pretty decent understanding about how our router is structured and what components were used. We also found two interesting debug pads that showed oscillating voltages during boot up. In this post, we will take a closer look at exactly these and try to get a Read-Write...
0x434b
This is a write-up for solving the devils-swapper RE challenge. It was mostly intended for my personal archive, but since it may be interesting to all of you. This especially applies if you're still rather new to the whole RE world, as the write-up turned out to be quite verbose. I hope it is...
0x434b
It has been a while since I did some hardware hacking, and this time I want to review the basics. The LinkSys EA6100 router intrigued me since I was only able to find encrypted firmware images (or updates). Known tools like binwalk were unable to unpack the system:$ file...
0x434b
Welcome back to part 2.2 of this series! If you have not yet checked out part 1 or part 2.1, please do so first as they highlight important reconnaissance steps as well as the first half of the disassembly analysis! Let's recall the current functionality we've encountered based on the developed...
0x434b
Welcome back to part 2 of this series! If you have not checked out part 1 yet, please do so first, as it highlights important reconnaissance steps!So let us dive right into the IDA adventure to get a better look at how imgdecrypt operates to secure firmware integrity of recent router models.We'll...
0x434b
Recently, we came across some firmware samples from D-Link routers that we were unable to unpack properly. Luckily, we got our hands on an older, cheaper but similar device (DIR882) that we could analyze more closely. The goal is to find a way to mitigate the firmware encryption that was put in...
0x434b
Note: This is a re-upload of an old write-up.This is another write-up from an interesting little challenge. The original forum post about it can be found here. To get your hands on the challenge I've prepared the base64 text representation of it once again below so you can try it yourself. base64...
0x434b
Note: Re-write/Re-upload due to dead linksThis write up are my thoughts and steps to statically analyze a given unknown binary. I want to understand the binary to a point where I can freely write about it. So here it is. I'm always open for you pointing out mistakes or giving feedback to me. The...