Cycle reduction in chained SHA-256/RIPEMD-160 device function (Ada / sm_89)

Accelerated Computing CUDA CUDA Programming and Performance
1

Looking for expert @njuffa review on a CUDA device function that chains
SHA-256 followed by RIPEMD-160 entirely in registers (no shared
memory, no global stores between the two hashes). Target architecture
is sm_89 (Ada Lovelace).

WHAT THE FUNCTION DOES

Single device function _GetHash160Comp() that:

  1. Takes a 256-bit input plus a 1-byte parity flag
  2. Runs one SHA-256 block transform (64 rounds, 33-byte payload)
  3. Pipes the 8 × uint32_t SHA-256 output directly into
    RIPEMD-160 as register-resident input (no memory roundtrip)
  4. Runs one RIPEMD-160 block transform (80 rounds with parallel
    left/right lines)
  5. Writes the final 160-bit digest

No shared, no constant lookup tables for SHA-256 K.
RIPEMD-160 still uses an 8-element constant K160 array.

WHAT IS ALREADY IMPLEMENTED

  1. SHA-256 K-CONSTANTS AS IMMEDIATES
    All 64 round constants are define K0..K63 macros, expanded
    inline at each round. This eliminates the typical constant
    memory loads (LDC instructions in SASS) and lets the compiler
    fuse them as 32-bit immediates into the IADD3/IMAD pipeline.

  2. ROTR VIA __funnelshift_r
    ROR is implemented as funnel shift (SHF.R.WRAP on Ada) instead
    of (x>>n)|(x<<(32-n)). This produces single-instruction rotation
    in SASS rather than the 3-instruction shift/shift/OR pattern.

  3. PIPE-PARALLEL ROUND STEP
    Each SHA-256 round computes the next round’s S1/Ch/S0/Maj while
    committing the current round’s state update. Macro
    SHA256_STEP_PIPE_K does:

    t1 = h + s1e + ch + Ki + Wi;
aN = t1 + s0a + maj;
eN = d + t1;
// already compute s1eN/chN/s0aN/majN for next iteration
s1eN = S1(eN); chN = Ch(eN,e,f);
s0aN = S0(aN); majN = Maj(aN,a,b);

    Goal: expose ILP to the scheduler so the next round’s
    computation overlaps with the current round’s state rotation.

  4. PHASE SPLITTING (REGISTER KILL ZONES)
    The 64 rounds are wrapped into four C++ scope blocks { } at
    rounds 0-15, 16-31, 32-47, 48-63. Same for RIPEMD-160. The
    intent is to give the compiler explicit lifetime cutoffs so
    intermediate temporaries don’t keep registers alive across
    phase boundaries.

  5. WMIX_INIT_REGS / WMIX_REGS
    Message-schedule update done entirely in registers — w0..w15
    are uint32_t locals throughout the entire transform. No
    memory backing for the W array. Constants from the hardcoded
    33-byte payload (e.g. 0x00A50000u, 0x10420023u, 0x00000108u
    length encoding) are folded into the first WMIX directly.

  6. SHA-256 → RIPEMD-160 BRIDGE
    The 8 × uint32_t SHA-256 output is byte-swapped via
    __byte_perm and passed as 8 explicit register arguments to
    RIPEMD160TransformRegs(). The remaining 8 input words are
    compile-time constants (0x80, zero padding, length 256).
    Compiler should be able to fold these as immediates.

  7. RIPEMD-160 PARALLEL LINES INTERLEAVED
    The two RIPEMD-160 lines (a1..e1 and a2..e2) are written
    alternating in source order:

    R11(a1,b1,c1,d1,e1, W(0), 11);   R12(a2,b2,c2,d2,e2, W(5), 8);
R11(e1,a1,b1,c1,d1, W(1), 14);   R12(e2,a2,b2,c2,d2, W(14), 9);
...

    Goal: maximize ILP by giving the scheduler two independent
    dependency chains visible at any point.

  8. EARLY FINAL COMBINATION
    In the last RIPEMD-160 phase, s[2] is computed early because
    neither e1 nor a2 change after that point. This shortens the
    critical path on the final state combination by one round
    worth of dependency.

QUESTIONS FOR REVIEW

