Diskuze: Jak upravit program pro hledání hesla ?

Caster

Člen

Caster:6.10.2022 13:09

Přestože program níže na Linuxu v pohodě funguje, netuším, jak ho upravit aby hledal heslo v testovacím souboru *.hc22000 Hash-Mode 22000 (WPA-PBKDF2-PMKID+EAPOL).

Zkusil jsem: main.cu

#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <cuda.h>
#include <sys/time.h>
#include <pthread.h>
#include <locale.h>
#include "sha256.cuh"

#define TEXT "Caster"
#define TEXT_LEN 13
#define THREADS 1500
#define BLOCKS 256
#define GPUS 4
#define DIFFICULTY 4
#define RANDOM_LEN 20

__constant__ BYTE characterSet[63] = {"abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ1234567890"};

__global__ void initSolutionMemory(int *blockContainsSolution) {
    *blockContainsSolution = -1;
}

__device__ unsigned long deviceRandomGen(unsigned long x) {
    x ^= (x << 21);
    x ^= (x >> 35);
    x ^= (x << 4);
    return x;
}

__global__ void sha256_cuda(BYTE *prefix, BYTE *solution, int *blockContainsSolution, unsigned long baseSeed) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    SHA256_CTX ctx;
    BYTE digest[32];
    BYTE random[RANDOM_LEN];
    unsigned long seed = baseSeed;
    seed += (unsigned long) i;
    for (int j = 0; j < RANDOM_LEN; j++) {
        seed = deviceRandomGen(seed);
        int randomIdx = (int) (seed % 62);
        random[j] = characterSet[randomIdx];
    }
    sha256_init(&ctx);
    sha256_update(&ctx, prefix, TEXT_LEN);
    sha256_update(&ctx, random, RANDOM_LEN);
    sha256_final(&ctx, digest);
    for (int j = 0; j < DIFFICULTY; j++)
        if (digest[j] > 0)
            return;
    if ((digest[DIFFICULTY] & 0xF0) > 0)
        return;
    if (*blockContainsSolution == 1)
        return;
    *blockContainsSolution = 1;
    for (int j = 0; j < RANDOM_LEN; j++)
        solution[j] = random[j];
}

void hostRandomGen(unsigned long *x) {
    *x ^= (*x << 21);
    *x ^= (*x >> 35);
    *x ^= (*x << 4);
}

void pre_sha256() {
    cudaMemcpyToSymbol(dev_k, host_k, sizeof(host_k), 0, cudaMemcpyHostToDevice);
}

long long timems() {
    struct timeval end;
    gettimeofday(&end, NULL);
    return end.tv_sec * 1000LL + end.tv_usec / 1000;
}

struct HandlerInput {
    int device;
    unsigned long hashesProcessed;
};
typedef struct HandlerInput HandlerInput;

pthread_mutex_t solutionLock;
BYTE *solution;

void *launchGPUHandlerThread(void *vargp) {
    HandlerInput *hi = (HandlerInput *) vargp;
    cudaSetDevice(hi->device);

    pre_sha256();

    BYTE cpuPrefix[] = {TEXT};
    BYTE *d_prefix;
    cudaMalloc(&d_prefix, TEXT_LEN);
    cudaMemcpy(d_prefix, cpuPrefix, TEXT_LEN, cudaMemcpyHostToDevice);

    BYTE *blockSolution = (BYTE *) malloc(sizeof(BYTE) * RANDOM_LEN);
    BYTE *d_solution;
    cudaMalloc(&d_solution, sizeof(BYTE) * RANDOM_LEN);

    int *blockContainsSolution = (int *) malloc(sizeof(int));
    int *d_blockContainsSolution;
    cudaMalloc(&d_blockContainsSolution, sizeof(int));

    unsigned long rngSeed = timems();

    initSolutionMemory<<<1, 1>>>(d_blockContainsSolution);

    while (1) {
        hostRandomGen(&rngSeed);

        hi->hashesProcessed += THREADS * BLOCKS;
        sha256_cuda<<<THREADS, BLOCKS>>>(d_prefix, d_solution, d_blockContainsSolution, rngSeed);
        cudaDeviceSynchronize();

        cudaMemcpy(blockContainsSolution, d_blockContainsSolution, sizeof(int), cudaMemcpyDeviceToHost);
        if (*blockContainsSolution == 1) {
            cudaMemcpy(blockSolution, d_solution, sizeof(BYTE) * RANDOM_LEN, cudaMemcpyDeviceToHost);
            solution = blockSolution;
            pthread_mutex_unlock(&solutionLock);
            break;
        }

        if (solution) {
            break;
        }
    }

    cudaDeviceReset();
    return NULL;
}

