Skip to content

Commit

Permalink
Add Cryptonight-fast
Browse files Browse the repository at this point in the history
  • Loading branch information
@gnock
gnock committed Jun 5, 2018
1 parent 4f34bd1 commit 81437eb
Show file tree
Hide file tree
Showing 8 changed files with 89 additions and 34 deletions.
2 changes: 1 addition & 1 deletion xmrstak/backend/amd/amd_gpu/gpu.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1004,7 +1004,7 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
return(ERR_OCL_API);
}

if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite)
if(miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon || miner_algo == cryptonight_ipbc || miner_algo == cryptonight_stellite || miner_algo == cryptonight_fast)
{
// Input
if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)
Expand Down
20 changes: 10 additions & 10 deletions xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -553,8 +553,8 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad,

__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states, ulong Threads
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7)
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
, __global ulong *input
#endif
)
Expand All @@ -574,8 +574,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
}

barrier(CLK_LOCAL_MEM_FENCE);
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7)
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
uint2 tweak1_2;
#endif
uint4 b_x;
Expand All @@ -599,8 +599,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
b[1] = states[3] ^ states[7];

b_x = ((uint4 *)b)[0];
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7)
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
tweak1_2 = as_uint2(input[4]);
tweak1_2.s0 >>= 24;
tweak1_2.s0 |= tweak1_2.s1 << 8;
Expand All @@ -627,8 +627,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
((uint4 *)c)[0] = AES_Round(AES0, AES1, AES2, AES3, ((uint4 *)c)[0], ((uint4 *)a)[0]);

b_x ^= ((uint4 *)c)[0];
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7)
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)
uint table = 0x75310U;
// cryptonight_stellite
# if(ALGO == 7)
Expand All @@ -646,8 +646,8 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
a[1] += c[0] * as_ulong2(tmp).s0;
a[0] += mul_hi(c[0], as_ulong2(tmp).s0);

// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7)
// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc || cryptonight_stellite || cryptonight_fast
#if(ALGO == 3 || ALGO == 5 || ALGO == 6 || ALGO == 7 || ALGO == 8)

# if(ALGO == 6)
uint2 ipbc_tmp = tweak1_2 ^ ((uint2 *)&(a[0]))[0];
Expand Down
32 changes: 16 additions & 16 deletions xmrstak/backend/cpu/crypto/cryptonight_aesni.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -432,7 +432,7 @@ inline void cryptonight_monero_tweak(uint64_t* mem_out, __m128i tmp)

