<![CDATA[hashcat Forum - hashcat]]> https://hashcat.net/forum/ Sun, 05 Apr 2026 00:15:59 +0000 MyBB <![CDATA[SNMPv3 engineID minimum length, am I missing something?]]> https://hashcat.net/forum/thread-13492.html Sat, 21 Feb 2026 17:45:28 +0000 c4p0ne]]> https://hashcat.net/forum/thread-13492.html
Code:
totoken.len_min[3] = 26

in module_25000.c (and its siblings). RFC 3411 says engineID is 5-32 bytes, so the minimum should be 10 hex chars, not 26. IPv4-based engine IDs are 9 bytes (18 hex) and MAC-based are 11 bytes (22 hex). Well both of these appear to be under the floor.

Also noticed m25000.pm uses

Code:
random_hex_string(26, 34)

so the test suite never generates a short enough engineID to catch it. Is this a known limitation or an actual bug? Seems like a one-line fix across the all the sibling modules. Will someone knowledgeable please come to the rescue here and tell me if I'm looking at this wrong?]]>
Code:
totoken.len_min[3] = 26

in module_25000.c (and its siblings). RFC 3411 says engineID is 5-32 bytes, so the minimum should be 10 hex chars, not 26. IPv4-based engine IDs are 9 bytes (18 hex) and MAC-based are 11 bytes (22 hex). Well both of these appear to be under the floor.

Also noticed m25000.pm uses

Code:
random_hex_string(26, 34)

so the test suite never generates a short enough engineID to catch it. Is this a known limitation or an actual bug? Seems like a one-line fix across the all the sibling modules. Will someone knowledgeable please come to the rescue here and tell me if I'm looking at this wrong?]]> <![CDATA[Token Length Exception]]> https://hashcat.net/forum/thread-13457.html Sun, 04 Jan 2026 16:51:10 +0000 Ricardo Alexis]]> https://hashcat.net/forum/thread-13457.html When I try to use 
Code:
hashcat -m 26620
, below error was occupied.

"Token length exception" in hashcat terminal.

How to resolve this issue?]]> When I try to use 
Code:
hashcat -m 26620
, below error was occupied.

"Token length exception" in hashcat terminal.

How to resolve this issue?]]> <![CDATA[Introduction & Tips for New Hashcat Users]]> https://hashcat.net/forum/thread-13448.html Fri, 26 Dec 2025 16:55:55 +0000 b15063368]]> https://hashcat.net/forum/thread-13448.html I’d like to share a few questions and tips for beginners:
  1. Best practices for getting started:
    What is the best way to begin with common hash types like MD5 or NTLM? Should I start with dictionary attacks or masks?
]]> I’d like to share a few questions and tips for beginners:
  1. Best practices for getting started:
    What is the best way to begin with common hash types like MD5 or NTLM? Should I start with dictionary attacks or masks?
]]> <![CDATA[looking for module]]> https://hashcat.net/forum/thread-13425.html Sun, 23 Nov 2025 01:20:49 +0000 misterapple282]]> https://hashcat.net/forum/thread-13425.html i’m looking for this module
sha1(salt.sha1(salt.sha1(password))
example 
п
salt: 01rIbTytUx300851
plaintext: romeoandjuliet
digest: 52cc8adbd7a00d40973c174f297833944162451d
]]> i’m looking for this module
sha1(salt.sha1(salt.sha1(password))
example 
п
salt: 01rIbTytUx300851
plaintext: romeoandjuliet
digest: 52cc8adbd7a00d40973c174f297833944162451d
]]> <![CDATA[Problem with multi keys in kernel]]> https://hashcat.net/forum/thread-13411.html Tue, 21 Oct 2025 08:21:04 +0000 zamgold]]> https://hashcat.net/forum/thread-13411.html I am developing a kernel - a modified RC4, as a result, in the kernel there are many candidates for the correct key that meet my criteria, which I want to throw out to potfile.
if you use the function mark_hash - some keys are not included in the potfile, so I made the following implementation:


KERNEL_FQ void m33545_comp (KERN_ATTR_TMPS(rc4_tmp_t))
{
  const u64 gid = get_global_id (0);
  const u64 lid = get_local_id (0);


  if (gid >= GID_CNT) return;



    LOCAL_VK u32 S[64 * FIXED_LOCAL_SIZE];


    //MY RC4_REALIZATION
    .....

    if ( value > CRITERION ) {
        u32 idx = hc_atomic_inc (d_return_buf);
        printf("\nkey=%02x%x value=%d hc_atomic_idx=%d\n", pws[gid].i[1], pws[gid].i[0],value, idx);

        plains_buf[idx].salt_pos  = SALT_POS_HOST;
        plains_buf[idx].digest_pos = 0;  // relative
        plains_buf[idx].hash_pos  = DIGESTS_OFFSET_HOST + 0;    // absolute
        plains_buf[idx].gidvid    = gid;
        plains_buf[idx].il_pos    = 0;
        plains_buf[idx].extra1    = 0;      // unused so far
        plains_buf[idx].extra2    = 0;      // unused so far
    }

}



here is part of my output to the screen:


[s]tatus [p]ause ypass [c]heckpoint [f]inish [q]uit =>
key=05dfb32e79 value=1064 hc_atomic_idx=0

key=0316b32e79 value=504 hc_atomic_idx=1

key=3594b32e79 value=520 hc_atomic_idx=2

key=4c9eb32e79 value=504 hc_atomic_idx=3

key=6a96b32e79 value=648 hc_atomic_idx=4

key=88e5b32e79 value=528 hc_atomic_idx=6

key=8fbab32e79 value=536 hc_atomic_idx=5

key=a1a5b32e79 value=520 hc_atomic_idx=7

key=add1b32e79 value=576 hc_atomic_idx=8

key=c13ab32e79 value=552 hc_atomic_idx=9

key=bae8b32e79 value=512 hc_atomic_idx=10

key=f6b6b32e79 value=520 hc_atomic_idx=11

key=f295b32e79 value=568 hc_atomic_idx=12

->check_crack2:13
cuMemcpyDtoHAsync(): invalid argument


<-check_crack2 return:-1
Session..........: hashcat
Status...........: Exhausted
Hash.Mode........: 33545 (RC4_MY)
Hash.Target......: xxxxxxxxxxxxxxxxxxxxxx
Time.Started.....: Tue Oct 21 12:56:06 2025 (0 secs)
Time.Estimated...: Tue Oct 21 12:56:06 2025 (0 secs)
Kernel.Feature...: Pure Kernel
Guess.Mask.......: 792eb3?b?b [5]
Guess.Queue......: 1/1 (100.00%)
Speed.#1.........: 11787.5 kH/s (0.00ms) @ Accel:1024 Loops:1024 Thr:32 Vec:1
Recovered........: 0/1 (0.00%) Digests (total), 0/1 (0.00%) Digests (new)
Progress.........: 65536/65536 (100.00%)
Rejected.........: 0/65536 (0.00%)
Restore.Point....: 65536/65536 (100.00%)
Restore.Sub.#1...: Salt:0 Amplifier:0-1 Iteration:0-1024
Candidate.Engine.: Device Generator
Candidates.#1....: $HEX[792eb3e0b8] -> $HEX[792eb3ffff]
Hardware.Mon.#1..: Temp: 65c Fan: 70% Util:100% Core:1860MHz Mem:9251MHz Bus:16

