diff --git a/README.md b/README.md index 046a930e136541792a7288c66e6eeb4ffc910218..e0b62503cc383bdf5f22c26e6cde3bcbcffdef45 100644 --- a/README.md +++ b/README.md @@ -48,12 +48,14 @@ Besides [Monero](https://getmonero.org), following coins can be mined using this - [QRL](https://theqrl.org) - **[Ryo](https://ryo-currency.com) - Upcoming xmr-stak-gui is sponsored by Ryo** - [TurtleCoin](https://turtlecoin.lol) +- [Plenteum](https://www.plenteum.com/) Ryo currency is a way for us to implement the ideas that we were unable to in Monero. See [here](https://github.com/fireice-uk/cryptonote-speedup-demo/) for details. If your prefered coin is not listed, you can choose one of the following algorithms: - +- 256Kib scratchpad memory + - cryptonight_turtle - 1MiB scratchpad memory - cryptonight_lite - cryptonight_lite_v7 diff --git a/xmrstak/backend/amd/amd_gpu/gpu.cpp b/xmrstak/backend/amd/amd_gpu/gpu.cpp index e0f3bcf1ff1def9b9d2e5fb040abd82c0d53a3c8..4ad4d59e944af5aceb97e9109c126d675744b7f7 100644 --- a/xmrstak/backend/amd/amd_gpu/gpu.cpp +++ b/xmrstak/backend/amd/amd_gpu/gpu.cpp @@ -406,7 +406,7 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_ * this is required if the dev pool is mining monero * but the user tuned there settings for another currency */ - if(miner_algo == cryptonight_monero_v8) + if(miner_algo == cryptonight_monero_v8 || miner_algo == cryptonight_turtle) { if(ctx->memChunk < 2) mem_chunk_exp = 1u << 2; diff --git a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl index 6a3def72c4dbb0c675ab206479d5013f6a1ec82a..f647bcafc1e7f051cd574814b2a9e71158efcbae 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl @@ -538,7 +538,7 @@ __kernel void JOIN(cn0,ALGO)(__global ulong *input, __global uint4 *Scratchpad, R"===( // cryptonight_monero_v8 && NVIDIA -#if(ALGO==11 && defined(__NV_CL_C_VERSION)) +#if((ALGO==11 || ALGO==13) && defined(__NV_CL_C_VERSION)) # define SCRATCHPAD_CHUNK(N) (*(__local uint4*)((__local uchar*)(scratchpad_line) + (idxS ^ (N << 4)))) # define SCRATCHPAD_CHUNK_GLOBAL (*((__global uint16*)(Scratchpad + (IDX((idx0 & 0x1FFFC0U) >> 4))))) #else @@ -556,7 +556,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states ulong a[2]; // cryptonight_monero_v8 -#if(ALGO==11) +#if(ALGO==11 || ALGO==13) ulong b[4]; uint4 b_x[2]; // NVIDIA @@ -571,7 +571,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states __local uint AES0[256], AES1[256]; // cryptonight_monero_v8 -#if(ALGO==11) +#if(ALGO==11 || ALGO==13) # if defined(__clang__) && !defined(__NV_CL_C_VERSION) __local uint RCP[256]; # endif @@ -587,7 +587,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states AES0[i] = tmp; AES1[i] = rotate(tmp, 8U); // cryptonight_monero_v8 -#if(ALGO==11 && (defined(__clang__) && !defined(__NV_CL_C_VERSION))) +#if((ALGO==11 || ALGO==13) && (defined(__clang__) && !defined(__NV_CL_C_VERSION))) RCP[i] = RCP_C[i]; #endif } @@ -622,7 +622,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states b_x[0] = ((uint4 *)b)[0]; // cryptonight_monero_v8 -#if(ALGO==11) +#if(ALGO==11 || ALGO==13) a[1] = states[1] ^ states[5]; b[2] = states[8] ^ states[10]; b[3] = states[9] ^ states[11]; @@ -654,7 +654,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states { ulong c[2]; // cryptonight_monero_v8 && NVIDIA -#if(ALGO==11 && defined(__NV_CL_C_VERSION)) +#if((ALGO==11 || ALGO==13) && defined(__NV_CL_C_VERSION)) uint idxS = idx0 & 0x30U; *scratchpad_line = SCRATCHPAD_CHUNK_GLOBAL; #endif @@ -668,7 +668,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states #endif // cryptonight_monero_v8 -#if(ALGO==11) +#if(ALGO==11 || ALGO==13) { ulong2 chunk1 = as_ulong2(SCRATCHPAD_CHUNK(1)); ulong2 chunk2 = as_ulong2(SCRATCHPAD_CHUNK(2)); @@ -693,7 +693,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states SCRATCHPAD_CHUNK(0) = b_x[0]; idx0 = as_uint2(c[0]).s0 & MASK; // cryptonight_monero_v8 -#elif(ALGO==11) +#elif(ALGO==11 || ALGO==13) SCRATCHPAD_CHUNK(0) = b_x[0] ^ ((uint4 *)c)[0]; # ifdef __NV_CL_C_VERSION // flush shuffled data @@ -712,7 +712,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states uint4 tmp; tmp = SCRATCHPAD_CHUNK(0); // cryptonight_monero_v8 -#if(ALGO==11) +#if(ALGO==11 || ALGO==13) // Use division and square root results from the _previous_ iteration to hide the latency tmp.s0 ^= division_result.s0; tmp.s1 ^= division_result.s1 ^ sqrt_result; @@ -770,7 +770,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states ((uint4 *)a)[0] ^= tmp; // cryptonight_monero_v8 -#if (ALGO == 11) +#if (ALGO == 11 || ALGO==13) # if defined(__NV_CL_C_VERSION) // flush shuffled data SCRATCHPAD_CHUNK_GLOBAL = *scratchpad_line; diff --git a/xmrstak/backend/amd/amd_gpu/opencl/fast_int_math_v2.cl b/xmrstak/backend/amd/amd_gpu/opencl/fast_int_math_v2.cl index c170387b49cfde725b06035fe85812a102c7f1e5..4205a67c36922baa4f3df914d12af8939e39d74c 100644 --- a/xmrstak/backend/amd/amd_gpu/opencl/fast_int_math_v2.cl +++ b/xmrstak/backend/amd/amd_gpu/opencl/fast_int_math_v2.cl @@ -4,7 +4,7 @@ R"===( */ // cryptonight_monero_v8 -#if(ALGO==11) +#if(ALGO==11 || ALGO==13) static const __constant uint RCP_C[256] = { diff --git a/xmrstak/backend/amd/autoAdjust.hpp b/xmrstak/backend/amd/autoAdjust.hpp index 0442fde917d90ba8a8ec72f042ce6c013227ea2f..f2dce7f90dbdec86e4030223830081082d60292f 100644 --- a/xmrstak/backend/amd/autoAdjust.hpp +++ b/xmrstak/backend/amd/autoAdjust.hpp @@ -131,7 +131,8 @@ private: } // check if cryptonight_monero_v8 is selected for the user or dev pool - bool useCryptonight_v8 = std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end(); + bool useCryptonight_v8 = (std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end() || + std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_turtle) != neededAlgorithms.end()); // true for all cryptonight_heavy derivates since we check the user and dev pool bool useCryptonight_heavy = std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_heavy) != neededAlgorithms.end(); @@ -153,6 +154,10 @@ private: if(hashMemSize <= CRYPTONIGHT_LITE_MEMORY) maxThreads *= 2u; + // increase all intensity limits by eight for turtle (*2u shadowed from lite) + if (hashMemSize <= CRYPTONIGHT_TURTLE_MEMORY) + maxThreads *= 4u; + // keep 128MiB memory free (value is randomly chosen) from the max available memory const size_t maxAvailableFreeMem = ctx.freeMem - minFreeMem; diff --git a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h index 06cbe87402b9f4fec445877f3b5adfc82b05c5e1..2218bf88a8700fd7e6fdc55e29a996027a2e7398 100644 --- a/xmrstak/backend/cpu/crypto/cryptonight_aesni.h +++ b/xmrstak/backend/cpu/crypto/cryptonight_aesni.h @@ -545,7 +545,7 @@ inline void set_float_rounding_mode() #define CN_MONERO_V8_SHUFFLE_0(n, l0, idx0, ax0, bx0, bx1) \ /* Shuffle the other 3x16 byte chunks in the current 64-byte cache line */ \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) \ { \ const uint64_t idx1 = idx0 & MASK; \ const __m128i chunk1 = _mm_load_si128((__m128i *)&l0[idx1 ^ 0x10]); \ @@ -558,7 +558,7 @@ inline void set_float_rounding_mode() #define CN_MONERO_V8_SHUFFLE_1(n, l0, idx0, ax0, bx0, bx1, lo, hi) \ /* Shuffle the other 3x16 byte chunks in the current 64-byte cache line */ \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) \ { \ const uint64_t idx1 = idx0 & MASK; \ const __m128i chunk1 = _mm_xor_si128(_mm_load_si128((__m128i *)&l0[idx1 ^ 0x10]), _mm_set_epi64x(lo, hi)); \ @@ -572,7 +572,7 @@ inline void set_float_rounding_mode() } #define CN_MONERO_V8_DIV(n, cx, sqrt_result, division_result_xmm, cl) \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) \ { \ uint64_t sqrt_result_tmp; \ assign(sqrt_result_tmp, sqrt_result); \ @@ -627,7 +627,7 @@ inline void set_float_rounding_mode() idx0 = h0[0] ^ h0[4]; \ ax0 = _mm_set_epi64x(h0[1] ^ h0[5], idx0); \ bx0 = _mm_set_epi64x(h0[3] ^ h0[7], h0[2] ^ h0[6]); \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) \ { \ bx1 = _mm_set_epi64x(h0[9] ^ h0[11], h0[8] ^ h0[10]); \ division_result_xmm = _mm_cvtsi64_si128(h0[12]); \ @@ -664,7 +664,7 @@ inline void set_float_rounding_mode() ptr0 = (__m128i *)&l0[idx0 & MASK]; \ if(PREFETCH) \ _mm_prefetch((const char*)ptr0, _MM_HINT_T0); \ - if(ALGO != cryptonight_monero_v8) \ + if(ALGO != cryptonight_monero_v8 && ALGO != cryptonight_turtle) \ bx0 = cx #define CN_STEP3(n, monero_const, l0, ax0, bx0, idx0, ptr0, lo, cl, ch, al0, ah0, cx, bx1, sqrt_result, division_result_xmm) \ @@ -681,7 +681,7 @@ inline void set_float_rounding_mode() ah0 += lo; \ al0 += hi; \ } \ - if(ALGO == cryptonight_monero_v8) \ + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) \ { \ bx1 = bx0; \ bx0 = cx; \ diff --git a/xmrstak/backend/cpu/minethd.cpp b/xmrstak/backend/cpu/minethd.cpp index 975f2d4a2a6de80a2abb1b8f69c27d811011455d..2327bed1d08294928f8bacbefc6f5852d0173ffe 100644 --- a/xmrstak/backend/cpu/minethd.cpp +++ b/xmrstak/backend/cpu/minethd.cpp @@ -403,6 +403,15 @@ bool minethd::self_test() hashf("\x03\x05\xa0\xdb\xd6\xbf\x05\xcf\x16\xe5\x03\xf3\xa6\x6f\x78\x00\x7c\xbf\x34\x14\x43\x32\xec\xbf\xc2\x2e\xd9\x5c\x87\x00\x38\x3b\x30\x9a\xce\x19\x23\xa0\x96\x4b\x00\x00\x00\x08\xba\x93\x9a\x62\x72\x4c\x0d\x75\x81\xfc\xe5\x76\x1e\x9d\x8a\x0e\x6a\x1c\x3f\x92\x4f\xdd\x84\x93\xd1\x11\x56\x49\xc0\x5e\xb6\x01", 76, out, ctx); bResult = bResult && memcmp(out, "\x40\x86\x5a\xa8\x87\x41\xec\x1d\xcc\xbd\x2b\xc6\xff\x36\xb9\x4d\x54\x71\x58\xdb\x94\x69\x8e\x3c\xa0\x3d\xe4\x81\x9a\x65\x9f\xef", 32) == 0; } + else if (algo == cryptonight_turtle) + { + hashf = func_selector(::jconf::inst()->HaveHardwareAes(), false, xmrstak_algo::cryptonight_turtle); + hashf("This is a test This is a test This is a test", 44, out, ctx); + bResult = bResult && memcmp(out, "\x30\x5f\x66\xfe\xbb\xf3\x60\x0e\xda\xbb\x60\xf7\xf1\xc9\xb9\x0a\x3a\xe8\x5a\x31\xd4\x76\xca\x38\x1d\x56\x18\xa6\xc6\x27\x60\xd7", 32) == 0; + hashf = func_selector(::jconf::inst()->HaveHardwareAes(), true, xmrstak_algo::cryptonight_turtle); + hashf("This is a test This is a test This is a test", 44, out, ctx); + bResult = bResult && memcmp(out, "\x30\x5f\x66\xfe\xbb\xf3\x60\x0e\xda\xbb\x60\xf7\xf1\xc9\xb9\x0a\x3a\xe8\x5a\x31\xd4\x76\xca\x38\x1d\x56\x18\xa6\xc6\x27\x60\xd7", 32) == 0; + } if(!bResult) @@ -532,6 +541,9 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc case cryptonight_superfast: algv = 11; break; + case cryptonight_turtle: + algv = 12; + break; default: algv = 2; break; @@ -596,7 +608,12 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc Cryptonight_hash<N>::template hash<cryptonight_superfast, false, false>, Cryptonight_hash<N>::template hash<cryptonight_superfast, true, false>, Cryptonight_hash<N>::template hash<cryptonight_superfast, false, true>, - Cryptonight_hash<N>::template hash<cryptonight_superfast, true, true> + Cryptonight_hash<N>::template hash<cryptonight_superfast, true, true>, + + Cryptonight_hash<N>::template hash<cryptonight_turtle, false, false>, + Cryptonight_hash<N>::template hash<cryptonight_turtle, true, false>, + Cryptonight_hash<N>::template hash<cryptonight_turtle, false, true>, + Cryptonight_hash<N>::template hash<cryptonight_turtle, true, true> }; std::bitset<2> digit; @@ -636,6 +653,35 @@ minethd::cn_hash_fun minethd::func_multi_selector(bool bHaveAes, bool bNoPrefetc } } + if (N <= 2 && (algo == cryptonight_turtle) && bHaveAes) + { + std::string selected_asm = asm_version_str; + if (selected_asm == "auto") + selected_asm = cpu::getAsmName(N); + + if (selected_asm != "off") + { + if (selected_asm == "intel_avx" && asm_version_str != "auto") + { + // Intel Ivy Bridge (Xeon v2, Core i7/i5/i3 3xxx, Pentium G2xxx, Celeron G1xxx) + if (N == 1) + selected_function = Cryptonight_hash_asm<1u, 0u>::template hash<cryptonight_turtle>; + else if (N == 2) + selected_function = Cryptonight_hash_asm<2u, 0u>::template hash<cryptonight_turtle>; + } + // supports only 1 thread per hash + if (N == 1 && selected_asm == "amd_avx") + { + // AMD Ryzen (1xxx and 2xxx series) + selected_function = Cryptonight_hash_asm<1u, 1u>::template hash<cryptonight_turtle>; + } + if (asm_version_str == "auto" && (selected_asm != "intel_avx" || selected_asm != "amd_avx")) + printer::inst()->print_msg(L3, "Switch to assembler version for '%s' cpu's", selected_asm.c_str()); + else if (selected_asm != "intel_avx" && selected_asm != "amd_avx") // unknown asm type + printer::inst()->print_msg(L1, "Assembler '%s' unknown, fallback to non asm version of cryptonight_v8", selected_asm.c_str()); + } + } + return selected_function; } diff --git a/xmrstak/backend/cryptonight.hpp b/xmrstak/backend/cryptonight.hpp index e905caa9f24ce264a000fdd54525c1ef70fd7873..2dd922f912970f40c6b8cc90cfb438bc2c3d1976 100644 --- a/xmrstak/backend/cryptonight.hpp +++ b/xmrstak/backend/cryptonight.hpp @@ -17,7 +17,8 @@ enum xmrstak_algo cryptonight_haven = 9, // equal to cryptonight_heavy with a small tweak cryptonight_bittube2 = 10, // derived from cryptonight_heavy with own aes-round implementation and minor other tweaks cryptonight_monero_v8 = 11, - cryptonight_superfast = 12 + cryptonight_superfast = 12, + cryptonight_turtle = 13 }; // define aeon settings @@ -37,6 +38,10 @@ constexpr uint32_t CRYPTONIGHT_MASARI_ITER = 0x40000; constexpr uint32_t CRYPTONIGHT_SUPERFAST_ITER = 0x20000; +constexpr size_t CRYPTONIGHT_TURTLE_MEMORY = 256 * 1024; +constexpr uint32_t CRYPTONIGHT_TURTLE_MASK = 0x1FFF0; +constexpr uint32_t CRYPTONIGHT_TURTLE_ITER = 0x10000; + template<xmrstak_algo ALGO> inline constexpr size_t cn_select_memory() { return 0; } @@ -76,6 +81,9 @@ inline constexpr size_t cn_select_memory<cryptonight_bittube2>() { return CRYPTO template<> inline constexpr size_t cn_select_memory<cryptonight_superfast>() { return CRYPTONIGHT_MEMORY; } +template<> +inline constexpr size_t cn_select_memory<cryptonight_turtle>() { return CRYPTONIGHT_TURTLE_MEMORY; } + inline size_t cn_select_memory(xmrstak_algo algo) { switch(algo) @@ -95,6 +103,8 @@ inline size_t cn_select_memory(xmrstak_algo algo) case cryptonight_haven: case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MEMORY; + case cryptonight_turtle: + return CRYPTONIGHT_TURTLE_MEMORY; default: return 0; } @@ -139,6 +149,9 @@ inline constexpr uint32_t cn_select_mask<cryptonight_bittube2>() { return CRYPTO template<> inline constexpr uint32_t cn_select_mask<cryptonight_superfast>() { return CRYPTONIGHT_MASK; } +template<> +inline constexpr uint32_t cn_select_mask<cryptonight_turtle>() { return CRYPTONIGHT_TURTLE_MASK; } + inline size_t cn_select_mask(xmrstak_algo algo) { switch(algo) @@ -158,6 +171,8 @@ inline size_t cn_select_mask(xmrstak_algo algo) case cryptonight_haven: case cryptonight_heavy: return CRYPTONIGHT_HEAVY_MASK; + case cryptonight_turtle: + return CRYPTONIGHT_TURTLE_MASK; default: return 0; } @@ -202,6 +217,9 @@ inline constexpr uint32_t cn_select_iter<cryptonight_bittube2>() { return CRYPTO template<> inline constexpr uint32_t cn_select_iter<cryptonight_superfast>() { return CRYPTONIGHT_SUPERFAST_ITER; } +template<> +inline constexpr uint32_t cn_select_iter<cryptonight_turtle>() { return CRYPTONIGHT_TURTLE_ITER; } + inline size_t cn_select_iter(xmrstak_algo algo) { switch(algo) @@ -223,6 +241,8 @@ inline size_t cn_select_iter(xmrstak_algo algo) return CRYPTONIGHT_MASARI_ITER; case cryptonight_superfast: return CRYPTONIGHT_SUPERFAST_ITER; + case cryptonight_turtle: + return CRYPTONIGHT_TURTLE_ITER; default: return 0; } diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu index 87c1befa8229a8f5234a5bdda534b4562371fce6..c3a97808c7f7a4f7633f91eed289cc118b423089 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_core.cu @@ -310,7 +310,7 @@ __global__ void cryptonight_core_gpu_phase2_double( int threads, int bfactor, in uint64_t bx1; uint32_t sqrt_result; uint64_t division_result; - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { bx0 = ((uint64_t*)(d_ctx_b + thread * 12))[sub]; bx1 = ((uint64_t*)(d_ctx_b + thread * 12 + 4))[sub]; @@ -350,7 +350,7 @@ __global__ void cryptonight_core_gpu_phase2_double( int threads, int bfactor, in t_fn0( cx.y & 0xff ) ^ t_fn1( (cx2.x >> 8) & 0xff ) ^ rotate16(t_fn0( (cx2.y >> 16) & 0xff ) ^ t_fn1( (cx.x >> 24 ) )) ); - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { const uint64_t chunk1 = myChunks[ idx1 ^ 2 + sub ]; @@ -393,14 +393,14 @@ __global__ void cryptonight_core_gpu_phase2_double( int threads, int bfactor, in else ((ulonglong4*)myChunks)[sub] = ((ulonglong4*)ptr0)[sub]; - if(ALGO != cryptonight_monero_v8) + if(ALGO != cryptonight_monero_v8 && ALGO != cryptonight_turtle) bx0 = cx_aes; uint64_t cx_mul; ((uint32_t*)&cx_mul)[0] = shuffle<2>(sPtr, sub, cx_aes.x , 0); ((uint32_t*)&cx_mul)[1] = shuffle<2>(sPtr, sub, cx_aes.y , 0); - if(ALGO == cryptonight_monero_v8 && sub == 1) + if((ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) && sub == 1) { // Use division and square root results from the _previous_ iteration to hide the latency ((uint32_t*)&division_result)[1] ^= sqrt_result; @@ -424,7 +424,7 @@ __global__ void cryptonight_core_gpu_phase2_double( int threads, int bfactor, in uint64_t cl = ((uint64_t*)myChunks)[ idx1 ]; // sub 0 -> hi, sub 1 -> lo uint64_t res = sub == 0 ? __umul64hi( cx_mul, cl ) : cx_mul * cl; - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { const uint64_t chunk1 = myChunks[ idx1 ^ 2 + sub ] ^ res; uint64_t chunk2 = myChunks[ idx1 ^ 4 + sub ]; @@ -441,7 +441,7 @@ __global__ void cryptonight_core_gpu_phase2_double( int threads, int bfactor, in } ax0 += res; } - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { bx1 = bx0; bx0 = cx_aes; @@ -464,7 +464,7 @@ __global__ void cryptonight_core_gpu_phase2_double( int threads, int bfactor, in if ( bfactor > 0 ) { ((uint64_t*)(d_ctx_a + thread * 4))[sub] = ax0; - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { ((uint64_t*)(d_ctx_b + thread * 12))[sub] = bx0; ((uint64_t*)(d_ctx_b + thread * 12 + 4))[sub] = bx1; @@ -771,7 +771,7 @@ void cryptonight_core_gpu_hash(nvid_ctx* ctx, uint32_t nonce) for ( int i = 0; i < partcount; i++ ) { - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { // two threads per block CUDA_CHECK_MSG_KERNEL( @@ -884,7 +884,10 @@ 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_monero_v8, 1>, cryptonight_core_gpu_hash<CRYPTONIGHT_SUPERFAST_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_superfast, 0>, - cryptonight_core_gpu_hash<CRYPTONIGHT_SUPERFAST_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_superfast, 1> + cryptonight_core_gpu_hash<CRYPTONIGHT_SUPERFAST_ITER, CRYPTONIGHT_MASK, CRYPTONIGHT_MEMORY/4, cryptonight_superfast, 1>, + + cryptonight_core_gpu_hash<CRYPTONIGHT_TURTLE_ITER, CRYPTONIGHT_TURTLE_MASK, CRYPTONIGHT_TURTLE_MEMORY/4, cryptonight_turtle, 0>, + cryptonight_core_gpu_hash<CRYPTONIGHT_TURTLE_ITER, CRYPTONIGHT_TURTLE_MASK, CRYPTONIGHT_TURTLE_MEMORY/4, cryptonight_turtle, 1> }; std::bitset<1> digit; diff --git a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu index f98de2c141fd821a12daade10b47d42e7e71240c..7149f37a661187a9dad8a3032cc55067be446c1e 100644 --- a/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu +++ b/xmrstak/backend/nvidia/nvcc_code/cuda_extra.cu @@ -142,7 +142,7 @@ __global__ void cryptonight_extra_gpu_prepare( int threads, uint32_t * __restric XOR_BLOCKS_DST( ctx_state, ctx_state + 8, ctx_a ); XOR_BLOCKS_DST( ctx_state + 4, ctx_state + 12, ctx_b ); memcpy( d_ctx_a + thread * 4, ctx_a, 4 * 4 ); - if(ALGO == cryptonight_monero_v8) + if(ALGO == cryptonight_monero_v8 || ALGO == cryptonight_turtle) { memcpy( d_ctx_b + thread * 12, ctx_b, 4 * 4 ); // bx1 @@ -310,7 +310,8 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx) // create a double buffer for the state to exchange the mixed state to phase1 CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state2, 50 * sizeof(uint32_t) * wsize)); } - else if(std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end()) + else if(std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end() || + std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_turtle) != neededAlgorithms.end() ) { // bx1 (16byte), division_result (8byte) and sqrt_result (8byte) ctx_b_size = 3 * 4 * sizeof(uint32_t) * wsize; @@ -362,11 +363,16 @@ extern "C" void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_bittube2><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce, ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); } - if(miner_algo == cryptonight_monero_v8) + else if(miner_algo == cryptonight_monero_v8) { CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_monero_v8><<<grid, block >>>( wsize, ctx->d_input, ctx->inputlen, startNonce, ctx->d_ctx_state,ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2 )); } + else if (miner_algo == cryptonight_turtle) + { + CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_extra_gpu_prepare<cryptonight_turtle> << <grid, block >> > (wsize, ctx->d_input, ctx->inputlen, startNonce, + ctx->d_ctx_state, ctx->d_ctx_state2, ctx->d_ctx_a, ctx->d_ctx_b, ctx->d_ctx_key1, ctx->d_ctx_key2)); + } else { /* pass two times d_ctx_state because the second state is used later in phase1, @@ -712,7 +718,8 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx) } // check if cryptonight_monero_v8 is selected for the user pool - bool useCryptonight_v8 = std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end(); + bool useCryptonight_v8 = (std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_monero_v8) != neededAlgorithms.end() || + std::find(neededAlgorithms.begin(), neededAlgorithms.end(), cryptonight_turtle) != neededAlgorithms.end()); // overwrite default config if cryptonight_monero_v8 is mined and GPU has at least compute capability 5.0 if(useCryptonight_v8 && gpuArch >= 50) diff --git a/xmrstak/jconf.cpp b/xmrstak/jconf.cpp index 2a2dc8dbc6580b455dac164db08e626a778f4a4d..165595c5f2c25599c51d21c117c813301ec64e20 100644 --- a/xmrstak/jconf.cpp +++ b/xmrstak/jconf.cpp @@ -99,6 +99,7 @@ xmrstak::coin_selection coins[] = { { "cryptonight_lite_v7", {cryptonight_aeon, cryptonight_aeon, 0u}, {cryptonight_aeon, cryptonight_aeon, 0u}, nullptr }, { "cryptonight_lite_v7_xor", {cryptonight_aeon, cryptonight_ipbc, 255u}, {cryptonight_aeon, cryptonight_aeon, 0u}, nullptr }, { "cryptonight_superfast", {cryptonight_heavy, cryptonight_superfast, 255u},{cryptonight_heavy, cryptonight_superfast, 0u}, nullptr }, + { "cryptonight_turtle", {cryptonight_turtle, cryptonight_turtle, 0u}, {cryptonight_turtle, cryptonight_turtle, 0u}, nullptr }, { "cryptonight_v7", {cryptonight_monero_v8, cryptonight_monero, 255u}, {cryptonight_monero_v8, cryptonight_monero_v8, 0u}, nullptr }, { "cryptonight_v8", {cryptonight_monero_v8, cryptonight_monero_v8, 255u}, {cryptonight_monero_v8, cryptonight_monero_v8, 0u}, nullptr }, { "cryptonight_v7_stellite", {cryptonight_monero_v8, cryptonight_stellite, 255u}, {cryptonight_monero_v8, cryptonight_monero_v8, 0u}, nullptr }, @@ -111,7 +112,8 @@ xmrstak::coin_selection coins[] = { { "qrl", {cryptonight_monero_v8, cryptonight_monero, 255u}, {cryptonight_monero_v8, cryptonight_monero_v8, 0u}, nullptr }, { "ryo", {cryptonight_heavy, cryptonight_heavy, 0u}, {cryptonight_heavy, cryptonight_heavy, 0u}, nullptr }, { "stellite", {cryptonight_monero_v8, cryptonight_stellite, 255u}, {cryptonight_monero_v8, cryptonight_monero_v8, 0u}, nullptr }, - { "turtlecoin", {cryptonight_aeon, cryptonight_aeon, 0u}, {cryptonight_aeon, cryptonight_aeon, 0u}, nullptr } + { "turtlecoin", {cryptonight_turtle, cryptonight_aeon, 5u}, {cryptonight_aeon, cryptonight_aeon, 0u}, nullptr }, + { "plenteum", {cryptonight_turtle, cryptonight_aeon, 5u}, {cryptonight_aeon, cryptonight_aeon, 0u}, nullptr } }; constexpr size_t coin_algo_size = (sizeof(coins)/sizeof(coins[0])); diff --git a/xmrstak/misc/executor.cpp b/xmrstak/misc/executor.cpp index a303b34cda51571b128feb72cd4842b7128f32b0..58d5b0e83ba15a569fbd5791b898e1e95f64e4b3 100644 --- a/xmrstak/misc/executor.cpp +++ b/xmrstak/misc/executor.cpp @@ -562,6 +562,7 @@ void executor::ex_main() break; case cryptonight_monero_v8: case cryptonight_monero: + case cryptonight_turtle: if(dev_tls) pools.emplace_front(0, "donate.xmr-stak.net:8800", "", "", "", 0.0, true, true, "", false); else diff --git a/xmrstak/net/jpsock.cpp b/xmrstak/net/jpsock.cpp index 406c535d2de65fb7dedd17da5339f50cb6f5562b..cbdf1d0c1271192ac79f48738660fbc81ff9f97f 100644 --- a/xmrstak/net/jpsock.cpp +++ b/xmrstak/net/jpsock.cpp @@ -709,6 +709,9 @@ bool jpsock::cmd_submit(const char* sJobId, uint32_t iNonce, const uint8_t* bRes case cryptonight_superfast: algo_name = "cryptonight_superfast"; break; + case cryptonight_turtle: + algo_name = "cryptonight_turtle"; + break; default: algo_name = "unknown"; break; diff --git a/xmrstak/pools.tpl b/xmrstak/pools.tpl index 58762de56f0db711313e788210e2d72f424d12ad..2019f2b86ec3549352808c846f28c79527613e1c 100644 --- a/xmrstak/pools.tpl +++ b/xmrstak/pools.tpl @@ -33,9 +33,12 @@ POOLCONF], * qrl - Quantum Resistant Ledger * ryo * turtlecoin + * plenteum * * Native algorithms which not depends on any block versions: * + * # 256KiB scratchpad memory + * cryptonight_turtle * # 1MiB scratchpad memory * cryptonight_lite * cryptonight_lite_v7