HMAC, SHA1 (C++/AVX2 VS2022)

Kevil

HMAC, SHA1 (C++/AVX2 VS2022)
« kdy: 05. 12. 2022, 16:37:39 »
Pomocí instrukcí AVX2 jsem naprogramoval funkci pro šifrování HMAC, 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

Výsledek by měl být:
effcdf6ae5eb2fa2d274­16d5f184df9c259a7c79

spočítá ale:
974020ac606d3­8554f91a4786d23a­ec3cf4d77de

viz HMAC generator zde
SHA1 generator tady

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ě.

Kód: [Vybrat]
#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);

 }


mhi

  • ****
  • 451
    • Zobrazit profil
Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #1 kdy: 06. 12. 2022, 11:39:43 »
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.

Kevil

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #2 kdy: 07. 12. 2022, 01:15:31 »
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 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ů):
Kód: [Vybrat]
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 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 ?




RDa

  • *****
  • 2 036
    • Zobrazit profil
    • E-mail
Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #3 kdy: 07. 12. 2022, 10:43:59 »
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)

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #4 kdy: 08. 12. 2022, 22:12:35 »
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


Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #5 kdy: 10. 12. 2022, 00:12:10 »
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]
Kód: [Vybrat]
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]
Kód: [Vybrat]
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++) {

Kevil

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #6 kdy: 31. 12. 2022, 02:01:43 »
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 ;).

_Jenda

  • *****
  • 1 266
    • Zobrazit profil
    • https://jenda.hrach.eu/
    • E-mail
Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #7 kdy: 31. 12. 2022, 04:19:49 »
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…
« Poslední změna: 31. 12. 2022, 04:21:56 od _Jenda »

Kevil

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #8 kdy: 31. 12. 2022, 09:40:03 »
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).

_Jenda

  • *****
  • 1 266
    • Zobrazit profil
    • https://jenda.hrach.eu/
    • E-mail
Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #9 kdy: 31. 12. 2022, 21:58:20 »
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.

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #10 kdy: 01. 01. 2023, 08:11:42 »
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.

Kevil

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #11 kdy: 01. 01. 2023, 11:58:48 »
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. 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.

Kevil

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #12 kdy: 01. 01. 2023, 12:00:44 »
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  ;).

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #13 kdy: 01. 01. 2023, 14:21:17 »
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ý?

Kevil

Re:HMAC, SHA1 (C++/AVX2 VS2022)
« Odpověď #14 kdy: 01. 01. 2023, 16:15:49 »
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).

Kód: [Vybrat]
#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;
}
« Poslední změna: 01. 01. 2023, 16:17:45 od Kevil »