diff options
author | psychocrypt <psychocryptHPC@gmail.com> | 2018-04-22 21:11:54 +0200 |
---|---|---|
committer | Timothy Pearson <tpearson@raptorengineering.com> | 2018-06-04 21:07:11 +0000 |
commit | f149b57458a05a9d5d67ad2834fa5e6c282fa76b (patch) | |
tree | d6d9562f7789770593d85ef10a3551fe299794f7 /xmrstak/backend/nvidia/nvcc_code | |
parent | 83db451571df19842d0791cfe11c98103543d84a (diff) | |
download | xmr-stak-f149b57458a05a9d5d67ad2834fa5e6c282fa76b.zip xmr-stak-f149b57458a05a9d5d67ad2834fa5e6c282fa76b.tar.gz |
add support for IPBC coin
- add algorithm `cryptonight_lite_v7_xor`
- update documentation
Diffstat (limited to 'xmrstak/backend/nvidia/nvcc_code')
-rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 19 |
1 files changed, 15 insertions, 4 deletions
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 43740d2..3bb910e 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -231,7 +231,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti uint32_t t1[2], t2[2], res; uint32_t tweak1_2[2]; - if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,7 +275,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti t1[0] = shuffle<4>(sPtr,sub, d[x], 0); const uint32_t z = d[0] ^ d[1]; - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) { const uint32_t table = 0x75310U; const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); @@ -304,10 +304,17 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti res = *( (uint64_t *) t2 ) >> ( sub & 1 ? 32 : 0 ); - if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; - const uint32_t long_state_update = sub2 ? tweaked_res : res; + uint32_t long_state_update = sub2 ? tweaked_res : res; + + if (ALGO == cryptonight_ipbc) + { + uint32_t value = shuffle<4>(sPtr,sub, long_state_update, sub & 1) ^ long_state_update; + long_state_update = sub >= 2 ? value : long_state_update; + } + storeGlobal32( long_state + j, long_state_update ); } else @@ -492,5 +499,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t { cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_aeon>(ctx, startNonce); } + else if(miner_algo == cryptonight_ipbc) + { + cryptonight_core_gpu_hash<CRYPTONIGHT_LITE_ITER, CRYPTONIGHT_LITE_MASK, CRYPTONIGHT_LITE_MEMORY/4, cryptonight_ipbc>(ctx, startNonce); + } } |