|
|
@ -1,4 +1,4 @@ |
|
|
|
// DiaKGCN 23-02-2012 - OpenCL kernel by Diapolo |
|
|
|
// DiaKGCN 24-02-2012 - OpenCL kernel by Diapolo |
|
|
|
// |
|
|
|
// |
|
|
|
// Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3. |
|
|
|
// Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3. |
|
|
|
// The kernel was rewritten by me (Diapolo) and is still public-domain! |
|
|
|
// The kernel was rewritten by me (Diapolo) and is still public-domain! |
|
|
@ -53,8 +53,12 @@ __kernel |
|
|
|
u V[8]; |
|
|
|
u V[8]; |
|
|
|
u W[16]; |
|
|
|
u W[16]; |
|
|
|
|
|
|
|
|
|
|
|
#if defined(VECTORS2) || defined(VECTORS4) || defined(VECTORS8) |
|
|
|
#ifdef VECTORS8 |
|
|
|
const u nonce = (uint)(get_local_id(0)) * (uint)(vec_step(u)) + (uint)(get_group_id(0)) * (uint)(WORKVEC) + base; |
|
|
|
const u nonce = (uint)(get_local_id(0)) * 8U + (uint)(get_group_id(0)) * (uint)(WORKVEC) + base; |
|
|
|
|
|
|
|
#elif defined VECTORS4 |
|
|
|
|
|
|
|
const u nonce = (uint)(get_local_id(0)) * 4U + (uint)(get_group_id(0)) * (uint)(WORKVEC) + base; |
|
|
|
|
|
|
|
#elif defined VECTORS2 |
|
|
|
|
|
|
|
const u nonce = (uint)(get_local_id(0)) * 2U + (uint)(get_group_id(0)) * (uint)(WORKVEC) + base; |
|
|
|
#else |
|
|
|
#else |
|
|
|
const u nonce = (uint)(get_local_id(0)) + (uint)(get_group_id(0)) * (uint)(WORKSIZE) + base; |
|
|
|
const u nonce = (uint)(get_local_id(0)) + (uint)(get_group_id(0)) * (uint)(WORKSIZE) + base; |
|
|
|
#endif |
|
|
|
#endif |
|
|
@ -113,8 +117,8 @@ __kernel |
|
|
|
//---------------------------------------------------------------------------------- |
|
|
|
//---------------------------------------------------------------------------------- |
|
|
|
|
|
|
|
|
|
|
|
#ifdef VECTORS8 |
|
|
|
#ifdef VECTORS8 |
|
|
|
W[0] = PreW18 + (u)( rotr25(nonce.s0), rotr25(nonce.s0) ^ 0x2004000U, rotr25(nonce.s0) ^ 0x4008000U, rotr25(nonce.s0) ^ 0x600c000U, |
|
|
|
W[0] = PreW18 + (u)( rotr25(nonce.s0), rotr25(nonce.s0) ^ 0x2004000U, rotr25(nonce.s0) ^ 0x4008000U, rotr25(nonce.s0) ^ 0x600c000U, |
|
|
|
rotr25(nonce.s0) ^ 0x8010000U, rotr25(nonce.s0) ^ 0xa014000U, rotr25(nonce.s0) ^ 0xc018000U, rotr25(nonce.s0) ^ 0xe01c000U); |
|
|
|
rotr25(nonce.s0) ^ 0x8010000U, rotr25(nonce.s0) ^ 0xa014000U, rotr25(nonce.s0) ^ 0xc018000U, rotr25(nonce.s0) ^ 0xe01c000U); |
|
|
|
#elif defined VECTORS4 |
|
|
|
#elif defined VECTORS4 |
|
|
|
W[0] = PreW18 + (u)(rotr25(nonce.x), rotr25(nonce.x) ^ 0x2004000U, rotr25(nonce.x) ^ 0x4008000U, rotr25(nonce.x) ^ 0x600c000U); |
|
|
|
W[0] = PreW18 + (u)(rotr25(nonce.x), rotr25(nonce.x) ^ 0x2004000U, rotr25(nonce.x) ^ 0x4008000U, rotr25(nonce.x) ^ 0x600c000U); |
|
|
|
#elif defined VECTORS2 |
|
|
|
#elif defined VECTORS2 |
|
|
|