ATTACK TYPE \n\n"
+ "Options:\n\n"
+ " -h"
+ "\t\tShow this help\n"
+ " -f"
+ "\t\tPath to your input hash file (HashExtractor output)\n"
+ " -d"
+ "\t\tPath to dictionary file\n", name);
+}
+
+int main (int argc, char **argv)
+{
+ std::chrono::steady_clock::time_point time_start;
+ std::chrono::steady_clock::time_point time_end;
+ double time_total = 0.0;
+ double time_total_ = 0.0;
+
+ TIMER_START()
+
+ int opt = 0;
+ int pass_batch_size = 60000;
+ char * input_hash = NULL;
+ char * input_dictionary = NULL;
+ unsigned char *nonce;
+ unsigned char *vmk;
+ unsigned char *mac;
+ uint32_t * d_w_words_uint32;
+
+ printf("\n---------> BitCracker: BitLocker password cracking tool <---------\n");
+
+ if (argc < 2) {
+ printf("Missing argument!\n");
+ usage(argv[0]);
+ exit(EXIT_FAILURE);
+ }
+
+ while (1) {
+ opt = getopt(argc, argv, "f:d:b:h");
+ if (opt == -1)
+ break;
+
+ switch (opt) {
+ case 'f':
+ if(strlen(optarg) >= INPUT_SIZE)
+ {
+ fprintf(stderr, "ERROR: Inut hash file path is bigger than %d\n", INPUT_SIZE);
+ exit(EXIT_FAILURE);
+ }
+ input_hash=(char *)Calloc(INPUT_SIZE, sizeof(char));
+ strncpy(input_hash, optarg, strlen(optarg)+1);
+ break;
+
+ case 'd':
+ if(strlen(optarg) >= INPUT_SIZE)
+ {
+ fprintf(stderr, "ERROR: Dictionary file path is bigger than %d\n", INPUT_SIZE);
+ exit(EXIT_FAILURE);
+ }
+ input_dictionary=(char *)Calloc(INPUT_SIZE, sizeof(char));
+ strncpy(input_dictionary,optarg, strlen(optarg)+1);
+ break;
+
+ case 'b':
+ pass_batch_size = atoi(optarg);
+ break;
+
+ case 'h':
+ usage(argv[0]);
+ exit(EXIT_FAILURE);
+ break;
+
+ default:
+ exit(EXIT_FAILURE);
+ }
+ }
+
+ if (optind < argc) {
+ printf ("non-option ARGV-elements: ");
+ while (optind < argc)
+ printf ("%s ", argv[optind++]);
+ putchar ('\n');
+ exit(EXIT_FAILURE);
+ }
+
+ if (input_dictionary == NULL){
+ printf("Missing dictionary file!\n");
+ usage(argv[0]);
+ exit(EXIT_FAILURE);
+ }
+
+ if (input_hash == NULL){
+ printf("Missing input hash file!\n");
+ usage(argv[0]);
+ exit(EXIT_FAILURE);
+ }
+
+ // max_num_pswd_per_read = gridBlocks * THREADS_PER_BLOCK * 4;
+ // max_num_pswd_per_read = 11520;//11520 passwords can be processed in one pass in PVC-B3
+ // max_num_pswd_per_read = 50000;//26624/2;//26624 passwords can be processed in one pass in PVC-B4
+ max_num_pswd_per_read = pass_batch_size;
+ // max_num_pswd_per_read = 27648; //27648 passwords can be processed in one pass in A100
+ // if(getGPUStats())
+ // {
+ // fprintf(stderr, "Device error... exit!\n");
+ // // goto cleanup;
+ // }
+
+ printf("\n\n==================================\n");
+ printf("Retrieving Info\n==================================\n\n");
+ if(parse_data(input_hash, &salt, &nonce, &vmk, &mac) == BIT_FAILURE)
+ {
+ fprintf(stderr, "Input hash format error... exit!\n");
+ // goto cleanup;
+ }
+
+ if(mac == NULL)
+ {
+ fprintf(stderr, "NULL MAC string error... exit!\n");
+ // goto cleanup;
+ }
+
+ double duration = 0.0;
+
+#ifdef DEBUG_TIME
+ auto time11 = std::chrono::steady_clock::now();
+#endif
+
+ // create sycl queue
+ sycl::queue qbc;
+
+#ifdef DEBUG_TIME
+ auto time12 = std::chrono::steady_clock::now();
+ double duration1 = std::chrono::duration(time12 - time11).count();
+ duration += duration1;
+ std::cout << "init: " << duration1 << " us\n\n";
+
+ auto time21 = std::chrono::steady_clock::now();
+#endif
+
+ // allocate memory
+ d_w_words_uint32 = (uint32_t *)sycl::malloc_device(NUM_HASH_BLOCKS * HASH_BLOCK_NUM_UINT32 * sizeof(uint32_t), qbc);
+
+#ifdef DEBUG_TIME
+ auto time22 = std::chrono::steady_clock::now();
+ double duration2 = std::chrono::duration(time22 - time21).count();
+ duration += duration2;
+ std::cout << "main() - alloc : duration2: " << duration2 << " us\n\n";
+#endif
+
+ if(evaluate_w_block(salt, d_w_words_uint32, duration, qbc) == BIT_FAILURE)
+ {
+ fprintf(stderr, "Words error... exit!\n");
+ goto cleanup;
+ }
+
+ std::cout << "================================================\n";
+ std::cout << " Attack\n";
+ std::cout << "================================================\n";
+
+ time_total_ = attack(input_dictionary, d_w_words_uint32, vmk, nonce, mac, pass_batch_size, duration, qbc);
+
+cleanup:
+ free(input_hash);
+ free(input_dictionary);
+
+ sycl::free(d_w_words_uint32, qbc);
+
+ // std::cout << "Total time for whole calculation: " << duration / 1e6 << " s\n\n";
+ TIMER_END()
+ TIMER_PRINT("bitcracker - total time for whole calculation")
+
+ return 0;
+}
diff --git a/SYCL/src/sha256.h b/SYCL/src/sha256.h
new file mode 100644
index 0000000..3e783fe
--- /dev/null
+++ b/SYCL/src/sha256.h
@@ -0,0 +1,363 @@
+/* Modifications Copyright (C) 2023 Intel Corporation
+ *
+ * This program is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2, as published
+ * by the Free Software Foundation.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, see .
+ *
+ *
+ * SPDX-License-Identifier: GPL-2.0-only
+ */
+
+/*
+ * BitCracker: BitLocker password cracking tool, CUDA version.
+ * Copyright (C) 2013-2017 Elena Ago
+ * Massimo Bernaschi
+ *
+ * This file is part of the BitCracker project: https://github.com/e-ago/bitcracker
+ *
+ * BitCracker is free software: you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation, either version 2 of the License, or
+ * (at your option) any later version.
+ *
+ * BitCracker is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with BitCracker. If not, see .
+ */
+
+#include
+
+#define ROR07(x) (((x) << 25) | ((x) >> 7))
+#define ROR18(x) (((x) << 14) | ((x) >> 18))
+
+#define ROR17(x) (((x) << 15) | ((x) >> 17))
+#define ROR19(x) (((x) << 13) | ((x) >> 19))
+
+// #define SWAP_UINT32(x) (((x) >> 24) | (((x) & 0x00FF0000) >> 8) | (((x) & 0x0000FF00) << 8) | ((x) << 24))
+#define __dpcpp_inline__ __inline__ __attribute__((always_inline))
+
+static __dpcpp_inline__ uint32_t LOP3LUT_XOR(uint32_t a, uint32_t b, uint32_t c) {
+ return a^b^c;
+}
+
+static __dpcpp_inline__ uint32_t LOP3LUT_XORAND(uint32_t g, uint32_t e, uint32_t f) {
+ return (g ^ (e & (f ^ g)));
+}
+
+static __dpcpp_inline__ uint32_t LOP3LUT_ANDOR(uint32_t a, uint32_t b, uint32_t c) {
+ return ((a & (b | c)) | (b & c));
+}
+
+#define SCHEDULE00() \
+ schedule00 = schedule16 + schedule25 \
+ + LOP3LUT_XOR(ROR07(schedule17) , ROR18(schedule17) , (schedule17 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule30) , ROR19(schedule30) , (schedule30 >> 10));
+
+#define SCHEDULE01() \
+ schedule01 = schedule17 + schedule26 \
+ + LOP3LUT_XOR(ROR07(schedule18) , ROR18(schedule18) , (schedule18 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule31) , ROR19(schedule31) , (schedule31 >> 10));
+
+#define SCHEDULE02() \
+ schedule02 = schedule18 + schedule27 \
+ + LOP3LUT_XOR(ROR07(schedule19) , ROR18(schedule19) , (schedule19 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule00) , ROR19(schedule00) , (schedule00 >> 10));
+
+#define SCHEDULE03() \
+ schedule03 = schedule19 + schedule28 \
+ + LOP3LUT_XOR(ROR07(schedule20) , ROR18(schedule20) , (schedule20 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule01) , ROR19(schedule01) , (schedule01 >> 10));
+
+#define SCHEDULE04() \
+ schedule04 = schedule20 + schedule29 \
+ + LOP3LUT_XOR(ROR07(schedule21) , ROR18(schedule21) , (schedule21 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule02) , ROR19(schedule02) , (schedule02 >> 10));
+
+#define SCHEDULE05() \
+ schedule05 = schedule21 + schedule30 \
+ + LOP3LUT_XOR(ROR07(schedule22) , ROR18(schedule22) , (schedule22 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule03) , ROR19(schedule03) , (schedule03 >> 10));
+
+#define SCHEDULE06() \
+ schedule06 = schedule22 + schedule31 \
+ + LOP3LUT_XOR(ROR07(schedule23) , ROR18(schedule23) , (schedule23 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule04) , ROR19(schedule04) , (schedule04 >> 10));
+
+#define SCHEDULE07() \
+ schedule07 = schedule23 + schedule00 \
+ + LOP3LUT_XOR(ROR07(schedule24) , ROR18(schedule24) , (schedule24 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule05) , ROR19(schedule05) , (schedule05 >> 10));
+
+#define SCHEDULE08() \
+ schedule08 = schedule24 + schedule01 \
+ + LOP3LUT_XOR(ROR07(schedule25) , ROR18(schedule25) , (schedule25 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule06) , ROR19(schedule06) , (schedule06 >> 10));
+
+#define SCHEDULE09() \
+ schedule09 = schedule25 + schedule02 \
+ + LOP3LUT_XOR(ROR07(schedule26) , ROR18(schedule26) , (schedule26 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule07) , ROR19(schedule07) , (schedule07 >> 10));
+
+#define SCHEDULE10() \
+ schedule10 = schedule26 + schedule03 \
+ + LOP3LUT_XOR(ROR07(schedule27) , ROR18(schedule27) , (schedule27 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule08) , ROR19(schedule08) , (schedule08 >> 10));
+
+#define SCHEDULE11() \
+ schedule11 = schedule27 + schedule04 \
+ + LOP3LUT_XOR(ROR07(schedule28) , ROR18(schedule28) , (schedule28 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule09) , ROR19(schedule09) , (schedule09 >> 10));
+
+#define SCHEDULE12() \
+ schedule12 = schedule28 + schedule05 \
+ + LOP3LUT_XOR(ROR07(schedule29) , ROR18(schedule29) , (schedule29 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule10) , ROR19(schedule10) , (schedule10 >> 10));
+
+#define SCHEDULE13() \
+ schedule13 = schedule29 + schedule06 \
+ + LOP3LUT_XOR(ROR07(schedule30) , ROR18(schedule30) , (schedule30 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule11) , ROR19(schedule11) , (schedule11 >> 10));
+
+#define SCHEDULE14() \
+ schedule14 = schedule30 + schedule07 \
+ + LOP3LUT_XOR(ROR07(schedule31) , ROR18(schedule31) , (schedule31 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule12) , ROR19(schedule12) , (schedule12 >> 10));
+
+#define SCHEDULE15() \
+ schedule15 = schedule31 + schedule08 \
+ + LOP3LUT_XOR(ROR07(schedule00) , ROR18(schedule00) , (schedule00 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule13) , ROR19(schedule13) , (schedule13 >> 10));
+
+#define SCHEDULE16() \
+ schedule16 = schedule00 + schedule09 \
+ + LOP3LUT_XOR( ROR07(schedule01), ROR18(schedule01), (schedule01 >> 3)) \
+ + LOP3LUT_XOR( ROR17(schedule14), ROR19(schedule14), (schedule14 >> 10));
+
+#define SCHEDULE17() \
+ schedule17 = schedule01 + schedule10 \
+ + LOP3LUT_XOR(ROR07(schedule02) , ROR18(schedule02) , (schedule02 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule15) , ROR19(schedule15) , (schedule15 >> 10));
+
+#define SCHEDULE18() \
+ schedule18 = schedule02 + schedule11 \
+ + LOP3LUT_XOR(ROR07(schedule03) ,ROR18(schedule03) ,(schedule03 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule16), ROR19(schedule16), (schedule16 >> 10));
+#define SCHEDULE19() \
+ schedule19 = schedule03 + schedule12 \
+ + LOP3LUT_XOR(ROR07(schedule04) , ROR18(schedule04) , (schedule04 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule17) , ROR19(schedule17) , (schedule17 >> 10));
+
+#define SCHEDULE20() \
+ schedule20 = schedule04 + schedule13 \
+ + LOP3LUT_XOR(ROR07(schedule05) , ROR18(schedule05) , (schedule05 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule18) , ROR19(schedule18) , (schedule18 >> 10));
+
+#define SCHEDULE21() \
+ schedule21 = schedule05 + schedule14 \
+ + LOP3LUT_XOR(ROR07(schedule06) , ROR18(schedule06) , (schedule06 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule19) , ROR19(schedule19) , (schedule19 >> 10));
+
+#define SCHEDULE22() \
+ schedule22 = schedule06 + schedule15 \
+ + LOP3LUT_XOR(ROR07(schedule07) , ROR18(schedule07) , (schedule07 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule20) , ROR19(schedule20) , (schedule20 >> 10));
+
+#define SCHEDULE23() \
+ schedule23 = schedule07 + schedule16 \
+ + LOP3LUT_XOR(ROR07(schedule08) , ROR18(schedule08) , (schedule08 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule21) , ROR19(schedule21) , (schedule21 >> 10));
+
+#define SCHEDULE24() \
+ schedule24 = schedule08 + schedule17 \
+ + LOP3LUT_XOR(ROR07(schedule09) , ROR18(schedule09) , (schedule09 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule22) , ROR19(schedule22) , (schedule22 >> 10));
+
+#define SCHEDULE25() \
+ schedule25 = schedule09 + schedule18 \
+ + LOP3LUT_XOR(ROR07(schedule10) , ROR18(schedule10) , (schedule10 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule23) , ROR19(schedule23) , (schedule23 >> 10));
+
+#define SCHEDULE26() \
+ schedule26 = schedule10 + schedule19 \
+ + LOP3LUT_XOR(ROR07(schedule11) , ROR18(schedule11) , (schedule11 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule24) , ROR19(schedule24) , (schedule24 >> 10));
+
+#define SCHEDULE27() \
+ schedule27 = schedule11 + schedule20 \
+ + LOP3LUT_XOR(ROR07(schedule12) , ROR18(schedule12) , (schedule12 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule25) , ROR19(schedule25) , (schedule25 >> 10));
+
+#define SCHEDULE28() \
+ schedule28 = schedule12 + schedule21 \
+ + LOP3LUT_XOR(ROR07(schedule13) , ROR18(schedule13) , (schedule13 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule26) , ROR19(schedule26) , (schedule26 >> 10));
+
+#define SCHEDULE29() \
+ schedule29 = schedule13 + schedule22 \
+ + LOP3LUT_XOR(ROR07(schedule14) , ROR18(schedule14) , (schedule14 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule27) , ROR19(schedule27) , (schedule27 >> 10));
+
+#define SCHEDULE30() \
+ schedule30 = schedule14 + schedule23 \
+ + LOP3LUT_XOR(ROR07(schedule15) , ROR18(schedule15) , (schedule15 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule28) , ROR19(schedule28) , (schedule28 >> 10));
+
+#define SCHEDULE31() \
+ schedule31 = schedule15 + schedule24 \
+ + LOP3LUT_XOR(ROR07(schedule16) , ROR18(schedule16) , (schedule16 >> 3)) \
+ + LOP3LUT_XOR(ROR17(schedule29) , ROR19(schedule29) , (schedule29 >> 10));
+
+#define ROR06(x) (((x) << 26) | ((x) >> 6))
+#define ROR11(x) (((x) << 21) | ((x) >> 11))
+#define ROR25(x) (((x) << 7) | ((x) >> 25))
+
+#define ROR02(x) (((x) << 30) | ((x) >> 2))
+#define ROR13(x) (((x) << 19) | ((x) >> 13))
+#define ROR22(x) (((x) << 10) | ((x) >> 22))
+
+#define ROUND(a, b, c, d, e, f, g, h, W, k) \
+ h += LOP3LUT_XOR(ROR06(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + W; \
+ d += h; \
+ h += LOP3LUT_XOR(ROR02(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
+
+#define ROUND_SECOND_BLOCK(a, b, c, d, e, f, g, h, i, k, indexW) \
+ h += LOP3LUT_XOR(ROR06(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + w_words_uint32[(indexW + i)]; \
+ d += h; \
+ h += LOP3LUT_XOR(ROR02(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
+
+#define ROUND_SECOND_BLOCK_CONST(a, b, c, d, e, f, g, h, i, k, w) \
+ h += LOP3LUT_XOR(ROR06(e), ROR11(e), ROR25(e)) + LOP3LUT_XORAND(g,e,f) + k + w; \
+ d += h; \
+ h += LOP3LUT_XOR(ROR02(a), ROR13(a), ROR22(a)) + LOP3LUT_ANDOR(a,b,c);
+
+#define ALL_SCHEDULE_LAST16() \
+ SCHEDULE16() \
+ SCHEDULE17() \
+ SCHEDULE18() \
+ SCHEDULE19() \
+ SCHEDULE20() \
+ SCHEDULE21() \
+ SCHEDULE22() \
+ SCHEDULE23() \
+ SCHEDULE24() \
+ SCHEDULE25() \
+ SCHEDULE26() \
+ SCHEDULE27() \
+ SCHEDULE28() \
+ SCHEDULE29() \
+ SCHEDULE30() \
+ SCHEDULE31()
+
+#define ALL_SCHEDULE32() \
+ SCHEDULE00() \
+ SCHEDULE01() \
+ SCHEDULE02() \
+ SCHEDULE03() \
+ SCHEDULE04() \
+ SCHEDULE05() \
+ SCHEDULE06() \
+ SCHEDULE07() \
+ SCHEDULE08() \
+ SCHEDULE09() \
+ SCHEDULE10() \
+ SCHEDULE11() \
+ SCHEDULE12() \
+ SCHEDULE13() \
+ SCHEDULE14() \
+ SCHEDULE15() \
+ SCHEDULE16() \
+ SCHEDULE17() \
+ SCHEDULE18() \
+ SCHEDULE19() \
+ SCHEDULE20() \
+ SCHEDULE21() \
+ SCHEDULE22() \
+ SCHEDULE23() \
+ SCHEDULE24() \
+ SCHEDULE25() \
+ SCHEDULE26() \
+ SCHEDULE27() \
+ SCHEDULE28() \
+ SCHEDULE29() \
+ SCHEDULE30() \
+ SCHEDULE31()
+
+#define ALL_ROUND_B1_1() \
+ ROUND(a, b, c, d, e, f, g, h, schedule00, 0x428A2F98) \
+ ROUND(h, a, b, c, d, e, f, g, schedule01, 0x71374491) \
+ ROUND(g, h, a, b, c, d, e, f, schedule02, 0xB5C0FBCF) \
+ ROUND(f, g, h, a, b, c, d, e, schedule03, 0xE9B5DBA5) \
+ ROUND(e, f, g, h, a, b, c, d, schedule04, 0x3956C25B) \
+ ROUND(d, e, f, g, h, a, b, c, schedule05, 0x59F111F1) \
+ ROUND(c, d, e, f, g, h, a, b, schedule06, 0x923F82A4) \
+ ROUND(b, c, d, e, f, g, h, a, schedule07, 0xAB1C5ED5) \
+ ROUND(a, b, c, d, e, f, g, h, schedule08, 0xD807AA98) \
+ ROUND(h, a, b, c, d, e, f, g, schedule09, 0x12835B01) \
+ ROUND(g, h, a, b, c, d, e, f, schedule10, 0x243185BE) \
+ ROUND(f, g, h, a, b, c, d, e, schedule11, 0x550C7DC3) \
+ ROUND(e, f, g, h, a, b, c, d, schedule12, 0x72BE5D74) \
+ ROUND(d, e, f, g, h, a, b, c, schedule13, 0x80DEB1FE) \
+ ROUND(c, d, e, f, g, h, a, b, schedule14, 0x9BDC06A7) \
+ ROUND(b, c, d, e, f, g, h, a, schedule15, 0xC19BF174) \
+ ROUND(a, b, c, d, e, f, g, h, schedule16, 0xE49B69C1) \
+ ROUND(h, a, b, c, d, e, f, g, schedule17, 0xEFBE4786) \
+ ROUND(g, h, a, b, c, d, e, f, schedule18, 0x0FC19DC6) \
+ ROUND(f, g, h, a, b, c, d, e, schedule19, 0x240CA1CC) \
+ ROUND(e, f, g, h, a, b, c, d, schedule20, 0x2DE92C6F) \
+ ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4A7484AA) \
+ ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5CB0A9DC) \
+ ROUND(b, c, d, e, f, g, h, a, schedule23, 0x76F988DA) \
+ ROUND(a, b, c, d, e, f, g, h, schedule24, 0x983E5152) \
+ ROUND(h, a, b, c, d, e, f, g, schedule25, 0xA831C66D) \
+ ROUND(g, h, a, b, c, d, e, f, schedule26, 0xB00327C8) \
+ ROUND(f, g, h, a, b, c, d, e, schedule27, 0xBF597FC7) \
+ ROUND(e, f, g, h, a, b, c, d, schedule28, 0xC6E00BF3) \
+ ROUND(d, e, f, g, h, a, b, c, schedule29, 0xD5A79147) \
+ ROUND(c, d, e, f, g, h, a, b, schedule30, 0x06CA6351) \
+ ROUND(b, c, d, e, f, g, h, a, schedule31, 0x14292967)
+
+#define ALL_ROUND_B1_2() \
+ ROUND(a, b, c, d, e, f, g, h, schedule00, 0x27B70A85) \
+ ROUND(h, a, b, c, d, e, f, g, schedule01, 0x2E1B2138) \
+ ROUND(g, h, a, b, c, d, e, f, schedule02, 0x4D2C6DFC) \
+ ROUND(f, g, h, a, b, c, d, e, schedule03, 0x53380D13) \
+ ROUND(e, f, g, h, a, b, c, d, schedule04, 0x650A7354) \
+ ROUND(d, e, f, g, h, a, b, c, schedule05, 0x766A0ABB) \
+ ROUND(c, d, e, f, g, h, a, b, schedule06, 0x81C2C92E) \
+ ROUND(b, c, d, e, f, g, h, a, schedule07, 0x92722C85) \
+ ROUND(a, b, c, d, e, f, g, h, schedule08, 0xA2BFE8A1) \
+ ROUND(h, a, b, c, d, e, f, g, schedule09, 0xA81A664B) \
+ ROUND(g, h, a, b, c, d, e, f, schedule10, 0xC24B8B70) \
+ ROUND(f, g, h, a, b, c, d, e, schedule11, 0xC76C51A3) \
+ ROUND(e, f, g, h, a, b, c, d, schedule12, 0xD192E819) \
+ ROUND(d, e, f, g, h, a, b, c, schedule13, 0xD6990624) \
+ ROUND(c, d, e, f, g, h, a, b, schedule14, 0xF40E3585) \
+ ROUND(b, c, d, e, f, g, h, a, schedule15, 0x106AA070) \
+ ROUND(a, b, c, d, e, f, g, h, schedule16, 0x19A4C116) \
+ ROUND(h, a, b, c, d, e, f, g, schedule17, 0x1E376C08) \
+ ROUND(g, h, a, b, c, d, e, f, schedule18, 0x2748774C) \
+ ROUND(f, g, h, a, b, c, d, e, schedule19, 0x34B0BCB5) \
+ ROUND(e, f, g, h, a, b, c, d, schedule20, 0x391C0CB3) \
+ ROUND(d, e, f, g, h, a, b, c, schedule21, 0x4ED8AA4A) \
+ ROUND(c, d, e, f, g, h, a, b, schedule22, 0x5B9CCA4F) \
+ ROUND(b, c, d, e, f, g, h, a, schedule23, 0x682E6FF3) \
+ ROUND(a, b, c, d, e, f, g, h, schedule24, 0x748F82EE) \
+ ROUND(h, a, b, c, d, e, f, g, schedule25, 0x78A5636F) \
+ ROUND(g, h, a, b, c, d, e, f, schedule26, 0x84C87814) \
+ ROUND(f, g, h, a, b, c, d, e, schedule27, 0x8CC70208) \
+ ROUND(e, f, g, h, a, b, c, d, schedule28, 0x90BEFFFA) \
+ ROUND(d, e, f, g, h, a, b, c, schedule29, 0xA4506CEB) \
+ ROUND(c, d, e, f, g, h, a, b, schedule30, 0xBEF9A3F7) \
+ ROUND(b, c, d, e, f, g, h, a, schedule31, 0xC67178F2)
diff --git a/SYCL/src/utils.cpp b/SYCL/src/utils.cpp
new file mode 100644
index 0000000..5d61313
--- /dev/null
+++ b/SYCL/src/utils.cpp
@@ -0,0 +1,293 @@
+/* Modifications Copyright (C) 2023 Intel Corporation
+ *
+ * This program is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2, as published
+ * by the Free Software Foundation.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, see .
+ *
+ *
+ * SPDX-License-Identifier: GPL-2.0-only
+ */
+
+/*
+ * BitCracker: BitLocker password cracking tool, CUDA version.
+ * Copyright (C) 2013-2017 Elena Ago
+ * Massimo Bernaschi
+ *
+ * This file is part of the BitCracker project: https://github.com/e-ago/bitcracker
+ *
+ * BitCracker is free software: you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation, either version 2 of the License, or
+ * (at your option) any later version.
+ *
+ * BitCracker is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with BitCracker. If not, see .
+ */
+
+#include "bitcracker.h"
+
+/* John The Ripper function */
+char *strtokm(char *s1, const char *delims)
+{
+ static char *last = NULL;
+ char *endp;
+
+ if (!s1)
+ s1 = last;
+ if (!s1 || *s1 == 0)
+ return last = NULL;
+ endp = strpbrk(s1, delims);
+ if (endp) {
+ *endp = '\0';
+ last = endp + 1;
+ } else
+ last = NULL;
+ return s1;
+}
+
+void * Calloc(size_t len, size_t size) {
+ void * ptr = NULL;
+ if( size <= 0)
+ {
+ fprintf(stderr, "Critical error: requested memory size is 0\n");
+ exit(EXIT_FAILURE);
+ }
+
+ ptr = (void *)calloc(len, size);
+ if( ptr == NULL )
+ {
+ fprintf(stderr, "Critical error: Memory allocation\n");
+ exit(EXIT_FAILURE);
+ }
+ return ptr;
+}
+
+int parse_data(char *input_hash, unsigned char ** salt, unsigned char ** nonce, unsigned char ** vmk, unsigned char ** mac)
+{
+ if(!input_hash)
+ {
+ fprintf(stderr, "No input hash provided\n");
+ return BIT_FAILURE;
+ }
+
+ FILE * fphash = nullptr;
+ fphash = fopen(input_hash, "r");
+ if (!fphash) {
+ fprintf(stderr, "! %s : %s\n", input_hash, strerror(errno));
+ return BIT_FAILURE;
+ }
+
+ char * hash;
+ char * p;
+ int i, j;
+ int salt_size, iterations, vmk_size, nonce_size;
+ char temp[3];
+ const char zero_string[17]="0000000000000000";
+
+ (*salt) = (unsigned char *) Calloc(SALT_SIZE, sizeof(unsigned char));
+ (*nonce) = (unsigned char *) Calloc(NONCE_SIZE, sizeof(unsigned char));
+ (*vmk) = (unsigned char *) Calloc(VMK_SIZE, sizeof(unsigned char));
+ (*mac) = (unsigned char *) Calloc(MAC_SIZE, sizeof(unsigned char));
+
+ hash = (char *) Calloc(INPUT_HASH_SIZE, sizeof(char));
+
+ if(fgets(hash, INPUT_HASH_SIZE, fphash) == NULL)
+ {
+ fprintf(stderr, "No correct input hash provided\n");
+ goto out;
+ }
+
+ // printf("Reading hash file \"%s\"\n%s", input_hash, hash);
+ printf("Reading hash file \"%s\"\n", input_hash);
+
+ if (strncmp(hash, HASH_TAG, HASH_TAG_LEN) != 0)
+ {
+ fprintf(stderr, "Wrong hash format\n");
+ goto out;
+ }
+
+ hash += HASH_TAG_LEN;
+ p = strtokm(hash, "$");
+
+ p = strtokm(NULL, "$"); // salt length
+ salt_size = atoi(p);
+ if(salt_size != SALT_SIZE)
+ {
+ fprintf(stderr, "Wrong Salt size\n");
+ goto out;
+ }
+
+ p = strtokm(NULL, "$"); // salt
+ for (i = 0, j = 0; i < salt_size * 2; i += 2, j++)
+ {
+ temp[0] = p[i];
+ temp[1] = p[i + 1];
+ temp[2] = '\0';
+ long int ret = strtol(temp, NULL, 16);
+ (*salt)[j] = (unsigned char)(ret);
+ }
+
+ p = strtokm(NULL, "$"); // iterations
+ iterations = atoi(p);
+ if(iterations != NUM_HASH_BLOCKS)
+ {
+ fprintf(stderr, "Wrong Iterations parameter\n");
+ goto out;
+ }
+
+ p = strtokm(NULL, "$"); // nonce length
+ nonce_size = atoi(p);
+ if(nonce_size != NONCE_SIZE)
+ {
+ fprintf(stderr, "Wrong Nonce size\n");
+ goto out;
+ }
+
+ p = strtokm(NULL, "$"); // nonce
+ for (i = 0, j = 0; i < nonce_size*2; i+=2, j++)
+ {
+ temp[0] = p[i];
+ temp[1] = p[i + 1];
+ temp[2] = '\0';
+ long int ret = strtol(temp, NULL, 16);
+ (*nonce)[j] = (unsigned char)(ret);
+ }
+
+ p = strtokm(NULL, "$"); // vmk size
+ vmk_size = atoi(p);
+ if(vmk_size != VMK_SIZE)
+ {
+ fprintf(stderr, "Wrong VMK size\n");
+ goto out;
+ }
+
+ p = strtokm(NULL, "$"); // mac
+ for (i = 0, j = 0; i < MAC_SIZE*2; i+=2, j++)
+ {
+ temp[0] = p[i];
+ temp[1] = p[i + 1];
+ temp[2] = '\0';
+ long int ret = strtol(temp, NULL, 16);
+ (*mac)[j] = (unsigned char)(ret);
+ }
+
+ if(!memcmp((*mac), zero_string, MAC_SIZE))
+ {
+ free(*mac);
+ (*mac)=NULL;
+ }
+
+ // vmk
+ for (j=0; i < vmk_size*2; i+=2, j++)
+ {
+ temp[0] = p[i];
+ temp[1] = p[i + 1];
+ temp[2] = '\0';
+ long int ret = strtol(temp, NULL, 16);
+ (*vmk)[j] = (unsigned char)(ret);
+ }
+
+ fclose(fphash);
+
+ return BIT_SUCCESS;
+
+ out:
+ fclose(fphash);
+
+ free(*salt);
+ free(*nonce);
+ free(*vmk);
+ free(*mac);
+ free(hash);
+
+ return BIT_FAILURE;
+}
+
+static int print_once = 0;
+uint32_t read_password(
+ uint32_t ** buf_uint32, // a 32 uint32 slot for each password
+ char ** buf_char, // a 64 char slot for each password
+ uint32_t max_num_pswd_per_read,
+ FILE *fp)
+{
+ int j, k, size;
+ uint32_t num_pswd = 0; // count of passwords
+ char this_pswd[PSWD_NUM_CHAR]; // temporary storage for current password
+
+ if (fp == NULL || feof(fp) || buf_uint32 == NULL || buf_char == NULL) {
+ return 0;
+ }
+
+ memset(this_pswd, 0, PSWD_NUM_CHAR); // clear this_pswd, then
+ while(fgets(this_pswd, PSWD_NUM_CHAR, fp)) { // read current password into this_pswd
+ size = strlen(this_pswd) - 1;
+
+ // print warning
+ if((size < MIN_INPUT_PASSWORD_LEN || size > SECOND_LENGHT) && print_once == 0) {
+ fprintf(stderr, "WARNING: During USER PASSWORD attack, "
+ "only passwords between %d and %d character are considered. "
+ "Passwords like %s will be ignored.\n",
+ MIN_INPUT_PASSWORD_LEN, SECOND_LENGHT, this_pswd);
+ print_once = 1;
+ }
+
+ // if not good password, continue to next
+ if(size < MIN_INPUT_PASSWORD_LEN || size > SECOND_LENGHT || this_pswd[0] == '\n') {
+ continue;
+ }
+
+ // save this password to buf_char
+ memset((*buf_char) + (num_pswd * PSWD_NUM_CHAR), 0, PSWD_NUM_CHAR);
+ memcpy((*buf_char) + (num_pswd * PSWD_NUM_CHAR), this_pswd, size);
+
+ this_pswd[size] = 0x80; // terminate this password with 0x80
+ j = 0; // buf_uint32 element position
+ k = 0; // this_pswd char position
+ // For each password, there 32 slots.
+ // Each slot is a uint32 and is filled up by 'transformed' two-consecutive-chars of the password.
+ // j is half of k
+ do {
+ ((*buf_uint32) + (num_pswd * PSWD_NUM_UINT32) + j)[0] = (((uint32_t)this_pswd[k]) << 24) & 0xFF000000;
+ k++;
+ if(k <= size) {
+ ((*buf_uint32) + (num_pswd * PSWD_NUM_UINT32) + j)[0] |= (((uint32_t)this_pswd[k]) << 8) & 0x0000FF00;
+ }
+ k++;
+ j++;
+ } while(k <= size);
+
+ // based on password size, fill up (14 and 15) or (30 and 31) positions
+ if(size <= FIRST_LENGHT)
+ {
+ ((*buf_uint32) + (num_pswd * PSWD_NUM_UINT32) + 14)[0] = 0xFFFFFFFF;
+ ((*buf_uint32) + (num_pswd * PSWD_NUM_UINT32) + 15)[0] = (((uint8_t)(((size * 2) << 3) >> 8)) << 8) | ((uint8_t)((size * 2) << 3));
+ }
+ else
+ {
+ ((*buf_uint32) + (num_pswd * PSWD_NUM_UINT32) + 30)[0] = 0;
+ ((*buf_uint32) + (num_pswd * PSWD_NUM_UINT32) + 31)[0] = (((uint8_t)(((size * 2) << 3) >> 8)) << 8) | ((uint8_t)((size * 2) << 3));
+ }
+
+ memset(this_pswd, 0, PSWD_NUM_CHAR); // clear this_pswd
+ num_pswd++;
+
+ if(num_pswd >= max_num_pswd_per_read) {
+ break;
+ }
+ }
+
+ return num_pswd;
+}
diff --git a/SYCL/src/w_blocks.cpp b/SYCL/src/w_blocks.cpp
new file mode 100644
index 0000000..48e145b
--- /dev/null
+++ b/SYCL/src/w_blocks.cpp
@@ -0,0 +1,237 @@
+/* Modifications Copyright (C) 2023 Intel Corporation
+ *
+ * This program is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2, as published
+ * by the Free Software Foundation.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, see .
+ *
+ *
+ * SPDX-License-Identifier: GPL-2.0-only
+ */
+
+/*
+ * BitCracker: BitLocker password cracking tool, CUDA version.
+ * Copyright (C) 2013-2017 Elena Ago
+ * Massimo Bernaschi
+ *
+ * This file is part of the BitCracker project: https://github.com/e-ago/bitcracker
+ *
+ * BitCracker is free software: you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License as published by
+ * the Free Software Foundation, either version 2 of the License, or
+ * (at your option) any later version.
+ *
+ * BitCracker is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with BitCracker. If not, see .
+ */
+
+#include
+#include "bitcracker.h"
+#include
+#include
+
+#define ROR(x, i) (((x) << (32 - (i))) | ((x) >> (i)))
+
+#define LOADSCHEDULE_WPRE(j, i) \
+ d_w_words_uint32[j] = \
+ (uint32_t)block[i * 4 + 0] << 24 \
+ | (uint32_t)block[i * 4 + 1] << 16 \
+ | (uint32_t)block[i * 4 + 2] << 8 \
+ | (uint32_t)block[i * 4 + 3];
+
+#define CALCSCHEDULE_WPRE(j) \
+ d_w_words_uint32[j] = d_w_words_uint32[j - 16] + d_w_words_uint32[j - 7] \
+ + (ROR(d_w_words_uint32[j - 15], 7) ^ ROR(d_w_words_uint32[j - 15], 18) ^ (d_w_words_uint32[j - 15] >> 3)) \
+ + (ROR(d_w_words_uint32[j - 2], 17) ^ ROR(d_w_words_uint32[j - 2], 19) ^ (d_w_words_uint32[j - 2] >> 10));
+
+void kernel_w_block(
+ unsigned char salt[SALT_SIZE],
+ unsigned char padding[40],
+ uint32_t * d_w_words_uint32,
+ sycl::nd_item<1> item) //, sycl::stream out)
+{
+ uint64_t tid = item.get_global_id(0);
+ if (tid >= NUM_HASH_BLOCKS) return;
+
+ uint64_t texBlockId;
+ unsigned char block[HASH_BLOCK_NUM_UINT32];
+
+ int i;
+ // index 0-15
+ for(i = 0; i < SALT_SIZE; i++){
+ block[i] = salt[i];
+ }
+ // index 24-63
+ for(i = 0; i < PADDING_SIZE; i++){
+ block[i + 24] = padding[i];
+ }
+
+ // while(tid < NUM_HASH_BLOCKS)
+ {
+ // index 16-23
+ block[16] = (unsigned char) (tid >> (0 * 8));
+ block[17] = (unsigned char) (tid >> (1 * 8));
+ block[18] = (unsigned char) (tid >> (2 * 8));
+ block[19] = (unsigned char) (tid >> (3 * 8));
+ block[20] = (unsigned char) (tid >> (4 * 8));
+ block[21] = (unsigned char) (tid >> (5 * 8));
+ block[22] = (unsigned char) (tid >> (6 * 8));
+ block[23] = (unsigned char) (tid >> (7 * 8));
+
+ texBlockId = HASH_BLOCK_NUM_UINT32 * tid;
+ LOADSCHEDULE_WPRE(texBlockId + 0, 0)
+ LOADSCHEDULE_WPRE(texBlockId + 1, 1)
+ LOADSCHEDULE_WPRE(texBlockId + 2, 2)
+ LOADSCHEDULE_WPRE(texBlockId + 3, 3)
+ LOADSCHEDULE_WPRE(texBlockId + 4, 4)
+ LOADSCHEDULE_WPRE(texBlockId + 5, 5)
+ LOADSCHEDULE_WPRE(texBlockId + 6, 6)
+ LOADSCHEDULE_WPRE(texBlockId + 7, 7)
+ LOADSCHEDULE_WPRE(texBlockId + 8, 8)
+ LOADSCHEDULE_WPRE(texBlockId + 9, 9)
+ LOADSCHEDULE_WPRE(texBlockId + 10, 10)
+ LOADSCHEDULE_WPRE(texBlockId + 11, 11)
+ LOADSCHEDULE_WPRE(texBlockId + 12, 12)
+ LOADSCHEDULE_WPRE(texBlockId + 13, 13)
+ LOADSCHEDULE_WPRE(texBlockId + 14, 14)
+ LOADSCHEDULE_WPRE(texBlockId + 15, 15)
+ CALCSCHEDULE_WPRE(texBlockId + 16)
+ CALCSCHEDULE_WPRE(texBlockId + 17)
+ CALCSCHEDULE_WPRE(texBlockId + 18)
+ CALCSCHEDULE_WPRE(texBlockId + 19)
+ CALCSCHEDULE_WPRE(texBlockId + 20)
+ CALCSCHEDULE_WPRE(texBlockId + 21)
+ CALCSCHEDULE_WPRE(texBlockId + 22)
+ CALCSCHEDULE_WPRE(texBlockId + 23)
+ CALCSCHEDULE_WPRE(texBlockId + 24)
+ CALCSCHEDULE_WPRE(texBlockId + 25)
+ CALCSCHEDULE_WPRE(texBlockId + 26)
+ CALCSCHEDULE_WPRE(texBlockId + 27)
+ CALCSCHEDULE_WPRE(texBlockId + 28)
+ CALCSCHEDULE_WPRE(texBlockId + 29)
+ CALCSCHEDULE_WPRE(texBlockId + 30)
+ CALCSCHEDULE_WPRE(texBlockId + 31)
+ CALCSCHEDULE_WPRE(texBlockId + 32)
+ CALCSCHEDULE_WPRE(texBlockId + 33)
+ CALCSCHEDULE_WPRE(texBlockId + 34)
+ CALCSCHEDULE_WPRE(texBlockId + 35)
+ CALCSCHEDULE_WPRE(texBlockId + 36)
+ CALCSCHEDULE_WPRE(texBlockId + 37)
+ CALCSCHEDULE_WPRE(texBlockId + 38)
+ CALCSCHEDULE_WPRE(texBlockId + 39)
+ CALCSCHEDULE_WPRE(texBlockId + 40)
+ CALCSCHEDULE_WPRE(texBlockId + 41)
+ CALCSCHEDULE_WPRE(texBlockId + 42)
+ CALCSCHEDULE_WPRE(texBlockId + 43)
+ CALCSCHEDULE_WPRE(texBlockId + 44)
+ CALCSCHEDULE_WPRE(texBlockId + 45)
+ CALCSCHEDULE_WPRE(texBlockId + 46)
+ CALCSCHEDULE_WPRE(texBlockId + 47)
+ CALCSCHEDULE_WPRE(texBlockId + 48)
+ CALCSCHEDULE_WPRE(texBlockId + 49)
+ CALCSCHEDULE_WPRE(texBlockId + 50)
+ CALCSCHEDULE_WPRE(texBlockId + 51)
+ CALCSCHEDULE_WPRE(texBlockId + 52)
+ CALCSCHEDULE_WPRE(texBlockId + 53)
+ CALCSCHEDULE_WPRE(texBlockId + 54)
+ CALCSCHEDULE_WPRE(texBlockId + 55)
+ CALCSCHEDULE_WPRE(texBlockId + 56)
+ CALCSCHEDULE_WPRE(texBlockId + 57)
+ CALCSCHEDULE_WPRE(texBlockId + 58)
+ CALCSCHEDULE_WPRE(texBlockId + 59)
+ CALCSCHEDULE_WPRE(texBlockId + 60)
+ CALCSCHEDULE_WPRE(texBlockId + 61)
+ CALCSCHEDULE_WPRE(texBlockId + 62)
+ CALCSCHEDULE_WPRE(texBlockId + 63)
+
+ // if (texBlockId == (64*10)) {
+ // out <<"hello from kernel_w_block\n\n";
+ // out << "NUM_HASH_BLOCKS: " << NUM_HASH_BLOCKS << "\n";
+ // out << "texBlockId: ";
+ // out << texBlockId << "\n";
+ // for (int j = 0; j < 64; j++) {
+ // out << "d_w_words_uint32[" << sycl::setw(2) << j << "] = " << d_w_words_uint32[texBlockId + j] << "\n";
+ // }
+ // out << "\n";
+ // }
+
+ // tid += (item.get_local_range().get(0) * item.get_group_range(0));
+ }
+}
+
+int evaluate_w_block(
+ unsigned char * salt,
+ uint32_t * d_w_words_uint32,
+ double& duration,
+ sycl::queue qbc)
+{
+ unsigned char * salt_d;
+ unsigned char * padding;
+ unsigned char * padding_d;
+ uint64_t msgLen;
+
+ if(salt == NULL || d_w_words_uint32 == NULL)
+ return BIT_FAILURE;
+
+ padding = (unsigned char *) Calloc(PADDING_SIZE, sizeof(unsigned char));
+ padding[0] = 0x80;
+ memset(padding + 1, 0, 31);
+ msgLen = (FIXED_PART_INPUT_CHAIN_HASH << 3);
+ for (int i = 0; i < 8; i++) {
+ padding[PADDING_SIZE - 1 - i] = (uint8_t)(msgLen >> (i * 8));
+ }
+
+#ifdef DEBUG_TIME
+ auto time11 = std::chrono::steady_clock::now();
+#endif
+
+ // allocate device memory
+ salt_d = (unsigned char *)sycl::malloc_device(SALT_SIZE * sizeof(unsigned char), qbc);
+ padding_d = (unsigned char *)sycl::malloc_device(PADDING_SIZE * sizeof(unsigned char), qbc);
+
+ // copy to device memory
+ auto e1 = qbc.memcpy(salt_d, salt, SALT_SIZE * sizeof(unsigned char));
+ auto e2 = qbc.memcpy(padding_d, padding, PADDING_SIZE * sizeof(unsigned char));
+
+ // auto time12 = std::chrono::steady_clock::now();
+ // auto duration1 = std::chrono::duration(time12 - time11).count();
+ // duration += duration1;
+ // std::cout << "evaluate_w_block() - alloc and memcpy, duration1: " << duration1 << " us\n\n";
+
+ // auto time21 = std::chrono::steady_clock::now();
+
+ // launch kernel
+ qbc.parallel_for(
+ sycl::nd_range<1>(NUM_HASH_BLOCKS, 32), {std::move(e1), std::move(e2)},
+ [=](sycl::nd_item<1> item) {
+ kernel_w_block(salt_d, padding_d, d_w_words_uint32, item);
+ }
+ );
+ qbc.wait();
+
+#ifdef DEBUG_TIME
+ auto time22 = std::chrono::steady_clock::now();
+ auto duration2 = std::chrono::duration(time22 - time11).count();
+ duration += duration2;
+ std::cout << "evaluate_w_block -> malloc + H2D + kernel_w_block, duration2: " << duration2 << " us\n\n";
+#endif
+
+ free(padding);
+
+ sycl::free(salt_d, qbc);
+ sycl::free(padding_d, qbc);
+
+ return BIT_SUCCESS;
+}