Started: Tue Oct 21 12:56:03 2025
Stopped: Tue Oct 21 12:56:07 2025



Th finction check_cracked  from hashes.c sees that I have 13 keys, but returns -1, although it should return 0! Then I have a error cuMemcpyDtoHAsync()
[b]What am I doing wrong, how do I implement this?[/b]]]> I am developing a kernel - a modified RC4, as a result, in the kernel there are many candidates for the correct key that meet my criteria, which I want to throw out to potfile.
if you use the function mark_hash - some keys are not included in the potfile, so I made the following implementation:


KERNEL_FQ void m33545_comp (KERN_ATTR_TMPS(rc4_tmp_t))
{
  const u64 gid = get_global_id (0);
  const u64 lid = get_local_id (0);


  if (gid >= GID_CNT) return;



    LOCAL_VK u32 S[64 * FIXED_LOCAL_SIZE];


    //MY RC4_REALIZATION
    .....

    if ( value > CRITERION ) {
        u32 idx = hc_atomic_inc (d_return_buf);
        printf("\nkey=%02x%x value=%d hc_atomic_idx=%d\n", pws[gid].i[1], pws[gid].i[0],value, idx);

        plains_buf[idx].salt_pos  = SALT_POS_HOST;
        plains_buf[idx].digest_pos = 0;  // relative
        plains_buf[idx].hash_pos  = DIGESTS_OFFSET_HOST + 0;    // absolute
        plains_buf[idx].gidvid    = gid;
        plains_buf[idx].il_pos    = 0;
        plains_buf[idx].extra1    = 0;      // unused so far
        plains_buf[idx].extra2    = 0;      // unused so far
    }

}



here is part of my output to the screen:


[s]tatus [p]ause ypass [c]heckpoint [f]inish [q]uit =>
key=05dfb32e79 value=1064 hc_atomic_idx=0

key=0316b32e79 value=504 hc_atomic_idx=1

key=3594b32e79 value=520 hc_atomic_idx=2

key=4c9eb32e79 value=504 hc_atomic_idx=3

key=6a96b32e79 value=648 hc_atomic_idx=4

key=88e5b32e79 value=528 hc_atomic_idx=6

key=8fbab32e79 value=536 hc_atomic_idx=5

key=a1a5b32e79 value=520 hc_atomic_idx=7

key=add1b32e79 value=576 hc_atomic_idx=8

key=c13ab32e79 value=552 hc_atomic_idx=9

key=bae8b32e79 value=512 hc_atomic_idx=10

key=f6b6b32e79 value=520 hc_atomic_idx=11

key=f295b32e79 value=568 hc_atomic_idx=12

->check_crack2:13
cuMemcpyDtoHAsync(): invalid argument


<-check_crack2 return:-1
Session..........: hashcat
Status...........: Exhausted
Hash.Mode........: 33545 (RC4_MY)
Hash.Target......: xxxxxxxxxxxxxxxxxxxxxx
Time.Started.....: Tue Oct 21 12:56:06 2025 (0 secs)
Time.Estimated...: Tue Oct 21 12:56:06 2025 (0 secs)
Kernel.Feature...: Pure Kernel
Guess.Mask.......: 792eb3?b?b [5]
Guess.Queue......: 1/1 (100.00%)
Speed.#1.........: 11787.5 kH/s (0.00ms) @ Accel:1024 Loops:1024 Thr:32 Vec:1
Recovered........: 0/1 (0.00%) Digests (total), 0/1 (0.00%) Digests (new)
Progress.........: 65536/65536 (100.00%)
Rejected.........: 0/65536 (0.00%)
Restore.Point....: 65536/65536 (100.00%)
Restore.Sub.#1...: Salt:0 Amplifier:0-1 Iteration:0-1024
Candidate.Engine.: Device Generator
Candidates.#1....: $HEX[792eb3e0b8] -> $HEX[792eb3ffff]
Hardware.Mon.#1..: Temp: 65c Fan: 70% Util:100% Core:1860MHz Mem:9251MHz Bus:16

Started: Tue Oct 21 12:56:03 2025
Stopped: Tue Oct 21 12:56:07 2025



Th finction check_cracked  from hashes.c sees that I have 13 keys, but returns -1, although it should return 0! Then I have a error cuMemcpyDtoHAsync()
[b]What am I doing wrong, how do I implement this?[/b]]]> <![CDATA[Help setting up Hashcat please]]> https://hashcat.net/forum/thread-13391.html Sun, 28 Sep 2025 19:18:48 +0000 kerberstui]]> https://hashcat.net/forum/thread-13391.html <![CDATA[$wp$2y$10$]]> https://hashcat.net/forum/thread-13369.html Sat, 06 Sep 2025 08:50:23 +0000 174region174]]> https://hashcat.net/forum/thread-13369.html $wp$2y$10$........................:example
We managed to parse this type of hash. This is not a simple bcrypt
This password is encrypted several times.
The approximate encryption algorithm is as follows.
bcrypt hmac-sha384 base64
 Next, I'll show you how to search for it. But we need to write a module for hashcat, because our search method is very inconvenient. but it's possible.
cat test.dic | ./wp-sha384-hmac.bin | ./hashcat -m 3200 -a 0  3200.hash

You need to compile from the code  wp-sha384-hmac.bin
gcc wp-sha384-hmac.bin.c -O3 -march=native -flto -ffast-math -funroll-loops -o wp-sha384-hmac.bin -lssl -lcrypto

Code:
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <ctype.h>

#include <openssl/hmac.h>
#include <openssl/evp.h>
#include <openssl/buffer.h>

#define HMAC_KEY "wp-sha384"
#define HMAC_KEY_LEN 10

char *trim_whitespace(char *str) {
    if (!str) return NULL;

    while (isspace((unsigned char)*str)) str++;

    if (*str == '\0')
        return str;

    char *end = str + strlen(str) - 1;
    while (end > str && isspace((unsigned char)*end)) end--;

    *(end + 1) = '\0';
    return str;
}

