summaryrefslogtreecommitdiffstats
path: root/xmrstak/backend/nvidia/nvcc_code
diff options
context:
space:
mode:
authorpsychocrypt <psychocryptHPC@gmail.com>2018-04-22 21:11:54 +0200
committerTimothy Pearson <tpearson@raptorengineering.com>2018-06-04 21:07:11 +0000
commitf149b57458a05a9d5d67ad2834fa5e6c282fa76b (patch)
treed6d9562f7789770593d85ef10a3551fe299794f7 /xmrstak/backend/nvidia/nvcc_code
parent83db451571df19842d0791cfe11c98103543d84a (diff)
downloadxmr-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.cu19
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);
+ }
}
OpenPOWER on IntegriCloud