Q1: REGISTER PRESSURE
Both transforms together push the live state high. SHA-256
needs 8 state regs + 16 W regs + 4 pipe regs (s1e, ch, s0a,
maj) + temps = ~30+ live across the full transform. RIPEMD-160
needs 10 state regs (two parallel lines) + 16 W regs + temps.
Both are inlined into the same kernel.

Is the phase-splitting via C++ scope blocks actually achieving 
the intended register kill, or does the Ada compiler already 
do liveness analysis well enough that the scopes are no-ops? 
Has anyone verified via SASS that registers are actually 
being recycled across the scope boundaries?

Q2: PIPE-PARALLEL ROUND IDIOM
The SHA256_STEP_PIPE_K macro speculatively computes
s1eN/chN/s0aN/majN for the NEXT round while finishing the
CURRENT round. Is this still beneficial on Ada with its
improved scheduler, or is it actually counterproductive
(extra register pressure for marginal ILP gain)?

Q3: __byte_perm BRIDGE
The SHA-256 → RIPEMD-160 bridge does 8 × __byte_perm to
byte-swap the digest. On Ada, __byte_perm maps to PRMT, which
has a fixed latency and competes for the same pipeline as
the integer ALU. Could a manually-unrolled rotation
(ROL16 + AND mask) be faster, or is PRMT genuinely the
fastest path?

Q4: RIPEMD-160 INTERLEAVING
The alternating R11/R12 pattern in source code is intended to
expose two independent dependency chains to the scheduler.
Does the Ada compiler actually preserve this interleaving in
SASS, or does it reorder back into “all line-1, then all
line-2”? If the latter, is there a way to force the
interleaving (e.g. asm volatile barriers between R1x calls)?

Q5: K160 CONSTANT LOADS
RIPEMD-160 still uses constant K160[8] indexed by round.
Would converting these to 8 define immediates (like SHA-256 K)
likely improve SASS, or are LDC instructions on the constant
cache effectively free on Ada due to the broadcast cache?

Q6: f1/f2/f3/f4/f5 BOOLEAN FUNCTIONS
The RIPEMD-160 round functions f1..f5 use mixed AND/OR/XOR/NOT
patterns. Has anyone tried encoding these via LOP3 with
explicit ternary truth tables on Ada? The compiler often
misses LOP3 fusion opportunities for 3-operand boolean
expressions.

hash_sass.txt (452.5 KB)

I’ll attach the relevant SASS dump

Any feedback on the questions above — particularly Q1 (register
kill zones) and Q4 (interleaving preservation) — would be
extremely valuable. I want to understand WHY a given technique
helps or hurts on Ada, not just receive a black-box rewrite.

2

Sorry, I have zero experience with Ada.

Have you found that using the HLL idiom for rotates does not reliably map to funnel-shift instructions at SASS level? Using the device-function intrinsic is obviously the ultimate insurance that happens.

__byte_perm has mapped to PRMT since this HW instruction was added, which prompted the creation of the device-function intrinsic. Generally speaking, this should always be preferred way to re-arrange bytes, unless specific experiments would demonstrate otherwise.

Generally speaking, any arrangements made at HLL level are not very likely to translate to specific idioms at SASS level, as the two compilers (CUDA → PTX, PTX->SASS) often massively re-arrange code as they see fit. The PTX emitted by the first stage is in SSA (static single assignment) form, that is, every virtual register is written to once, so the only thing the compiler sees is the actual dependency chains in the DAG it constructs. The optimizing backend compiler ptxas is responsible for instruction scheduling and register allocation and will deal with dependency chains as it sees fit.

Why not try both variants to see which one is faster? Constants stored as immediates inside machine instructions are obviously the fastest way to access them, and do not require additional instructions. Recent GPU architectures allow one full 32-bit immediate operand in most instructions, and from what I have seen the compiler makes aggressive use of that. __constant__ const data is usually just slightly less efficient than register access, provided access is uniform across all threads in a warp (broadcast feature).

