1/* 2 * This software is Copyright (c) 2015 Sayantan Datta <std2048 at gmail dot com> 3 * and it is hereby released to the general public under the following terms: 4 * Redistribution and use in source and binary forms, with or without 5 * modification, are permitted. 6 * Based on Solar Designer implementation of DES_bs_b.c in jtr-v1.7.9 7 */ 8 9#include "opencl_DES_kernel_params.h" 10 11#define GET_HASH_0(hash, x, k, bits) \ 12 for (bit = bits; bit < k; bit++) \ 13 hash |= ((((uint)B[bit]) >> x) & 1) << bit; 14 15#define GET_HASH_1(hash, x, k, bits) \ 16 for (bit = bits; bit < k; bit++) \ 17 hash |= ((((uint)B[32 + bit]) >> x) & 1) << bit; 18 19#define OFFSET_TABLE_SIZE hash_chk_params.offset_table_size 20#define HASH_TABLE_SIZE hash_chk_params.hash_table_size 21 22inline void cmp_final(__private unsigned DES_bs_vector *B, 23 __private unsigned DES_bs_vector *binary, 24 __global unsigned int *offset_table, 25 __global unsigned int *hash_table, 26 DES_hash_check_params hash_chk_params, 27 volatile __global uint *hash_ids, 28 volatile __global uint *bitmap_dupe, 29 unsigned int section, 30 unsigned int depth, 31 unsigned int start_bit) 32{ 33 unsigned long hash; 34 unsigned int hash_table_index, t, bit; 35 36 GET_HASH_0(binary[0], depth, 32, start_bit); 37 GET_HASH_1(binary[1], depth, 32, start_bit); 38 39 hash = ((unsigned long)binary[1] << 32) | (unsigned long)binary[0]; 40 hash += (unsigned long)offset_table[hash % OFFSET_TABLE_SIZE]; 41 hash_table_index = hash % HASH_TABLE_SIZE; 42 43 if (hash_table[hash_table_index + HASH_TABLE_SIZE] == binary[1]) 44 if (hash_table[hash_table_index] == binary[0]) 45 if (!(atomic_or(&bitmap_dupe[hash_table_index/32], (1U << (hash_table_index % 32))) & (1U << (hash_table_index % 32)))) { 46 t = atomic_inc(&hash_ids[0]); 47 hash_ids[1 + 2 * t] = (section * 32) + depth; 48 hash_ids[2 + 2 * t] = hash_table_index; 49 } 50} 51 52#define BITMAP_SIZE_BITS hash_chk_params.bitmap_size_bits 53#define BITMAP_SIZE_BITS_LESS_ONE (BITMAP_SIZE_BITS - 1) 54 55__kernel void DES_bs_cmp_high(__global unsigned DES_bs_vector *unchecked_hashes, 56 __global unsigned int *offset_table, 57 __global unsigned int *hash_table, 58 DES_hash_check_params hash_chk_params, 59 volatile __global uint *hash_ids, 60 volatile __global uint *bitmap_dupe, 61 __global uint *bitmaps) { 62 63 int i; 64 unsigned DES_bs_vector B[64]; 65 int section = get_global_id(0); 66 int gws = get_global_size(0); 67 unsigned int value[2] , bit, bitmap_index; 68 69 for (i = 0; i < 64; i++) 70 B[i] = unchecked_hashes[section + i * gws]; 71 72 for (i = 0; i < 32; i++) { 73 value[0] = 0; 74 value[1] = 0; 75 GET_HASH_1(value[1], i, hash_chk_params.cmp_bits, 0); 76 bitmap_index = value[1] & BITMAP_SIZE_BITS_LESS_ONE; 77 bit = (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U; 78 if (bit) 79 cmp_final(B, value, offset_table, hash_table, hash_chk_params, hash_ids, bitmap_dupe, section, i, 0); 80 } 81} 82 83#define num_uncracked_hashes hash_chk_params.num_uncracked_hashes 84 85__kernel void DES_bs_cmp(__global unsigned DES_bs_vector *unchecked_hashes, 86 __global unsigned int *offset_table, 87 __global unsigned int *hash_table, 88 DES_hash_check_params hash_chk_params, 89 volatile __global uint *hash_ids, 90 volatile __global uint *bitmap_dupe, 91 __global int *uncracked_hashes) { 92 93 unsigned DES_bs_vector value[2] , mask, i, bit; 94 unsigned DES_bs_vector B[64]; 95 int section = get_global_id(0); 96 int gws = get_global_size(0); 97 98 for (i = 0; i < 64; i++) 99 B[i] = unchecked_hashes[section + i * gws]; 100 101 for (i = 0; i < num_uncracked_hashes; i++) { 102 103 value[0] = uncracked_hashes[i]; 104 value[1] = uncracked_hashes[i + num_uncracked_hashes]; 105 106 mask = B[0] ^ -(value[0] & 1); 107 108 for (bit = 1; bit < 32; bit++) 109 mask |= B[bit] ^ -((value[0] >> bit) & 1); 110 111 for (; bit < 64; bit += 2) { 112 mask |= B[bit] ^ -((value[1] >> (bit & 0x1F)) & 1); 113 mask |= B[bit + 1] ^ -((value[1] >> ((bit + 1) & 0x1F)) & 1); 114 } 115 116 if (mask != ~0U) { 117 for (mask = 0; mask < 32; mask++) { 118 value[0] = value[1] = 0; 119 cmp_final(B, value, offset_table, hash_table, hash_chk_params, hash_ids, bitmap_dupe, section, mask, 0); 120 } 121 } 122 } 123} 124