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