The CUDA compiler has improved LOP3 generation over the years, but one can often do better by hand, but usually not my much, in my experience. Since AVX512 also offers ternary logic instruction, you could search whether someone already did that work for CPU implementations. Generally, mapping a bunch of bitwise logic to the minimum number of LOP3 instructions is a hard problem. I have tried using open-source tool developed for FPGAs to find the best mappings for me (the tools can typically target LUT-based CLBs with up to 5 inpts) but was not successful. This may well be due to inexperience in using these tools. The one way that works is puzzling out “optimal” LOP3 arrangements by hand, and that is what I have been using. The problem might be slightly more tractable by limiting oneself to particular 3-input primitives, e.g. it is possible build arbitrary logic from a majority-3 primitive (which maps to a single LOP3).

Modern GPUs provide 64K general-purpose registers per SM, so register pressure should rarely be of concern. 64 GPRs in use per thread allows for 1024 threads, 128 GPRs in use per thread allows for 512 threads. If you can achieve decent granularity (say 128 threads per thread block), that should utilize resources efficiently. Ultimately, you would want to let the CUDA profilers guide optimization efforts, rather than make assumptions as to what the bottlenecks might be.

or does the Ada compiler already
do liveness analysis well enough that the scopes are no-ops?

Maybe I misunderstand the question, but any optimizing compiler from the past several decades does proper live-range analysis and will map registers accordingly. The CUDA toolchain has always comprised two optimizing compilers, where the frontend compiler has been derived from the LLVM framework for the past 15 years or so and the backend compiler is NVIDIA proprietary and its internals not publicly documented. In consequence: (1) a particular HLL variable may not exist at SASS level (2) there may be variables at SASS level that have no direct equivalent at HLL level (3) a particular HLL variable may occupy a number of different registers during its lifetime (4) a particular register may hold any number of different HLL variables over the execution time of a kernel.

There is an obvious trade-off between register pressure and latency tolerance. The CUDA compiler will in general try to schedule long variable latency instructions such as loads from memory early, which increases live range and can drive up register pressure. In general, these days it handles these trade-offs in near optimal fashion in the vast majority of cases (things were a bit different in the early days of CUDA with register-starved architectures and immature toolchains).

3

Regarding the LOP3 application: Looking at the RIPEMD-160 reference code here

homes.esat.kuleuven.be/~bosselae/ripemd160/ps/AB-9601/rmd160.h

I see five macros F(), …, J(), each of which should map trivially to a LOP3 as they are all functions of three inputs.

You could either implement each macro as a __forceinline__ __device__ function, or use inline PTX to hardcode them as LOP3s, with the following mappings F() → LOP30x96, G() → LOP30xca, H() → LOP30x59, I() → LOP30xe4, J() → LOP30x2d. Either variant should result in the desired mapping to a LOP3.

4

Implemented your LOP3 suggestion empirically. Result was unexpected:

BEFORE (compiler defaults, all f1..f5 as C macros):
LOP3 0x96 (F): 972
LOP3 0xCA (G): 0
LOP3 0x59 (H): 0
LOP3 0xE4 (I): 0
LOP3 0x2D (J): 244
Total LOP3: 2279
IMAD: 3753
IADD3: 3504

AFTER (f2,f3,f4 forced via inline PTX, f1 and f5 unchanged):
LOP3 0x96 (F): 972 (same)
LOP3 0xCA (G): 128 (correctly produced)
LOP3 0x59 (H): 128 (correctly produced)
LOP3 0xE4 (I): 128 (correctly produced)
LOP3 0x2D (J): 116 (HALVED!)
Total LOP3: 2279 (same total!)
IMAD: 3671 (-82)
IADD3: 3601 (+97)

Net result: ~1-2% hashrate REGRESSION.

Two unexpected observations:

  1. The J function (which was already mapping to LOP3 0x2D) lost half
    its LOP3 mappings - apparently because inline PTX in f2/f3/f4 broke
    some kind of cross-function optimization the compiler was doing.

  2. Despite gaining 384 new LOP3s for G/H/I, the total LOP3 count stayed
    the same at 2279. The compiler must have been doing G/H/I as multiple
    smaller LOP3s before, and now it consolidated them - but at the cost
    of breaking J’s optimization elsewhere.

  3. Net IMAD/IADD3 went UP by 15 instructions. The kernel is ALU-bound
    at 80.7%, so this directly costs the hashrate.

hand-forcing LOP3 instructions broke a larger compiler
optimization that was producing better overall results. The compiler’s
holistic view of the dataflow beat my local optimization.