int main() {
    setlocale(LC_NUMERIC, "");

    pthread_mutex_init(&solutionLock, NULL);
    pthread_mutex_lock(&solutionLock);

    unsigned long **processedPtrs = (unsigned long **) malloc(sizeof(unsigned long *) * GPUS);
    pthread_t *tids = (pthread_t *) malloc(sizeof(pthread_t) * GPUS);
    long long start = timems();
    for (int i = 0; i < GPUS; i++) {
        HandlerInput *hi = (HandlerInput *) malloc(sizeof(HandlerInput));
        hi->device = i;
        hi->hashesProcessed = 0;
        processedPtrs[i] = &hi->hashesProcessed;
        pthread_create(tids + i, NULL, launchGPUHandlerThread, hi);
        usleep(10);
    }

    while (1) {
        unsigned long totalProcessed = 0;
        for (int i = 0; i < GPUS; i++) {
            totalProcessed += *(processedPtrs[i]);
        }
        long long elapsed = timems() - start;
        printf("Hashes (%'lu) Seconds (%'f) Hashes/sec (%'lu)\r", totalProcessed, ((float) elapsed) / 1000.0, (unsigned long) ((double) totalProcessed / (double) elapsed) * 1000);
        if (solution) {
            break;
        }
    }
    printf("\n");

    pthread_mutex_lock(&solutionLock);
    long long end = timems();
    long long elapsed = end - start;

    for (int i = 0; i < GPUS; i++) {
        pthread_join(tids[i], NULL);
    }

    unsigned long totalProcessed = 0;
    for (int i = 0; i < GPUS; i++) {
        totalProcessed += *(processedPtrs[i]);
    }

    printf("Solution: %.20s\n", solution);
    printf("Hashes processed: %'lu\n", totalProcessed);
    printf("Time: %llu\n", elapsed);
    printf("Hashes/sec: %'lu\n", (unsigned long) ((double) totalProcessed / (double) elapsed) * 1000);

    return 0;
}

sha256.cuh

#ifndef SHA256_H
#define SHA256_H


/****************************** MACROS ******************************/
#define SHA256_BLOCK_SIZE 32            // SHA256 outputs a 32 byte digest

#define ROTLEFT(a,b) (((a) << (b)) | ((a) >> (32-(b))))
#define ROTRIGHT(a,b) (((a) >> (b)) | ((a) << (32-(b))))

#define CH(x,y,z) (((x) & (y)) ^ (~(x) & (z)))
#define MAJ(x,y,z) (((x) & (y)) ^ ((x) & (z)) ^ ((y) & (z)))
#define EP0(x) (ROTRIGHT(x,2) ^ ROTRIGHT(x,13) ^ ROTRIGHT(x,22))
#define EP1(x) (ROTRIGHT(x,6) ^ ROTRIGHT(x,11) ^ ROTRIGHT(x,25))
#define SIG0(x) (ROTRIGHT(x,7) ^ ROTRIGHT(x,18) ^ ((x) >> 3))
#define SIG1(x) (ROTRIGHT(x,17) ^ ROTRIGHT(x,19) ^ ((x) >> 10))

#define checkCudaErrors(x) \
{ \
    cudaGetLastError(); \
    x; \
    cudaError_t err = cudaGetLastError(); \
    if (err != cudaSuccess) \
        printf("GPU: cudaError %d (%s)\n", err, cudaGetErrorString(err)); \
}
/**************************** DATA TYPES ****************************/
typedef unsigned char BYTE;             // 8-bit byte
typedef uint32_t  WORD;             // 32-bit word, change to "long" for 16-bit machines

typedef struct {
        BYTE data[64];
        WORD datalen;
        unsigned long long bitlen;
        WORD state[8];
} SHA256_CTX;

__constant__ WORD dev_k[64];

static const WORD host_k[64] = {
        0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
        0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174,
        0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da,
        0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967,
        0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85,
        0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070,
        0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3,
        0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2
};

/*********************** FUNCTION DECLARATIONS **********************/
char * print_sha(BYTE * buff);
__device__ void sha256_init(SHA256_CTX *ctx);
__device__ void sha256_update(SHA256_CTX *ctx, const BYTE data[], size_t len);
__device__ void sha256_final(SHA256_CTX *ctx, BYTE hash[]);


char * hash_to_string(BYTE * buff) {
        char * string = (char *)malloc(70);
        int k, i;
        for (i = 0, k = 0; i < 32; i++, k+= 2)
        {
                sprintf(string + k, "%.2x", buff[i]);
                //printf("%02x", buff[i]);
        }
        string[64] = 0;
        return string;
}

