diff options
Diffstat (limited to 'xmrstak/backend/nvidia/nvcc_code/cuda_core.cu')
| -rw-r--r-- | xmrstak/backend/nvidia/nvcc_code/cuda_core.cu | 25 |
1 files changed, 19 insertions, 6 deletions
diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 3bb910e..4eacfb6 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 || ALGO == cryptonight_ipbc) + if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { uint32_t * state = d_ctx_state + thread * 50; tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8); @@ -275,12 +275,21 @@ __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 || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { const uint32_t table = 0x75310U; - const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); - const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; - storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) + { + const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2); + const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; + storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + } + else if(ALGO == cryptonight_stellite) + { + const uint32_t index = ((z >> 27) & 12) | ((z >> 23) & 2); + const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24; + storeGlobal32( long_state + j, sub == 2 ? fork_7 : z ); + } } else storeGlobal32( long_state + j, z ); @@ -304,7 +313,7 @@ __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 || ALGO == cryptonight_ipbc) + if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) { const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res; uint32_t long_state_update = sub2 ? tweaked_res : res; @@ -503,5 +512,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_ipbc>(ctx, startNonce); } + else if(miner_algo == cryptonight_stellite) + { + cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_stellite>(ctx, startNonce); + } } |
