1/*
2 * RIPEMD-160 implementation. Copyright (c) 2015, magnum
3 * This software is hereby released to the general public under
4 * the following terms: Redistribution and use in source and binary
5 * forms, with or without modification, are permitted.
6 *
7 * KEYLEN  should be PLAINTEXT_LENGTH for passwords or 20 for hash
8 * OUTLEN  should be sizeof(outbuffer->v)
9 * SALTLEN should be sizeof(currentsalt.salt)
10 */
11
12#include "opencl_misc.h"
13#include "opencl_ripemd.h"
14#define AES_SRC_TYPE __constant
15#define AES_DST_TYPE __global
16#include "opencl_aes.h"
17
18#define ITERATIONS 2000
19
20typedef struct {
21	uint length;
22	uchar v[KEYLEN];
23} pbkdf2_password;
24
25typedef struct {
26	uint v[16 / 4];
27} tc_hash;
28
29typedef struct {
30	uint salt[SALTLEN / 4];
31	uint bin[(512 - 64) / 4];
32} tc_salt;
33
34#define RIPEMD160_DIGEST_LENGTH 20
35
36inline void preproc(__global const uchar *key, uint keylen, uint *state,
37                    uint padding)
38{
39	uint i;
40	uint W[16];
41
42	for (i = 0; i < 16; i++)
43		W[i] = padding;
44
45	for (i = 0; i < keylen; i++)
46		XORCHAR(W, i, key[i]);
47
48	state[0] = INIT_A;
49	state[1] = INIT_B;
50	state[2] = INIT_C;
51	state[3] = INIT_D;
52	state[4] = INIT_E;
53
54	ripemd160(W, state);
55}
56
57inline void hmac_ripemd160(uint *output, uint *ipad_state, uint *opad_state,
58                           __constant uint *salt, uchar add)
59{
60	uint i;
61	uint W[16] = { 0 };
62
63	for (i = 0; i < 5; i++)
64		output[i] = ipad_state[i];
65
66	for (i = 0; i < 16; i++)
67		W[i] = salt[i];
68
69	ripemd160(W, output);
70
71	W[0] = add << 24;
72	W[1] = 0x80;
73	for (i = 2; i < 14; i++)
74		W[i] = 0;
75	W[14] = (64 + SALTLEN + 4) << 3;
76	W[15] = 0;
77
78	ripemd160(W, output);
79
80	for (i = 0; i < 5; i++)
81		W[i] = output[i];
82
83	for (i = 0; i < 5; i++)
84		output[i] = opad_state[i];
85
86	ripemd160_160Z(W, output);
87}
88
89inline void big_hmac_ripemd160(uint *input, uint inputlen, uint *ipad_state,
90                               uint *opad_state, uint *tmp_out)
91{
92	uint i;
93	uint W[5];
94
95	for (i = 0; i < 5; i++)
96		W[i] = input[i];
97
98	for (i = 1; i < ITERATIONS; i++) {
99		uint ctx[5];
100		uint j;
101
102		for (j = 0; j < 5; j++)
103			ctx[j] = ipad_state[j];
104
105		ripemd160_160Z(W, ctx);
106
107		for (j = 0; j < 5; j++)
108			W[j] = ctx[j];
109
110		for (j = 0; j < 5; j++)
111			ctx[j] = opad_state[j];
112
113		ripemd160_160Z(W, ctx);
114
115		for (j = 0; j < 5; j++)
116			W[j] = ctx[j];
117
118		for (j = 0; j < 5; j++)
119			tmp_out[j] ^= ctx[j];
120	}
121}
122
123inline void pbkdf2(__global const uchar *pass, uint passlen,
124                   __constant uint *salt, uint *out)
125{
126	uint ipad_state[5];
127	uint opad_state[5];
128	uint r, t = 0;
129
130	preproc(pass, passlen, ipad_state, 0x36363636);
131	preproc(pass, passlen, opad_state, 0x5c5c5c5c);
132
133	for (r = 1; r <= (OUTLEN + 19) / 20; r++) {
134		uint tmp_out[5];
135		uint i;
136
137		hmac_ripemd160(tmp_out, ipad_state, opad_state, salt, r);
138
139		big_hmac_ripemd160(tmp_out, RIPEMD160_DIGEST_LENGTH,
140		                   ipad_state, opad_state,
141		                   tmp_out);
142
143		for (i = 0; i < 20 && t < (OUTLEN + 3) / 4 * 4; i++, t++)
144			PUTCHAR(out, t, ((uchar*)tmp_out)[i]);
145	}
146}
147
148__kernel void tc_ripemd_aesxts(__global const pbkdf2_password *inbuffer,
149                               __global tc_hash *outbuffer,
150                               __constant tc_salt *salt)
151{
152	uint idx = get_global_id(0);
153	uint key[64 / 4];
154
155	pbkdf2(inbuffer[idx].v, inbuffer[idx].length, salt->salt, key);
156
157	AES_256_XTS_first_sector(salt->bin, outbuffer[idx].v, (uchar*)key);
158}
159