__device__ void mycpy12(uint32_t *d, const uint32_t *s) {
#pragma unroll 3
    for (int k=0; k < 3; k++) d[k] = s[k];
}

__device__ void mycpy16(uint32_t *d, const uint32_t *s) {
#pragma unroll 4
    for (int k=0; k < 4; k++) d[k] = s[k];
}

__device__ void mycpy32(uint32_t *d, const uint32_t *s) {
#pragma unroll 8
    for (int k=0; k < 8; k++) d[k] = s[k];
}

__device__ void mycpy44(uint32_t *d, const uint32_t *s) {
#pragma unroll 11
    for (int k=0; k < 11; k++) d[k] = s[k];
}

__device__ void mycpy48(uint32_t *d, const uint32_t *s) {
#pragma unroll 12
    for (int k=0; k < 12; k++) d[k] = s[k];
}

__device__ void mycpy64(uint32_t *d, const uint32_t *s) {
#pragma unroll 16
    for (int k=0; k < 16; k++) d[k] = s[k];
}

__device__ void sha256_transform(SHA256_CTX *ctx, const BYTE data[])
{
        WORD a, b, c, d, e, f, g, h, i, j, t1, t2, m[64];
    WORD S[8];

    //mycpy32(S, ctx->state);

    #pragma unroll 16
        for (i = 0, j = 0; i < 16; ++i, j += 4)
                m[i] = (data[j] << 24) | (data[j + 1] << 16) | (data[j + 2] << 8) | (data[j + 3]);

    #pragma unroll 64
        for (; i < 64; ++i)
                m[i] = SIG1(m[i - 2]) + m[i - 7] + SIG0(m[i - 15]) + m[i - 16];

        a = ctx->state[0];
        b = ctx->state[1];
        c = ctx->state[2];
        d = ctx->state[3];
        e = ctx->state[4];
        f = ctx->state[5];
        g = ctx->state[6];
        h = ctx->state[7];

    #pragma unroll 64
        for (i = 0; i < 64; ++i) {
                t1 = h + EP1(e) + CH(e, f, g) + dev_k[i] + m[i];
                t2 = EP0(a) + MAJ(a, b, c);
                h = g;
                g = f;
                f = e;
                e = d + t1;
                d = c;
                c = b;
                b = a;
                a = t1 + t2;
        }

        ctx->state[0] += a;
        ctx->state[1] += b;
        ctx->state[2] += c;
        ctx->state[3] += d;
        ctx->state[4] += e;
        ctx->state[5] += f;
        ctx->state[6] += g;
        ctx->state[7] += h;
}

__device__ void sha256_init(SHA256_CTX *ctx)
{
        ctx->datalen = 0;
        ctx->bitlen = 0;
        ctx->state[0] = 0x6a09e667;
        ctx->state[1] = 0xbb67ae85;
        ctx->state[2] = 0x3c6ef372;
        ctx->state[3] = 0xa54ff53a;
        ctx->state[4] = 0x510e527f;
        ctx->state[5] = 0x9b05688c;
        ctx->state[6] = 0x1f83d9ab;
        ctx->state[7] = 0x5be0cd19;
}

__device__ void sha256_update(SHA256_CTX *ctx, const BYTE data[], size_t len)
{
        WORD i;

        // for each byte in message
        for (i = 0; i < len; ++i) {
                // ctx->data == message 512 bit chunk
                ctx->data[ctx->datalen] = data[i];
                ctx->datalen++;
                if (ctx->datalen == 64) {
                        sha256_transform(ctx, ctx->data);
                        ctx->bitlen += 512;
                        ctx->datalen = 0;
                }
        }
}