uint8_t x = static_cast<uint8_t>(vh >> 24);
static const uint16_t table = 0x7531;
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_fast)
{
const uint8_t index = (((x >> 3) & 6) | (x & 1)) << 1;
vh ^= ((table >> index) & 0x3) << 28;
Expand All @@ -456,7 +456,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();

if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43)
if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43)
{
memset(output, 0, 32);
return;
Expand All @@ -465,7 +465,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
keccak((const uint8_t *)input, len, ctx0->hash_state, 200);

uint64_t monero_const;
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
monero_const = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
monero_const ^= *(reinterpret_cast<const uint64_t*>(ctx0->hash_state) + 24);
Expand Down Expand Up @@ -494,7 +494,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
else
cx = _mm_aesenc_si128(cx, _mm_set_epi64x(ah0, al0));

if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
cryptonight_monero_tweak<ALGO>((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
else
_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
Expand All @@ -518,7 +518,7 @@ void cryptonight_hash(const void* input, size_t len, void* output, cryptonight_c
_mm_prefetch((const char*)&l0[al0 & MASK], _MM_HINT_T0);
ah0 += lo;

if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
if(ALGO == cryptonight_ipbc)
((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const ^ ((uint64_t*)&l0[idx0 & MASK])[0];
Expand Down Expand Up @@ -561,7 +561,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();

if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43)
if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43)
{
memset(output, 0, 64);
return;
Expand All @@ -571,7 +571,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
keccak((const uint8_t *)input+len, len, ctx[1]->hash_state, 200);

uint64_t monero_const_0, monero_const_1;
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
monero_const_0 = *reinterpret_cast<const uint64_t*>(reinterpret_cast<const uint8_t*>(input) + 35);
monero_const_0 ^= *(reinterpret_cast<const uint64_t*>(ctx[0]->hash_state) + 24);
Expand Down Expand Up @@ -609,7 +609,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
else
cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh0, axl0));

if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
cryptonight_monero_tweak<ALGO>((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
else
_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
Expand All @@ -627,7 +627,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
else
cx = _mm_aesenc_si128(cx, _mm_set_epi64x(axh1, axl1));

if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
cryptonight_monero_tweak<ALGO>((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
else
_mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
Expand All @@ -648,7 +648,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
axh0 += lo;
((uint64_t*)&l0[idx0 & MASK])[0] = axl0;

if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
if(ALGO == cryptonight_ipbc)
((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0 ^ ((uint64_t*)&l0[idx0 & MASK])[0];
Expand Down Expand Up @@ -684,7 +684,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
axh1 += lo;
((uint64_t*)&l1[idx1 & MASK])[0] = axl1;

if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
if(ALGO == cryptonight_ipbc)
((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1 ^ ((uint64_t*)&l1[idx1 & MASK])[0];
Expand Down Expand Up @@ -736,7 +736,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
else \
c = _mm_aesenc_si128(c, a); \
b = _mm_xor_si128(b, c); \
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) \
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) \
cryptonight_monero_tweak<ALGO>((uint64_t*)ptr, b); \
else \
_mm_store_si128(ptr, b);\
Expand All @@ -751,7 +751,7 @@ void cryptonight_double_hash(const void* input, size_t len, void* output, crypto
#define CN_STEP4(a, b, c, l, mc, ptr, idx) \
lo = _umul128(idx, _mm_cvtsi128_si64(b), &hi); \
a = _mm_add_epi64(a, _mm_set_epi64x(lo, hi)); \
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) \
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) \
{ \
_mm_store_si128(ptr, _mm_xor_si128(a, mc)); \
if (ALGO == cryptonight_ipbc) \
Expand Down Expand Up @@ -782,7 +782,7 @@ void cryptonight_triple_hash(const void* input, size_t len, void* output, crypto
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();

if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43)
if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43)
{
memset(output, 0, 32 * 3);
return;
Expand Down Expand Up @@ -876,7 +876,7 @@ void cryptonight_quad_hash(const void* input, size_t len, void* output, cryptoni
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();

if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43)
if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43)
{
memset(output, 0, 32 * 4);
return;
Expand Down Expand Up @@ -985,7 +985,7 @@ void cryptonight_penta_hash(const void* input, size_t len, void* output, crypton
constexpr size_t ITERATIONS = cn_select_iter<ALGO>();
constexpr size_t MEM = cn_select_memory<ALGO>();

if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite) && len < 43)
if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast) && len < 43)
{
memset(output, 0, 32 * 5);
return;
Expand Down
33 changes: 32 additions & 1 deletion xmrstak/backend/cpu/minethd.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -292,6 +292,9 @@ bool minethd::self_test()
else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_stellite)
{
}
else if(::jconf::inst()->GetCurrentCoinSelection().GetDescription(1).GetMiningAlgo() == cryptonight_fast)
{
}
for (int i = 0; i < MAX_N; i++)
cryptonight_free_ctx(ctx[i]);

Expand Down Expand Up @@ -377,6 +380,9 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
case cryptonight_stellite:
algv = 6;
break;
case cryptonight_fast:
algv = 7;
break;
default:
algv = 2;
break;
Expand Down Expand Up @@ -410,7 +416,11 @@ minethd::cn_hash_fun minethd::func_selector(bool bHaveAes, bool bNoPrefetch, xmr
cryptonight_hash<cryptonight_stellite, false, false>,
cryptonight_hash<cryptonight_stellite, true, false>,
cryptonight_hash<cryptonight_stellite, false, true>,
cryptonight_hash<cryptonight_stellite, true, true>
cryptonight_hash<cryptonight_stellite, true, true>,
cryptonight_hash<cryptonight_fast, false, false>,
cryptonight_hash<cryptonight_fast, true, false>,
cryptonight_hash<cryptonight_fast, false, true>,
cryptonight_hash<cryptonight_fast, true, true>
};

std::bitset<2> digit;
Expand Down Expand Up @@ -555,6 +565,9 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
case cryptonight_stellite:
algv = 6;
break;
case cryptonight_fast:
algv = 7;
break;
default:
algv = 2;
break;
Expand Down Expand Up @@ -679,6 +692,24 @@ minethd::cn_hash_fun_multi minethd::func_multi_selector(size_t N, bool bHaveAes,
cryptonight_penta_hash<cryptonight_stellite, true, false>,
cryptonight_penta_hash<cryptonight_stellite, false, true>,
cryptonight_penta_hash<cryptonight_stellite, true, true>,


cryptonight_double_hash<cryptonight_fast, false, false>,
cryptonight_double_hash<cryptonight_fast, true, false>,
cryptonight_double_hash<cryptonight_fast, false, true>,
cryptonight_double_hash<cryptonight_fast, true, true>,
cryptonight_triple_hash<cryptonight_fast, false, false>,
cryptonight_triple_hash<cryptonight_fast, true, false>,
cryptonight_triple_hash<cryptonight_fast, false, true>,
cryptonight_triple_hash<cryptonight_fast, true, true>,
cryptonight_quad_hash<cryptonight_fast, false, false>,
cryptonight_quad_hash<cryptonight_fast, true, false>,
cryptonight_quad_hash<cryptonight_fast, false, true>,
cryptonight_quad_hash<cryptonight_fast, true, true>,
cryptonight_penta_hash<cryptonight_fast, false, false>,
cryptonight_penta_hash<cryptonight_fast, true, false>,
cryptonight_penta_hash<cryptonight_fast, false, true>,
cryptonight_penta_hash<cryptonight_fast, true, true>,
};

std::bitset<2> digit;
Expand Down
18 changes: 17 additions & 1 deletion xmrstak/backend/cryptonight.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,8 @@ enum xmrstak_algo
cryptonight_heavy = 4,
cryptonight_aeon = 5,
cryptonight_ipbc = 6, // equal to cryptonight_aeon with a small tweak in the miner code
cryptonight_stellite = 7 //equal to cryptonight_monero but with one tiny change
cryptonight_stellite = 7, //equal to cryptonight_monero but with one tiny change
cryptonight_fast = 8 //equal to cryptonight_monero but with less iterations, used by masari
};

// define aeon settings
Expand All @@ -28,6 +29,8 @@ constexpr size_t CRYPTONIGHT_HEAVY_MEMORY = 4 * 1024 * 1024;
constexpr uint32_t CRYPTONIGHT_HEAVY_MASK = 0x3FFFF0;
constexpr uint32_t CRYPTONIGHT_HEAVY_ITER = 0x40000;

constexpr uint32_t CRYPTONIGHT_FAST_ITER = 0x40000;

template<xmrstak_algo ALGO>
inline constexpr size_t cn_select_memory() { return 0; }

Expand All @@ -52,13 +55,17 @@ inline constexpr size_t cn_select_memory<cryptonight_ipbc>() { return CRYPTONIGH
template<>
inline constexpr size_t cn_select_memory<cryptonight_stellite>() { return CRYPTONIGHT_MEMORY; }

template<>
inline constexpr size_t cn_select_memory<cryptonight_fast>() { return CRYPTONIGHT_MEMORY; }


inline size_t cn_select_memory(xmrstak_algo algo)
{
switch(algo)
{
case cryptonight_stellite:
case cryptonight_monero:
case cryptonight_fast:
case cryptonight:
return CRYPTONIGHT_MEMORY;
case cryptonight_ipbc:
Expand Down Expand Up @@ -96,12 +103,16 @@ inline constexpr uint32_t cn_select_mask<cryptonight_ipbc>() { return CRYPTONIGH
template<>
inline constexpr uint32_t cn_select_mask<cryptonight_stellite>() { return CRYPTONIGHT_MASK; }

template<>
inline constexpr uint32_t cn_select_mask<cryptonight_fast>() { return CRYPTONIGHT_MASK; }

inline size_t cn_select_mask(xmrstak_algo algo)
{
switch(algo)
{
case cryptonight_stellite:
case cryptonight_monero:
case cryptonight_fast:
case cryptonight:
return CRYPTONIGHT_MASK;
case cryptonight_ipbc:
Expand Down Expand Up @@ -139,6 +150,9 @@ inline constexpr uint32_t cn_select_iter<cryptonight_ipbc>() { return CRYPTONIGH
template<>
inline constexpr uint32_t cn_select_iter<cryptonight_stellite>() { return CRYPTONIGHT_ITER; }

template<>
inline constexpr uint32_t cn_select_iter<cryptonight_fast>() { return CRYPTONIGHT_FAST_ITER; }

inline size_t cn_select_iter(xmrstak_algo algo)
{
switch(algo)
Expand All @@ -153,6 +167,8 @@ inline size_t cn_select_iter(xmrstak_algo algo)
return CRYPTONIGHT_LITE_ITER;
case cryptonight_heavy:
return CRYPTONIGHT_HEAVY_ITER;
case cryptonight_fast:
return CRYPTONIGHT_FAST_ITER;
default:
return 0;
}
Expand Down
12 changes: 8 additions & 4 deletions xmrstak/backend/nvidia/nvcc_code/cuda_core.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 || ALGO == cryptonight_stellite)
if (ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
uint32_t * state = d_ctx_state + thread * 50;
tweak1_2[0] = (d_input[8] >> 24) | (d_input[9] << 8);
Expand Down Expand Up @@ -275,10 +275,10 @@ __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 || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
const uint32_t table = 0x75310U;
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_fast)
{
const uint32_t index = ((z >> 26) & 12) | ((z >> 23) & 2);
const uint32_t fork_7 = z ^ ((table >> index) & 0x30U) << 24;
Expand Down Expand Up @@ -312,7 +312,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 || ALGO == cryptonight_stellite)
if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc || ALGO == cryptonight_stellite || ALGO == cryptonight_fast)
{
const uint32_t tweaked_res = tweak1_2[sub & 1] ^ res;
uint32_t long_state_update = sub2 ? tweaked_res : res;
Expand Down Expand Up @@ -515,5 +515,9 @@ void cryptonight_core_cpu_hash(nvid_ctx* ctx, xmrstak_algo miner_algo, uint32_t
{
cryptonight_core_gpu_hash<CRYPTONIGHT_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_stellite>(ctx, startNonce);
}
else if(miner_algo == cryptonight_fast)
{
cryptonight_core_gpu_hash<CRYPTONIGHT_FAST_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_fast>(ctx, startNonce);
}

}
Loading

0 comments on commit 81437eb

Please sign in to comment.