bitmsghash.cl
1 /* 2 * This is based on the John The Ripper SHA512 code, modified for double SHA512 and for use as a miner in Bitmessage. 3 * This software is originally Copyright (c) 2012 Myrice <qqlddg at gmail dot com> 4 * and it is hereby released to the general public under the following terms: 5 * Redistribution and use in source and binary forms, with or without modification, are permitted. 6 */ 7 8 #ifdef cl_khr_byte_addressable_store 9 #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : disable 10 #endif 11 12 #define uint8_t unsigned char 13 #define uint32_t unsigned int 14 #define uint64_t unsigned long 15 #define SALT_SIZE 0 16 17 #define BINARY_SIZE 8 18 #define FULL_BINARY_SIZE 64 19 20 21 #define PLAINTEXT_LENGTH 72 22 23 #define CIPHERTEXT_LENGTH 128 24 25 26 /// Warning: This version of SWAP64(n) is slow and avoid bugs on AMD GPUs(7970) 27 // #define SWAP64(n) as_ulong(as_uchar8(n).s76543210) 28 29 #define SWAP64(n) \ 30 (((n) << 56) \ 31 | (((n) & 0xff00) << 40) \ 32 | (((n) & 0xff0000) << 24) \ 33 | (((n) & 0xff000000) << 8) \ 34 | (((n) >> 8) & 0xff000000) \ 35 | (((n) >> 24) & 0xff0000) \ 36 | (((n) >> 40) & 0xff00) \ 37 | ((n) >> 56)) 38 39 40 41 #define rol(x,n) ((x << n) | (x >> (64-n))) 42 #define ror(x,n) ((x >> n) | (x << (64-n))) 43 #define Ch(x,y,z) ((x & y) ^ ( (~x) & z)) 44 #define Maj(x,y,z) ((x & y) ^ (x & z) ^ (y & z)) 45 #define Sigma0(x) ((ror(x,28)) ^ (ror(x,34)) ^ (ror(x,39))) 46 #define Sigma1(x) ((ror(x,14)) ^ (ror(x,18)) ^ (ror(x,41))) 47 #define sigma0(x) ((ror(x,1)) ^ (ror(x,8)) ^(x>>7)) 48 #define sigma1(x) ((ror(x,19)) ^ (ror(x,61)) ^(x>>6)) 49 50 51 52 typedef struct { // notice memory align problem 53 uint64_t H[8]; 54 uint32_t buffer[32]; //1024 bits 55 uint32_t buflen; 56 } sha512_ctx; 57 58 typedef struct { 59 uint64_t target; 60 char v[PLAINTEXT_LENGTH+1]; 61 } sha512_key; 62 63 64 /* Macros for reading/writing chars from int32's */ 65 #define PUTCHAR(buf, index, val) (buf)[(index)>>2] = ((buf)[(index)>>2] & ~(0xffU << (((index) & 3) << 3))) + ((val) << (((index) & 3) << 3)) 66 67 68 __constant uint64_t k[] = { 69 0x428a2f98d728ae22UL, 0x7137449123ef65cdUL, 0xb5c0fbcfec4d3b2fUL, 70 0xe9b5dba58189dbbcUL, 71 0x3956c25bf348b538UL, 0x59f111f1b605d019UL, 0x923f82a4af194f9bUL, 72 0xab1c5ed5da6d8118UL, 73 0xd807aa98a3030242UL, 0x12835b0145706fbeUL, 0x243185be4ee4b28cUL, 74 0x550c7dc3d5ffb4e2UL, 75 0x72be5d74f27b896fUL, 0x80deb1fe3b1696b1UL, 0x9bdc06a725c71235UL, 76 0xc19bf174cf692694UL, 77 0xe49b69c19ef14ad2UL, 0xefbe4786384f25e3UL, 0x0fc19dc68b8cd5b5UL, 78 0x240ca1cc77ac9c65UL, 79 0x2de92c6f592b0275UL, 0x4a7484aa6ea6e483UL, 0x5cb0a9dcbd41fbd4UL, 80 0x76f988da831153b5UL, 81 0x983e5152ee66dfabUL, 0xa831c66d2db43210UL, 0xb00327c898fb213fUL, 82 0xbf597fc7beef0ee4UL, 83 0xc6e00bf33da88fc2UL, 0xd5a79147930aa725UL, 0x06ca6351e003826fUL, 84 0x142929670a0e6e70UL, 85 0x27b70a8546d22ffcUL, 0x2e1b21385c26c926UL, 0x4d2c6dfc5ac42aedUL, 86 0x53380d139d95b3dfUL, 87 0x650a73548baf63deUL, 0x766a0abb3c77b2a8UL, 0x81c2c92e47edaee6UL, 88 0x92722c851482353bUL, 89 0xa2bfe8a14cf10364UL, 0xa81a664bbc423001UL, 0xc24b8b70d0f89791UL, 90 0xc76c51a30654be30UL, 91 0xd192e819d6ef5218UL, 0xd69906245565a910UL, 0xf40e35855771202aUL, 92 0x106aa07032bbd1b8UL, 93 0x19a4c116b8d2d0c8UL, 0x1e376c085141ab53UL, 0x2748774cdf8eeb99UL, 94 0x34b0bcb5e19b48a8UL, 95 0x391c0cb3c5c95a63UL, 0x4ed8aa4ae3418acbUL, 0x5b9cca4f7763e373UL, 96 0x682e6ff3d6b2b8a3UL, 97 0x748f82ee5defb2fcUL, 0x78a5636f43172f60UL, 0x84c87814a1f0ab72UL, 98 0x8cc702081a6439ecUL, 99 0x90befffa23631e28UL, 0xa4506cebde82bde9UL, 0xbef9a3f7b2c67915UL, 100 0xc67178f2e372532bUL, 101 0xca273eceea26619cUL, 0xd186b8c721c0c207UL, 0xeada7dd6cde0eb1eUL, 102 0xf57d4f7fee6ed178UL, 103 0x06f067aa72176fbaUL, 0x0a637dc5a2c898a6UL, 0x113f9804bef90daeUL, 104 0x1b710b35131c471bUL, 105 0x28db77f523047d84UL, 0x32caab7b40c72493UL, 0x3c9ebe0a15c9bebcUL, 106 0x431d67c49c100d4cUL, 107 0x4cc5d4becb3e42b6UL, 0x597f299cfc657e2aUL, 0x5fcb6fab3ad6faecUL, 108 0x6c44198c4a475817UL, 109 }; 110 111 112 113 static void setup_ctx(sha512_ctx* ctx, const char * password, uint8_t pass_len) 114 { 115 uint32_t* b32 = ctx->buffer; 116 117 //set password to buffer 118 for (uint32_t i = 0; i < pass_len; i++) { 119 PUTCHAR(b32,i,password[i]); 120 } 121 ctx->buflen = pass_len; 122 123 //append 1 to ctx buffer 124 uint32_t length = ctx->buflen; 125 PUTCHAR(b32, length, 0x80); 126 while((++length & 3) != 0) { 127 PUTCHAR(b32, length, 0); 128 } 129 130 uint32_t* buffer32 = b32+(length>>2); 131 for(uint32_t i = length; i < 128; i+=4) {// append 0 to 128 132 *buffer32++=0; 133 } 134 135 //append length to buffer 136 uint64_t *buffer64 = (uint64_t *)ctx->buffer; 137 buffer64[15] = SWAP64(((uint64_t) ctx->buflen) * 8); 138 } 139 140 inline uint64_t sha512(char* password) 141 { 142 __private sha512_ctx ctx; 143 setup_ctx(&ctx, password, 72); 144 // sha512 main` 145 int i; 146 147 uint64_t a = 0x6a09e667f3bcc908UL; 148 uint64_t b = 0xbb67ae8584caa73bUL; 149 uint64_t c = 0x3c6ef372fe94f82bUL; 150 uint64_t d = 0xa54ff53a5f1d36f1UL; 151 uint64_t e = 0x510e527fade682d1UL; 152 uint64_t f = 0x9b05688c2b3e6c1fUL; 153 uint64_t g = 0x1f83d9abfb41bd6bUL; 154 uint64_t h = 0x5be0cd19137e2179UL; 155 156 __private uint64_t w[16]; 157 158 uint64_t *data = (uint64_t *) ctx.buffer; 159 160 for (i = 0; i < 16; i++) 161 w[i] = SWAP64(data[i]); 162 163 uint64_t t1, t2; 164 for (i = 0; i < 16; i++) { 165 t1 = k[i] + w[i] + h + Sigma1(e) + Ch(e, f, g); 166 t2 = Maj(a, b, c) + Sigma0(a); 167 168 h = g; 169 g = f; 170 f = e; 171 e = d + t1; 172 d = c; 173 c = b; 174 b = a; 175 a = t1 + t2; 176 } 177 178 for (i = 16; i < 80; i++) { 179 180 w[i & 15] =sigma1(w[(i - 2) & 15]) + sigma0(w[(i - 15) & 15]) + w[(i -16) & 15] + w[(i - 7) & 15]; 181 t1 = k[i] + w[i & 15] + h + Sigma1(e) + Ch(e, f, g); 182 t2 = Maj(a, b, c) + Sigma0(a); 183 184 h = g; 185 g = f; 186 f = e; 187 e = d + t1; 188 d = c; 189 c = b; 190 b = a; 191 a = t1 + t2; 192 } 193 194 uint64_t finalhash[8]; 195 196 finalhash[0] = SWAP64(a + 0x6a09e667f3bcc908UL); 197 finalhash[1] = SWAP64(b + 0xbb67ae8584caa73bUL); 198 finalhash[2] = SWAP64(c + 0x3c6ef372fe94f82bUL); 199 finalhash[3] = SWAP64(d + 0xa54ff53a5f1d36f1UL); 200 finalhash[4] = SWAP64(e + 0x510e527fade682d1UL); 201 finalhash[5] = SWAP64(f + 0x9b05688c2b3e6c1fUL); 202 finalhash[6] = SWAP64(g + 0x1f83d9abfb41bd6bUL); 203 finalhash[7] = SWAP64(h + 0x5be0cd19137e2179UL); 204 205 setup_ctx(&ctx, (char*) finalhash, 64); 206 207 a = 0x6a09e667f3bcc908UL; 208 b = 0xbb67ae8584caa73bUL; 209 c = 0x3c6ef372fe94f82bUL; 210 d = 0xa54ff53a5f1d36f1UL; 211 e = 0x510e527fade682d1UL; 212 f = 0x9b05688c2b3e6c1fUL; 213 g = 0x1f83d9abfb41bd6bUL; 214 h = 0x5be0cd19137e2179UL; 215 216 data = (uint64_t *) ctx.buffer; 217 //((uint64_t*)ctx.buffer)[8] = SWAP64((uint64_t)0x80); 218 219 for (i = 0; i < 16; i++) 220 w[i] = SWAP64(data[i]); 221 222 for (i = 0; i < 16; i++) { 223 t1 = k[i] + w[i] + h + Sigma1(e) + Ch(e, f, g); 224 t2 = Maj(a, b, c) + Sigma0(a); 225 226 h = g; 227 g = f; 228 f = e; 229 e = d + t1; 230 d = c; 231 c = b; 232 b = a; 233 a = t1 + t2; 234 } 235 236 for (i = 16; i < 80; i++) { 237 238 w[i & 15] =sigma1(w[(i - 2) & 15]) + sigma0(w[(i - 15) & 15]) + w[(i -16) & 15] + w[(i - 7) & 15]; 239 t1 = k[i] + w[i & 15] + h + Sigma1(e) + Ch(e, f, g); 240 t2 = Maj(a, b, c) + Sigma0(a); 241 242 h = g; 243 g = f; 244 f = e; 245 e = d + t1; 246 d = c; 247 c = b; 248 b = a; 249 a = t1 + t2; 250 } 251 return SWAP64(a + 0x6a09e667f3bcc908UL); 252 } 253 254 __kernel void kernel_sha512(__global const sha512_key *password,__global uint64_t *hash, uint64_t start) 255 { 256 uint64_t idx = get_global_id(0); 257 if (idx == 0 && start == 0) { 258 *hash = 0; 259 } 260 uint64_t winval; 261 262 uint64_t junk[9]; 263 264 __global uint64_t * source = (__global uint64_t*) password->v; 265 for (int i = 1; i < 9; i++) { 266 junk[i] = source[i]; 267 } 268 269 junk[0] = SWAP64(idx + (start)); 270 271 winval = sha512((char*)junk); 272 if (SWAP64(winval) < password->target) { 273 *hash = SWAP64(junk[0]); 274 } 275 } 276