void process_stream(FILE *fp) {
    char *line = NULL;
    size_t len = 0;

    while (getline(&line, &len, fp) != -1) {
        char *clean = trim_whitespace(line);
        size_t linelen = strlen(clean);

        unsigned char hmac[EVP_MAX_MD_SIZE];
        unsigned int hmac_len = 0;

        HMAC(EVP_sha384(),
            HMAC_KEY, HMAC_KEY_LEN,
            (unsigned char *)clean, linelen,
            hmac, &hmac_len);

        BIO *b64 = BIO_new(BIO_f_base64());
        BIO_set_flags(b64, BIO_FLAGS_BASE64_NO_NL);
        BIO *bio = BIO_new(BIO_s_mem());
        b64 = BIO_push(b64, bio);

        BIO_write(b64, hmac, hmac_len);
        BIO_flush(b64);

        BUF_MEM *bptr;
        BIO_get_mem_ptr(b64, &bptr);
        fwrite(bptr->data, 1, bptr->length, stdout);
        putchar('\n');

        BIO_free_all(b64);
    }

    free(line);
}

int main(int argc, char *argv[]) {
    if (argc < 2) {
        process_stream(stdin);
        return 0;
    }

    for (int i = 1; i < argc; ++i) {
        if (strcmp(argv[i], "-") == 0) {
            process_stream(stdin);
            continue;
        }

        FILE *fp = fopen(argv[i], "r");
        if (!fp) {
            perror(argv[i]);
            continue;
        }

        process_stream(fp);
        fclose(fp);
    }

    return 0;
}
If your dictionary contains the password to your hash. Then you will have to search for it again in some way in your dictionary. That's the whole problem. We need your help with the revision.
If necessary, I can post a hash with a known password for your experiments in this thread.]]> $wp$2y$10$........................:example
We managed to parse this type of hash. This is not a simple bcrypt
This password is encrypted several times.
The approximate encryption algorithm is as follows.
bcrypt hmac-sha384 base64
 Next, I'll show you how to search for it. But we need to write a module for hashcat, because our search method is very inconvenient. but it's possible.
cat test.dic | ./wp-sha384-hmac.bin | ./hashcat -m 3200 -a 0  3200.hash

You need to compile from the code  wp-sha384-hmac.bin
gcc wp-sha384-hmac.bin.c -O3 -march=native -flto -ffast-math -funroll-loops -o wp-sha384-hmac.bin -lssl -lcrypto

Code:
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <ctype.h>

#include <openssl/hmac.h>
#include <openssl/evp.h>
#include <openssl/buffer.h>

#define HMAC_KEY "wp-sha384"
#define HMAC_KEY_LEN 10

char *trim_whitespace(char *str) {
    if (!str) return NULL;

    while (isspace((unsigned char)*str)) str++;

    if (*str == '\0')
        return str;

    char *end = str + strlen(str) - 1;
    while (end > str && isspace((unsigned char)*end)) end--;

    *(end + 1) = '\0';
    return str;
}

void process_stream(FILE *fp) {
    char *line = NULL;
    size_t len = 0;

    while (getline(&line, &len, fp) != -1) {
        char *clean = trim_whitespace(line);
        size_t linelen = strlen(clean);

        unsigned char hmac[EVP_MAX_MD_SIZE];
        unsigned int hmac_len = 0;

        HMAC(EVP_sha384(),
            HMAC_KEY, HMAC_KEY_LEN,
            (unsigned char *)clean, linelen,
            hmac, &hmac_len);

        BIO *b64 = BIO_new(BIO_f_base64());
        BIO_set_flags(b64, BIO_FLAGS_BASE64_NO_NL);
        BIO *bio = BIO_new(BIO_s_mem());
        b64 = BIO_push(b64, bio);

        BIO_write(b64, hmac, hmac_len);
        BIO_flush(b64);

        BUF_MEM *bptr;
        BIO_get_mem_ptr(b64, &bptr);
        fwrite(bptr->data, 1, bptr->length, stdout);
        putchar('\n');

        BIO_free_all(b64);
    }

    free(line);
}

int main(int argc, char *argv[]) {
    if (argc < 2) {
        process_stream(stdin);
        return 0;
    }

    for (int i = 1; i < argc; ++i) {
        if (strcmp(argv[i], "-") == 0) {
            process_stream(stdin);
            continue;
        }

        FILE *fp = fopen(argv[i], "r");
        if (!fp) {
            perror(argv[i]);
            continue;
        }

        process_stream(fp);
        fclose(fp);
    }

    return 0;
}
If your dictionary contains the password to your hash. Then you will have to search for it again in some way in your dictionary. That's the whole problem. We need your help with the revision.
If necessary, I can post a hash with a known password for your experiments in this thread.]]> <![CDATA[[Discussion] Ideas for (at least one/two) new rules]]> https://hashcat.net/forum/thread-13360.html Mon, 25 Aug 2025 10:32:49 +0000 Snoopy]]> https://hashcat.net/forum/thread-13360.html
After all these new updates coming with the release of HC 7.0, how about brainstorming for some ideas for new rules?

After years of using rules, or answering questions in the forum, there are some things which pop up frequently. So here are at least one, two suggestions for rules which came to my mind

TOGGLE every xth char: i know this one can be replaced by T@ and maybe can be misleading because of, start with first char of the pass or not, so it think this is not a good example, but the next one could be

Title: Capitalize every first letter in a pass, resetting trigger at non letter chars like space, colons or digits etc.
example: "my first pass" -> " My First Pass", my1first2pass -> My1First2Pass"

Untitle: reverse logic from Title

Most other things frequently asked would involve UTF-8 support or multibyte rules, and i think this is not possible with the actual rules engine, am i right?

Like deleting given words, replacing given words, replacing single chars with words and so on.]]>
After all these new updates coming with the release of HC 7.0, how about brainstorming for some ideas for new rules?

After years of using rules, or answering questions in the forum, there are some things which pop up frequently. So here are at least one, two suggestions for rules which came to my mind

TOGGLE every xth char: i know this one can be replaced by T@ and maybe can be misleading because of, start with first char of the pass or not, so it think this is not a good example, but the next one could be

Title: Capitalize every first letter in a pass, resetting trigger at non letter chars like space, colons or digits etc.
example: "my first pass" -> " My First Pass", my1first2pass -> My1First2Pass"

Untitle: reverse logic from Title

Most other things frequently asked would involve UTF-8 support or multibyte rules, and i think this is not possible with the actual rules engine, am i right?

Like deleting given words, replacing given words, replacing single chars with words and so on.]]> <![CDATA[Md5+ pair sum + base62 best approach to custom mode]]> https://hashcat.net/forum/thread-13348.html Thu, 14 Aug 2025 07:52:10 +0000 Sasquatch]]> https://hashcat.net/forum/thread-13348.html
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.]]>
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.]]> <![CDATA[Hashcat to RX480]]> https://hashcat.net/forum/thread-13347.html Tue, 12 Aug 2025 20:15:01 +0000 Haswl]]> https://hashcat.net/forum/thread-13347.html <![CDATA[Using the Assimilation Bridge (Python Plugin) for Rapid Prototyping]]> https://hashcat.net/forum/thread-13346.html Mon, 11 Aug 2025 09:11:11 +0000 atom]]> https://hashcat.net/forum/thread-13346.html
I originally added this plugin for situations where I had reversed some algorithm from an application or library, but usually it was a one-hash scenario. In most cases, I would simply write a quick proof of concept in, crack the single hash, and move on. The Python plugin is for such cases where the hashes are more resistant, for example when they are not based on a weak password. In that case you might need both computational power and some smart password guessing. You may want to take advantage of hashcat features such as the rule engine, hybrid attack modes, restart functionality, or built-in multithreading.

