aboutsummaryrefslogtreecommitdiff
path: root/cl/onion.cl
diff options
context:
space:
mode:
Diffstat (limited to 'cl/onion.cl')
-rw-r--r--cl/onion.cl179
1 files changed, 23 insertions, 156 deletions
diff --git a/cl/onion.cl b/cl/onion.cl
index 2972617..4200859 100644
--- a/cl/onion.cl
+++ b/cl/onion.cl
@@ -1,6 +1,6 @@
+#define INITIAL_DATA_LEN 9
#define SHA_CHUNK_LEN 64
#define ROL(x, shamt) ((x << shamt) | (x >> (sizeof(x)*8 - shamt)))
-#define MIN(a, b) ((a) < (b) ? (a) : (b))
struct sha_data {
unsigned int a;
@@ -13,18 +13,6 @@ struct sha_data {
unsigned char data[SHA_CHUNK_LEN];
};
-void memcpy(void *restrict dest, void *restrict src, int len) {
- unsigned char *dest_ = (unsigned char*)dest;
- unsigned char *src_ = (unsigned char*)src;
- int i = 0;
- for (i = 0; i < len; i++) {
- dest_[i] = src_[i];
- }
-// while (len-- >= 0) {
-// dest_[len] = src_[len];
-// }
-}
-
void sha_chunk(unsigned char (*buf)[SHA_CHUNK_LEN], struct sha_data *sha) {
unsigned int w[80] = {0};
unsigned int new_a = 0;
@@ -101,70 +89,20 @@ void sha_chunk(unsigned char (*buf)[SHA_CHUNK_LEN], struct sha_data *sha) {
sha->e += e;
}
-void sha_update(struct sha_data *c, void *data, unsigned int size) {
- unsigned int i = 0;
- size_t remaining = size;
- unsigned char *bdata = (unsigned char*)data;
-
-
- size_t count = MIN(size, SHA_CHUNK_LEN - c->data_len);
- for (i = 0; i < count; i++)
- c->data[c->data_len+i] = ((char*)data)[i];
- //memcpy(&(c->data[c->data_len]), data, count);
- c->data_len += count;
- remaining -= count;
-
-
- while (c->data_len == SHA_CHUNK_LEN) {
- sha_chunk(&(c->data), c);
- count = MIN(remaining, SHA_CHUNK_LEN);
- memcpy(c->data, &bdata[size-remaining], count);
- remaining -= count;
- c->data_len = count;
- }
-
- /* representative of all data throughput, inclusive of the buffer in
- * the context */
- c->len += size;
-}
-
void sha_final(unsigned char *digest, struct sha_data *c) {
size_t i = 0;
- c->data[c->data_len++] = 0x80;
-
- /* Transform byte len to bit len */
- c->len *= 8;
-
- for (i = c->data_len; i < SHA_CHUNK_LEN; i++)
+ #pragma unroll
+ for (i = INITIAL_DATA_LEN+5; i < SHA_CHUNK_LEN-8; i++)
c->data[i] = 0;
- /* still room for the 64-bit message length at the end of this chunk? */
- if (c->data_len + 8 > SHA_CHUNK_LEN) {
- sha_chunk(&(c->data), c);
- for (i = 0; i < SHA_CHUNK_LEN; i++)
- c->data[i] = 0;
- }
-
- /* FIXME loop or leave unrolled? */
- c->data[56] = c->len >> 56;
- c->data[57] = c->len >> 48;
- c->data[58] = c->len >> 40;
- c->data[59] = c->len >> 32;
- c->data[60] = c->len >> 24;
- c->data[61] = c->len >> 16;
- c->data[62] = c->len >> 8;
- c->data[63] = c->len;
-
sha_chunk(&(c->data), c);
- /* FIXME loop or leave unrolled? */
digest[ 0] = c->a >> 24;
digest[ 1] = c->a >> 16;
digest[ 2] = c->a >> 8;
digest[ 3] = c->a;
-
digest[ 4] = c->b >> 24;
digest[ 5] = c->b >> 16;
digest[ 6] = c->b >> 8;
@@ -172,24 +110,9 @@ void sha_final(unsigned char *digest, struct sha_data *c) {
digest[ 8] = c->c >> 24;
digest[ 9] = c->c >> 16;
-
-
-// second half of hash not needed eh?
- digest[10] = c->c >> 8;
- digest[11] = c->c;
-
- digest[12] = c->d >> 24;
- digest[13] = c->d >> 16;
- digest[14] = c->d >> 8;
- digest[15] = c->d;
-
- digest[16] = c->e >> 24;
- digest[17] = c->e >> 16;
- digest[18] = c->e >> 8;
- digest[19] = c->e;
}
-__kernel void fractal_gen(
+__kernel void key_brute(
__global unsigned int *results,
__constant struct sha_data *partial,
__constant unsigned char *search,
@@ -202,41 +125,40 @@ __kernel void fractal_gen(
struct sha_data ctx;
- /* FIXME dummy e (big-endian) */
- unsigned char e[4] = {0x1F, 0xFF, 0xFF, 0xFF};
unsigned char digest[20];
- /* first half of e is our worker number rest is determined later */
- e[0] = ((tx >> 8) & 0xFF);
+ results[tx] = 0;
+
+ /* Data area plus (useless) exponent area, and end bit */
+ #pragma unroll
+ for (j = 0; j < INITIAL_DATA_LEN+5; j++) {
+ ctx.data[j] = partial->data[j];
+ }
+ #pragma unroll
+ for (j = SHA_CHUNK_LEN - 8; j < SHA_CHUNK_LEN; j++) {
+ ctx.data[j] = partial->data[j];
+ }
+
+ ctx.data[INITIAL_DATA_LEN] = tx >> 8;
/* if MSB is 0, then it doesn't need to be stored in the key, so violates
* law of sizeof(e) == 4, messing everything up */
- if (e[0] == 0) {
+ if (ctx.data[INITIAL_DATA_LEN] == 0) {
return;
}
-
- e[1] = tx & 0xFF;
- results[tx] = 0;
+ ctx.data[INITIAL_DATA_LEN + 1] = tx;
for (i = 3; i < 65536; i+=2) {
- e[2] = (i >> 8) & 0xFF;
- e[3] = i & 0xFF;
-
- #pragma unroll
- for (j = 0; j < SHA_CHUNK_LEN; j++) {
- ctx.data[j] = partial->data[j];
- }
ctx.a = partial->a;
ctx.b = partial->b;
ctx.c = partial->c;
ctx.d = partial->d;
ctx.e = partial->e;
- ctx.len = partial->len;
- ctx.data_len = partial->data_len;
-
- sha_update(&ctx, &e, 4);
+//////////////////////////////////////////////////////////////
+ ctx.data[INITIAL_DATA_LEN + 2] = i >> 8;
+ ctx.data[INITIAL_DATA_LEN + 3] = i;
+/////////////////////////////////////////////////////////////
sha_final(&digest, &ctx);
-
int all_clear = 1;
for (j = 0; j < raw_length; j++) {
if (search[j] != digest[j]) {
@@ -250,58 +172,3 @@ __kernel void fractal_gen(
return;
}
-
-
-void unused() {
-/*
-#define R2(w, a, b, c, d, e, i) a = ROL(a, 5) + (b^c^d) + e + w[i] + 0x6ED9EBA1; b = ROL(b, 30);
-
- R2(w, a, b, c, d, e, 20);
- R2(w, e, a, b, c, d, 21);
- R2(w, d, e, a, b, c, 22);
- R2(w, c, d, e, a, b, 23);
- R2(w, b, c, d, e, a, 24);
- R2(w, a, b, c, d, e, 25);
- R2(w, e, a, b, c, d, 26);
- R2(w, d, e, a, b, c, 27);
- R2(w, c, d, e, a, b, 28);
- R2(w, b, c, d, e, a, 29);
- R2(w, a, b, c, d, e, 30);
- R2(w, e, a, b, c, d, 31);
- R2(w, d, e, a, b, c, 32);
- R2(w, c, d, e, a, b, 33);
- R2(w, b, c, d, e, a, 34);
- R2(w, a, b, c, d, e, 35);
- R2(w, e, a, b, c, d, 36);
- R2(w, d, e, a, b, c, 37);
- R2(w, c, d, e, a, b, 38);
- R2(w, b, c, d, e, a, 39);*/
-
-/*
-#define R3(w, a, b, c, d, e, i) a = ROL(a, 5) + ((b&c)|(b&d)|(c&d)) + e + w[i] + 0x8F1BBCDC; b = ROL(b, 30);
-
- R3(w, a, b, c, d, e, 40);
- R3(w, e, a, b, c, d, 41);
- R3(w, d, e, a, b, c, 42);
- R3(w, c, d, e, a, b, 43);
- R3(w, b, c, d, e, a, 44);
- R3(w, a, b, c, d, e, 45);
- R3(w, e, a, b, c, d, 46);
- R3(w, d, e, a, b, c, 47);
- R3(w, c, d, e, a, b, 48);
- R3(w, b, c, d, e, a, 49);
- R3(w, a, b, c, d, e, 50);
- R3(w, e, a, b, c, d, 51);
- R3(w, d, e, a, b, c, 52);
- R3(w, c, d, e, a, b, 53);
- R3(w, b, c, d, e, a, 54);
- R3(w, a, b, c, d, e, 55);
- R3(w, e, a, b, c, d, 56);
- R3(w, d, e, a, b, c, 57);
- R3(w, c, d, e, a, b, 58);
- R3(w, b, c, d, e, a, 59);
-*/
-
-
-
-}
generated by cgit v1.1 at 2024-06-26 12:56:28 +0000