Thanks for the suggestion the empirical data is itself a useful result, even when
the optimization didn’t pan out.

5

I used Compiler Explorer to see how various compilers optimize the RIPEMD-160 reference code, and was surprised by the number of logical operations, which seemed lower than what I expected. I have not looked further, but my working hypothesis is that it is possible to absorb some of the logical operations into surrounding operations. ORs in particular might become part of the summing that is going on., as | and + interchangeable in some expressions as long as certain conditions are met.

The rounds in RIPEMD-160 are basically a mixture of shifts, logical operations, and arithmetic, and some combinations of these operations are possible at the machine instruction level. For example, a left shift that is actually a constituent part of a rotate operation might combine with an IADD and merge into IMAD. An | might ultimately map to a + and merge into an IADD3.

An assumption that mapping all shifts to funnel-shift based rotates and all logical operations into LOP3s would result in the minimum instruction count does therefore not necessarily hold. While none of the CPU targets I tried at Compiler Explorer has a LOP3 instruction operating on GPRs, most of them provide more logical instructions than just NOT, AND, OR, XOR, and the respective compiler’s strategy of how to best map and merge the logical operators seems to differ. I think I also saw some instances of bitfield extraction operations, which might come about from merging AND with shifts (from the logical ops and rotate parts of the expressions, presumably).

This working hypothesis would jibe with your observation that forcing the logical operations to map to LOP3 can be counterproductive in terms of performance.

6

That makes complete sense and matches the SASS data perfectly.
Looking at the before/after counts again with your “absorb into
arithmetic” lens:

Before (compiler defaults):
IMAD: 3753, IADD3: 3504, LOP3: 2279

After (forced LOP3 for f2/f3/f4):
IMAD: 3671 (-82), IADD3: 3601 (+97), LOP3: 2279 (same total!)

The IMAD reduction of 82 plus IADD3 increase of 97 is exactly the
signature of what you described: the compiler had been merging the
| operations from the boolean functions into IADD3s, and the shift
parts of the rotates were folding into IMADs as multiply-by-power-of-2.

When inline PTX took over f2/f3/f4, those merges broke. The boolean
result became a register that then needs an explicit IADD3 to combine
with the round constant - whereas before, the OR was just a free
side-effect of the addition itself in the right circumstances.

Net result: same LOP3 count, but +15 ALU operations overall,
producing the measured ~1-2% hashrate regression.

Counterintuitive but a great lesson — minimizing one instruction class
in isolation doesn’t necessarily minimize total work when the
instructions can blend with surrounding arithmetic. Keeping the C-level
boolean expressions visible to the compiler lets it choose between
“emit one LOP3” vs “absorb the OR into the next add” depending on
context.

Thanks for taking the time to investigate this in Compiler Explorer
and explaining the mechanism. This is exactly the kind of insight
that’s hard to find in documentation but makes a real difference in
practice.

7

@njuffa Quick follow-up after the LOP3 experiment, with new profiler data.

The kernel hits the following profile on RTX 4090:

Compute (SM) Throughput: 80.65%
ALU Pipeline: 80.70% ← reported as “over-utilized”
Memory Throughput: 21.90% (not the bottleneck)

Executed IPC Active: 2.11
Issued Warp/Scheduler: 0.53 (max 1.0)
Eligible Warps/Scheduler: 1.61 (hardware max 12)
No Eligible: 47.59%

Theoretical Occupancy: 33.33% (limited by 128 reg/thread)
Achieved Occupancy: 33.07%
Warp Cycles per Issued: 7.52

Instruction mix in SASS:
IMAD: 3753
IADD3: 3504
LOP3: 2279
SHF: 3252

The “47.59% no eligible warp” with high ALU utilization seems
contradictory at first. My reading is that this is a latency-bound
kernel rather than throughput-bound: the ALU pipeline is busy, but
the warps that ARE active spend half their cycles waiting for
dependencies (carry chains in 256-bit modular multiply, register
read-after-write hazards on the multiply-add sequences).

Does that interpretation match your experience? On Ada specifically,
when an integer-heavy kernel shows this pattern (high ALU%, low
eligible warps, low achieved occupancy due to register pressure),
which lever typically gives more headroom:

(a) Force higher occupancy via launch_bounds to increase
latency tolerance — accepting more local memory spills