__device__ void sha256_final(SHA256_CTX *ctx, BYTE hash[])
{
        WORD i;

        i = ctx->datalen;

        // Pad whatever data is left in the buffer.
        if (ctx->datalen < 56) {
                ctx->data[i++] = 0x80;
                while (i < 56)
                        ctx->data[i++] = 0x00;
        }
        else {
                ctx->data[i++] = 0x80;
                while (i < 64)
                        ctx->data[i++] = 0x00;
                sha256_transform(ctx, ctx->data);
                memset(ctx->data, 0, 56);
        }

        // Append to the padding the total message's length in bits and transform.
        ctx->bitlen += ctx->datalen * 8;
        ctx->data[63] = ctx->bitlen;
        ctx->data[62] = ctx->bitlen >> 8;
        ctx->data[61] = ctx->bitlen >> 16;
        ctx->data[60] = ctx->bitlen >> 24;
        ctx->data[59] = ctx->bitlen >> 32;
        ctx->data[58] = ctx->bitlen >> 40;
        ctx->data[57] = ctx->bitlen >> 48;
        ctx->data[56] = ctx->bitlen >> 56;
        sha256_transform(ctx, ctx->data);

        // Since this implementation uses little endian byte ordering and SHA uses big endian,
        // reverse all the bytes when copying the final state to the output hash.
        for (i = 0; i < 4; ++i) {
                hash[i] = (ctx->state[0] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 4] = (ctx->state[1] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 8] = (ctx->state[2] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 12] = (ctx->state[3] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 16] = (ctx->state[4] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 20] = (ctx->state[5] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 24] = (ctx->state[6] >> (24 - i * 8)) & 0x000000ff;
                hash[i + 28] = (ctx->state[7] >> (24 - i * 8)) & 0x000000ff;
        }
}

#endif   // SHA256_H

\---

Chci docílit: Upravit program, aby hledal heslo (8 znaků - malá a velká písmena + číslice) v testovacím souboru *.hc22000, případně data zadat rovnou do programu:

WPA01e4e144106e26778d0b2349f00014b90b2c4d549675942446c84d075c*5374757a6b615f3547***
WPA0283e42ba03461533b2960cd1cc4fb28ab2c4d549675942446c84d075c5374757a6b615f35478b2ed778f07caa662b67cb07ad1a3bb29fa8014284b0081df4982b5ed703e1b90103007502010a0000000000000000000047b0cf74c4e4884e8e3e38c3a14949503e8c654fec91ec11f459f9a8837176eb000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000001630140100000fac040100000fac040100000fac02000002

Odpovědět

6.10.2022 13:09

Caster

Člen

Caster:9.10.2022 13:18

Po úspěšném spuštění prvního testovacího programu na GPU (Visual Studio 2022 a CUDA Toolkit) se pokouším převést program výše z Linuxu do Visual Studia 2022 (VS 2022).

Potřeboval bych poradit, jakou knihovnou ve VS 2022 nahradit Linuxovou knihovnu:

#include <sys/time.h>

Při měření času je v programu použita funkce:

long long timems() {
    struct timeval end;
    gettimeofday(&end, NULL);
    return end.tv_sec * 1000LL + end.tv_usec / 1000;
}

Nahoru Odpovědět

9.10.2022 13:18

DarkCoder

Člen

DarkCoder:9.10.2022 14:39

#include <stdio.h>
#include <winsock2.h> // struct timeval
#include <windows.h> // timeGetTime()

#pragma comment( lib, "winmm.lib") // timeGetTime()

typedef unsigned long DWORD;

int gettimeofday(struct timeval* tp, void* tzp);
long long timems(void);

int main(void) {
        long long val = timems();
        printf("%lld\n", val);

        return 0;
}

int gettimeofday(struct timeval* tp, void* tzp) {
        DWORD t;

        t = timeGetTime();
        tp->tv_sec = t / 1000;
        tp->tv_usec = t % 1000;

        return 0;
}

long long timems(void) {
        struct timeval end;
        gettimeofday(&end, NULL);
        return end.tv_sec * 1000LL + end.tv_usec / 1000;
}

Nahoru Odpovědět

9.10.2022 14:39

"I ta nejlepší poučka postrádá na významu, není-li patřičně předána." - DarkCoder

Caster

Člen

Caster:9.10.2022 15:03

Super skvělý, díky, na to bych nepřišel. Chyby u času již zmizely ;-) .

Poslední problém je jak nahradit funkce (ovládání vláken):

pthread_mutex_init()
pthread_mutex_lock()
pthread_create()
pthread_mutex_unlock()

Linux knihovny:

#include <pthread.h>

část programu viz výše:

pthread_mutex_init(&solutionLock, NULL);
pthread_mutex_lock(&solutionLock);

unsigned long** processedPtrs = (unsigned long**)malloc(sizeof(unsigned long*) * GPUS);
pthread_t* tids = (pthread_t*)malloc(sizeof(pthread_t) * GPUS);
long long start = timems();
for (int i = 0; i < GPUS; i++) {
    HandlerInput* hi = (HandlerInput*)malloc(sizeof(HandlerInput));
    hi->device = i;
    hi->hashesProcessed = 0;
    processedPtrs[i] = &hi->hashesProcessed;
    pthread_create(tids + i, NULL, launchGPUHandlerThread, hi);
    usleep(10);
}

funkcemi CUDA Toolkit. Dal jsem dotaz do fóra developer.nvidia.com

