I have removed some codes that are doing support work (which are important in production code) outside the kernel and improved occupancy from 18.75% to 25% (launch_bounds(64, 6) to launch_bounds(64, 8)) to test whether it can hide instruction stall effectively to support my hypothesis in #29.
It does. With only 33% increase in occupancy, I am now seeing ~70% performance boost (though probably some of it are due to the reduced register pressure, and it’s still ~29% slower than the second test in post #29). Warp issue efficiency is also up from 37% to 49%. Actual stall reasons are now 90% since I’ve removed support code, and execution dependency stall is zero, which really shows how well Maxwell can hide execution dependency at low occupancy when it is all compute.
Which part in particular do you consider “extremely relevant”? I scanned through the entire clip but have been unable to locate it. In particular, is there any solution being detailed that has a gate count lower than what JanetYellen achieved according to the data given earlier in this thread?
I’d be happy to take a look at your sass if you like.
Or you might want to read this document carefully yourself:
You can use maxas to output a kernel into my format.. though it currently requires you overwrite the cuda7.5 nvdisasm with the one from cuda6.5. Nvidia seems to have changed the output format somewhat and I haven’t had time to update my parser.
In your case you’re probably mainly concerned with the stall counts of the instructions, which is going to be determined by the dependencies present.
But if you’re spending this much time messing with ptx and fighting the compiler, just hand assemble the sass yourself :) It’s actually much easier than programming in ptx with tools I provide. It looks like much of your code could be generated from meta programming.
and got a 2% performance boost and saved a few registers. Considering switching my 18 gates s4 for Roman’s 17 gates one granted less than 1% extra performance, 2% is actually a lot. Little things like this always help.
scottgray, I’m definitely bookmarking that and I’ll see what I can do with it. You are right the code could be generated from meta-programming, though the performance code will need different reordering for different salt value (4096 variations), since the salt alters the data expansion function and subsequently the dependency between instructions. Automatically generating performance code for all 4096 salts might not be trivial.
Once I’m done with this round of optimization in the near future I’ll share the sass or the code (it’ll be open-source in the end anyways).
I really wish I’d written the assembler in python so it could be better integrated with pycuda. But even still, it should be possible to dynamically construct a cubin on the fly and pass it straight to cuModuleLoadData without even writing it to disk. You would just need to embed the perl interpreter which I think it fairly easy to do. Unlike nvcc, there is a trivial amount of overhead in generating a cubin from a sass file.
Then you wouldn’t need to hand map registers to variables (unless you wanted to override that and apply a complicated register coloring to reduce bank conflicts). I’ll get to that eventually.
Which reminds me: LOP.LUT3 takes 3 input registers and there’s only 4 banks so you likely have a lot of opportunity for register bank conflicts in your code. My tool will measure that for you (while also factoring in reuse codes).
So I thought up a clever way to do this mentioned in #10:
By turning this:
d = d ^ x_1; // xor in the SBOXes
s1(d ^ k_1, ...); // xor outside the SBOXes
d = d ^ x_2;
s1(d ^ k_2, ...);
Into this:
d = d ^ x_1 ^ k_1; // lop3(d, d, x_1, k_1, 0x96);
s1(d, ...); // saved one instruction here.
d = d ^ x_2 ^ k_1; // lop3(d, d, x_2, k_1, 0x96);
s1(d ^ k_2, ...);
Unfortunately I can only do this for four out of 16 DES rounds since there’s a “cooldown round” after each use, and there are conditional key swaps. Additionally, only two outputs per SBOX can be optimized this way, since the expansion function duplicates the other two outputs and their duplication are xored with different keys. But even so, by doing this I got another 2% boost in performance.
All codes are now open-sourced here as part of my personal fork of Meriken’s Tripcode Engine, a third-party tripcode finder for 10 character (DES crypt(3)) and 12 character (SHA-1) tripcodes.
Grab the following files if you are interested in:
UNIX DES crypt(3) with salts:
DES_Crypt_##.h in /MerikensTripcodeEngine/Source Files/out/, where ## ranges from 0 to 63. The 64 files contain 4096 functions DES_Crypt_###(), where ### ranges from 0 to 4095 and corresponds with the 12-bit salt value used to perturb the DES expansion function.
Note: initial permutation and final permutation are not included in the function; you need to do it yourself. Normally they’re part of the pre-processing. Both have minimal performance impact.
Suggested register usage or launch bound is 168 regs, or__launch_bounds__(128, 3), even when you wrap things around it.
With this version, I get a performance of 950 MH/s for UNIX DES crypt(3) (or equivalently 23750 MH/s for 1 round of DES) on my reference Gigabyte GTX 980 Ti (+270 MHz). Considering hashcat’s implementation gets 165.5 MH/s on a GTX Titan X (+225 MHz), it’s a great improvement. Even my naive implementation bounded by shared memory/synchronization with old SBOXes from JtR is faster (300 MH/s on 980 Ti +300 MHz).
This further dispelled the myth that Nvidia cards are still bad at DES crypt(3); while it may hold some truth before Maxwell (or even sm_32, since it’s the first version that lift the register limit from 63 to 255) when compared against GCN, it’s no longer the case, and it all depends on the implementation.
What an awesome job! I’m quite impressed with your ability to find effective transformations for complex formulas. Your new S-boxes with the LOP3.LUT instruction are absolutely remarkable. I am so glad I open-sourced my tripcode finder. This kind of work is exactly what I hoped for.
As for your Nvidia/AMD comparison, I am currently getting 800MH/s on my 7990 with OpenCL and rewriting my implementation with a GCN assembler. We will see how that goes :)
That’s right. What a bummer. Oh well. Seriously, you deserve all the praise people would give you there. I would have never imagined I would see contributions from others comparable to the legendary Toriya/chapuni of the mty fame as far as tripcodes are concerned. Now it’s time for me to grab a new 980Ti…
When you’ve got the time you might want to revisit the CUDA tripcode matching routine (DES_GetTripcodeChunks etc, which I haven’t changed); in it’s current form it’s subject to (slight) hit rate loss. This is partly due to its multi-pass nature which allows potential matches to be overwritten by other potential matches, and partly because tripcodes matching as little as five characters are marked as potential matches. The hit rate loss is negligible at ~10k chunks and on MTY for all practical purposes, but grows with both number of chunks and number of passes. I’ve measured the hit rate loss to be around ~0.3% per pass at 262,144 chunks, and around ~95% at 64 passes and 16,777,216 chunks.
Here’s my hit rate benchmark at 32 passes. A pass count between 16 and 32 (with 1% to 3% less performance compared to 64 passes) seems to be a good compromise between performance and hit rate for high-end Nvidia GPUs.
You are quite right about that. I thought about that possibility when I started working on MTE, but I brushed it off as being negligible. However, the ideal length of chunks is hard to determine because, as they get longer, the number of chunks grows exponentially when complex regex patterns are expanded. I will defnitey revisit the issue, though.
On a related note, Toriya/chapuni implemented a pretty neat tripcode matching routine in mty in which Quine-McCluskey is used to find potential matches without the final permutation. With this routine, there would be no need for DES_GetTripcodeChunks(). I will try to incorporate this routine into MTE.
Thank you for sharing the benchmark. I have been playing with your codes on my new 980 Ti for a few days now, and, I must say, your data swapping macros are absolutely amazing. It was quite an experience to see a 100+% speed gain immediately after I applied them to the main loop.
Hi I’m new to CUDA programming. I was interested in your implementation of the DES algorithm, but I could not find it at the link https://github.com/DeepLearningJohnDoe/SLUT/blob/master/sbox.h Could you please drop a link to the source code or help me and answer some questions about how to revise this algorithm. Thank you in advance. Here is my mail - alipov09@mail.ru