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