Instead of just writing your own standalone cracking tool, you can later move your PoC logic into the Python bridge and directly benefit from hashcat features.

We recently participated in the Jabbercracky password contest (DEFCON 33) where we faced a set of completely unknown hashes with no context at all.

When you do not know the hash type and there is no obvious signature or pattern in the hashes, a great tool to try is mdxfind. It is multithreaded, has many useful features, and is ideal for such situations. It supports iterations, external salt or username inputs, and more. In this contest, mdxfind could identify some of the algorithms, but some others the players had to figure out manually.

That is a perfect example of when the Python bridge is useful because it gives you a programming language with access to a large pool of third-party cryptographic libraries, and using a programming language gives you maximum flexibility.

Note!!! that in order to use this you need a version higher than v7.0.0. For instance new version v7.1.0 or the beta binary from https://hashcat.net/beta or by compiling from GitHub master sources.


0. Set up the Python bridge

If hashcat -b -m 73000 runs successfully, you are ready for step 1. If you have never used it before, check out the Quickstart Guide.

1. Download the contest hash list

Here is a mirror.

2. Convert it for use with the Python bridge

Code:
$ cat hash_list_3.txt | grep -v '^\$2a\$' | sed -e 's/$/*/g' > hashes73000.txt

Mode 73000 expects the format HASH*SALT. In this contest there were no salts, only peppers, so we leave the salt field empty and handle the pepper in our Python code.  We skip the bcrypts intentionally. You can also implement bcrypt in the Python bridge, but since bcrypt has a recognizable signature and is supported directly with mode 3200, it is better to handle those natively.

We use mode 73000 instead of 72000 for better cross-platform compatibility.

3. Edit Python/generic_hash_mp.py

Code:
import sys
import struct
import hashlib
import hcshared
import hcmp
import hmac

# SELF TEST pair
ST_HASH = "9c016f79fdcdc7422af93fcc203b0eda43278d32e3a6edd607d35c63*"
ST_PASS = "Location"

def calc_hash(password: bytes, salt: dict) -> list:
  results = []

  # SHA2-224 (20 percent)
  results.append(hashlib.sha224(password).hexdigest())

  # SHA2-224 with constant prefix (10 percent)
  results.append(hashlib.sha224(b"d3fc0n" + password).hexdigest())

  # SHA2-224 HMAC (10 percent)
  results.append(hmac.new(b"DEFCON", password, hashlib.sha224).hexdigest())

  # SHA2-224 of SHA2-224 (15 percent total with shuck bonus)
  results.append(hashlib.sha224(hashlib.sha224(password).hexdigest().encode()).hexdigest())

  # SHA2-224 of SHA1 (15 percent total with shuck bonus)
  results.append(hashlib.sha224(hashlib.sha1(password).hexdigest().encode()).hexdigest())

  # SHA2-384 (100k Flat)
  results.append(hashlib.sha384(password).hexdigest())

  return results

4. Run hashcat normally

Code:
$ ./hashcat -m 73000 hashes73000.txt example.dict
hashcat (v7.0.0-117-ge6758bf60+) starting

Initializing bridges. Please be patient...
Loaded python library from: /home/atom/.pyenv/versions/3.13.3/lib/libpython3.13.so

Assimilation Bridge
===================
* Unit #01 -> #01: Python Interpreter (3.13.3 (main, Jun  2 2025, 22:02:22) [GCC 13.3.0])
...
Hashes: 275002 digests; 275002 unique digests, 1 unique salts
...
2397b137125c9057e24395767580e9766a340bc13447b5dca4f240d7*:CHOCOLATE
b6a567591e336b320e866187f7993de073589e8518e0b24a6902bcf9*:Learning
bb838248f9023a39bfa873d982aed3260279286ba51975684c42f0ec*:Property
f1cc31f19daa3f07b0278cada69ddfce3619ccbcaf33dd584739e0d8*:Universal
...

- atom]]>
I originally added this plugin for situations where I had reversed some algorithm from an application or library, but usually it was a one-hash scenario. In most cases, I would simply write a quick proof of concept in, crack the single hash, and move on. The Python plugin is for such cases where the hashes are more resistant, for example when they are not based on a weak password. In that case you might need both computational power and some smart password guessing. You may want to take advantage of hashcat features such as the rule engine, hybrid attack modes, restart functionality, or built-in multithreading.

Instead of just writing your own standalone cracking tool, you can later move your PoC logic into the Python bridge and directly benefit from hashcat features.

We recently participated in the Jabbercracky password contest (DEFCON 33) where we faced a set of completely unknown hashes with no context at all.

When you do not know the hash type and there is no obvious signature or pattern in the hashes, a great tool to try is mdxfind. It is multithreaded, has many useful features, and is ideal for such situations. It supports iterations, external salt or username inputs, and more. In this contest, mdxfind could identify some of the algorithms, but some others the players had to figure out manually.

That is a perfect example of when the Python bridge is useful because it gives you a programming language with access to a large pool of third-party cryptographic libraries, and using a programming language gives you maximum flexibility.

Note!!! that in order to use this you need a version higher than v7.0.0. For instance new version v7.1.0 or the beta binary from https://hashcat.net/beta or by compiling from GitHub master sources.


0. Set up the Python bridge

If hashcat -b -m 73000 runs successfully, you are ready for step 1. If you have never used it before, check out the Quickstart Guide.

1. Download the contest hash list

Here is a mirror.

2. Convert it for use with the Python bridge

Code:
$ cat hash_list_3.txt | grep -v '^\$2a\$' | sed -e 's/$/*/g' > hashes73000.txt

Mode 73000 expects the format HASH*SALT. In this contest there were no salts, only peppers, so we leave the salt field empty and handle the pepper in our Python code.  We skip the bcrypts intentionally. You can also implement bcrypt in the Python bridge, but since bcrypt has a recognizable signature and is supported directly with mode 3200, it is better to handle those natively.

We use mode 73000 instead of 72000 for better cross-platform compatibility.

3. Edit Python/generic_hash_mp.py

Code:
import sys
import struct
import hashlib
import hcshared
import hcmp
import hmac

# SELF TEST pair
ST_HASH = "9c016f79fdcdc7422af93fcc203b0eda43278d32e3a6edd607d35c63*"
ST_PASS = "Location"

