Názor k článku Použití instrukcí SSE a AVX pro zrychlení bitových operací od Kevil - Pěkný článek. Pro daný účel bych zkusil výpočet...

Pěkný článek. Pro daný účel bych zkusil výpočet na GPU ;-). Rád programuji v C++, x64 asm a AVX2.

Aktuálně se pokouším naprogramovat prolomení WiFi hesla o délce 8 znaků (malá/velká písmena a číslice) hrubou silou, což je 62^8 tj. cca 200 bilionů kombinací. Zkušebně v C++ CUDA programu (Visual Studio 2022) všechna hesla vygeneruji za cca 40 ms (notebook má NVIDIA GeForce GTX 960M). Umím už v C++ spočítat SHA-1 a chystám se doplnit HMAC a PBKDF2, abych našel správné heslo v rozumném čase.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <iostream>
#include <chrono>

using namespace std;
using namespace std::chrono;

# define blocks 4
# define threads 992
# define characters 8

cudaError_t cudaStatus;

// "abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789"
__constant__ char1 charset[] = { 0x61, 0x62, 0x63, 0x64, 0x65, 0x66, 0x67, 0x68, 0x69, 0x6a, 0x6b, 0x6c, 0x6d, 0x6e, 0x6f, 0x70, 0x71, 0x72, 0x73, 0x74, 0x75, 0x76, 0x77, 0x78, 0x79, 0x7a, 0x41, 0x42, 0x43, 0x44, 0x45, 0x46, 0x47, 0x48, 0x49, 0x4a, 0x4b, 0x4c, 0x4d, 0x4e, 0x4f, 0x50, 0x51, 0x52, 0x53, 0x54, 0x55, 0x56, 0x57, 0x58, 0x59, 0x5a, 0x30, 0x31, 0x32, 0x33, 0x34, 0x35, 0x36, 0x37, 0x38, 0x39 };

__global__ void blackcat(void) {

        char1 password[characters];

        uint8_t counters[characters];
        uint64_t n = (pow(62, characters) / threads);   // Number of search cycles per thread

        // Nastavení počátečních hodnot hesla pro každé vlákno, odkud se mají začít generovat
        for (int i = characters - 1; i >= 0; i--) {
                counters[i] = (n * threadIdx.x / (uint64_t)pow(62, characters - 1 - i) % 62);
        }

        while (n > 0) {

                bool flag = false;
                for (int i = characters - 1; i >= 0; i--) {
                        password[i] = charset[counters[i]];
                        if (i == characters - 1) {
                                counters[i]++;
                                if (counters[i] > 61) {
                                        counters[i] = (uint8_t)0;
                                        flag = true;
                                }
                        }
                        else {
                                if (flag) {
                                        counters[i]++;
                                        if (counters[i] > 61) {
                                                counters[i] = (uint8_t)0;
                                        }
                                        else {
                                                flag = false;
                                        }
                                }
                        }
                }
                // Po odkomentování vypíše poslední 3 generované hesla
                //if (threadIdx.x == threads - 1 && blockIdx.x == blocks - 1 && n < 4) {
                //      printf("Thread[%d]",threadIdx.x);
                //      for (int i = 0; i < characters; i++) {
                //              printf(" %c", password[i]);
                //      }
                //      printf("\n");
                //}

                /* Test zda jsme našli password,
                pokud ano vypíšeme password, ukončíme všechna vlákna a předčasně se vrátíme z funkce,
                možná bude dobré občas vypsat čas běhu, abychom věděli, že program stále běží */
                n--;
        }
}

int main() {

        auto start = high_resolution_clock::now();

        cudaSetDevice(0);
        cudaStatus = cudaGetLastError();
                if (cudaStatus != cudaSuccess) {
                fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
         }
        blackcat << <blocks, threads >> > ();
        cudaStatus = cudaGetLastError();
                if (cudaStatus != cudaSuccess) {
                fprintf(stderr, "Kernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        }
        cudaDeviceSynchronize();
        cudaStatus = cudaGetLastError();
                if (cudaStatus != cudaSuccess) {
                fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
         }

        auto stop = high_resolution_clock::now();
        auto duration = duration_cast<microseconds>(stop - start);
        printf("\nTime  = %llx (HEX)\n", duration.count());

        return 0;
}