Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions .clang-format
Original file line number Diff line number Diff line change
@@ -0,0 +1,4 @@
---
BasedOnStyle: WebKit
...

37 changes: 21 additions & 16 deletions src/work.cl
Original file line number Diff line number Diff line change
Expand Up @@ -10,9 +10,9 @@ enum Blake2b_IV {
};

enum IV_Derived {
nano_xor_iv0 = 0x6a09e667f2bdc900UL, // iv1 ^ 0x1010000 ^ outlen
nano_xor_iv4 = 0x510e527fade682f9UL, // iv4 ^ inbytes
nano_xor_iv6 = 0xe07c265404be4294UL, // iv6 ^ ~0
nano_xor_iv0 = 0x6a09e667f2bdc900UL, // iv1 ^ 0x1010000 ^ outlen
nano_xor_iv4 = 0x510e527fade682f9UL, // iv4 ^ inbytes
nano_xor_iv6 = 0xe07c265404be4294UL, // iv6 ^ ~0
};

#ifdef cl_amd_media_ops
Expand Down Expand Up @@ -55,22 +55,27 @@ static inline ulong rotr64(ulong x, int shift)
G32(m0, m1, m2, m3, vv[a / 2], vb1, vb2, vv[c / 2], vd1, vd2)

#define ROUND(m0, m1, m2, m3, m4, m5, m6, m7, m8, m9, m10, m11, m12, m13, m14, \
m15) \
m15) \
do { \
G2v(m0, m1, m2, m3, 0, 4, 8, 12); \
G2v(m4, m5, m6, m7, 2, 6, 10, 14); \
G2v_split(m8, m9, m10, m11, 0, vv[5 / 2].s1, vv[6 / 2].s0, 10, \
vv[15 / 2].s1, vv[12 / 2].s0); \
vv[15 / 2].s1, vv[12 / 2].s0); \
G2v_split(m12, m13, m14, m15, 2, vv[7 / 2].s1, vv[4 / 2].s0, 8, \
vv[13 / 2].s1, vv[14 / 2].s0); \
vv[13 / 2].s1, vv[14 / 2].s0); \
} while (0)

static inline ulong blake2b(ulong const nonce, __constant ulong *h)
static inline ulong blake2b(ulong const nonce, __constant ulong* h)
{
ulong2 vv[8] = {
{nano_xor_iv0, iv1}, {iv2, iv3}, {iv4, iv5},
{iv6, iv7}, {iv0, iv1}, {iv2, iv3},
{nano_xor_iv4, iv5}, {nano_xor_iv6, iv7},
{ nano_xor_iv0, iv1 },
{ iv2, iv3 },
{ iv4, iv5 },
{ iv6, iv7 },
{ iv0, iv1 },
{ iv2, iv3 },
{ nano_xor_iv4, iv5 },
{ nano_xor_iv6, iv7 },
};

ROUND(nonce, h[0], h[1], h[2], h[3], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
Expand All @@ -93,12 +98,12 @@ static inline ulong blake2b(ulong const nonce, __constant ulong *h)
#undef G2v_split
#undef ROUND

__kernel void nano_work(__constant uchar *attempt,
__global uchar *result_a,
__constant uchar *item_a,
const ulong difficulty)
__kernel void nano_work(__constant uchar* attempt,
__global uchar* result_a,
__constant uchar* item_a,
const ulong difficulty)
{
const ulong attempt_l = *((__constant ulong *) attempt) + get_global_id(0);
const ulong attempt_l = *((__constant ulong*)attempt) + get_global_id(0);
if (blake2b(attempt_l, item_a) >= difficulty)
*((__global ulong *) result_a) = attempt_l;
*((__global ulong*)result_a) = attempt_l;
}