SHA256 How to get digest value in `unsigned char` format?

[LeMoussel]

Hi,

First of all: thank you very much for open sourcing OpenCL hashcat!

For a custom application based on sha256, I must get the digest value in `unsigned char` format.
It's OK with CPU implementation with deps\LZMA-SDK\CSha256.c

Code:

void Sha256_Final(CSha256 *p, Byte *digest)
{
.....
  for (i = 0; i < 8; i += 2)
  {
    UInt32 v0 = p->state[i];
    UInt32 v1 = p->state[i + 1];
    SetBe32(digest    , v0);
    SetBe32(digest + 4, v1);
    digest += 8;
  }
}

How to do with OpenCL implementation? (OpenCL\inc_hash_sha1.cl)

Code:

__kernel void ocl_sha256(__global input_t *input, const unsigned int gid_max)
{
    unsigned int lid = get_local_id(0);
    unsigned int gid = get_global_id(0);

    if (gid >= gid_max) return;

    // TODO
    unsigned int test[64] = { 0 };

    sha1_ctx_t ctx0;
    sha1_init(&ctx0);
    sha1_update(&ctx0, test, 64);
    sha1_final(&ctx0);

/*

unsigned int digest32[4] = { ctx0.h[0], ctx0.h[1], ctx0.h[2], ctx0.h[3] }

unsigned char digest8[16] = ????

*/

}

Thanks,

[philsmd]

I think it makes sense to explain the problem with some exampes too.

just an example, maybe not applicable to your problem at all:
"I want to convert 0x01020304 to 0x04030201"

"I want to convert raw/binary data \x65\x66\x67\x68 to hexadecimal characters 65666768"

This "Byte" type from lzma does nothing special, it's just a custom variable type in the lzma code.

it's actually more clever to just try to understand the problem and see what is "wrong" and fix it.
Maybe your input is different etc. You should defintely show us some data too, e.g. I got this xyz, but want to have this instead: abc

btw: we normally do not use "int" array type but u32 arrays instead (in theory also unsigned char should work, but in hashcat we normally just use the unsigned u32 type instead)

note2: what I totally forgot to mention, but I already saw it before... you are calling sha1 but your function says sha256 ("ocl_sha256"). this is very, very weird . Maybe this is the main reason for the confusion/wrong results ?

[LeMoussel]

Note2: rectified, change sha1 to sha256

Example :

Code:

    unsigned int r0 = 3369586927, r1 = 250477186, r2 = 3538557605, r3 = 2458147738;
    const unsigned int digest32[4] = { r0, r1, r2, r3 };
    unsigned char digest8[16];

I want to store digest32 in digest8.
In result, in hexadecimal, the values in digest8 will be: `0XEF 0XD0 0XD7 0XC8 0X82 0XFA 0XED 0X0E 0XA5 0X1A 0XEA 0XD2 0X9A 0X5B 0X84 0X92`

[philsmd]

the u32 and u8 are primitive types, you can simply cast an array of 4-byte integers to an array with 1-byte chars/bytes.

Code:

u8 * bytes = (u8 *) digest32;

or copy them directly

Code:

u8 bytes[16];

bytes[0] = (digest32[0] >>  0) & 0xff;
bytes[1] = (digest32[0] >>  8) & 0xff;
bytes[2] = (digest32[0] >> 16) & 0xff;
bytes[3] = (digest32[0] >> 24) & 0xff;

bytes[4] = (digest32[1] >>  0) & 0xff;
bytes[5] = (digest32[1] >>  8) & 0xff;
bytes[6] = (digest32[1] >> 16) & 0xff;
bytes[7] = (digest32[1] >> 24) & 0xff;

....

btw: if you need the "reverse order", you just use >> 24, >> 16, >> 8 and finally >> 0 instead of the reverse order.

While it is often easier to think in bytes (or unsigned chars), the compilers/runtimes normally prefer to use the full 4-byte integers... therefore hashcat actually uses u32 wherever possible. There are of course exceptions, because compilers (unless you use volatile) are most of the time clever enough to just optimize the code from the single chars to the more easier to use u32/integer type (it's more native/easier for most OpenCL device architectures to just use u32).

btw: the output of sha256 is 32 bytes, not 16 bytes: I see that you are using 16 bytes output, but sha256 outputs 32 raw bytes. This might be fine, depending on what you need to do with the output. I just want to mention it here, because it's very confusing to use 16 bytes with sha256.

[LeMoussel]

Thank you very much for your help.

In hashcat file: inc_common.cl, I find `DECLSPEC u8 v8a_from_v32_S (const u32 v32)` function.
From what I understand (I'm starting in OpenCL), it is also possible to do :

Code:

    const u32 r0 = ctx0.h[0];
    const u32 r1 = ctx0.h[1];
    const u32 r2 = ctx0.h[2];
    const u32 r3 = ctx0.h[3];
    
    const u32 digest32[4] = { r0, r1, r2, r3 };
    u8 digest8[16];
    vconv32_t v;

    for (unsigned int i = 0; i < 4; i++)
    {
        v.v32 = digest32[i];
        digest8[i * 4] = v.v8a;
        digest8[i * 4 + 1] = v.v8b;
        digest8[i * 4 + 2] = v.v8c;
        digest8[i * 4 + 3] = v.v8d;
    }

What do you mean by "the output of sha256 is 32 bytes, not 16 bytes: I see that you are using 16 bytes output, but sha256 outputs 32 raw bytes"?  With OpenCL u32 is an unsigned int ?

[philsmd]

The length of a normal sha256 hash is 64 hexadecimal characters and therefore 32 raw bytes. You are trying to use 4 ints of 4 bytes = 4 * 4 = 16 bytes. That's only half the bytes of the full hash.

[LeMoussel]

Ah OK!
But then how to get full hash?

Code:

4 ints of 8 bytes = 4 * 8 = 32 bytes => u8 digest8[32]

[philsmd]

no, it's 8 ints (all of which are 4 bytes), see https://github.com/hashcat/hashcat/blob/...a256.h#L64

[LeMoussel]

Oups!
the code is then

Code:

const u32 digest32[8] = { ctx0.h[0], ctx0.h[1], ctx0.h[2], ctx0.h[3], ctx0.h[4], ctx0.h[5], ctx0.h[6], ctx0.h[7] }    
u8 digest8[32];

for (unsigned int i = 0; i < 8; i++)
{
    v.v32 = digest32[i];
    digest8[i * 4] = v.v8a;
    digest8[i * 4 + 1] = v.v8b;
    digest8[i * 4 + 2] = v.v8c;
    digest8[i * 4 + 3] = v.v8d;
}