diff options
| author | fireice-uk <fireice-uk@users.noreply.github.com> | 2018-05-03 10:40:16 +0100 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2018-05-03 10:40:16 +0100 |
| commit | 478a023ff62d1ebccb9547c1a152b7d296c11fb2 (patch) | |
| tree | 3282b5b36f412381a9c3eec11333ccdc766ef355 /xmrstak/backend/nvidia/nvcc_code | |
| parent | 3eaf54b2fad71cad23c304335dc63024d9939171 (diff) | |
| parent | 624b4ca870ef184fa3947c8368964e89647d9447 (diff) | |
| download | xmr-stak-478a023ff62d1ebccb9547c1a152b7d296c11fb2.zip xmr-stak-478a023ff62d1ebccb9547c1a152b7d296c11fb2.tar.gz | |
Merge pull request #1512 from psychocrypt/topic-stelliteFork
support stellite v4 fork
Diffstat (limited to 'xmrstak/backend/nvidia/nvcc_code')
| -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); + } } |