Nahoru Odpovědět

9.10.2022 15:03

Peter Mlich

Člen

Peter Mlich:9.10.2022 16:36

Nebylo by rychlejsi misto volani cyklu pouzit primo instrukci?

for (int k=0; k < 3; k++) d[k] = s[k];

d[0] = s[0];
d[1] = s[1];
d[2] = s[2];

Zvlast, kdyz to cele provadis opakovane 1000x, tak by to mohlo mit vyznam

Nahoru Odpovědět

9.10.2022 16:36

Caster

Člen

Caster:9.10.2022 17:25

@Peter Mlich

Nerozumím tvému komentáři ohledně volání cyklu. Jaký cyklus myslíš ?

P.S. Program v úvodu běží na Linuxu (KALI) na mém notebooku s GPU NVIDIA GeForce 960M v pohodě rychlostí cca 32 Gh/s ;-) . Upravuji ho, abych ho mohl spustit pod Visual Studio 2022 s NVIDIA Toolkitem. Hledám také nějaký příklad, jak naprogramovat funkci PBKDF2(HMAC−SHA1, passphrase, ssid, 4096, 256) aby běžela co nejrychleji (s využitím nejnovějších instrukcí CPU a během ve více warpech na GPU).

Nahoru Odpovědět

9.10.2022 17:25

Peter Mlich

Člen

Peter Mlich:9.10.2022 20:20

for (int k=0; k < 3; k++) d[k] = s[k]; // tohle tam mas, a je to preci totez jako:

d[0] = s[0];
d[1] = s[1];
d[2] = s[2];

Az na to, ze, bez cyklu to pobezi rychleji a spotrebuje mene pameti.
Ale, nevim, jaky zpusobem funguji moderni urychlovace cyklu. Je mozne, ze to cpu dokaze zpracovat paralelne a ten rozdil je pak minimalni.

Tez je zajimave, ze mas cykly:

  for (i = 0; i < 64; ++i) // 1
    for (int k=0; k < 11; k++) // 2
// kde ++i by melo byt rychlejsi jez i++

Proste, to vypada jako slepenec nekolika kodu, kde vlastne ani nerozumis, co to dela a proc a to se pokousis prepsat na neco jineho, cemu take nerozumis? Jen mi to prislo zvlastni, jako.

Nahoru Odpovědět

9.10.2022 20:20

Caster

Člen

Caster:9.10.2022 20:33

Díky, jen upřesňuji, že kód je jiného autora viz odkaz zde. Prvotní odkaz na video programu v chodu je tady.

Také mi připadá, že je to celé nějaké zmatečné. Snažím se pochopit co to vlastně dělá a udělat vlastní program od začátku. Nejdříve to chci rozchodit ve Windows (Visual Studio 2022 a CUDA Toolkit), pochopit co to dělá (např. debugem proměnných) a pak to zkusit udělat od začátku nově pro funkci PBKDF2(HMAC−SHA1, passphrase, ssid, 4096, 256), aby to i na obyčejné grafické kartě běželo s maximální rychlostí.

Editováno 9.10.2022 20:35

Nahoru Odpovědět

9.10.2022 20:33

DarkCoder

Člen

DarkCoder:9.10.2022 22:01

Ano, takovéto kopírování je rychlejší, ale hodí se pro velmi malý počet kopírování. V drtivé většině se používá traversování pomocí cyklu.

Kopírování prvků pole musí mít význam, v opačném případě stačí pouze předat ukazatel na pole.

Co se týká inkrementace v prefixové podobě ++i proti postfixové podobě i++, již jsem to zde na fóru psal. Pro vestavěné typy zde není žádný rozdíl. Drobný rozdíl je u tzv. compound typů, tedy typů co vytváří programátor, struktury, objekty, unie, tam je prefixová podoba zápisu nepatrně rychlejší.

Překladače ale za dobu své existence jistě neusnuly na vavřínech a dost pravděpodobně dokáží rozpoznat kód a optimalizovat jej tak, že nebude rozdíl mezi prefixovou a postfixovou formou zápisu.

Nahoru Odpovědět

9.10.2022 22:01

"I ta nejlepší poučka postrádá na významu, není-li patřičně předána." - DarkCoder

Peter Mlich

Člen

Peter Mlich:10.10.2022 8:33

Take tam vidim parkrat cislo * 8, to by slo tez resit shiftovanim, ne? cislo << 3

Nahoru Odpovědět

10.10.2022 8:33

Nejčastěji vyhledáváné

Diskuze: Jak upravit program pro hledání hesla ?