diff --git a/fw/Core/Hitcon/Secret/.gitignore b/fw/Core/Hitcon/Secret/.gitignore index 77b1560..68862d3 100644 --- a/fw/Core/Hitcon/Secret/.gitignore +++ b/fw/Core/Hitcon/Secret/.gitignore @@ -1,3 +1 @@ -*.log -cache secret-out*.h diff --git a/fw/Core/Hitcon/Secret/replace.py b/fw/Core/Hitcon/Secret/replace.py new file mode 100644 index 0000000..57f24d3 --- /dev/null +++ b/fw/Core/Hitcon/Secret/replace.py @@ -0,0 +1,127 @@ +from pathlib import Path +import sys +import re +import random + +random.seed(0x1337) + +def load_log(): + if Path('cache').exists(): + with open('cache', 'r') as f: + return eval(f.read()) + + total = {} + + for file in Path('.').rglob('log-*.log'): + with open(file, 'r') as f: + lines = f.readlines() + + for line in lines: + try: + column_id, num_of_0, num = map(int, line.strip().split()) + total[column_id] = total.get(column_id, {}) + total[column_id][num_of_0] = total[column_id].get(num_of_0, []) + if len(total[column_id][num_of_0]) < 10: + total[column_id][num_of_0].append(num) + except ValueError as e: + print(f'Warning {file}: {e}', file=sys.stderr) + print(f' Line: {line.strip()}', file=sys.stderr) + + with open('cache', 'w') as f: + f.write(str(total)) + + return total + + +def main_old(): + total = load_log() + + with open('secret.h') as f: + content_original = f.read() + + pattern = re.compile(r'(0x[0-9a-fA-F]{2}, ){9}// (\d+)') + + has_38 = [0x0f, 0x0f, 0x01, 0x0e] + random.shuffle(has_38) + + for i in range(4): + start = 0 + content = content_original + column_ids = list(range(16)) * 3 + this_38 = has_38.pop() + column_ids.remove(this_38) # one column that has 38 is reserved + random.shuffle(column_ids) + + while True: + match = pattern.search(content, start) + if match is None: + break + + # column_id = int(match.group().split(', ')[0], 16) + try: + column_id = column_ids.pop(0) + except IndexError: + column_id = this_38 + count = int(match.group().split('//')[1]) + + try: + data, *total[column_id][count] = total[column_id][count] + except (KeyError, IndexError, ValueError) as e: + column_ids.append(column_id) + # print(f'Warning: {e}', file=sys.stderr) + # print(f' column_id: {column_id}', file=sys.stderr) + # print(f' count: {count}', file=sys.stderr) + # print(f' total: {total[column_id].get(count, [])}', file=sys.stderr) + continue + + # print(f'column_id: {column_id}, count: {count}, data: {data}') + + start = match.end() + content = content[:match.start()] +\ + f'0x{column_id:02x}, ' +\ + ', '.join(f'0x{d:02x}' for d in int.to_bytes(data, 8, 'big')) +\ + f', // {count}' +\ + content[start:] + + # print(content) + + with open(f'secret-out-{i}.h', 'w') as f: + f.write(content) + + +def main(): + total = load_log() + + with open('secret.h') as f: + content_original = f.read() + + pattern = re.compile(r'DATA\(0x([0-9a-fA-F]{2})\), // (\d+)') + + start = 0 + content = content_original + + while True: + match = pattern.search(content, start) + if match is None: + break + + column_id = int(match.group(1), 16) + count = int(match.group(2)) + + data, *total[column_id][count] = total[column_id][count] + + start = match.end() + content = content[:match.start()] +\ + f'0x{column_id:02x}, ' +\ + ', '.join(f'0x{d:02x}' for d in int.to_bytes(data, 8, 'big')) +\ + f', // {count}' +\ + content[start:] + + print(content) + + # with open(f'secret-out-{i}.h', 'w') as f: + # f.write(content) + + +if __name__ == '__main__': + main() diff --git a/fw/Core/Hitcon/Secret/secret.h b/fw/Core/Hitcon/Secret/secret.h index 29271b1..0619c0f 100644 --- a/fw/Core/Hitcon/Secret/secret.h +++ b/fw/Core/Hitcon/Secret/secret.h @@ -26,58 +26,58 @@ constexpr uint8_t kGameAchievementData[] = { // clang-format off // Tetris - 0x00, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 15 - 0x00, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 16 - 0x01, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 17 - 0x01, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 18 - 0x02, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 20 - 0x02, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 22 - 0x03, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 24 - 0x03, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 26 - 0x04, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 28 - 0x04, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 30 - 0x05, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 31 - 0x05, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 32 - 0x06, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 33 - 0x06, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 34 - 0x07, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 35 - 0x07, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 37 + DATA(0x00), // 15 + DATA(0x00), // 16 + DATA(0x01), // 17 + DATA(0x01), // 18 + DATA(0x02), // 20 + DATA(0x02), // 22 + DATA(0x03), // 24 + DATA(0x03), // 26 + DATA(0x04), // 28 + DATA(0x04), // 30 + DATA(0x05), // 31 + DATA(0x05), // 32 + DATA(0x06), // 33 + DATA(0x06), // 34 + DATA(0x07), // 35 + DATA(0x07), // 37 // Dino - 0x08, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 15 - 0x08, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 16 - 0x09, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 17 - 0x09, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 18 - 0x0a, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 20 - 0x0a, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 22 - 0x0b, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 24 - 0x0b, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 26 - 0x0c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 28 - 0x0c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 30 - 0x0d, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 31 - 0x0d, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 32 - 0x0e, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 33 - 0x0e, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 34 - 0x0f, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 35 - 0x0f, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 37 + DATA(0x08), // 15 + DATA(0x08), // 16 + DATA(0x09), // 17 + DATA(0x09), // 18 + DATA(0x0a), // 20 + DATA(0x0a), // 22 + DATA(0x0b), // 24 + DATA(0x0b), // 26 + DATA(0x0c), // 28 + DATA(0x0c), // 30 + DATA(0x0d), // 31 + DATA(0x0d), // 32 + DATA(0x0e), // 33 + DATA(0x0e), // 35 + DATA(0x0f), // 37 + DATA(0x0f), // 38 // Snake - 0x00, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 10 - 0x01, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 10 - 0x02, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 12 - 0x03, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 12 - 0x04, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 14 - 0x05, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 15 - 0x06, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 16 - 0x07, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 18 - 0x08, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 20 - 0x09, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 22 - 0x0a, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 24 - 0x0b, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 26 - 0x0c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 28 - 0x0d, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 30 - 0x0e, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 32 - 0x0f, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, // 38 + DATA(0x00), // 10 + DATA(0x01), // 10 + DATA(0x02), // 12 + DATA(0x03), // 12 + DATA(0x04), // 14 + DATA(0x05), // 15 + DATA(0x06), // 16 + DATA(0x07), // 18 + DATA(0x08), // 20 + DATA(0x09), // 22 + DATA(0x0a), // 24 + DATA(0x0b), // 26 + DATA(0x0c), // 28 + DATA(0x0d), // 30 + DATA(0x0e), // 32 + DATA(0x0f), // 38 // clang-format on }; constexpr size_t kGameAchievementDataSize = @@ -85,8 +85,10 @@ constexpr size_t kGameAchievementDataSize = static_assert(kGameAchievementDataSize % 9 == 0); constexpr size_t kGameAchievementDataCount = kGameAchievementDataSize / 9; -constexpr uint32_t kPerDataPeriod = 5 * 60; -constexpr uint8_t kXBoardWindowSize = 4; +constexpr uint8_t PARTITION_LEN = 4; +constexpr uint32_t PER_DATA_PERIOD = 5 * 60; +constexpr uint8_t XBOARD_WINDOW_SIZE = 4; +#if defined SPONSOR constexpr uint8_t kStrayIRData[] = { // First byte is the column @@ -95,24 +97,42 @@ constexpr uint8_t kStrayIRData[] = { // clang-format off /* PARTITION 1 */ - DATA(0x0a), // 15 - DATA(0x0b), // 15 - DATA(0x0c), // 15 - DATA(0x0f), // 15 - DATA(0x0a), // 15 - DATA(0x0b), // 15 - DATA(0x0c), // 15 - DATA(0x0f), // 15 - DATA(0x0a), // 16 - DATA(0x0b), // 16 - DATA(0x0c), // 16 - DATA(0x0f), // 16 - DATA(0x0a), // 16 - DATA(0x0b), // 16 - DATA(0x0c), // 16 - DATA(0x0f), // 16 + DATA(0x0a), // 9 + DATA(0x0b), // 9 + DATA(0x0c), // 9 + DATA(0x0f), // 9 + DATA(0x0a), // 9 + DATA(0x0b), // 9 + DATA(0x0c), // 9 + DATA(0x0f), // 9 + DATA(0x0a), // 10 + DATA(0x0b), // 10 + DATA(0x0c), // 10 + DATA(0x0f), // 10 + DATA(0x0a), // 10 + DATA(0x0b), // 10 + DATA(0x0c), // 10 + DATA(0x0f), // 10 /* PARTITION 2 */ + DATA(0x0a), // 13 + DATA(0x0b), // 13 + DATA(0x0c), // 13 + DATA(0x0f), // 13 + DATA(0x0a), // 13 + DATA(0x0b), // 13 + DATA(0x0c), // 13 + DATA(0x0f), // 13 + DATA(0x0a), // 14 + DATA(0x0b), // 14 + DATA(0x0c), // 14 + DATA(0x0f), // 14 + DATA(0x0a), // 14 + DATA(0x0b), // 14 + DATA(0x0c), // 14 + DATA(0x0f), // 14 + + /* PARTITION 3 */ DATA(0x0a), // 17 DATA(0x0b), // 17 DATA(0x0c), // 17 @@ -130,24 +150,6 @@ constexpr uint8_t kStrayIRData[] = { DATA(0x0c), // 18 DATA(0x0f), // 18 - /* PARTITION 3 */ - DATA(0x0a), // 19 - DATA(0x0b), // 19 - DATA(0x0c), // 19 - DATA(0x0f), // 19 - DATA(0x0a), // 19 - DATA(0x0b), // 19 - DATA(0x0c), // 19 - DATA(0x0f), // 19 - DATA(0x0a), // 20 - DATA(0x0b), // 20 - DATA(0x0c), // 20 - DATA(0x0f), // 20 - DATA(0x0a), // 20 - DATA(0x0b), // 20 - DATA(0x0c), // 20 - DATA(0x0f), // 20 - /* PARTITION 4 */ DATA(0x0a), // 21 DATA(0x0b), // 21 @@ -175,58 +177,58 @@ constexpr uint8_t kStrayXBoardData[] = { // clang-format off /* PARTITION 1 */ - DATA(0x0a), // 17 - DATA(0x0b), // 17 - DATA(0x0c), // 17 - DATA(0x0d), // 17 - DATA(0x0e), // 17 - DATA(0x0f), // 17 - DATA(0x0a), // 17 - DATA(0x0b), // 17 - DATA(0x0c), // 18 - DATA(0x0d), // 18 - DATA(0x0e), // 18 - DATA(0x0f), // 18 - DATA(0x0a), // 18 - DATA(0x0b), // 18 - DATA(0x0c), // 18 - DATA(0x0d), // 18 + DATA(0x0a), // 11 + DATA(0x0b), // 11 + DATA(0x0c), // 11 + DATA(0x0d), // 11 + DATA(0x0e), // 11 + DATA(0x0f), // 11 + DATA(0x0a), // 11 + DATA(0x0b), // 11 + DATA(0x0c), // 12 + DATA(0x0d), // 12 + DATA(0x0e), // 12 + DATA(0x0f), // 12 + DATA(0x0a), // 12 + DATA(0x0b), // 12 + DATA(0x0c), // 12 + DATA(0x0d), // 12 /* PARTITION 2 */ + DATA(0x0e), // 15 + DATA(0x0f), // 15 + DATA(0x0a), // 15 + DATA(0x0b), // 15 + DATA(0x0c), // 15 + DATA(0x0d), // 15 + DATA(0x0e), // 15 + DATA(0x0f), // 15 + DATA(0x0a), // 16 + DATA(0x0b), // 16 + DATA(0x0c), // 16 + DATA(0x0d), // 16 + DATA(0x0e), // 16 + DATA(0x0f), // 16 + DATA(0x0a), // 16 + DATA(0x0b), // 16 + + /* PARTITION 3 */ + DATA(0x0c), // 19 + DATA(0x0d), // 19 DATA(0x0e), // 19 DATA(0x0f), // 19 DATA(0x0a), // 19 DATA(0x0b), // 19 DATA(0x0c), // 19 DATA(0x0d), // 19 - DATA(0x0e), // 19 - DATA(0x0f), // 19 + DATA(0x0e), // 20 + DATA(0x0f), // 20 DATA(0x0a), // 20 DATA(0x0b), // 20 DATA(0x0c), // 20 DATA(0x0d), // 20 DATA(0x0e), // 20 DATA(0x0f), // 20 - DATA(0x0a), // 20 - DATA(0x0b), // 20 - - /* PARTITION 3 */ - DATA(0x0c), // 21 - DATA(0x0d), // 21 - DATA(0x0e), // 21 - DATA(0x0f), // 21 - DATA(0x0a), // 21 - DATA(0x0b), // 21 - DATA(0x0c), // 21 - DATA(0x0d), // 21 - DATA(0x0e), // 22 - DATA(0x0f), // 22 - DATA(0x0a), // 22 - DATA(0x0b), // 22 - DATA(0x0c), // 22 - DATA(0x0d), // 22 - DATA(0x0e), // 22 - DATA(0x0f), // 22 /* PARTITION 4 */ DATA(0x0a), // 23 @@ -247,27 +249,539 @@ constexpr uint8_t kStrayXBoardData[] = { DATA(0x0d), // 24 // clang-format on }; +constexpr size_t PER_IR_PARTITION_SIZE = + sizeof(kStrayIRData) / sizeof(uint8_t) / PARTITION_LEN / 9; +constexpr size_t PER_XBoard_PARTITION_SIZE = + sizeof(kStrayXBoardData) / sizeof(uint8_t) / PARTITION_LEN / 9; -#ifdef DATA -#undef DATA -#endif +// defined SPONSOR +#elif defined SPEAKER + +constexpr uint8_t kStrayIRData[] = { + // First byte is the column + // 0x00, + // Repeat... + + // clang-format off + /* PARTITION 1 */ + DATA(0x00), // 9 + DATA(0x02), // 9 + DATA(0x04), // 9 + DATA(0x06), // 9 + DATA(0x08), // 9 + DATA(0x0a), // 9 + DATA(0x0c), // 9 + DATA(0x0e), // 9 + DATA(0x00), // 10 + DATA(0x02), // 10 + DATA(0x04), // 10 + DATA(0x06), // 10 + DATA(0x08), // 10 + DATA(0x0a), // 10 + DATA(0x0c), // 10 + DATA(0x0e), // 10 + + /* PARTITION 2 */ + DATA(0x00), // 13 + DATA(0x02), // 13 + DATA(0x04), // 13 + DATA(0x06), // 13 + DATA(0x08), // 13 + DATA(0x0a), // 13 + DATA(0x0c), // 13 + DATA(0x0e), // 13 + DATA(0x00), // 14 + DATA(0x02), // 14 + DATA(0x04), // 14 + DATA(0x06), // 14 + DATA(0x08), // 14 + DATA(0x0a), // 14 + DATA(0x0c), // 14 + DATA(0x0e), // 14 + + /* PARTITION 3 */ + DATA(0x00), // 17 + DATA(0x02), // 17 + DATA(0x04), // 17 + DATA(0x06), // 17 + DATA(0x08), // 17 + DATA(0x0a), // 17 + DATA(0x0c), // 17 + DATA(0x0e), // 17 + DATA(0x00), // 18 + DATA(0x02), // 18 + DATA(0x04), // 18 + DATA(0x06), // 18 + DATA(0x08), // 18 + DATA(0x0a), // 18 + DATA(0x0c), // 18 + DATA(0x0e), // 18 + + /* PARTITION 4 */ + DATA(0x00), // 21 + DATA(0x02), // 21 + DATA(0x04), // 21 + DATA(0x06), // 21 + DATA(0x08), // 21 + DATA(0x0a), // 21 + DATA(0x0c), // 21 + DATA(0x0e), // 21 + DATA(0x00), // 22 + DATA(0x02), // 22 + DATA(0x04), // 22 + DATA(0x06), // 22 + DATA(0x08), // 22 + DATA(0x0a), // 22 + DATA(0x0c), // 22 + DATA(0x0e), // 22 + // clang-format on +}; + +constexpr uint8_t kStrayXBoardData[] = { + // First byte is the column + // 0x00, + // Repeat... + + // clang-format off + /* PARTITION 1 */ + DATA(0x00), // 11 + DATA(0x01), // 11 + DATA(0x02), // 11 + DATA(0x03), // 11 + DATA(0x04), // 11 + DATA(0x05), // 11 + DATA(0x06), // 11 + DATA(0x07), // 11 + DATA(0x08), // 12 + DATA(0x09), // 12 + DATA(0x0a), // 12 + DATA(0x0b), // 12 + DATA(0x0c), // 12 + DATA(0x0d), // 12 + DATA(0x0e), // 12 + DATA(0x0f), // 12 + + /* PARTITION 2 */ + DATA(0x00), // 15 + DATA(0x01), // 15 + DATA(0x02), // 15 + DATA(0x03), // 15 + DATA(0x04), // 15 + DATA(0x05), // 15 + DATA(0x06), // 15 + DATA(0x07), // 15 + DATA(0x08), // 16 + DATA(0x09), // 16 + DATA(0x0a), // 16 + DATA(0x0b), // 16 + DATA(0x0c), // 16 + DATA(0x0d), // 16 + DATA(0x0e), // 16 + DATA(0x0f), // 16 + + /* PARTITION 3 */ + DATA(0x00), // 19 + DATA(0x01), // 19 + DATA(0x02), // 19 + DATA(0x03), // 19 + DATA(0x04), // 19 + DATA(0x05), // 19 + DATA(0x06), // 19 + DATA(0x07), // 19 + DATA(0x08), // 20 + DATA(0x09), // 20 + DATA(0x0a), // 20 + DATA(0x0b), // 20 + DATA(0x0c), // 20 + DATA(0x0d), // 20 + DATA(0x0e), // 20 + DATA(0x0f), // 20 + + /* PARTITION 4 */ + DATA(0x00), // 23 + DATA(0x01), // 23 + DATA(0x02), // 23 + DATA(0x03), // 23 + DATA(0x04), // 23 + DATA(0x05), // 23 + DATA(0x06), // 23 + DATA(0x07), // 23 + DATA(0x08), // 24 + DATA(0x09), // 24 + DATA(0x0a), // 24 + DATA(0x0b), // 24 + DATA(0x0c), // 24 + DATA(0x0d), // 24 + DATA(0x0e), // 24 + DATA(0x0f), // 24 + // clang-format on +}; +constexpr size_t PER_IR_PARTITION_SIZE = + sizeof(kStrayIRData) / sizeof(uint8_t) / PARTITION_LEN / 9; +constexpr size_t PER_XBoard_PARTITION_SIZE = + sizeof(kStrayXBoardData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +// defined SPEAKER +#elif defined HITCON_R1 + +constexpr uint8_t kStrayIRData[] = { + // First byte is the column + // 0x00, + // Repeat... + + // clang-format off + /* PARTITION 1 */ + DATA(0x00), // 8 + DATA(0x01), // 8 + DATA(0x02), // 8 + DATA(0x00), // 8 + DATA(0x01), // 8 + DATA(0x02), // 8 + DATA(0x00), // 8 + DATA(0x01), // 8 + DATA(0x02), // 9 + DATA(0x00), // 9 + DATA(0x01), // 9 + DATA(0x02), // 9 + DATA(0x00), // 9 + DATA(0x01), // 9 + DATA(0x02), // 9 + DATA(0x00), // 9 + + /* PARTITION 2 */ + DATA(0x01), // 12 + DATA(0x02), // 12 + DATA(0x00), // 12 + DATA(0x01), // 12 + DATA(0x02), // 12 + DATA(0x00), // 12 + DATA(0x01), // 12 + DATA(0x02), // 12 + DATA(0x00), // 13 + DATA(0x01), // 13 + DATA(0x02), // 13 + DATA(0x00), // 13 + DATA(0x01), // 13 + DATA(0x02), // 13 + DATA(0x00), // 13 + DATA(0x01), // 13 + + /* PARTITION 3 */ + DATA(0x02), // 16 + DATA(0x00), // 16 + DATA(0x01), // 16 + DATA(0x02), // 16 + DATA(0x00), // 16 + DATA(0x01), // 16 + DATA(0x02), // 16 + DATA(0x00), // 16 + DATA(0x01), // 17 + DATA(0x02), // 17 + DATA(0x00), // 17 + DATA(0x01), // 17 + DATA(0x02), // 17 + DATA(0x00), // 17 + DATA(0x01), // 17 + DATA(0x02), // 17 + + /* PARTITION 4 */ + DATA(0x00), // 20 + DATA(0x01), // 20 + DATA(0x02), // 20 + DATA(0x00), // 20 + DATA(0x01), // 20 + DATA(0x02), // 20 + DATA(0x00), // 20 + DATA(0x01), // 20 + DATA(0x02), // 21 + DATA(0x00), // 21 + DATA(0x01), // 21 + DATA(0x02), // 21 + DATA(0x00), // 21 + DATA(0x01), // 21 + DATA(0x02), // 21 + DATA(0x00), // 21 + // clang-format on +}; +constexpr size_t PER_IR_PARTITION_SIZE = + sizeof(kStrayIRData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +constexpr uint8_t kStrayXBoardData[] = {}; +constexpr size_t PER_XBoard_PARTITION_SIZE = + sizeof(kStrayXBoardData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +// defined HITCON_R1 +#elif defined HITCON_R2 + +constexpr uint8_t kStrayIRData[] = { + // First byte is the column + // 0x00, + // Repeat... + + // clang-format off + /* PARTITION 1 */ + DATA(0x02), // 8 + DATA(0x03), // 8 + DATA(0x04), // 8 + DATA(0x02), // 8 + DATA(0x03), // 8 + DATA(0x04), // 8 + DATA(0x02), // 8 + DATA(0x03), // 8 + DATA(0x04), // 9 + DATA(0x02), // 9 + DATA(0x03), // 9 + DATA(0x04), // 9 + DATA(0x02), // 9 + DATA(0x03), // 9 + DATA(0x04), // 9 + DATA(0x02), // 9 + + /* PARTITION 2 */ + DATA(0x03), // 12 + DATA(0x04), // 12 + DATA(0x02), // 12 + DATA(0x03), // 12 + DATA(0x04), // 12 + DATA(0x02), // 12 + DATA(0x03), // 12 + DATA(0x04), // 12 + DATA(0x02), // 13 + DATA(0x03), // 13 + DATA(0x04), // 13 + DATA(0x02), // 13 + DATA(0x03), // 13 + DATA(0x04), // 13 + DATA(0x02), // 13 + DATA(0x03), // 13 + + /* PARTITION 3 */ + DATA(0x04), // 16 + DATA(0x02), // 16 + DATA(0x03), // 16 + DATA(0x04), // 16 + DATA(0x02), // 16 + DATA(0x03), // 16 + DATA(0x04), // 16 + DATA(0x02), // 16 + DATA(0x03), // 17 + DATA(0x04), // 17 + DATA(0x02), // 17 + DATA(0x03), // 17 + DATA(0x04), // 17 + DATA(0x02), // 17 + DATA(0x03), // 17 + DATA(0x04), // 17 + + /* PARTITION 4 */ + DATA(0x02), // 20 + DATA(0x03), // 20 + DATA(0x04), // 20 + DATA(0x02), // 20 + DATA(0x03), // 20 + DATA(0x04), // 20 + DATA(0x02), // 20 + DATA(0x03), // 20 + DATA(0x04), // 21 + DATA(0x02), // 21 + DATA(0x03), // 21 + DATA(0x04), // 21 + DATA(0x02), // 21 + DATA(0x03), // 21 + DATA(0x04), // 21 + DATA(0x02), // 21 + // clang-format on +}; +constexpr size_t PER_IR_PARTITION_SIZE = + sizeof(kStrayIRData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +constexpr uint8_t kStrayXBoardData[] = {}; +constexpr size_t PER_XBoard_PARTITION_SIZE = + sizeof(kStrayXBoardData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +// defined HITCON_R2 +#elif defined HITCON_R3 + +constexpr uint8_t kStrayIRData[] = { + // First byte is the column + // 0x00, + // Repeat... + + // clang-format off + /* PARTITION 1 */ + DATA(0x04), // 8 + DATA(0x05), // 8 + DATA(0x06), // 8 + DATA(0x04), // 8 + DATA(0x05), // 8 + DATA(0x06), // 8 + DATA(0x04), // 8 + DATA(0x05), // 8 + DATA(0x06), // 9 + DATA(0x04), // 9 + DATA(0x05), // 9 + DATA(0x06), // 9 + DATA(0x04), // 9 + DATA(0x05), // 9 + DATA(0x06), // 9 + DATA(0x04), // 9 + + /* PARTITION 2 */ + DATA(0x05), // 12 + DATA(0x06), // 12 + DATA(0x04), // 12 + DATA(0x05), // 12 + DATA(0x06), // 12 + DATA(0x04), // 12 + DATA(0x05), // 12 + DATA(0x06), // 12 + DATA(0x04), // 13 + DATA(0x05), // 13 + DATA(0x06), // 13 + DATA(0x04), // 13 + DATA(0x05), // 13 + DATA(0x06), // 13 + DATA(0x04), // 13 + DATA(0x05), // 13 + + /* PARTITION 3 */ + DATA(0x06), // 16 + DATA(0x04), // 16 + DATA(0x05), // 16 + DATA(0x06), // 16 + DATA(0x04), // 16 + DATA(0x05), // 16 + DATA(0x06), // 16 + DATA(0x04), // 16 + DATA(0x05), // 17 + DATA(0x06), // 17 + DATA(0x04), // 17 + DATA(0x05), // 17 + DATA(0x06), // 17 + DATA(0x04), // 17 + DATA(0x05), // 17 + DATA(0x06), // 17 + + /* PARTITION 4 */ + DATA(0x04), // 20 + DATA(0x05), // 20 + DATA(0x06), // 20 + DATA(0x04), // 20 + DATA(0x05), // 20 + DATA(0x06), // 20 + DATA(0x04), // 20 + DATA(0x05), // 20 + DATA(0x06), // 21 + DATA(0x04), // 21 + DATA(0x05), // 21 + DATA(0x06), // 21 + DATA(0x04), // 21 + DATA(0x05), // 21 + DATA(0x06), // 21 + DATA(0x04), // 21 + // clang-format on +}; +constexpr size_t PER_IR_PARTITION_SIZE = + sizeof(kStrayIRData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +constexpr uint8_t kStrayXBoardData[] = {}; +constexpr size_t PER_XBoard_PARTITION_SIZE = + sizeof(kStrayXBoardData) / sizeof(uint8_t) / PARTITION_LEN / 9; + +// defined HITCONR3 +#elif defined HITCON_R4 + +constexpr uint8_t kStrayIRData[] = { + // First byte is the column + // 0x00, + // Repeat... + + // clang-format off + /* PARTITION 1 */ + DATA(0x06), // 8 + DATA(0x07), // 8 + DATA(0x08), // 8 + DATA(0x06), // 8 + DATA(0x07), // 8 + DATA(0x08), // 8 + DATA(0x06), // 8 + DATA(0x07), // 8 + DATA(0x08), // 9 + DATA(0x06), // 9 + DATA(0x07), // 9 + DATA(0x08), // 9 + DATA(0x06), // 9 + DATA(0x07), // 9 + DATA(0x08), // 9 + DATA(0x06), // 9 + + /* PARTITION 2 */ + DATA(0x07), // 12 + DATA(0x08), // 12 + DATA(0x06), // 12 + DATA(0x07), // 12 + DATA(0x08), // 12 + DATA(0x06), // 12 + DATA(0x07), // 12 + DATA(0x08), // 12 + DATA(0x06), // 13 + DATA(0x07), // 13 + DATA(0x08), // 13 + DATA(0x06), // 13 + DATA(0x07), // 13 + DATA(0x08), // 13 + DATA(0x06), // 13 + DATA(0x07), // 13 -constexpr int kPartitionOffset = 0x0; + /* PARTITION 3 */ + DATA(0x08), // 16 + DATA(0x06), // 16 + DATA(0x07), // 16 + DATA(0x08), // 16 + DATA(0x06), // 16 + DATA(0x07), // 16 + DATA(0x08), // 16 + DATA(0x06), // 16 + DATA(0x07), // 17 + DATA(0x08), // 17 + DATA(0x06), // 17 + DATA(0x07), // 17 + DATA(0x08), // 17 + DATA(0x06), // 17 + DATA(0x07), // 17 + DATA(0x08), // 17 -constexpr uint8_t kPartitionPacketId = 0x4; + /* PARTITION 4 */ + DATA(0x06), // 20 + DATA(0x07), // 20 + DATA(0x08), // 20 + DATA(0x06), // 20 + DATA(0x07), // 20 + DATA(0x08), // 20 + DATA(0x06), // 20 + DATA(0x07), // 20 + DATA(0x08), // 21 + DATA(0x06), // 21 + DATA(0x07), // 21 + DATA(0x08), // 21 + DATA(0x06), // 21 + DATA(0x07), // 21 + DATA(0x08), // 21 + DATA(0x06), // 21 + // clang-format on +}; +constexpr size_t PER_IR_PARTITION_SIZE = + sizeof(kStrayIRData) / sizeof(uint8_t) / PARTITION_LEN / 9; -constexpr size_t kPartitionCount = 4; +constexpr uint8_t kStrayXBoardData[] = {}; +constexpr size_t PER_XBoard_PARTITION_SIZE = + sizeof(kStrayXBoardData) / sizeof(uint8_t) / PARTITION_LEN / 9; -constexpr size_t kIrPartitionSize = - sizeof(kStrayIRData) / sizeof(kStrayIRData[0]) / kPartitionCount / 9; -constexpr size_t kXBoardPartitionSize = sizeof(kStrayXBoardData) / - sizeof(kStrayXBoardData[0]) / - kPartitionCount / 9; +#endif // HITCON_R4 -// From 0 to 0x10000, 0x10000 is using it 100% of the time. -constexpr int kIrUsePreparedChance = 0x8000; -// From 0 to 0x10000, 0x10000 is using it 100% of the time. -constexpr int kXBoardUsePreparedChance = 0x8000; +#ifdef DATA +#undef DATA +#endif } // namespace hitcon #endif // HITCON_SECRET_SECRET_H diff --git a/sw/secret-mining/.gitignore b/sw/secret-mining/.gitignore new file mode 100644 index 0000000..d5de9f2 --- /dev/null +++ b/sw/secret-mining/.gitignore @@ -0,0 +1,2 @@ +*.log +cache diff --git a/sw/secret-mining/Makefile b/sw/secret-mining/Makefile new file mode 100644 index 0000000..6c1a2d8 --- /dev/null +++ b/sw/secret-mining/Makefile @@ -0,0 +1,8 @@ +CXX = g++ +CXXFLAGS = -std=c++17 -Wall -Wextra -pedantic -O2 -fopenmp + +brute-cpu: *.cc *.h + $(CXX) $(CXXFLAGS) -o $@ brute-cpu.cc keccak.cc sha3_cpu.cc + +format: + clang-format -i *.cc *.h diff --git a/sw/secret-mining/brute-cpu.cc b/sw/secret-mining/brute-cpu.cc new file mode 100644 index 0000000..10e21a3 --- /dev/null +++ b/sw/secret-mining/brute-cpu.cc @@ -0,0 +1,176 @@ +#include +#include +#include +#include +#include +#include +#include + +#include "keccak.h" +#include "sha3_cpu.h" + +struct bdata_t { + union { + uint64_t u64[2]; + uint8_t u8[16]; + } u; +}; + +namespace { +// Look-up Table for the number of leading zero bits in a nibble +constexpr int LEADING_ZERO_BITS_LUT[16] = {4, 3, 2, 2, 1, 1, 1, 1, + 0, 0, 0, 0, 0, 0, 0, 0}; +} // namespace + +// Computes the number of leading zero bits in the given binary hash +int ComputePrefixZero(const uint8_t *bin_hash) { + int count = 0; + for (size_t i = 0; i < SHA3_256_HASH_SIZE; i++) { + uint8_t byte = bin_hash[i]; + uint8_t high_nibble = (byte & 0xF0) >> 4; + if (high_nibble != 0) { + return count + LEADING_ZERO_BITS_LUT[high_nibble]; + } + count += 4; + + uint8_t low_nibble = byte & 0x0F; + if (low_nibble != 0) { + return count + LEADING_ZERO_BITS_LUT[low_nibble]; + } + count += 4; + } + // All bytes are zero, so return the total number of bits in the hash. + // return SHA3_256_HASH_SIZE * 8; + // Ah screw it, nobody gets here, might as well **** around. + return *reinterpret_cast(0); +} + +bool check_compatibility_and_speed() { + bdata_t d1; + + // Test compatibility + for (int i = 0; i < 100000; ++i) { + d1.u.u64[0] = i; + d1.u.u64[1] = i; + + int cnt1, cnt2; + { + sha3_context c; + sha3_Init256(&c); + sha3_UpdateWord(&c, &d1.u.u64[0]); + sha3_UpdateWord(&c, &d1.u.u64[1]); + const uint8_t *hash = + reinterpret_cast(sha3_Finalize(&c)); + cnt1 = ComputePrefixZero(hash); + } + { + SHA3_cpu sha3_cpu(256); + sha3_cpu.init(); + sha3_cpu.add(reinterpret_cast(&d1.u.u64[0]), 8); + sha3_cpu.add(reinterpret_cast(&d1.u.u64[1]), 8); + std::vector hash = sha3_cpu.digest(); + cnt2 = ComputePrefixZero(hash.data()); + } + + if (cnt1 != cnt2) { + printf("Mismatch at %d: %d vs %d\n", i, cnt1, cnt2); + return false; + } + } + + // Test speed + { + auto start = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < 1000000; ++i) { + d1.u.u64[0] = i; + d1.u.u64[1] = i; + sha3_context c; + sha3_Init256(&c); + sha3_UpdateWord(&c, &d1.u.u64[0]); + sha3_UpdateWord(&c, &d1.u.u64[1]); + sha3_Finalize(&c); + } + auto end = std::chrono::high_resolution_clock::now(); + auto duration = + std::chrono::duration_cast(end - start) + .count(); + printf("keccak: %ld ms\n", duration); + } + { + auto start = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < 1000000; ++i) { + d1.u.u64[0] = i; + d1.u.u64[1] = i; + SHA3_cpu sha3_cpu(256); + sha3_cpu.init(); + sha3_cpu.add(reinterpret_cast(&d1.u.u64[0]), 8); + sha3_cpu.add(reinterpret_cast(&d1.u.u64[1]), 8); + sha3_cpu.digest(); + } + auto end = std::chrono::high_resolution_clock::now(); + auto duration = + std::chrono::duration_cast(end - start) + .count(); + printf("SHA3_cpu: %ld ms\n", duration); + } + + return true; +} + +int john_brute_use_sha3_cpu(int col, uint64_t start) { + struct bdata_t d1; + + memset(&d1.u.u8[0], 0, 16); + memcpy(&d1.u.u8[0], "HITCON", 6); + d1.u.u8[7] = col & 0xFF; + + SHA3_cpu c(256); + std::vector> res(256); + for (uint64_t i = start;; i++) { + c.init(); + uint64_t i2 = i; + for (int j = 0; j < 8; j++) { + d1.u.u8[15 - j] = i2 & 0x0FF; + i2 = i2 >> 8; + } + c.add(reinterpret_cast(&d1.u.u64[0]), 8); + c.add(reinterpret_cast(&d1.u.u64[1]), 8); + std::vector hash = c.digest(); + // printf("%d %d %d\n", hash[0], hash[1], hash[2]); + int cnt = ComputePrefixZero(hash.data()); + // printf("%llu - %d\n", i, cnt); + auto &r = res[cnt]; + if (r.size() < 65536) { + r.insert(i); + + // currently, print it to stdout + // TODO: send to DB + printf("%d %d %lu\n", col, cnt, i); + } + } + return 0; +} + +int main() { + // if (!check_compatibility_and_speed()) { + // return 1; + // } + + // get environment variable + const char *col_str = getenv("COL"); + if (!col_str) { + printf("COL not set\n"); + return 1; + } + int col = atoi(col_str); + + const char *start_str = getenv("START"); + uint64_t start = 0; + if (start_str) { + start = strtoull(start_str, nullptr, 10); + } + + return john_brute_use_sha3_cpu(col, start); + + return 0; +} diff --git a/sw/secret-mining/calc-num.py b/sw/secret-mining/calc-num.py new file mode 100644 index 0000000..e72a02c --- /dev/null +++ b/sw/secret-mining/calc-num.py @@ -0,0 +1,61 @@ +from pathlib import Path +import sys +import random +import math + + +def load_log(): + if Path('cache').exists(): + with open('cache', 'r') as f: + return eval(f.read()) + + total = {} + + for file in Path('.').rglob('log-*.log'): + with open(file, 'r') as f: + lines = f.readlines() + + for line in lines: + try: + column_id, num_of_0, num = map(int, line.strip().split()) + total[column_id] = total.get(column_id, {}) + total[column_id][num_of_0] = total[column_id].get(num_of_0, []) + if len(total[column_id][num_of_0]) < 10: + total[column_id][num_of_0].append(num) + except ValueError as e: + print(f'Warning {file}: {e}', file=sys.stderr) + print(f' Line: {line.strip()}', file=sys.stderr) + + with open('cache', 'w') as f: + f.write(str(total)) + + return total + + +def main(): + if len(sys.argv) != 2: + print(f'Usage: {sys.argv[0]} ', file=sys.stderr) + sys.exit(1) + + target_score = int(sys.argv[1]) + + total = load_log() + + num_grids = 16 * 8 + avg_score_per_grid = round(math.sqrt((target_score * 16 / num_grids))) + print(f'{avg_score_per_grid = }') + print(f'target score passed by user = {target_score}') + print(f'expected target score = {(avg_score_per_grid ** 2) * num_grids // 16 + 1}') + + for col in range(16): + for _ in range(8): + candidates = total[col][avg_score_per_grid] + num = random.choice(candidates) + candidates.remove(num) + + # print(f'{col} {num} (num_of_0 = {avg_score_per_grid})') + print(f'cargo run -- --send-game-data --col {col} --data-int {num}') + + +if __name__ == '__main__': + main() diff --git a/sw/secret-mining/common.h b/sw/secret-mining/common.h new file mode 100644 index 0000000..bbb55df --- /dev/null +++ b/sw/secret-mining/common.h @@ -0,0 +1,36 @@ +#pragma once +#include +#include +#ifdef __linux__ +#include +#endif // __linux__ + +constexpr bool isLittleEndian() { return __BYTE_ORDER == __LITTLE_ENDIAN; } + +constexpr uint64_t rotateLeft(uint64_t x, unsigned n) { + assert(n < 64); + int reverse = 64 - n; + return (x << n) | (x >> reverse); +} + +constexpr void _swap(uint8_t &v1, uint8_t &v2) { + // std::swap becomes constexpr only in C++20. + uint8_t tmp = v1; + v1 = v2; + v2 = tmp; +} + +constexpr uint64_t toLittleEndian(uint64_t val) { + // use such coding to omit NVCC warning +#if __BYTE_ORDER == __LITTLE_ENDIAN + return val; +#else + uint8_t *ptr = reinterpret_cast(&val); + + _swap(ptr[0], ptr[7]); + _swap(ptr[1], ptr[6]); + _swap(ptr[2], ptr[5]); + _swap(ptr[3], ptr[4]); + return val; +#endif +} diff --git a/sw/secret-mining/gen-cache.py b/sw/secret-mining/gen-cache.py new file mode 100644 index 0000000..0731c49 --- /dev/null +++ b/sw/secret-mining/gen-cache.py @@ -0,0 +1,34 @@ +from pathlib import Path +import sys + + +def load_log(): + # if Path('cache').exists(): + # with open('cache', 'r') as f: + # return eval(f.read()) + + total = {} + + for file in Path('.').rglob('log-*.log'): + with open(file, 'r') as f: + lines = f.readlines() + + for line in lines: + try: + column_id, num_of_0, num = map(int, line.strip().split()) + total[column_id] = total.get(column_id, {}) + total[column_id][num_of_0] = total[column_id].get(num_of_0, []) + if len(total[column_id][num_of_0]) < 10: + total[column_id][num_of_0].append(num) + except ValueError as e: + print(f'Warning {file}: {e}', file=sys.stderr) + print(f' Line: {line.strip()}', file=sys.stderr) + + with open('cache', 'w') as f: + f.write(str(total)) + + return total + + +if __name__ == '__main__': + load_log() diff --git a/sw/secret-mining/keccak.cc b/sw/secret-mining/keccak.cc new file mode 100644 index 0000000..f5817c0 --- /dev/null +++ b/sw/secret-mining/keccak.cc @@ -0,0 +1,399 @@ +// Taken from https://github.com/brainhub/SHA3IUF +// Slightly modified for this project. +// See license below. + +/* ------------------------------------------------------------------------- + * Works when compiled for either 32-bit or 64-bit targets, optimized for + * 64 bit. + * + * Canonical implementation of Init/Update/Finalize for SHA-3 byte input. + * + * SHA3-256, SHA3-384, SHA-512 are implemented. SHA-224 can easily be added. + * + * Based on code from http://keccak.noekeon.org/ . + * + * I place the code that I wrote into public domain, free to use. + * + * I would appreciate if you give credits to this work if you used it to + * write or test * your code. + * + * Aug 2015. Andrey Jivsov. crypto@brainhub.org + * ---------------------------------------------------------------------- */ + +#include "keccak.h" + +#include +#include +#include + +#define SHA3_ASSERT(x) +#define SHA3_TRACE(format, ...) +#define SHA3_TRACE_BUF(format, buf, l) + +/* + * This flag is used to configure "pure" Keccak, as opposed to NIST SHA3. + */ +#define SHA3_USE_KECCAK_FLAG 0x80000000 +#define SHA3_CW(x) ((x) & (~SHA3_USE_KECCAK_FLAG)) + +#if defined(_MSC_VER) +#define SHA3_CONST(x) x +#else +#define SHA3_CONST(x) x##L +#endif + +#ifndef SHA3_ROTL64 +#define SHA3_ROTL64(x, y) \ + (((x) << (y)) | ((x) >> ((sizeof(uint64_t) * 8) - (y)))) +#endif + +static const uint64_t keccakf_rndc[24] = { + SHA3_CONST(0x0000000000000001UL), SHA3_CONST(0x0000000000008082UL), + SHA3_CONST(0x800000000000808aUL), SHA3_CONST(0x8000000080008000UL), + SHA3_CONST(0x000000000000808bUL), SHA3_CONST(0x0000000080000001UL), + SHA3_CONST(0x8000000080008081UL), SHA3_CONST(0x8000000000008009UL), + SHA3_CONST(0x000000000000008aUL), SHA3_CONST(0x0000000000000088UL), + SHA3_CONST(0x0000000080008009UL), SHA3_CONST(0x000000008000000aUL), + SHA3_CONST(0x000000008000808bUL), SHA3_CONST(0x800000000000008bUL), + SHA3_CONST(0x8000000000008089UL), SHA3_CONST(0x8000000000008003UL), + SHA3_CONST(0x8000000000008002UL), SHA3_CONST(0x8000000000000080UL), + SHA3_CONST(0x000000000000800aUL), SHA3_CONST(0x800000008000000aUL), + SHA3_CONST(0x8000000080008081UL), SHA3_CONST(0x8000000000008080UL), + SHA3_CONST(0x0000000080000001UL), SHA3_CONST(0x8000000080008008UL)}; + +static const unsigned keccakf_rotc[24] = {1, 3, 6, 10, 15, 21, 28, 36, + 45, 55, 2, 14, 27, 41, 56, 8, + 25, 43, 62, 18, 39, 61, 20, 44}; + +static const unsigned keccakf_piln[24] = {10, 7, 11, 17, 18, 3, 5, 16, + 8, 21, 24, 4, 15, 23, 19, 13, + 12, 2, 20, 14, 22, 9, 6, 1}; + +/* generally called after SHA3_KECCAK_SPONGE_WORDS-ctx->capacityWords words + * are XORed into the state s + */ +void keccakf(uint64_t s[25]) { + size_t i, j, round; + uint64_t t, bc[5]; + + for (round = 0; round < KECCAK_ROUNDS; round++) { + /* Theta */ + for (i = 0; i < 5; i++) + bc[i] = s[i] ^ s[i + 5] ^ s[i + 10] ^ s[i + 15] ^ s[i + 20]; + + for (i = 0; i < 5; i++) { + t = bc[(i + 4) % 5] ^ SHA3_ROTL64(bc[(i + 1) % 5], 1); + for (j = 0; j < 25; j += 5) s[j + i] ^= t; + } + + /* Rho Pi */ + t = s[1]; + for (i = 0; i < 24; i++) { + j = keccakf_piln[i]; + bc[0] = s[j]; + s[j] = SHA3_ROTL64(t, keccakf_rotc[i]); + t = bc[0]; + } + + /* Chi */ + for (j = 0; j < 25; j += 5) { + for (i = 0; i < 5; i++) bc[i] = s[j + i]; + for (i = 0; i < 5; i++) s[j + i] ^= (~bc[(i + 1) % 5]) & bc[(i + 2) % 5]; + } + + /* Iota */ + s[0] ^= keccakf_rndc[round]; + } +} + +// This is exactly the same as keccakf except that the caller is expected to +// call keccakf_split() variant exactly KECCAK_ROUNDS time, with round = 0 to +// KECCAK_ROUNDS-1. +void keccakf_split(uint64_t s[25], int round) { + int i, j; + uint64_t t, bc[5]; + /* Theta */ + for (i = 0; i < 5; i++) + bc[i] = s[i] ^ s[i + 5] ^ s[i + 10] ^ s[i + 15] ^ s[i + 20]; + + for (i = 0; i < 5; i++) { + t = bc[(i + 4) % 5] ^ SHA3_ROTL64(bc[(i + 1) % 5], 1); + for (j = 0; j < 25; j += 5) s[j + i] ^= t; + } + + /* Rho Pi */ + t = s[1]; + for (i = 0; i < 24; i++) { + j = keccakf_piln[i]; + bc[0] = s[j]; + s[j] = SHA3_ROTL64(t, keccakf_rotc[i]); + t = bc[0]; + } + + /* Chi */ + for (j = 0; j < 25; j += 5) { + for (i = 0; i < 5; i++) bc[i] = s[j + i]; + for (i = 0; i < 5; i++) s[j + i] ^= (~bc[(i + 1) % 5]) & bc[(i + 2) % 5]; + } + + /* Iota */ + s[0] ^= keccakf_rndc[round]; +} + +/* *************************** Public Inteface ************************ */ + +/* For Init or Reset call these: */ +sha3_return_t sha3_Init(void *priv, unsigned bitSize) { + sha3_context *ctx = (sha3_context *)priv; + if (bitSize != 256 && bitSize != 384 && bitSize != 512) + return SHA3_RETURN_BAD_PARAMS; + memset(ctx, 0, sizeof(*ctx)); + ctx->capacityWords = 2 * bitSize / (8 * sizeof(uint64_t)); + return SHA3_RETURN_OK; +} + +void sha3_Init256(void *priv) { sha3_Init(priv, 256); } + +void sha3_Init384(void *priv) { sha3_Init(priv, 384); } + +void sha3_Init512(void *priv) { sha3_Init(priv, 512); } + +enum SHA3_FLAGS sha3_SetFlags(void *priv, enum SHA3_FLAGS flags) { + sha3_context *ctx = (sha3_context *)priv; + flags = static_cast(flags & SHA3_FLAGS_KECCAK); + ctx->capacityWords |= (flags == SHA3_FLAGS_KECCAK ? SHA3_USE_KECCAK_FLAG : 0); + return flags; +} + +void sha3_UpdateWord(void *priv, void const *bufIn) { + sha3_context *ctx = (sha3_context *)priv; + + /* Ensure that byteIndex is reset to 0 (since we're processing a full word) */ + if (ctx->byteIndex != 0) { + return; // Error: byteIndex should be 0 when updating a full word + } + + const uint64_t *t = reinterpret_cast(bufIn); + + ctx->u.s[ctx->wordIndex] ^= *t; + if (++ctx->wordIndex == + (SHA3_KECCAK_SPONGE_WORDS - SHA3_CW(ctx->capacityWords))) { + keccakf(ctx->u.s); + ctx->wordIndex = 0; + } +} + +void sha3_Update(void *priv, void const *bufIn, size_t len) { + sha3_context *ctx = (sha3_context *)priv; + + if (len % 8 == 0 && ctx->byteIndex == 0) { + // The input length is a multiple of 8 bytes and byteIndex is 0, + // which means we are processing complete words from bufIn. + size_t words = len / sizeof(uint64_t); + const uint64_t *buf = reinterpret_cast(bufIn); + + for (size_t i = 0; i < words; i++, buf++) { + sha3_UpdateWord(ctx, buf); + } + return; + } + + /* 0...7 -- how much is needed to have a word */ + unsigned old_tail = (8 - ctx->byteIndex) & 7; + + size_t words; + unsigned tail; + size_t i; + + const uint8_t *buf = reinterpret_cast(bufIn); + + SHA3_TRACE_BUF("called to update with:", buf, len); + + SHA3_ASSERT(ctx->byteIndex < 8); + SHA3_ASSERT(ctx->wordIndex < sizeof(ctx->u.s) / sizeof(ctx->u.s[0])); + + if (len < old_tail) { /* have no complete word or haven't started + * the word yet */ + SHA3_TRACE("because %d<%d, store it and return", (unsigned)len, + (unsigned)old_tail); + /* endian-independent code follows: */ + while (len--) + ctx->saved |= (uint64_t)(*(buf++)) << ((ctx->byteIndex++) * 8); + SHA3_ASSERT(ctx->byteIndex < 8); + return; + } + + if (old_tail) { /* will have one word to process */ + SHA3_TRACE("completing one word with %d bytes", (unsigned)old_tail); + /* endian-independent code follows: */ + len -= old_tail; + while (old_tail--) + ctx->saved |= (uint64_t)(*(buf++)) << ((ctx->byteIndex++) * 8); + + /* now ready to add saved to the sponge */ + ctx->u.s[ctx->wordIndex] ^= ctx->saved; + SHA3_ASSERT(ctx->byteIndex == 8); + ctx->byteIndex = 0; + ctx->saved = 0; + if (++ctx->wordIndex == + (SHA3_KECCAK_SPONGE_WORDS - SHA3_CW(ctx->capacityWords))) { + keccakf(ctx->u.s); + ctx->wordIndex = 0; + } + } + + /* now work in full words directly from input */ + + SHA3_ASSERT(ctx->byteIndex == 0); + + words = len / sizeof(uint64_t); + tail = len - words * sizeof(uint64_t); + + SHA3_TRACE("have %d full words to process", (unsigned)words); + + for (i = 0; i < words; i++, buf += sizeof(uint64_t)) { + const uint64_t t = + (uint64_t)(buf[0]) | ((uint64_t)(buf[1]) << 8 * 1) | + ((uint64_t)(buf[2]) << 8 * 2) | ((uint64_t)(buf[3]) << 8 * 3) | + ((uint64_t)(buf[4]) << 8 * 4) | ((uint64_t)(buf[5]) << 8 * 5) | + ((uint64_t)(buf[6]) << 8 * 6) | ((uint64_t)(buf[7]) << 8 * 7); +#if defined(__x86_64__) || defined(__i386__) + SHA3_ASSERT(memcmp(&t, buf, 8) == 0); +#endif + ctx->u.s[ctx->wordIndex] ^= t; + if (++ctx->wordIndex == + (SHA3_KECCAK_SPONGE_WORDS - SHA3_CW(ctx->capacityWords))) { + keccakf(ctx->u.s); + ctx->wordIndex = 0; + } + } + + SHA3_TRACE("have %d bytes left to process, save them", (unsigned)tail); + + /* finally, save the partial word */ + SHA3_ASSERT(ctx->byteIndex == 0 && tail < 8); + while (tail--) { + SHA3_TRACE("Store byte %02x '%c'", *buf, *buf); + ctx->saved |= (uint64_t)(*(buf++)) << ((ctx->byteIndex++) * 8); + } + SHA3_ASSERT(ctx->byteIndex < 8); + SHA3_TRACE("Have saved=0x%016" PRIx64 " at the end", ctx->saved); +} + +void const *sha3_Finalize_split(void *priv, size_t round) { + sha3_context *ctx = (sha3_context *)priv; + + if (round == 0) { + /* Append 2-bit suffix 01, per SHA-3 spec. */ + uint64_t t; + + if (ctx->capacityWords & SHA3_USE_KECCAK_FLAG) { + /* Keccak version */ + t = (uint64_t)(((uint64_t)1) << (ctx->byteIndex * 8)); + } else { + /* SHA3 version */ + t = (uint64_t)(((uint64_t)(0x02 | (1 << 2))) << ((ctx->byteIndex) * 8)); + } + + ctx->u.s[ctx->wordIndex] ^= ctx->saved ^ t; + + /* Prepare for the final round */ + ctx->u.s[SHA3_KECCAK_SPONGE_WORDS - SHA3_CW(ctx->capacityWords) - 1] ^= + SHA3_CONST(0x8000000000000000UL); + } else if (round <= KECCAK_ROUNDS) { + /* Perform KECCAK_ROUNDS of keccakf */ + keccakf_split(ctx->u.s, round - 1); + } else { /* round = KECCAK_ROUNDS + 1 */ + /* Convert the context state to bytes after the final round */ + for (unsigned i = 0; i < SHA3_KECCAK_SPONGE_WORDS; i++) { + const unsigned t1 = (uint32_t)ctx->u.s[i]; + const unsigned t2 = (uint32_t)((ctx->u.s[i] >> 16) >> 16); + ctx->u.sb[i * 8 + 0] = (uint8_t)(t1); + ctx->u.sb[i * 8 + 1] = (uint8_t)(t1 >> 8); + ctx->u.sb[i * 8 + 2] = (uint8_t)(t1 >> 16); + ctx->u.sb[i * 8 + 3] = (uint8_t)(t1 >> 24); + ctx->u.sb[i * 8 + 4] = (uint8_t)(t2); + ctx->u.sb[i * 8 + 5] = (uint8_t)(t2 >> 8); + ctx->u.sb[i * 8 + 6] = (uint8_t)(t2 >> 16); + ctx->u.sb[i * 8 + 7] = (uint8_t)(t2 >> 24); + } + } + + return (round == KECCAK_ROUNDS + 1) ? ctx->u.sb : nullptr; +} + +/* This is simply the 'update' with the padding block. + * The padding block is 0x01 || 0x00* || 0x80. First 0x01 and last 0x80 + * bytes are always present, but they can be the same byte. + */ +void const *sha3_Finalize(void *priv) { + sha3_context *ctx = (sha3_context *)priv; + + SHA3_TRACE("called with %d bytes in the buffer", ctx->byteIndex); + + /* Append 2-bit suffix 01, per SHA-3 spec. Instead of 1 for padding we + * use 1<<2 below. The 0x02 below corresponds to the suffix 01. + * Overall, we feed 0, then 1, and finally 1 to start padding. Without + * M || 01, we would simply use 1 to start padding. */ + + uint64_t t; + + if (ctx->capacityWords & SHA3_USE_KECCAK_FLAG) { + /* Keccak version */ + t = (uint64_t)(((uint64_t)1) << (ctx->byteIndex * 8)); + } else { + /* SHA3 version */ + t = (uint64_t)(((uint64_t)(0x02 | (1 << 2))) << ((ctx->byteIndex) * 8)); + } + + ctx->u.s[ctx->wordIndex] ^= ctx->saved ^ t; + + ctx->u.s[SHA3_KECCAK_SPONGE_WORDS - SHA3_CW(ctx->capacityWords) - 1] ^= + SHA3_CONST(0x8000000000000000UL); + keccakf(ctx->u.s); + + /* Return first bytes of the ctx->s. This conversion is not needed for + * little-endian platforms e.g. wrap with #if !defined(__BYTE_ORDER__) + * || !defined(__ORDER_LITTLE_ENDIAN__) || + * __BYTE_ORDER__!=__ORDER_LITTLE_ENDIAN__ + * ... the conversion below ... + * #endif */ + { + unsigned i; + for (i = 0; i < SHA3_KECCAK_SPONGE_WORDS; i++) { + const unsigned t1 = (uint32_t)ctx->u.s[i]; + const unsigned t2 = (uint32_t)((ctx->u.s[i] >> 16) >> 16); + ctx->u.sb[i * 8 + 0] = (uint8_t)(t1); + ctx->u.sb[i * 8 + 1] = (uint8_t)(t1 >> 8); + ctx->u.sb[i * 8 + 2] = (uint8_t)(t1 >> 16); + ctx->u.sb[i * 8 + 3] = (uint8_t)(t1 >> 24); + ctx->u.sb[i * 8 + 4] = (uint8_t)(t2); + ctx->u.sb[i * 8 + 5] = (uint8_t)(t2 >> 8); + ctx->u.sb[i * 8 + 6] = (uint8_t)(t2 >> 16); + ctx->u.sb[i * 8 + 7] = (uint8_t)(t2 >> 24); + } + } + + SHA3_TRACE_BUF("Hash: (first 32 bytes)", ctx->u.sb, 256 / 8); + + return (ctx->u.sb); +} + +sha3_return_t sha3_HashBuffer(unsigned bitSize, enum SHA3_FLAGS flags, + const void *in, unsigned inBytes, void *out, + unsigned outBytes) { + sha3_return_t err; + sha3_context c; + + err = sha3_Init(&c, bitSize); + if (err != SHA3_RETURN_OK) return err; + if (sha3_SetFlags(&c, flags) != flags) { + return SHA3_RETURN_BAD_PARAMS; + } + sha3_Update(&c, in, inBytes); + const void *h = sha3_Finalize(&c); + + if (outBytes > bitSize / 8) outBytes = bitSize / 8; + memcpy(out, h, outBytes); + return SHA3_RETURN_OK; +} diff --git a/sw/secret-mining/keccak.h b/sw/secret-mining/keccak.h new file mode 100644 index 0000000..59842bc --- /dev/null +++ b/sw/secret-mining/keccak.h @@ -0,0 +1,95 @@ +#ifndef SHA3_H +#define SHA3_H + +// Taken from https://github.com/brainhub/SHA3IUF +// Slightly modified for this project. +// See license below. + +#include +#include + +/* ------------------------------------------------------------------------- + * Works when compiled for either 32-bit or 64-bit targets, optimized for + * 64 bit. + * + * Canonical implementation of Init/Update/Finalize for SHA-3 byte input. + * + * SHA3-256, SHA3-384, SHA-512 are implemented. SHA-224 can easily be added. + * + * Based on code from http://keccak.noekeon.org/ . + * + * I place the code that I wrote into public domain, free to use. + * + * I would appreciate if you give credits to this work if you used it to + * write or test * your code. + * + * Aug 2015. Andrey Jivsov. crypto@brainhub.org + * ---------------------------------------------------------------------- */ + +/* 'Words' here refers to uint64_t */ +#define SHA3_KECCAK_SPONGE_WORDS \ + (((1600) / 8 /*bits to byte*/) / sizeof(uint64_t)) +typedef struct sha3_context_ { + uint64_t saved; /* the portion of the input message that we + * didn't consume yet */ + union { /* Keccak's state */ + uint64_t s[SHA3_KECCAK_SPONGE_WORDS]; + uint8_t sb[SHA3_KECCAK_SPONGE_WORDS * 8]; + } u; + unsigned byteIndex; /* 0..7--the next byte after the set one + * (starts from 0; 0--none are buffered) */ + unsigned wordIndex; /* 0..24--the next word to integrate input + * (starts from 0) */ + unsigned capacityWords; /* the double size of the hash output in + * words (e.g. 16 for Keccak 512) */ +} sha3_context; + +enum SHA3_FLAGS { SHA3_FLAGS_NONE = 0, SHA3_FLAGS_KECCAK = 1 }; + +enum SHA3_RETURN { SHA3_RETURN_OK = 0, SHA3_RETURN_BAD_PARAMS = 1 }; +typedef enum SHA3_RETURN sha3_return_t; + +/* For Init or Reset call these: */ +sha3_return_t sha3_Init(void *priv, unsigned bitSize); + +// Takes 16ms on STM32@12MHz, should be split. +void keccakf(uint64_t s[25]); +// Takes 0.7ms on STM32@12MHz, should only be called once per task. +// round should be sequentially called with [0, KECCAK_ROUNDS-1], for a total of +// KECCAK_ROUNDS times. +void keccakf_split(uint64_t s[25], int round); + +void sha3_Init256(void *priv); +void sha3_Init384(void *priv); +void sha3_Init512(void *priv); + +enum SHA3_FLAGS sha3_SetFlags(void *priv, enum SHA3_FLAGS); + +// Takes up to 1 keccakf() call. +void sha3_UpdateWord(void *priv, void const *bufIn); + +// Takes multiple keccakf() call and should not be used on STM32. +void sha3_Update(void *priv, void const *bufIn, size_t len); + +// Takes 1 keccakf() call. +void const *sha3_Finalize(void *priv); + +// This is exactly the same as sha3_Finalize() but the caller is expected +// to call this KECCAK_ROUNDS+2 times, each time giving round = 0 to +// KECCAK_ROUNDS+1. +// For round=0, it should handle anything before keccakf(). +// For round=1 to KECCAK_ROUNDS, it should call keccakf(round-1). +// For round=KECCAK_ROUNDS+1, it should handle anything after keccakf(). +void const *sha3_Finalize_split(void *priv, int round); + +/* Single-call hashing */ +sha3_return_t sha3_HashBuffer( + unsigned bitSize, /* 256, 384, 512 */ + enum SHA3_FLAGS flags, /* SHA3_FLAGS_NONE or SHA3_FLAGS_KECCAK */ + const void *in, unsigned inBytes, void *out, + unsigned outBytes); /* up to bitSize/8; truncation OK */ + +constexpr size_t SHA3_256_HASH_SIZE = 256 / 8; +constexpr size_t KECCAK_ROUNDS = 24; + +#endif diff --git a/sw/secret-mining/sha3_cpu.cc b/sw/secret-mining/sha3_cpu.cc new file mode 100644 index 0000000..4f771a5 --- /dev/null +++ b/sw/secret-mining/sha3_cpu.cc @@ -0,0 +1,217 @@ +#include "sha3_cpu.h" + +#include + +#include +#include + +#include "common.h" + +namespace { +constexpr size_t idx(size_t x) { return x % 5; } + +constexpr size_t idx(size_t x, size_t y) { return idx(x) + 5 * idx(y); } + +// Array of indicies and rotation for P and Pi phases. +constexpr std::array, 24> g_ppi_aux = { + {{6, 44}, {12, 43}, {18, 21}, {24, 14}, {3, 28}, {9, 20}, + {10, 3}, {16, 45}, {22, 61}, {1, 1}, {7, 6}, {13, 25}, + {19, 8}, {20, 18}, {4, 27}, {5, 36}, {11, 10}, {17, 15}, + {23, 56}, {2, 62}, {8, 55}, {14, 39}, {15, 41}, {21, 2}}}; + +// Values, required for the last iota phase. +constexpr std::array g_iota_aux = { + 0x0000000000000001L, 0x0000000000008082L, 0x800000000000808aL, + 0x8000000080008000L, 0x000000000000808bL, 0x0000000080000001L, + 0x8000000080008081L, 0x8000000000008009L, 0x000000000000008aL, + 0x0000000000000088L, 0x0000000080008009L, 0x000000008000000aL, + 0x000000008000808bL, 0x800000000000008bL, 0x8000000000008089L, + 0x8000000000008003L, 0x8000000000008002L, 0x8000000000000080L, + 0x000000000000800aL, 0x800000008000000aL, 0x8000000080008081L, + 0x8000000000008080L, 0x0000000080000001L, 0x8000000080008008L}; + +void updateState(uint64_t A[25]) { + for (int round = 0; round < 24; ++round) { + // Thetta phase + uint64_t C[25]; + for (size_t x = 0; x < 5; x++) { + C[x] = A[idx(x, 0)] ^ A[idx(x, 1)] ^ A[idx(x, 2)] ^ A[idx(x, 3)] ^ + A[idx(x, 4)]; + } + + for (size_t x = 0; x < 5; ++x) { + uint64_t D = C[idx(x + 5 - 1)] ^ rotateLeft(C[idx(x + 1)], 1); + for (int y = 0; y < 5; ++y) { + A[idx(x, y)] ^= D; + } + } + + // P and Pi phases + // First element remains the same. + C[0] = A[0]; + for (size_t i = 0; i < 24; ++i) { + C[i + 1] = rotateLeft(A[g_ppi_aux[i].first], g_ppi_aux[i].second); + } + + // Ksi phase + for (size_t x = 0; x < 5; ++x) { + for (size_t y = 0; y < 5; ++y) { + A[idx(x, y)] = C[idx(x, y)] ^ (~C[idx(x + 1, y)] & C[idx(x + 2, y)]); + } + } + + // Iota phase + A[0] ^= g_iota_aux[round]; + } +} + +void processSingleBlock(uint64_t A[25], const uint8_t *data, size_t size) { + assert(size % 8 == 0); + for (unsigned int i = 0, ei = size / 8; i < ei; ++i) { + A[i] ^= toLittleEndian(reinterpret_cast(data)[i]); + } + + updateState(A); +} + +void addPadding(uint8_t *begin, uint8_t *end) { + if (std::next(begin) == end) { + *begin = 0x86; + return; + } + + *begin++ = 0x06; + *--end = 0x80; + std::fill(begin, end, 0); +} + +void copyLittleEndian64(const uint64_t A[25], uint8_t *data, size_t size) { + if (isLittleEndian()) { + // Help the compiler to recognize a simple memcpy. + const uint8_t *A8 = reinterpret_cast(A); + std::copy(A8, A8 + size, data); + return; + } + + uint64_t *data64 = reinterpret_cast(data); + for (; size >= 8; size -= 8, ++A, ++data64) { + *data64 = toLittleEndian(*A); + } + const uint8_t *A8 = reinterpret_cast(A); + uint8_t *data8 = reinterpret_cast(data64); + std::copy(A8, A8 + size, data8); +} + +} // namespace + +SHA3_cpu::SHA3_cpu(size_t block) + : m_digestSize(block / 8), m_bufferSize(200 - 2 * m_digestSize), + m_blockBuffer(static_cast( + aligned_alloc(sizeof(uint64_t), m_bufferSize))) { + assert(m_digestSize * 8 == block); + init(); +} + +void SHA3_cpu::init() { + std::fill(std::begin(m_A), std::end(m_A), uint64_t(0)); + m_bufferOffset = 0; + + m_finished = false; +} + +void SHA3_cpu::add(const uint8_t *data, size_t sz) { + assert(!m_finished && "Init should be called"); + while (sz != 0) { + if (sz < m_bufferSize - m_bufferOffset) { + std::copy(data, data + sz, m_blockBuffer.get() + m_bufferOffset); + m_bufferOffset += sz; + return; + } + + if (m_bufferOffset == 0) { + processBlock(data); + sz -= m_bufferSize; + data += m_bufferSize; + continue; + } + + size_t dataSize = m_bufferSize - m_bufferOffset; + std::copy(data, data + dataSize, m_blockBuffer.get() + m_bufferOffset); + processBlock(m_blockBuffer.get()); + m_bufferOffset = 0; + sz -= dataSize; + data += dataSize; + } +} + +void SHA3_cpu::finish() { + addPadding(m_blockBuffer.get() + m_bufferOffset, + m_blockBuffer.get() + m_bufferSize); + processBlock(m_blockBuffer.get()); + m_bufferOffset = 0; +} + +std::vector SHA3_cpu::digest() { + if (!m_finished) { + finish(); + m_finished = true; + } + std::vector result(m_digestSize); + copyLittleEndian64(m_A, result.data(), result.size()); + return result; +} + +void SHA3_cpu::processBlock(const uint8_t *buf) { + processSingleBlock(m_A, buf, m_bufferSize); +} + +SHA3_cpu_batch::SHA3_cpu_batch(size_t block) : m_digestSize(block / 8) { + assert(m_digestSize * 8 == block); + unsigned threads = omp_get_num_procs(); + threads = threads == 0 ? 2 : threads; + m_states.resize(threads); + for (auto &val : m_states) { + val.blockBuffer.reset(new uint8_t[200 - 2 * m_digestSize]); + } +} + +std::vector SHA3_cpu_batch::calculate( + const std::vector> &datas) { + auto result = prepareResult(datas.size()); +#pragma omp parallel num_threads(m_states.size()) + { + int tid = omp_get_thread_num(); + auto &state = m_states[tid]; + size_t blockSize = 200 - 2 * m_digestSize; + int nthreads = omp_get_num_threads(); + for (size_t i = tid; i < datas.size(); i += nthreads) { + std::fill(std::begin(state.A), std::end(state.A), uint64_t(0)); + size_t sizeLeft = datas[i].second; + const uint8_t *data = datas[i].first; + + while (true) { + if (sizeLeft < blockSize) { + std::copy(data, data + sizeLeft, state.blockBuffer.get()); + addPadding(state.blockBuffer.get() + sizeLeft, + state.blockBuffer.get() + blockSize); + processSingleBlock(state.A, state.blockBuffer.get(), blockSize); + copyLittleEndian64(state.A, result[i].data(), m_digestSize); + break; + } + processSingleBlock(state.A, data, blockSize); + data += blockSize; + sizeLeft -= blockSize; + } + } + } + return result; +} + +std::vector SHA3_cpu_batch::prepareResult(size_t size) { + std::vector result; + result.reserve(size); + for (size_t i = 0; i < size; ++i) { + result.emplace_back(m_digestSize); + } + return result; +} diff --git a/sw/secret-mining/sha3_cpu.h b/sw/secret-mining/sha3_cpu.h new file mode 100644 index 0000000..116c61f --- /dev/null +++ b/sw/secret-mining/sha3_cpu.h @@ -0,0 +1,52 @@ +#pragma once +#include +#include +#include +#include + +class SHA3_cpu { + public: + SHA3_cpu(size_t block); + void init(); + void add(const uint8_t *data, size_t sz); + + std::vector digest(); + + private: + // Argument buf should be at least m_buffer_size. + void processBlock(const uint8_t *buf); + + void finish(); + + private: + uint64_t m_A[25]; // State array. + size_t m_digestSize = 0; + + size_t m_bufferSize = 0; + std::unique_ptr m_blockBuffer; + size_t m_bufferOffset = 0; + + bool m_finished = false; +}; + +class SHA3_cpu_batch { + public: + using Digest = std::vector; + + SHA3_cpu_batch(size_t block); + + std::vector calculate( + const std::vector> &datas); + size_t batchSize() const { return m_states.size(); } + + private: + std::vector prepareResult(size_t size); + + private: + size_t m_digestSize = 0; + struct State { + uint64_t A[25]; + std::unique_ptr blockBuffer; + }; + std::vector m_states; +}; diff --git a/sw/secret-mining/trim.py b/sw/secret-mining/trim.py new file mode 100644 index 0000000..719e912 --- /dev/null +++ b/sw/secret-mining/trim.py @@ -0,0 +1,54 @@ +from collections import Counter, defaultdict +from pathlib import Path +import sys + +def main(): + if len(sys.argv) < 2: + print(f'Usage: {sys.argv[0]} files', file=sys.stderr) + sys.exit(1) + + total = defaultdict(Counter) + + for file in Path('.').rglob('log-*.log'): + with open(file, 'r') as f: + lines = f.readlines() + + for line in lines: + try: + column_id, num_of_0, num = map(int, line.strip().split()) + total[column_id][num_of_0] += 1 + except ValueError as e: + print(f'Warning {file}: {e}', file=sys.stderr) + print(f' Line: {line.strip()}', file=sys.stderr) + + for file in sys.argv[1:]: + with open(file, 'r') as f: + lines = f.readlines() + + out = [] + + for line in lines: + try: + column_id, num_of_0, num = map(int, line.strip().split()) + + if total[column_id][num_of_0] > 65536: + total[column_id][num_of_0] -= 1 + else: + out.append(line) + + except ValueError as e: + print(f'Warning {file}: {e}', file=sys.stderr) + print(f' Line: {line.strip()}', file=sys.stderr) + + with open(file, 'w') as f: + f.writelines(out) + + # print the result as a table + print('', 'col', *range(16), sep='\t') + print('count') + for num_of_0 in range(256): + print(f'{num_of_0}', '', *(total[column_id][num_of_0] for column_id in range(16)), sep='\t') + + +if __name__ == '__main__': + main()