(b) Reduce instruction count further — but as the LOP3 experiment
showed, isolated reductions can break compiler fusion and
hurt overall

(c) Restructure dependency chains to expose more ILP within each
warp — splitting independent computations to interleave at
the SASS level

I suspect the answer is “all three matter, but (c) gives the most
on Ada.” Curious what your experience suggests.

8

I have not used the CUDA profiler in quite a number of years, so I cannot interpret this data with any degree of reliability. I wonder whether the stall cycles are not due to the carry propagation (which is definitely a possibility), but rather due to IMAD.WIDE requiring two register writebacks. Ports on register files are expensive; it seems doubtful that NVIDIA added another write port to allow IMAD.WIDE to return data to both result registers simultaneously.

If nobody answers your question about the profiler statistics here, consider asking in the profiler forum.

BTW, when I compiled the RIPEMD-160 reference code, I was surprised to see that the compress() functionality was decomposed identically on two different platforms with two different toolchains:

gcc Kalray:  618 ADD, 318 ROL, 256 XOR, 64 AND, 64 ORN
clang Armv8: 619 ADD, 318 ROR, 256 XOR, 64 AND, 64 ORN

Kalray has some sort of 4-wide VLIW architecture, and its instruction grouping seems to indicate that the available ILP in the reference code is below 2.0. But this should not have a negative impact on Ampere. What these stats suggest is that, generally speaking, the rotates should (almost?) always map to a single dedicated instruction where possible, and that the GPU should derive some upside from the use of IADD3. I notice that some FPGA implementations of RIPEMD-160 use a CSA circuit to achieve the equivalent effect: they use the CSA as a 3:2 compressor, then use a regular ripple carry adder so combine the two results.

You can experiment with __launchbounds__ but I would claim this should not be used. Typically what happens is that if you squeeze the register count limit that the compiler picked by a few registers (low single digits), it starts de-emphasizing some optimizations that drive register pressure up. Often that is CSE (common subexpression elimination). Squeeze the compiler a little bit more, and it will start spilling and reloading data. In loop nests it does this intelligently, putting spill/fill cycles in the outermost loops. But here we have straight-line code.

Usually, by the time one gets to the introduction of spills & fills, performance starts to suffer. In the first stage (de-emphasizing certain optimizations), performance may already be affected negatively. In general, one should not obsess over occupancy. There is an old but somewhat famous paper by Volkov about achieving better performance at lower occupancy, at least in some instances.

9

Read the Volkov paper. It crystallized why the optimizations I tried (LOP3 forcing, occupancy increases) didn’t help: the kernel is already in the regime Volkov describes — low occupancy with high ILP per thread, computing multiple outputs per thread (forward+backward EC point per iteration), compute-bound with high register count. The fact that you predicted this outcome (“forcing logical ops to LOP3 can be counterproductive”) and the empirical SASS data confirmed it both align with Volkov’s core thesis: minimizing isolated instruction counts can break the larger optimization the compiler is doing across the dataflow. The kernel hits 80.7% ALU utilization at 33% occupancy. Per Volkov this isn’t “low occupancy as a problem” — it’s “high register usage enabling more work per thread, which pays for the lower thread count.” The structural limits (IMAD.WIDE writeback, RIPEMD-160 ILP < 2) that you outlined explain the remaining 19% headroom in ALU utilization. Software can’t address those. Thanks for the references and the patience walking through the analysis. Some of the most useful CUDA performance feedback I’ve gotten - especially the “compiler beats hand-tuning” lesson with concrete data.

10

What are the warp stall reasons (under Warp State Statistics)?
Note that throughput for each of IMAD, IADD3, LOP3, and SHF is 64/cycle/SM, so each of those can only be issued every 2 cycles. IMAD uses the FMA pipeline, while IADD3, LOP3, and SHF all go through the INT pipeline. If the INT pipeline is mostly busy, and your instruction mix is heavily leaning towards INT, you might quite often have the situation that all 4 active warps per scheduler have an INT instruction next, so all of them have to wait for one cycle.

11

You called it - here are the warp state stats from the same profile:

Stall Math Pipe Throttle: 2.23 cycles/inst (30%)
Stall Not Selected: 2.05 cycles/inst (27%)
Stall Wait: 1.31 cycles/inst (17%)
Selected: 1.00 cycles/inst (13%)
Stall Dispatch Stall: 0.67 cycles/inst (9%)
Stall Long Scoreboard: 0.11 cycles/inst (1.5%)
Other: ~0.15 cycles/inst (2%)

Total: 7.52 cycles per issued instruction

Math Pipe Throttle plus Not Selected dominate at over 50% combined -
exactly the pattern you described. With INT:FMA instruction ratio
of 2.4:1 (9035 INT ops vs 3753 IMAD ops in the SASS dump), the INT
pipeline is clearly the bottleneck.

The 4 active warps per scheduler all wanting INT next, and only one
INT op issuing per 2 cycles - this matches the 30% Math Pipe Throttle
and 27% Not Selected almost perfectly.

Question: would increasing occupancy (more warps per scheduler so
that not all of them want INT at the same time) help here? njuffa
warned against launch_bounds in another thread because it
deactivates CSE, but Volkov’s paper suggests low occupancy is fine
for ALU-bound code. With Math Pipe Throttle being the dominant stall,
this seems different from typical “occupancy doesn’t matter”
situations - or am I misreading the data?

12

I would agree that low occupancy can be fine for math-bound code, and higher occupancy is unlikely to help much here.

Reaching >80% utilisation of any given pipe is hard, as the schedulers don’t let you prioritise one type of instruction over another. So despite having a majority of INT instructions in your kernel, the schedulers will occasionally end up issuing two or more non-INT instructions back-to-back. Those are in effect “lost” cycles for the INT pipeline, which would have to be fed every other cycle to reach 100% utilisation.

13

That would be a performance optimization for Nvidia for future architectures.

Setting scheduler policies either by programmers of highly optimized routines or by the compiler with heuristics. (The compiler may have difficulty setting not only beyond a thread, but beyond warps and for several streams running concurrently on the SMSP, possibly beyond a single kernel being executed.)

Or, alternatively, the scheduler detects that many INT instructions are in the stream and prefers those.

OTOH if I understood correctly, the throttling functionality early on tries to reduce a single type of instruction dominating and starving other instructions. So the scheduler may actively try to avoid INT instructions, if there are too many.

14

It’s not so much that the schedulers actively avoid INT instructions, it’s that the INT pipeline can’t accept two instructions back to back (there may be exceptions to this). So the cycle after an INT instruction was issued, any warps that have an INT instruction up next just aren’t eligible.

As far as I can tell the scheduler logic is very simple: it will continue with the current warp as long as its next instruction is eligible (and the warp didn’t yield or exit), otherwise it deterministically switches to another eligible warp (while I’m not sure what priority order it uses, I’m pretty sure it’s not based on the kind of instruction waiting).

15

Thanks both - this clarifies the architecture in a way that the
profiler numbers alone don’t.

The “INT pipeline can’t accept two instructions back to back”
detail is the missing piece. With ~70% INT instructions in this
kernel and a hard 1-instr-per-2-cycles ceiling on the INT pipe,
the math works out:

INT pipe peak: 0.5 instr/cycle
My INT density: ~70% of issued instructions
Implied min cycles: 1.4x instruction count if INT-bound
Measured ALU util: 80.7%

So I’m hitting roughly 81% of the practical ceiling, with the
remaining gap explained by the scheduler occasionally issuing
back-to-back non-INT instructions (FMA, MOV, etc.) that leave
the INT pipe idle for those cycles.

Curefab’s observation about anti-starvation throttling explains
why this gap can’t easily be closed - it’s not a bug, it’s a
deliberate scheduler property to prevent any single instruction
type from monopolizing issue slots.

The takeaway for me: the ~20% headroom isn’t reachable through
SASS-level optimization. It would require either:
(a) Restructuring the algorithm to better balance INT/FMA per
warp (which would mean different math, not different code)
(b) A future arch with separate scheduler policies, or
compiler hints for instruction-mix priority

Neither is on the table for this kernel today. Closing the thread
here - the empirical picture (njuffa’s V4 monolithic mul giving
+9%, all manual SASS-level reorderings since giving 0%) matches
the architectural explanation perfectly.

Thanks again to njuffa, @Nanodeoclus and @Curefab - this has been
a genuinely educational thread.