ROR 7 on Fermi should be possible with [font=“Courier New”]vmad.u32.u32.u32.shr7 r0, r0, 128, r0[/font] although that consumes an extra register to hold the constant 128. Similarly al other may be handled with LSR + MAD or with LSR + BFI, although at the expense of extra registers to hold the constants. I’m not sure though whether this can lead to an overall speedup.
If you compile ROL or ROR idioms where shift count is a compile-time constant on Fermi, you will likely find that each ROR or ROL is respresented by a pair of SHR + ISCADD. Of course ROR and ROL are freely interchangeable, just the shift count changes.
Thank you for the information. I have examined the assembly output for
rotateright(x,bits) (((x & 0xffffffff) >> bits) | (x << (32 - bits)))
and it does indeed become SHR + ISCADD. Impressive!
Unfortunately, that probably also means that there’s little room for reducing AMD’s 3x speed advantage in this area.
Excuse my naivité (I am by no means a crypto expert). Is there a good reason why the bit slice method - which has proven to work greatly on DES encryption for example) is not applied in implementations of the SHA2 hash family? Essentially it interprets the 32 bit registers of the CPU (or GPU) as 32 independent single bit registers.
Wouldn’t be faster with the SHA series, since hashes mix integer addition with logical operations. The combined state of 32 SHA256 operations would also require 32*16 registers, which cannot all be addressed by PTX instructions.
Actually, SHA256 mixes 512 bits of state into 256 bits of hash, so you would need 768 registers to do 32 of them at once without spilling to memory.
Bumping this thread because of sm_30 and sm_35.
A couple thoughts:
sm_30: It’s probably not a win at all but you could scatter a 32-bit word’s bits across 32 lanes and use SHFL to perform rotations. XOR’s, AND’s and NOT’s would remain the same and be performed in parallel. It sounds like a good idea but the scattering before rotation and gathering before the additions would probably wipe out any benefit.
sm_35: CUDA finally has a ROL/ROR operation with the SHF.L/R opcode. It’s probably worth seeing what kind of performance gain this provides. If the Elcomsoft results are any indicator, there should be a solid improvement.
Last night after dinner I banged out a fully unrolled SHA256 kernel using macro expansion hackery. Unfortunately, NVCC 5.0 is stack dump’ing on the SHF.R operation on the ‘Live Variable Analysis’ pass.
I’ve filed a bug but my curiosity won’t be satisfied until it’s fixed. :)
Did you code inline PTX or simple C code? The compiler should recognize standard source code idioms for rotates and use the funnel shift instructions for them. It all worked fine when I tried it. The NOTs should get absorbed into modifiers in the LOP.AND and LOP.XOR instructions, although I recently discovered that this is not always the case in CUDA 5.0 (bug filed). With SHA1 one could reformulate the boolean logic slightly for an incremental performance increase, not sure whether that applies to SHA256. The goal would be to increase ILP and minimize operation count, keeping in mind that ANDN as well as XORN have single-instruction representations at SASS level.
@njuffa, You beat me to it. I was just about to write that I dumped the SASS and … the compiler spotted the standard “(a >> n) | (a << (32 - n))” idiom and has peppered the code with SHF’s. :)
The NVCC bug (#235811) is still nasty though. The crash occurs on the inlined SHF PTX. No need to use it at this point!
It’s very impressive – sm_35 only requires 37 registers in my SHA256 routine. sm_30 requires 62 while sm_21 and sm_12 use 61.
Not sure what that compiler bug is about, I have not encountered that. Thanks for filing a bug report. Did you check whether all the NOTs have been absorbed into subsequent LOP instructions?
I see the LOP.PASSB with a suspicious twiddle:
LOP.PASS_B R16, RZ, ~R7;
Is that what you were expecting? I assume this means the NOT wasn’t absorbed?
There are 59 of them so that looks good to me! I guess it’s not 64 since the first chunk starts with constant hashes of which a few are probably NOT’d at compile time before being mixed.
That is SASS for “NOT R7”, a “logic operation, pass-through source operand B” with a negation modifier on operand B. So it looks like this is one of those cases where the CUDA 5.0 compiler does not absorb the NOT into the dependant LOP instruction. For my code I have worked around this by use of inline PTX so you might want to give that a try if you want to achieve speed-of-light performance.
Are you saying that I can apply a twiddle to an argument in ‘and.b32’ and ‘xor.b32’ (which is not in the PTX manual) or are you saying that an inlined PTX “not,and” sequence should coax the current compiler into generating the proper SASS?
The negation modifier (~) is a feature of the hardware and thus SASS, not PTX. Same approach as is used with the negation and absolute value modifiers on the input to floating-point instructions.
In PTX, you will have to code “NOT” followed by separate “AND”, “OR”, or “XOR”. PTXAS then knows how to merge the “NOT” into the negation modifier of a dependent logical operation. The reason the CUDA 5.0 toolchain does not discover all the opportunities for merging NOT is that the frontend sometimes transforms Boolean expressions in such a way that there is no standalone NOT left for PTXAS to merge (this is easily seen by examining the PTX). By coding the “NOT” directly via inline PTX, PTXAS now has the opportunity to apply the optimization and save an instruction.
I have used this inline PTX approach in various places of the SIMD-in-a-word emulation paths. e.g.:
asm ("not.b32 %0,%0;" : "+r"(b));
Hmm… no joy in PTX-ville. I tried several variations and verified the PTX results in NOT + AND being adjacent. The SASS still appears to performing all NOT’s via LOP.PASS_B.
That is very odd and unexpected. Can you post your code? The NOT and AND do not need to be physically adjacent for the optimization to occur. It suffices that the result of the NOT is fed to the input of the AND, as the compiler works off the dependency graph. It may be the case that the result of the NOT is not only needed as an input to the AND but also has additional uses and merging is not appropriate as it would just cause redudant computation. In any event I would like to take a look at the code to make sure one way or the other. If you prefer not to post the code to this thread, a PM through the forum is an alternative.
I tried to paste a screen-capture but it didn’t work. Here is the “notand” routine with several variants – all appear to work but none appear to absorb the NOT:
//
// NOT AND
//
DEVICE_FUNCTION_QUALIFIERS
beu32
notand(beu32 a, const beu32 b)
{
#if __CUDA_ARCH__ >= 999
beu32 d;
asm ("not.b32 %0, %0;" : "+r"(a));
asm ("and.b32 %0, %1, %2;" : "=r"(d) : "r"(a), "r"(b));
return d;
#elif __CUDA_ARCH__ >= 999
beu32 d;
asm("{ "
" .reg .u32 t; "
" not.b32 t, %1; "
" and.b32 %0, t, %2; "
"} "
: "=r"(d) : "r"(a), "r"(b));
return d;
#elif __CUDA_ARCH__ >= 100
beu32 d;
asm("not.b32 %1, %1; "
"and.b32 %0, %1, %2; " : "=r"(d), "+r"(a) : "r"(b));
return d;
#else
return ~a & b;
#endif
}
I confirmed your observations with regard to CUDA 5.0. The merging happens for sm_1x, but not sm_2x or sm_3x (I tried various different idioms). I checked my old notes and merging of NOT into a dependent AND / OR/ XOR operation was working as desired in mid 2012, with the then current internal toolchain. I will follow up on this with the compiler team. Thanks for pointing this out.
Ah, I see the merged “LOP.AND D,~A,B” on sm_12. Very cool:
LOP.AND R45, ~R44, R46;
Thanks!
Has anyone confirmed whether sm_35 provides significant performance benefits in SHA-256 hashing?
Christian