def calc_hash(password: bytes, salt: dict) -> list:
  results = []

  # SHA2-224 (20 percent)
  results.append(hashlib.sha224(password).hexdigest())

  # SHA2-224 with constant prefix (10 percent)
  results.append(hashlib.sha224(b"d3fc0n" + password).hexdigest())

  # SHA2-224 HMAC (10 percent)
  results.append(hmac.new(b"DEFCON", password, hashlib.sha224).hexdigest())

  # SHA2-224 of SHA2-224 (15 percent total with shuck bonus)
  results.append(hashlib.sha224(hashlib.sha224(password).hexdigest().encode()).hexdigest())

  # SHA2-224 of SHA1 (15 percent total with shuck bonus)
  results.append(hashlib.sha224(hashlib.sha1(password).hexdigest().encode()).hexdigest())

  # SHA2-384 (100k Flat)
  results.append(hashlib.sha384(password).hexdigest())

  return results

4. Run hashcat normally

Code:
$ ./hashcat -m 73000 hashes73000.txt example.dict
hashcat (v7.0.0-117-ge6758bf60+) starting

Initializing bridges. Please be patient...
Loaded python library from: /home/atom/.pyenv/versions/3.13.3/lib/libpython3.13.so

Assimilation Bridge
===================
* Unit #01 -> #01: Python Interpreter (3.13.3 (main, Jun  2 2025, 22:02:22) [GCC 13.3.0])
...
Hashes: 275002 digests; 275002 unique digests, 1 unique salts
...
2397b137125c9057e24395767580e9766a340bc13447b5dca4f240d7*:CHOCOLATE
b6a567591e336b320e866187f7993de073589e8518e0b24a6902bcf9*:Learning
bb838248f9023a39bfa873d982aed3260279286ba51975684c42f0ec*:Property
f1cc31f19daa3f07b0278cada69ddfce3619ccbcaf33dd584739e0d8*:Universal
...

- atom]]> <![CDATA[Wrong password written to the POT file and output file]]> https://hashcat.net/forum/thread-13323.html Fri, 25 Jul 2025 18:06:52 +0000 Dr-Tibetor]]> https://hashcat.net/forum/thread-13323.html
29c37a5bfc5bcd22f185f1dcc5abd86f:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
f4661dc96fb199341e1d170582c0d28a:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
149c937814608b36cacf0824dfa1a817:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
198ebcb6b0c4cf477410fa72dee74716:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
e8a810c553c4e247c172dca674785626:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
49a93b0ef986ab6386c93392da128654:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
0ed82038ed3ddf6874afa080ddd28943:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
1a9b93d008375fcca9bdb60a62af31a6:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
52b00e4ebb1fb9bcbd5604bac23987ba:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
64bc754c198340891a4b23ef5be23d2b:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
f16a739b1354618a9222567a308dd720:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
4e370d11be575098e9eb0ce1f06172e8:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
2e4a4d75e8c16b8f6b07c84fb83e3c8b:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
de810a323d7d5eb8ffd247fc2b1f0289:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
c983b5bd585ad803ccceaa6ebe6821fb:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
b156548bd18ea0c72dfbfd120ff66a1b:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
c207fcc5c7b2d0687cadfc67d622ec63:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
9cce330bd69c57be2798da52e3646178:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…]]>
29c37a5bfc5bcd22f185f1dcc5abd86f:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
f4661dc96fb199341e1d170582c0d28a:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
149c937814608b36cacf0824dfa1a817:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
198ebcb6b0c4cf477410fa72dee74716:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
e8a810c553c4e247c172dca674785626:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
49a93b0ef986ab6386c93392da128654:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
0ed82038ed3ddf6874afa080ddd28943:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
1a9b93d008375fcca9bdb60a62af31a6:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
52b00e4ebb1fb9bcbd5604bac23987ba:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
64bc754c198340891a4b23ef5be23d2b:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
f16a739b1354618a9222567a308dd720:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
4e370d11be575098e9eb0ce1f06172e8:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
2e4a4d75e8c16b8f6b07c84fb83e3c8b:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
de810a323d7d5eb8ffd247fc2b1f0289:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
c983b5bd585ad803ccceaa6ebe6821fb:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
b156548bd18ea0c72dfbfd120ff66a1b:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
c207fcc5c7b2d0687cadfc67d622ec63:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…
9cce330bd69c57be2798da52e3646178:РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…РїС—Р…]]> <![CDATA[New algo module]]> https://hashcat.net/forum/thread-13318.html Thu, 17 Jul 2025 07:08:31 +0000 xray013]]> https://hashcat.net/forum/thread-13318.html
I need new hashcat module with algo sha256(sha256($salt.$pass).$pass) in OPTIMIZED mode (i created pure mode, but optimized very hard for me) - so i need only m0xxxxx-a0-optimized.cl and module_0xxxxx.c files

I will pay for this work.

Please PM me.I will pay forI wi\\ll pay for this workыы\\\\\\\
this work]]>
I need new hashcat module with algo sha256(sha256($salt.$pass).$pass) in OPTIMIZED mode (i created pure mode, but optimized very hard for me) - so i need only m0xxxxx-a0-optimized.cl and module_0xxxxx.c files

I will pay for this work.

Please PM me.I will pay forI wi\\ll pay for this workыы\\\\\\\
this work]]> <![CDATA[Alternative PBKDF1-SHA1 plugin]]> https://hashcat.net/forum/thread-13315.html Wed, 16 Jul 2025 19:22:59 +0000 penguinkeeper]]> https://hashcat.net/forum/thread-13315.html this pull request, adding PBKDF1-SHA1, there are 2 ways to tackle it. The faster but less stable way (exec inside), or the slower but more stable way (exec outside). In this case, I developed both but the slower but more stable way ended up being the one chosen by Atom to use in public Hashcat. To avoid developing that second plugin being a waste of time, it's being posted here so hopefully it can help someone in the future get a few more percent performance Smile

Performance increases are not guaranteed, especially for very weak hardware and support for this plugin is *not* official, please don't make a github issue if you have issues and instead just comment here.

Links:
https://mega.nz/file/fQIFkLgD#3uoyDNv6dc...vhqgRGMW9M
https://drive.google.com/file/d/17_S6JyO...sp=sharing]]> this pull request, adding PBKDF1-SHA1, there are 2 ways to tackle it. The faster but less stable way (exec inside), or the slower but more stable way (exec outside). In this case, I developed both but the slower but more stable way ended up being the one chosen by Atom to use in public Hashcat. To avoid developing that second plugin being a waste of time, it's being posted here so hopefully it can help someone in the future get a few more percent performance Smile

Performance increases are not guaranteed, especially for very weak hardware and support for this plugin is *not* official, please don't make a github issue if you have issues and instead just comment here.

