Fórum Root.cz
Hlavní témata => Vývoj => Téma založeno: Kevil 05. 12. 2022, 16:37:39
-
Pomocí instrukcí AVX2 jsem naprogramoval funkci pro šifrování HMAC (https://en.wikipedia.org/wiki/HMAC), SHA-1 (https://en.wikipedia.org/wiki/SHA-1). Funkci SHA1 mi funguje pro vstupní text "Ahoj1234", který je menší než 64 bytů. Po doplnění funkce HMAC mi ale výpočet s testovacími daty nefunguje správně a nemohu přijít na to kde mám chybu.
Testovací data pro HMAC, SHA1 jsou:
key = "Jefe"
message = "what do ya want for nothing?"
viz Test Cases for HMAC-MD5 and HMAC-SHA-1 (https://datatracker.ietf.org/doc/html/rfc2202)
Výsledek by měl být:
effcdf6ae5eb2fa2d27416d5f184df9c259a7c79
spočítá ale:
974020ac606d38554f91a4786d23aec3cf4d77de
viz HMAC generator zde (https://cryptii.com/pipes/hmac)
SHA1 generator tady
(https://emn178.github.io/online-tools/sha1.html)
Vlastní kód (Visual Studio 2022, C++ a AVX2)
32 bitová data ukládá Visual Studio 2022 ve tvaru little endian tj. string "Ahoj1234" je uložen jako "johA" a "4321" (lsb na první místo v paměti).
Specifikace HMAC požaduje uložení počtu bitů šifrované zprávy jako 64bitové big-endian číslo, pracuji s ním ale jako s 32 bitovým little-endian číslem (posledních 32 bitů v druhé AVX2 části 128 bitů a vyzkoušel jsem, že pro test SHA1 "Ahoj1234" to funguje správně.
#include <iostream>
#include <immintrin.h>
#include <chrono>
#include <stdint.h>
#include <stdio.h>
#include <intrin.h> // Byte swap, unsigned long _byteswap_ulong(unsigned long value);
#ifndef __cplusplus
#include <stdalign.h> // C11 defines _Alignas(). This header defines alignas()
#endif
using namespace std;
using namespace std::chrono;
uint32_t result;
__m256i sha1res; // výsledek SHA1
__m256i key;
__m256i indata[4]; //2x 512 bitů pro vstup HMAC
// HMAC block size block size of the SHA1 hash function is 64 bytes, output size pro SHA1 is 20 bytes
// https://en.m.wikipedia.org/wiki/HMAC
// SHA1
// https://en.wikipedia.org/wiki/SHA-1
void p256_hex_u32(__m256i in) {
alignas(32) uint32_t v[8];
_mm256_maskstore_epi32((int*)v, _mm256_setr_epi32(-1, -1, -1, -1, -1, 0, 0, 0), in);
printf("v8_u32: %x %x %x %x %x\n", v[0], v[1], v[2], v[3], v[4]);
}
inline uint32_t rotl32_30(uint32_t value) {
return value << 30 | value >> 2;
}
inline uint32_t rotl32_5(uint32_t value) {
return value << 5 | value >> 27;
}
inline uint32_t rotl32_1(uint32_t value) {
return value << 1 | value >> 31;
}
uint32_t rotl32(uint32_t value, unsigned int count) {
return value << count | value >> (32 - count);
}
void SHA1(__m256i* indata) {
uint32_t pole[80];
__m256i data;
uint32_t h0 = 0x67452301, a = h0;
uint32_t h1 = 0xEFCDAB89, b = h1;
uint32_t h2 = 0x98BADCFE, c = h2;
uint32_t h3 = 0x10325476, d = h3;
uint32_t h4 = 0xC3D2E1F0, e = h4;
for (int k = 0; k < 2; k++) {
data = _mm256_load_si256(indata + k * 2);
_mm256_store_si256((__m256i*)pole, data);
data = _mm256_load_si256(indata + k * 2 + 1);
_mm256_store_si256((__m256i*)pole + 1, data );
uint32_t temp;
for (int i = 0; i < 80; i++) {
if (i > 15) {
pole[i] = rotl32_1((pole[i - 3] ^ pole[i - 8] ^ pole[i - 14] ^ pole[i - 16]));
}
temp = rotl32_5(a) + e + pole[i];
if (i < 20) {
temp += ((b & c) | ((~b) & d)) + 0x5A827999;
}
else if (i < 40) {
temp += (b ^ c ^ d) + 0x6ED9EBA1;
}
else if (i < 60) {
temp += ((b & c) | (b & d) | (c & d)) + 0x8F1BBCDC;
}
else {
temp += (b ^ c ^ d) + 0xCA62C1D6;
}
e = d;
d = c;
c = rotl32_30(b);
b = a;
a = temp;
}
h0 += a;
h1 += b;
h2 += c;
h3 += d;
h4 += e;
}
sha1res = _mm256_setr_epi32(h0, h1, h2, h3, h4, 0, 0, 0);
}
int main() {
__m256i tmp;
auto start = high_resolution_clock::now();
//Pro test funkce SHA1 lze odkomentovat (a komentovat přiřazení pro Jefe..., nutno také změnit parametr pro načítání k, místo k < 2 na k < 1 a dát stopku za sha1res = ...
//sha1res = 8162b79671480fc441e2d54ee85dea77f43b5bc2
//viz odkaz na generátor SHA1 v úvodu
//key = _mm256_setr_epi32(0x41686f6a, 0x31323334, 0x80000000, 0, 0, 0, 0, 0); // "Ahoj1234"
//indata[0] = key;
//indata[1] = _mm256_setr_epi32(0, 0, 0, 0, 0, 0, 0, 0x40);
//key = _mm256_setr_epi32(_byteswap_ulong(0x4A656665), 0, 0, 0, 0, 0, 0, 0); // Jefe
key = _mm256_setr_epi32(0x4A656665, 0, 0, 0, 0, 0, 0, 0); // Jefe
indata[0] = _mm256_xor_si256(key, _mm256_set1_epi8(0x36));
indata[1] = _mm256_set1_epi8(0x36); // 0 XOR 0x36 = 0x36
//tmp = _mm256_setr_epi32(_byteswap_ulong(0x77686174), _byteswap_ulong(0x20646F20), _byteswap_ulong(0x79612077), _byteswap_ulong(0x616E7420), _byteswap_ulong(0x666F7220), _byteswap_ulong(0x6E6F7468), _byteswap_ulong(0x696E673F), 0); // "what do ya want for nothing?" 28 znaků
tmp = _mm256_setr_epi32(0x77686174, 0x20646F20, 0x79612077, 0x616E7420, 0x666F7220, 0x6E6F7468, 0x696E673F, 0); // "what do ya want for nothing?" 28 znaků
indata[2] = _mm256_or_si256(tmp, _mm256_setr_epi32(0, 0, 0, 0, 0, 0, 0, 0x80000000));
indata[3] = _mm256_setr_epi32(0, 0, 0, 0, 0, 0, 0, 736); //počet bitů pri ipad hashování = (64 + 28) * 8
SHA1(indata);
indata[0] = _mm256_xor_si256(key, _mm256_set1_epi8(0x5c));
indata[1] = _mm256_set1_epi8(0x5c); // 0 XOR 0x5c = 0x5c
indata[2] = _mm256_or_si256(sha1res, _mm256_setr_epi32(0, 0, 0, 0, 0, 0x80000000, 0, 0));
indata[3] = _mm256_setr_epi32(0, 0, 0, 0, 0, 0, 0, 672); //počet bitů pro opad hashování = (64 + 20) * 8
SHA1(indata);
auto stop = high_resolution_clock::now();
auto duration = duration_cast<microseconds>(stop - start);
printf("Time = %lli (us DEC)\n", duration.count());
p256_hex_u32(sha1res);
}
-
Postup ktery volim ja je nasledujici:
1) pridam si debug dump vsech prislusnych promennych po kazdem kroku do obou algoritmu
2) vydumpuju
3) hledam rozdil a tim i chybu
Neresim tedy typove uplne stejne problemy, ale kdyz neni chyba na prvni pohled videt, je to reseni. Treba otocena endianita bude videt v nejakem kroku ihned, ale v druhe runde to uz udela uplne jiny vysledek.
-
Myslím, že jsou výpočty v pořádku jen mám problém, že si nejsem zcela jistý, že počítám SHA-1 správně, pokud je délka zprávy delší než 64 znaků a překročí tedy 512 bitů.
Pokud je zpráva např. "Ahoj1234" uložím ji na začátek 512 bitů, za poslední znak přidám 1 dodatečný bit tj 0x80, zbytek doplním "0" bity a na posledních 32 bitů zadám původní délku zprávy v bitech tj. 8 znaků * 8 = 64 (jde o délku zprávy, která musí být ve tvaru 64 bitů, horních 32 bitů ale nikdy nepoužiju, nastavím je na 0). U této zprávy program spočítá výsledek SHA-1 správně.
U zprávy o délce 80 znaků, která přesáhne 512 bitů postupuji v programu takto (a dělám to určitě chybně). Do prvního pole 512 bitů uložím prvních 64 znaků zprávy. Do druhého navazující pole 512 bitů pak zbylých 80-64=16 znaků, za poslední znak přidám 1 bit tj. 0x80, zbylé bity do konce 512 bitů vynuluji a do posledních 32 bitů zadám celkovou délku původní zprávy tj. 80 znaků * 8 = 640.
SHA-1 (https://emn178.github.io/online-tools/sha1.html) pak je
50abf570 6a150990 a08b2c5e a40fa0e5 85554732 (správně)
14a5f567 f881dad3 721e1fa9 a09cde30 6efd7632 (vyjde mi chybně)
Vstupní zpráva je 80 znaků "12345678901234567890123456789012345678901234567890123456789012345678901234567890"
V programu zadám vstupní údaje takto (můj CPU i7-7500U umí AVX2 tj. max 256 bitů):
indata[0] = _mm256_setr_epi32(0x31323334, 0x35363738, 0x39303132, 0x33343536, 0x37383930, 0x31323334, 0x35363738, 0x39303132);
indata[1] = _mm256_setr_epi32(0x33343536, 0x37383930, 0x31323334, 0x35363738, 0x39303132, 0x33343536, 0x37383930, 0x31323334);
indata[2] = _mm256_setr_epi32(0x35363738, 0x39303132, 0x33343536, 0x37383930, 0x80000000, 0, 0, 0);
indata[3] = _mm256_setr_epi32(0, 0, 0, 0, 0, 0, 0, 640);
V popisu SHA-1 Breaking Down: SHA-1 Algorithm (https://infosecwriteups.com/breaking-down-sha-1-algorithm-c152ed353de2) se píše, že by se snad delší zpráva měla dělit na části po 448 bitech. To by ale znamenalo, že délka zprávy by měla být v každém 512 bitovém bloku na konci (posledních 64 bitů)?
Jak mám tedy správně zadat hodnoty do bloků 512 bitů jedna a dva ?
-
Postup ktery volim ja je nasledujici:
1) pridam si debug dump vsech prislusnych promennych po kazdem kroku do obou algoritmu
2) vydumpuju
3) hledam rozdil a tim i chybu
Neresim tedy typove uplne stejne problemy, ale kdyz neni chyba na prvni pohled videt, je to reseni. Treba otocena endianita bude videt v nejakem kroku ihned, ale v druhe runde to uz udela uplne jiny vysledek.
+1
Takto jsem posledne vyresil sifrovani ve fotakovem RAWu, ktere jede v uint32, ale v 32bit PHP je vsechno signed. Samozrejme bitove operace, pripadne nasobeni a maskovani nefungovalo.. ale mit trace z C implementace a PHP implementace, odhalil presne radek kde nastala chyba. (nakonec vznikla 32 i 64bit verze v PHP, ale lisi se od C implementace, aby se simulovali efekty uint32 spravne)
-
Ahoj, nedávno jsem implementoval (mimo jiné) jak SHA-1 tak HMAC v čistém C (C89/C90). Mám to na GitHibu a běží to i ve Visual Studiu. Je to otestované a funguje to zaručeně správně (použil jsem fuzzing a jako oracle Windows crypto API). Takže do mé implementace můžeš hodit svůj input a prokrokovat si to v debuggeru. Nepoužívám ovšem ani AVX ani SSE ani AES-NI. Je to k dispozici na https://github.com/MarekKnapek/mk_crypto/blob/refactor/fuzz/mk_mac_hmac_fuzz.c
-
Tak jsem se na to podíval a určitě tam máš chybu ve výpočtu SHA-1, ve druhém a dalším bloku, kód [1] nahraď kódem [2].
[1]
void SHA1(__m256i* indata) {
uint32_t pole[80];
__m256i data;
uint32_t h0 = 0x67452301, a = h0;
uint32_t h1 = 0xEFCDAB89, b = h1;
uint32_t h2 = 0x98BADCFE, c = h2;
uint32_t h3 = 0x10325476, d = h3;
uint32_t h4 = 0xC3D2E1F0, e = h4;
for (int k = 0; k < 2; k++) {
data = _mm256_load_si256(indata + k * 2);
_mm256_store_si256((__m256i*)pole, data);
data = _mm256_load_si256(indata + k * 2 + 1);
_mm256_store_si256((__m256i*)pole + 1, data );
uint32_t temp;
for (int i = 0; i < 80; i++) {
[2]
void SHA1(__m256i* indata) {
uint32_t pole[80];
__m256i data;
uint32_t h0 = 0x67452301;
uint32_t h1 = 0xEFCDAB89;
uint32_t h2 = 0x98BADCFE;
uint32_t h3 = 0x10325476;
uint32_t h4 = 0xC3D2E1F0;
for (int k = 0; k < 2; k++) {
data = _mm256_load_si256(indata + k * 2);
_mm256_store_si256((__m256i*)pole, data);
data = _mm256_load_si256(indata + k * 2 + 1);
_mm256_store_si256((__m256i*)pole + 1, data );
uint32_t temp;
a = h0;
b = h1;
c = h2;
d = h3;
e = h4;
for (int i = 0; i < 80; i++) {
-
Díky, tvé odpovědi jsem si všiml až teď, když jsem na chybu přišel důkladným laděním ;). Ještě to znovu zkontroluji a doplním funkci PBKDF2.
Mám už skoro hotový program na prolomení WiFi hesla o délce 8 znaků (malá a velká písmena, číslice), který na GPU vygeneruje 200 bilionů (62^8) možných kombinací hesla za 40 ms 8). Potřebuji do něho doplnit jen to šifrování a uvidím, zda se podaří heslo najit v rozumném čase ;).
-
To nějak nechápu. Generuješ na GPU, ale počítáš na CPU? Generování všech kombinací je jednoduché, počítání je to složité, protože u wifi je to PBKDF s 4095 iteracemi a to na CPU hodně štěstí. Ale když máš GPU, proč nepoužiješ hotový hashcat?
Našel jsem že na high-end GPU dává hashcat kolem 1 MH/s (tj. 4G iterací/s což by tak mohlo sedět a na CPU to imho rychlejc nedáš), takže projít 62**8 je pořád 7 let jestli dobře počítám…
-
Kód nejdříve odladím na CPU a pak ho upravím pro GPU.
Hashcat je bohužel pomalý. Na mém notebooku s NVIDIA GeForce GTX 960M by výpočet trval cca 130 let... :-\
Rád mám vše pod vlastní kontrolou. Vtip je v tom, že program poběží pouze na GPU, nepotřebuje vůbec žádné přesuny nějakých dat mezi RAM a GPU. Jak jsem již psal, generování všech 200 bilionů možných hesel a to včetně prázdné testovací smyčky 4096 iterací trvá těch 40 ms ;).
Zkoušel jsem samozřejmě i OpenSSL, funkce ale nelze volat z GPU.
Potřebuji tam jen doplnit ty šifry a porovnat hash, kterou jsem odchytil z WiFi signálu, abych našel správné heslo 8).
-
Kód nejdříve odladím na CPU a pak ho upravím pro GPU.
OpenCL/CUDA má nějaký ekvivalent těch použitých intrinsics?
Jak jsem již psal, generování všech 200 bilionů možných hesel a to včetně prázdné testovací smyčky 4096 iterací trvá těch 40 ms ;).
Generování je irelevantní když používáš vyčerpávající hledání nebo třeba nějaký jednoduchý slovník a pravidla, generování trvá dlouho jenom když děláš nějaké Markovovy řetězce nebo co se teď používá. Prázdnou smyčku ti to vyoptimalizovalo, není reálné abys udělal za 40 ms 62**8*4096 operací, to by bylo 22 exa operací/s.
-
ale v heslu muzou byt i jine znaky takze je to spis 70 na 8.
kdyby jeden vypocet trval milisekundu tak
projit vsecky moznosti trva 18 tisic let.
musis mit aspon nejaky odhad jak heslo vypada.
-
Kód nejdříve odladím na CPU a pak ho upravím pro GPU.
OpenCL/CUDA má nějaký ekvivalent těch použitých intrinsics?
Jak jsem již psal, generování všech 200 bilionů možných hesel a to včetně prázdné testovací smyčky 4096 iterací trvá těch 40 ms ;).
Generování je irelevantní když používáš vyčerpávající hledání nebo třeba nějaký jednoduchý slovník a pravidla, generování trvá dlouho jenom když děláš nějaké Markovovy řetězce nebo co se teď používá. Prázdnou smyčku ti to vyoptimalizovalo, není reálné abys udělal za 40 ms 62**8*4096 operací, to by bylo 22 exa operací/s.
Na ty intrinsics se ještě podívám (https://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__SIMD.html). I kdyby je nešlo v CUDA kódu volat nejde o žádný zásadní problém. Stejně je používám jen k přípravě nastavení hodnot pro hashování. Jádro hashe běží na uint32_t a to v CUDA kódu není žádný problém.
-
ale v heslu muzou byt i jine znaky takze je to spis 70 na 8.
kdyby jeden vypocet trval milisekundu tak
projit vsecky moznosti trva 18 tisic let.
musis mit aspon nejaky odhad jak heslo vypada.
U hesla bráchy znám jeho délku 8 znaků a že se skládá z malých a velkých písmen + číslic ;).
-
Nechci ti do toho kecat, ale opravdu tvrdíš, že optimalizovanýmu hashcatu by to trvalo 130 let, zatímco tvůj na koleně splácanej kód to zvládne za 40 ms, a není ti to ani trochu divný?
-
Nechci ti do toho kecat, ale opravdu tvrdíš, že optimalizovanýmu hashcatu by to trvalo 130 let, zatímco tvůj na koleně splácanej kód to zvládne za 40 ms, a není ti to ani trochu divný?
Samozřejmě uvidím, až tam přidám to hashování. Věřím ale, že program ušitý na míru bez jakýchkoliv přesunů dat mezi RAM a GPU, bez zbytečných testování pro "x" různých variant hashe bude podstatně rychlejší.
Program CUDA, který vygeneruje 200 bilionů všech kombinací hesla (62^8) za 40 ms + 0,5 s start kernelu 8).
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
#include <chrono>
#include <immintrin.h>
using namespace std;
using namespace std::chrono;
// NVIDIA GeForce GTX 960M v NB má 5 Multiprocessors
# define blocks 4
# define threads 961
# define characters 6
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
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;
}
}
}
}
// 960, zkušební výpis posledních 3 vygenerovaných hesel o pár sekund zpozdí výpočet
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 (PBKDF2 a HMAC SHA1), musím doplnit
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 iter = 4096;
//while (iter) {
// iter--;
//}
}
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;
}
-
i kdyby jedno heslo z 62^8 trvalo vypocitat miliontinu sekundy, tak je to vypocet na 7 let behu.
-
i kdyby jedno heslo z 62^8 trvalo vypocitat miliontinu sekundy, tak je to vypocet na 7 let behu.
Generování jednoho hesla trvá 1,83 E-16 s 8). Nemyslím si, že hashování bude více než miliardkrát pomalejší E-16 vs. E-6;)
-
Generování jednoho hesla trvá 1,83 E-16 s 8). Nemyslím si, že hashování bude více než miliardkrát pomalejší E-16 vs. E-6;)
checht. V tom výpočtu bude asi nějaká chyba. Při paralelizaci 10000 by to znamenalo cca 546E9 hesel za sekundu.
Vzhledem k tomu, že GTX 960M jede na cca 1GHz (1E9) to je cca 546 hesel na takt na jednu jednotku, což je evidentní nesmysl.
Jestli on nebude zakopaný pes v tom, že v ukázkovém programu je 6 znaků a ne 8 ... To pak vychází drobet jiná a dokonce i uvěřitelná čísla.
-
Generování jednoho hesla trvá 1,83 E-16 s 8). Nemyslím si, že hashování bude více než miliardkrát pomalejší E-16 vs. E-6;)
checht. V tom výpočtu bude asi nějaká chyba. Při paralelizaci 10000 by to znamenalo cca 546E9 hesel za sekundu.
Vzhledem k tomu, že GTX 960M jede na cca 1GHz (1E9) to je cca 546 hesel na takt na jednu jednotku, což je evidentní nesmysl.
Jestli on nebude zakopaný pes v tom, že v ukázkovém programu je 6 znaků a ne 8 ... To pak vychází drobet jiná a dokonce i uvěřitelná čísla.
To se omlouvám, přehlédl jsem těch 6 znaků při zkoušení. Pro 8 stačí jen změnit definice na začátku programu na:
# define blocks 4
# define threads 992
# define characters 8
Okomentoval jsem jen test na výpis posledních 3 hesel, přeložil a spustil bez ladění Ctrl + F5 s časem:
Time = 7e072 (HEX)
jde o mikrosekundy tj. 516210 DEC = 516,210 ms (včetně spuštění CUDA jádra cca 500 ms)
-
Je zajímavé, že pokud odkomentuji výpis posledních tří generovaných hesel, čas výpočtu se zásadně prodlouží. Pro 6 znaků na 21,5 s viz:
Thread[960] 9 9 9 9 9 7
Thread[960] 9 9 9 9 9 8
Thread[960] 9 9 9 9 9 9
Time = 148da70 (HEX)
Pro 8 znaků pak myslím až na 35 minut. Musím se na to podívat v CUDA debugeru.
-
Je zajímavé, že pokud odkomentuji výpis posledních tří generovaných hesel, čas výpočtu se zásadně prodlouží.
Ale o tom zde už dobrou čtvrthodinu hovoříme, Mlho!
Pokud s těmi vygenerovanými hesly nic neděláš, tak ten program nemá žádný side effect a překladač všechny výpočty vyhodí pryč, protože proč by je tam nechával, když se výsledek na nic nepoužije...
-
Je zajímavé, že pokud odkomentuji výpis posledních tří generovaných hesel, čas výpočtu se zásadně prodlouží.
Ale o tom zde už dobrou čtvrthodinu hovoříme, Mlho!
Pokud s těmi vygenerovanými hesly nic neděláš, tak ten program nemá žádný side effect a překladač všechny výpočty vyhodí pryč, protože proč by je tam nechával, když se výsledek na nic nepoužije...
az budu tesne pred smrti, tak se mozna dozvim vysledek :-D
-
Je zajímavé, že pokud odkomentuji výpis posledních tří generovaných hesel, čas výpočtu se zásadně prodlouží.
Ale o tom zde už dobrou čtvrthodinu hovoříme, Mlho!
Pokud s těmi vygenerovanými hesly nic neděláš, tak ten program nemá žádný side effect a překladač všechny výpočty vyhodí pryč, protože proč by je tam nechával, když se výsledek na nic nepoužije...
To je nesmysl. Překladač nevyhodí funkční kód 8).
-
Samozřejmě uvidím, až tam přidám to hashování. Věřím ale, že program ušitý na míru bez jakýchkoliv přesunů dat mezi RAM a GPU, bez zbytečných testování pro "x" různých variant hashe bude podstatně rychlejší.
Trochu jo, ale zejména u pomalých hashovacích funkcí (což 4095krát iterovaná HMAC-SHA je) je ta režie zanedbatelná. Kdyby tomu bylo naopak, tak by rychlé hashovací funkce (třeba čistá MD5) nebyly o moc rychlejší - ale ony jsou, a to skoro tak jak by teoreticky být měly (e.g. zpomaluje se to přesně podle počtu iterací).
Pokud s těmi vygenerovanými hesly nic neděláš, tak ten program nemá žádný side effect a překladač všechny výpočty vyhodí pryč, protože proč by je tam nechával, když se výsledek na nic nepoužije...
+1
Mimochodem tohle je taky divné:
counters[i] = (n * threadIdx.x / (uint64_t)pow(62, characters - 1 - i) % 62);crackujeme hesla s floatovými funkcemi ::). Když potřebuješ rozhodit prohledávaný prostor mezi kernely (nepotřebuješ, dokud si neobenchmarkuješ, že režie na kernel je dost velká na to abys to musel řešit) tak to udělej tak, že kernelu přiřadíš třeba prvních 5 znaků a zbylé 3 znaky se proiterují uvnitř kernelu.
-
Ahoj Kevile, matematika ti moc nejde, co? Říkáš, že zvládneš spočítat 62^8 netriviálních operací za 40 ms? To máme 62^8/40 netriviálních operací za ms, neboli 62^8/40/1000 netriviálních operací za μs, neboli 62^8/40/1000/1000 netriviálních operací za ns. Furt je to zhruba 5.5 miliónu netriviálních operací každou nanosekundu. Při rychlosti 5GHz, to je stále milión operací za takt. To asi ne, co?
-
Mimochodem tohle je taky divné:
counters[i] = (n * threadIdx.x / (uint64_t)pow(62, characters - 1 - i) % 62);crackujeme hesla s floatovými funkcemi ::). Když potřebuješ rozhodit prohledávaný prostor mezi kernely (nepotřebuješ, dokud si neobenchmarkuješ, že režie na kernel je dost velká na to abys to musel řešit) tak to udělej tak, že kernelu přiřadíš třeba prvních 5 znaků a zbylé 3 znaky se proiterují uvnitř kernelu.
Rád si nechám poradit, jak skupiny generovaných hesel rovnoměrně rozložit na jednotlivá vlákna bez floatu ;). Funkčnost svého kódu jsem si nejdříve otestoval v Excelu.
Programuji mj. také 8 bitové MCU, např. ATtiny202 (https://www.microchip.com/en-us/product/ATTINY202) a tam je nesmysl, že by překladač vyhodil kód, který údajně nic nedělá 8).
S odladěním na GPU mám problém. VS2022 a nejnovější NVIDIA Nsight Visual Studio Edition Debugger (režim Next-Gen) nepodporuje starší grafické karty architektury Maxwell (CUDA verze 5.0) což je má karta NVIDIA GeForce GTX 960M, kterou mám v netebooku. Ladění by mělo jít s Nsight v Legacy režimu ve VS2017. To jsem zkoušel, ale debugger se k aplikaci připojí jen na zlomek sekundy a pak se hned odpojí s chybovou hláškou, že nesouhlasí security nastavení v parametrech laděného programu a Nsight Monitoru. Defaultně je ale u obou nastaveno security False. I když jsem oboje změnil na True nic se nezměnilo. Možná je zádrhelem to, že v seznamu podporovaných grafik je u Nsight Legacy uvedena karta GeForce GTX 960 a ne GeForce GTX 960M (má mobilní). Hlavu si z toho ale moc nedělám a nový notebook s výkonější GPU kartou nebudu kvůli tomu kupovat. Zkusím do programu doplnit to hashování a odzkouším rychlost nalezení hesla měřením času.
Pokud má někdo na svém PC novější GPU může zkusit jak se můj program na generování hesel chová v debuggeru ;).
-
To je nesmysl. Překladač nevyhodí funkční kód 8).
stane se... ...jednou jsem taky dal data jako konstanty do programu,. rýpal se v unrollování loopů a pragmách a najednou mi z icc vyšla tak super binárka, že jsem už skoro pálil rachejtle..
No ale bylo to tím, že "výpočet" se provedl už při překladu.. bylo to vidět ve vygenerovaných instrukcích...
co jsem si z toho odnesl.. vstupy brát z externího souboru..
-
[
Takto jsem posledne vyresil sifrovani ve fotakovem RAWu,
co proboha šifruje foťák? (když nepočítám wpa nebo připojení k internetu)? Nebo je to nějaký foťák s funkcí content authenticity? (nějaká featura za kterou stojí adobe, že fotky budou podepsané jak vylezou z foťáku a má to dokazovat, že s nimi nrbylo manipulováno)