Md5+ pair sum + base62 best approach to custom mode

[Sasquatch]

Hi I have a  CCTV camera I got myself locked out of. I managed to get the stored hash and  password check algorithm out thanks to power of ghira and chatgpt.

It is standard MD5 followed by pair sum and base62 conversion to 8 character hash.

I have opencl kernel and host code that run 360MH/s on p400. And only 5300MH/s on rtx3070.

I'm thinking, since hashcat hits ~40GH/s on 3070 can I add extra steps for my use case to existing module and how much will it slow down?

Or should I just get vast.ai instance with 4090? And pay for the 2 weeks it would take at 6GH/s?

Here is revelant opencl rernel section:

Code:

that fixed it put my hasrate dropped from 5130 to 5100Mh/s what are sane values of local and -D STEPS_PER_THREAD for rtx3070
it was 64 and 1 for p400(giving 360Mh/s

my kernel snippet for refernce: 
[Code]

__kernel void crack_kernel(
const ulong base_idx,
const ulong total,
const uint charset_len,
const int L,
__constant const uchar *charset,
__constant const uchar *target8,
__global uint *found_count,
__global ulong *found_index,
__global uchar *found_digest
) {
ulong gid = (ulong)get_global_id(0);
ulong idx = base_idx + gid;
if (idx >= total) return;

uchar plain[16]; 
index_to_plain(idx, charset, charset_len, plain, L); 

uchar digest[16]; 
md5_singleblock(plain, L, digest); 

uchar out8[8]; 
// unrolled loop for 8 steps: 
{ 
    uint v0 = (uint)digest[0] + (uint)digest[1]; 
    uint v1 = (uint)digest[2] + (uint)digest[3]; 
    uint v2 = (uint)digest[4] + (uint)digest[5]; 
    uint v3 = (uint)digest[6] + (uint)digest[7]; 
    uint v4 = (uint)digest[8] + (uint)digest[9]; 
    uint v5 = (uint)digest[10] + (uint)digest[11]; 
    uint v6 = (uint)digest[12] + (uint)digest[13]; 
    uint v7 = (uint)digest[14] + (uint)digest[15]; 

     
out8[0] = map62[mod62_lut[v0 & 0xFFu]]; 
    out8[1] = map62[mod62_lut[v1 & 0xFFu]]; 
    out8[2] = map62[mod62_lut[v2 & 0xFFu]]; 
    out8[3] = map62[mod62_lut[v3 & 0xFFu]]; 
    out8[4] = map62[mod62_lut[v4 & 0xFFu]]; 
    out8[5] = map62[mod62_lut[v5 & 0xFFu]]; 
    out8[6] = map62[mod62_lut[v6 & 0xFFu]]; 
    out8[7] = map62[mod62_lut[v7 & 0xFFu]]; 
} 

// compare with target8 without branches 
int ok = 1; 
for (int i=0; i<8; i++) { 
    if (out8[i] != target8[i]) { 
        ok = 0; break; 
    } 
} 

if (ok) { 
    uint old = atomic_inc(found_count); 
    if (old == 0u) { 
        found_index[0] = idx; 
        for (int i=0; i<16; i++) found_digest[i] = digest[i]; 
    } 
}

}

P.s I cannot change salt for known one, as I'm getting CRC error on boot and I can't find where it's stored(possibly in  protected CPU memory)

And yes I realised(now) that I can gain some speed by  calculating salt indices on in host code  and comparing integers in kernel cutting out char conversion.

[Sasquatch]

m24900 seems to be closest but not finding abc12 from D1H3aIcP, differen algo, obviously, but afaik only mode using "compatible" hash.

hits 480MH/s on p400

[Sasquatch]

i foound the difference!
my camera's firmwre does:

Code:

LDRB   r0, [digest, #X]   ; load byteA
LDRB   r1, [digest, #Y]   ; load byteB
ADD    r0, r0, r1         ; sum into r0 (0–510 possible)
AND    r0, r0, #0xFF      ; mask to 8 bits
MOV    r1, #62
BL     __aeabi_idivmod    ; integer division/modulo

m24900 skips masking to 8 bit value before %62
without masking abc12 hash is D1P3aIkX

simple fix, thanks hashcat team

P.s it's not dahua camera so I can't claim this is m24900 bug...

[atom]

This could be interesting from forensic perspective. Would you like to share the details? Like Vendor/Model and the algorithm difference to 24900. Then we could make this a real hash mode for hashcat. Maybe you can generate a hash on that device with password "hashcat" so we can use this as a self-test.

[Sasquatch]

This could be interesting from forensic perspective. Would you like to share the details? Like Vendor/Model and the algorithm difference to 24900. Then we could make this a real hash mode for hashcat. Maybe you can generate a hash on that device with password "hashcat" so we can use this as a self-test.

Only difference I found is the masking.
Camera is Besder AliExpress special, using xmey firmware. 
The hash in question is U-Boot pre Linux Kernel load  password.
So technically that is U-Boot 4.something hash, not necessarily Besder/xmey specific
I tried skipping password verification by replacing funcion call with r0 load - CRC error on boot
Same for injected hash.
But CRC compare function grabs it from address before mapped firmware area...

[atom]

If you can provide a proof of concept in any language and a description of how to extract the required key material (or a extraction tool), I can turn this into a real hash mode and you will get the best performance.

[Sasquatch]

If you can provide a proof of concept in any language and a description of how to extract the required key material (or a extraction tool), I can turn this into a real hash mode and you will get the best performance.

All this needed to work with my camera's hashing is changing this:

Code:

a0 = (((a >> 0) & 0xff) + ((a >> 8) & 0xff)) % 62;
....
...
..

To this:

Code:

a0 = ((((a >> 0) & 0xff) + ((a >> 8) & 0xff)) & 0xFF)% 62;

....
...
..

In m24900_ax-optimized.cl

Subject to testing, whi will commit as soon as I get home from work.

[Sasquatch]

i cooked up module 24901 from 24900 tested hashes generated with mask.py (attached too) 

both are hitting 740MH/s on p400,