Links:
https://mega.nz/file/fQIFkLgD#3uoyDNv6dc...vhqgRGMW9M
https://drive.google.com/file/d/17_S6JyO...sp=sharing]]> <![CDATA[Extend ethereum wallet, PBKDF2-HMAC-SHA256 to allow arbitrary cipher length]]> https://hashcat.net/forum/thread-12520.html Thu, 06 Mar 2025 07:04:32 +0000 bobby_newport]]> https://hashcat.net/forum/thread-12520.html
I am trying to develop my own module that is very similar to module 15600.  It is a crypto wallet that has a keystore very similar to Ethereum Wallet. It uses PBKDF2-HMAC-SHA256, the correct password can be confirmed by comparing a mac to a generated mac which is the keccak hash of the last 16 bytes of the derived key and cipher text.  I have tested the logic and it does work. 

Module 15600 expects the cipher to be 64 characters long, but the wallet I am trying to crack can be much longer. The current cipher text I have is 156 characters long. 

I have been going at it for several days and struggling. I have also looked at module 16300 for guidance. I am hoping to get some direction on the changes that I need to make. Within the module c file I am replacing this part

  ethereum_pbkdf2->ciphertext[0] = hex_to_u32 (&ciphertext_pos[ 0]);
  ethereum_pbkdf2->ciphertext[1] = hex_to_u32 (&ciphertext_pos[ 8]);
  ethereum_pbkdf2->ciphertext[2] = hex_to_u32 (&ciphertext_pos[16]);
  ethereum_pbkdf2->ciphertext[3] = hex_to_u32 (&ciphertext_pos[24]);
  ethereum_pbkdf2->ciphertext[4] = hex_to_u32 (&ciphertext_pos[32]);
  ethereum_pbkdf2->ciphertext[5] = hex_to_u32 (&ciphertext_pos[40]);
  ethereum_pbkdf2->ciphertext[6] = hex_to_u32 (&ciphertext_pos[48]);
  ethereum_pbkdf2->ciphertext[7] = hex_to_u32 (&ciphertext_pos[56]);

with this part

  crypto_pbkdf2->ciphertext[0] = hex_to_u32 (&ciphertext_pos[ 0]);
  crypto_pbkdf2->ciphertext[1] = hex_to_u32 (&ciphertext_pos[ 8]);
  crypto_pbkdf2->ciphertext[2] = hex_to_u32 (&ciphertext_pos[16]);
  crypto_pbkdf2->ciphertext[3] = hex_to_u32 (&ciphertext_pos[24]);
  crypto_pbkdf2->ciphertext[4] = hex_to_u32 (&ciphertext_pos[32]);
  crypto_pbkdf2->ciphertext[5] = hex_to_u32 (&ciphertext_pos[40]);
  crypto_pbkdf2->ciphertext[6] = hex_to_u32 (&ciphertext_pos[48]);
  crypto_pbkdf2->ciphertext[7] = hex_to_u32 (&ciphertext_pos[56]);
  crypto_pbkdf2->ciphertext[8] = hex_to_u32 (&ciphertext_pos[64]);
  crypto_pbkdf2->ciphertext[9] = hex_to_u32 (&ciphertext_pos[72]);
  crypto_pbkdf2->ciphertext[10] = hex_to_u32 (&ciphertext_pos[80]);
  crypto_pbkdf2->ciphertext[11] = hex_to_u32 (&ciphertext_pos[88]);
  crypto_pbkdf2->ciphertext[12] = hex_to_u32 (&ciphertext_pos[96]);
  crypto_pbkdf2->ciphertext[13] = hex_to_u32 (&ciphertext_pos[104]);
  crypto_pbkdf2->ciphertext[14] = hex_to_u32 (&ciphertext_pos[112]);
  crypto_pbkdf2->ciphertext[15] = hex_to_u32 (&ciphertext_pos[120]);
  crypto_pbkdf2->ciphertext[16] = hex_to_u32 (&ciphertext_pos[128]);
  crypto_pbkdf2->ciphertext[17] = hex_to_u32 (&ciphertext_pos[136]);
  crypto_pbkdf2->ciphertext[18] = hex_to_u32 (&ciphertext_pos[144]);
  crypto_pbkdf2->ciphertext[19] = hex_to_u32 (&ciphertext_pos[152]);


With the encoder I have replaced 

const int line_len = snprintf (line_buf, line_size, "%s*%u*%s*%08x%08x%08x%08x%08x%08x%08x%08x*%08x%08x%08x%08x%08x%08x%08x%08x",
    SIGNATURE_ETHEREUM_PBKDF2,
    salt->salt_iter + 1,
    tmp_salt,
    byte_swap_32 (ethereum_pbkdf2->ciphertext[0]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[1]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[2]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[3]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[4]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[5]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[6]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[7]),
    byte_swap_32 (digest[0]),
    byte_swap_32 (digest[1]),
    byte_swap_32 (digest[2]),
    byte_swap_32 (digest[3]),
    byte_swap_32 (digest[4]),
    byte_swap_32 (digest[5]),
    byte_swap_32 (digest[6]),
    byte_swap_32 (digest[7])
  );

with this part

const int line_len = snprintf (line_buf, line_size, "%s*%u*%s*%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x*%08x%08x%08x%08x%08x%08x%08x%08x",
    SIGNATURE_CRYPTO_PBKDF2,
    salt->salt_iter + 1,
    tmp_salt,
    byte_swap_32 (crypto_pbkdf2->ciphertext[0]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[1]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[2]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[3]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[4]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[5]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[6]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[7]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[8]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[9]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[10]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[11]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[12]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[13]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[14]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[15]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[16]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[17]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[18]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[19]),
    byte_swap_32 (digest[0]),
    byte_swap_32 (digest[1]),
    byte_swap_32 (digest[2]),
    byte_swap_32 (digest[3]),
    byte_swap_32 (digest[4]),
    byte_swap_32 (digest[5]),
    byte_swap_32 (digest[6]),
    byte_swap_32 (digest[7])
  );


In the kernel I have replaced this

/**
   * keccak
   */
  u32 ciphertext[8];
  ciphertext[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[0];
  ciphertext[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[1];
  ciphertext[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[2];
  ciphertext[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[3];
  ciphertext[4] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[4];
  ciphertext[5] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[5];
  ciphertext[6] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[6];
  ciphertext[7] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[7];
  u32 key[4];
  key[0] = hc_swap32_S (tmps[gid].out[4]);
  key[1] = hc_swap32_S (tmps[gid].out[5]);
  key[2] = hc_swap32_S (tmps[gid].out[6]);
  key[3] = hc_swap32_S (tmps[gid].out[7]);
  u64 st[25];
  st[ 0] = hl32_to_64_S (key[1], key[0]);
  st[ 1] = hl32_to_64_S (key[3], key[2]);
  st[ 2] = hl32_to_64_S (ciphertext[1], ciphertext[0]);
  st[ 3] = hl32_to_64_S (ciphertext[3], ciphertext[2]);
  st[ 4] = hl32_to_64_S (ciphertext[5], ciphertext[4]);
  st[ 5] = hl32_to_64_S (ciphertext[7], ciphertext[6]);
  st[ 6] = 0x01;
  st[ 7] = 0;
  st[ 8] = 0;
  st[ 9] = 0;
  st[10] = 0;
  st[11] = 0;
  st[12] = 0;
  st[13] = 0;
  st[14] = 0;
  st[15] = 0;
  st[16] = 0;
  st[17] = 0;
  st[18] = 0;
  st[19] = 0;
  st[20] = 0;
  st[21] = 0;
  st[22] = 0;
  st[23] = 0;
  st[24] = 0;


with 

u32 ciphertext[20];
  ciphertext[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[0];
  ciphertext[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[1];
  ciphertext[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[2];
  ciphertext[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[3];
  ciphertext[4] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[4];
  ciphertext[5] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[5];
  ciphertext[6] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[6];
  ciphertext[7] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[7];
  ciphertext[8] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[8];
  ciphertext[9] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[9];
  ciphertext[10] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[10];
  ciphertext[11] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[11];
  ciphertext[12] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[12];
  ciphertext[13] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[13];
  ciphertext[14] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[14];
  ciphertext[15] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[15];
  ciphertext[16] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[16];
  ciphertext[17] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[17];
  ciphertext[18] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[18];
  ciphertext[19] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[19];
  u32 key[4];
  key[0] = hc_swap32_S (tmps[gid].out[4]);
  key[1] = hc_swap32_S (tmps[gid].out[5]);
  key[2] = hc_swap32_S (tmps[gid].out[6]);
  key[3] = hc_swap32_S (tmps[gid].out[7]);
  u64 st[25];
  st[ 0] = hl32_to_64_S (key[1], key[0]);
  st[ 1] = hl32_to_64_S (key[3], key[2]);
  st[ 2] = hl32_to_64_S (ciphertext[1], ciphertext[0]);
  st[ 3] = hl32_to_64_S (ciphertext[3], ciphertext[2]);
  st[ 4] = hl32_to_64_S (ciphertext[5], ciphertext[4]);
  st[ 5] = hl32_to_64_S (ciphertext[7], ciphertext[6]);
  st[ 6] = hl32_to_64_S (ciphertext[9], ciphertext[8]);
  st[ 7] = hl32_to_64_S (ciphertext[11], ciphertext[10]);
  st[ 8] = hl32_to_64_S (ciphertext[13], ciphertext[12]);
  st[ 9] = hl32_to_64_S (ciphertext[15], ciphertext[14]);
  st[10] = hl32_to_64_S (ciphertext[17], ciphertext[16]);
  st[11] = hl32_to_64_S (ciphertext[19], ciphertext[18]);
  st[12] = 0x01;
  st[13] = 0;
  st[14] = 0;
  st[15] = 0;
  st[16] = 0;
  st[17] = 0;
  st[18] = 0;
  st[19] = 0;
  st[20] = 0;
  st[21] = 0;
  st[22] = 0;
  st[23] = 0;
  st[24] = 0;

I am clearly missing something as the test fails and when I ignore the test, it doesn't return the true password that is in the list. The error when running the module is 

* Device #1: ATTENTION! OpenCL kernel self-test failed.


Your device driver installation is probably broken.
See also: https://hashcat.net/faq/wrongdriver

Aborting session due to kernel self-test failure.

You can use --self-test-disable to override, but do not report related errors.

Module 15600, along with many others work without an issue so I know it is a problem with my code. 


In an ideal world the module would take a range of cipher lengths, which would open it to many more wallets. Full disclosure, I have minimal understanding of the C language (mainly Python and Java). 

Any help will be appreciated.]]>
I am trying to develop my own module that is very similar to module 15600.  It is a crypto wallet that has a keystore very similar to Ethereum Wallet. It uses PBKDF2-HMAC-SHA256, the correct password can be confirmed by comparing a mac to a generated mac which is the keccak hash of the last 16 bytes of the derived key and cipher text.  I have tested the logic and it does work. 

Module 15600 expects the cipher to be 64 characters long, but the wallet I am trying to crack can be much longer. The current cipher text I have is 156 characters long. 

I have been going at it for several days and struggling. I have also looked at module 16300 for guidance. I am hoping to get some direction on the changes that I need to make. Within the module c file I am replacing this part

  ethereum_pbkdf2->ciphertext[0] = hex_to_u32 (&ciphertext_pos[ 0]);
  ethereum_pbkdf2->ciphertext[1] = hex_to_u32 (&ciphertext_pos[ 8]);
  ethereum_pbkdf2->ciphertext[2] = hex_to_u32 (&ciphertext_pos[16]);
  ethereum_pbkdf2->ciphertext[3] = hex_to_u32 (&ciphertext_pos[24]);
  ethereum_pbkdf2->ciphertext[4] = hex_to_u32 (&ciphertext_pos[32]);
  ethereum_pbkdf2->ciphertext[5] = hex_to_u32 (&ciphertext_pos[40]);
  ethereum_pbkdf2->ciphertext[6] = hex_to_u32 (&ciphertext_pos[48]);
  ethereum_pbkdf2->ciphertext[7] = hex_to_u32 (&ciphertext_pos[56]);

with this part

  crypto_pbkdf2->ciphertext[0] = hex_to_u32 (&ciphertext_pos[ 0]);
  crypto_pbkdf2->ciphertext[1] = hex_to_u32 (&ciphertext_pos[ 8]);
  crypto_pbkdf2->ciphertext[2] = hex_to_u32 (&ciphertext_pos[16]);
  crypto_pbkdf2->ciphertext[3] = hex_to_u32 (&ciphertext_pos[24]);
  crypto_pbkdf2->ciphertext[4] = hex_to_u32 (&ciphertext_pos[32]);
  crypto_pbkdf2->ciphertext[5] = hex_to_u32 (&ciphertext_pos[40]);
  crypto_pbkdf2->ciphertext[6] = hex_to_u32 (&ciphertext_pos[48]);
  crypto_pbkdf2->ciphertext[7] = hex_to_u32 (&ciphertext_pos[56]);
  crypto_pbkdf2->ciphertext[8] = hex_to_u32 (&ciphertext_pos[64]);
  crypto_pbkdf2->ciphertext[9] = hex_to_u32 (&ciphertext_pos[72]);
  crypto_pbkdf2->ciphertext[10] = hex_to_u32 (&ciphertext_pos[80]);
  crypto_pbkdf2->ciphertext[11] = hex_to_u32 (&ciphertext_pos[88]);
  crypto_pbkdf2->ciphertext[12] = hex_to_u32 (&ciphertext_pos[96]);
  crypto_pbkdf2->ciphertext[13] = hex_to_u32 (&ciphertext_pos[104]);
  crypto_pbkdf2->ciphertext[14] = hex_to_u32 (&ciphertext_pos[112]);
  crypto_pbkdf2->ciphertext[15] = hex_to_u32 (&ciphertext_pos[120]);
  crypto_pbkdf2->ciphertext[16] = hex_to_u32 (&ciphertext_pos[128]);
  crypto_pbkdf2->ciphertext[17] = hex_to_u32 (&ciphertext_pos[136]);
  crypto_pbkdf2->ciphertext[18] = hex_to_u32 (&ciphertext_pos[144]);
  crypto_pbkdf2->ciphertext[19] = hex_to_u32 (&ciphertext_pos[152]);


With the encoder I have replaced 

const int line_len = snprintf (line_buf, line_size, "%s*%u*%s*%08x%08x%08x%08x%08x%08x%08x%08x*%08x%08x%08x%08x%08x%08x%08x%08x",
    SIGNATURE_ETHEREUM_PBKDF2,
    salt->salt_iter + 1,
    tmp_salt,
    byte_swap_32 (ethereum_pbkdf2->ciphertext[0]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[1]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[2]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[3]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[4]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[5]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[6]),
    byte_swap_32 (ethereum_pbkdf2->ciphertext[7]),
    byte_swap_32 (digest[0]),
    byte_swap_32 (digest[1]),
    byte_swap_32 (digest[2]),
    byte_swap_32 (digest[3]),
    byte_swap_32 (digest[4]),
    byte_swap_32 (digest[5]),
    byte_swap_32 (digest[6]),
    byte_swap_32 (digest[7])
  );

with this part

const int line_len = snprintf (line_buf, line_size, "%s*%u*%s*%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x*%08x%08x%08x%08x%08x%08x%08x%08x",
    SIGNATURE_CRYPTO_PBKDF2,
    salt->salt_iter + 1,
    tmp_salt,
    byte_swap_32 (crypto_pbkdf2->ciphertext[0]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[1]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[2]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[3]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[4]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[5]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[6]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[7]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[8]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[9]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[10]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[11]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[12]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[13]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[14]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[15]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[16]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[17]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[18]),
    byte_swap_32 (crypto_pbkdf2->ciphertext[19]),
    byte_swap_32 (digest[0]),
    byte_swap_32 (digest[1]),
    byte_swap_32 (digest[2]),
    byte_swap_32 (digest[3]),
    byte_swap_32 (digest[4]),
    byte_swap_32 (digest[5]),
    byte_swap_32 (digest[6]),
    byte_swap_32 (digest[7])
  );


In the kernel I have replaced this

/**
   * keccak
   */
  u32 ciphertext[8];
  ciphertext[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[0];
  ciphertext[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[1];
  ciphertext[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[2];
  ciphertext[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[3];
  ciphertext[4] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[4];
  ciphertext[5] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[5];
  ciphertext[6] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[6];
  ciphertext[7] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[7];
  u32 key[4];
  key[0] = hc_swap32_S (tmps[gid].out[4]);
  key[1] = hc_swap32_S (tmps[gid].out[5]);
  key[2] = hc_swap32_S (tmps[gid].out[6]);
  key[3] = hc_swap32_S (tmps[gid].out[7]);
  u64 st[25];
  st[ 0] = hl32_to_64_S (key[1], key[0]);
  st[ 1] = hl32_to_64_S (key[3], key[2]);
  st[ 2] = hl32_to_64_S (ciphertext[1], ciphertext[0]);
  st[ 3] = hl32_to_64_S (ciphertext[3], ciphertext[2]);
  st[ 4] = hl32_to_64_S (ciphertext[5], ciphertext[4]);
  st[ 5] = hl32_to_64_S (ciphertext[7], ciphertext[6]);
  st[ 6] = 0x01;
  st[ 7] = 0;
  st[ 8] = 0;
  st[ 9] = 0;
  st[10] = 0;
  st[11] = 0;
  st[12] = 0;
  st[13] = 0;
  st[14] = 0;
  st[15] = 0;
  st[16] = 0;
  st[17] = 0;
  st[18] = 0;
  st[19] = 0;
  st[20] = 0;
  st[21] = 0;
  st[22] = 0;
  st[23] = 0;
  st[24] = 0;


with 

u32 ciphertext[20];
  ciphertext[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[0];
  ciphertext[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[1];
  ciphertext[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[2];
  ciphertext[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[3];
  ciphertext[4] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[4];
  ciphertext[5] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[5];
  ciphertext[6] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[6];
  ciphertext[7] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[7];
  ciphertext[8] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[8];
  ciphertext[9] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[9];
  ciphertext[10] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[10];
  ciphertext[11] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[11];
  ciphertext[12] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[12];
  ciphertext[13] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[13];
  ciphertext[14] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[14];
  ciphertext[15] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[15];
  ciphertext[16] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[16];
  ciphertext[17] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[17];
  ciphertext[18] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[18];
  ciphertext[19] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[19];
  u32 key[4];
  key[0] = hc_swap32_S (tmps[gid].out[4]);
  key[1] = hc_swap32_S (tmps[gid].out[5]);
  key[2] = hc_swap32_S (tmps[gid].out[6]);
  key[3] = hc_swap32_S (tmps[gid].out[7]);
  u64 st[25];
  st[ 0] = hl32_to_64_S (key[1], key[0]);
  st[ 1] = hl32_to_64_S (key[3], key[2]);
  st[ 2] = hl32_to_64_S (ciphertext[1], ciphertext[0]);
  st[ 3] = hl32_to_64_S (ciphertext[3], ciphertext[2]);
  st[ 4] = hl32_to_64_S (ciphertext[5], ciphertext[4]);
  st[ 5] = hl32_to_64_S (ciphertext[7], ciphertext[6]);
  st[ 6] = hl32_to_64_S (ciphertext[9], ciphertext[8]);
  st[ 7] = hl32_to_64_S (ciphertext[11], ciphertext[10]);
  st[ 8] = hl32_to_64_S (ciphertext[13], ciphertext[12]);
  st[ 9] = hl32_to_64_S (ciphertext[15], ciphertext[14]);
  st[10] = hl32_to_64_S (ciphertext[17], ciphertext[16]);
  st[11] = hl32_to_64_S (ciphertext[19], ciphertext[18]);
  st[12] = 0x01;
  st[13] = 0;
  st[14] = 0;
  st[15] = 0;
  st[16] = 0;
  st[17] = 0;
  st[18] = 0;
  st[19] = 0;
  st[20] = 0;
  st[21] = 0;
  st[22] = 0;
  st[23] = 0;
  st[24] = 0;

I am clearly missing something as the test fails and when I ignore the test, it doesn't return the true password that is in the list. The error when running the module is 

* Device #1: ATTENTION! OpenCL kernel self-test failed.


Your device driver installation is probably broken.
See also: https://hashcat.net/faq/wrongdriver

Aborting session due to kernel self-test failure.

You can use --self-test-disable to override, but do not report related errors.

Module 15600, along with many others work without an issue so I know it is a problem with my code. 


In an ideal world the module would take a range of cipher lengths, which would open it to many more wallets. Full disclosure, I have minimal understanding of the C language (mainly Python and Java). 

Any help will be appreciated.]]>