IPBC and Graft POW

xmrstak/backend/amd/amd_gpu/gpu.cpp

@@ -308,9 +308,10 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 		return ERR_OCL_API;
 	}
 
-	size_t scratchPadSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
-	int threadMemMask = cn_select_mask(::jconf::inst()->GetMiningAlgo());
-	int hashIterations = cn_select_iter(::jconf::inst()->GetMiningAlgo());
+	size_t scratchPadSize = std::max(
+		cn_select_memory(::jconf::inst()->GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetMiningAlgoRoot())
+	);
 
 	size_t g_thd = ctx->rawIntensity;
 	ctx->ExtraBuffers[0] = clCreateBuffer(opencl_ctx, CL_MEM_READ_WRITE, scratchPadSize * g_thd, NULL, &ret);
@@ -382,6 +383,7 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 
 	for(int ii = 0; ii < num_algos; ++ii)
 	{
+		// scratchpad size for the selected mining algorithm
 		size_t hashMemSize = cn_select_memory(miner_algo[ii]);
 		int threadMemMask = cn_select_mask(miner_algo[ii]);
 		int hashIterations = cn_select_iter(miner_algo[ii]);
@@ -493,7 +495,7 @@ size_t InitOpenCLGpu(cl_context opencl_ctx, GpuContext* ctx, const char* source_
 					p_id++;
 				}
 
-			if((ret = clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS)
+				if((ret = clGetProgramInfo(ctx->Program[ii], CL_PROGRAM_BINARIES, num_devices * sizeof(char*), all_programs.data(),NULL)) != CL_SUCCESS)
 				{
 					printer::inst()->print_msg(L1,"Error %s when calling clGetProgramInfo.", err_to_str(ret));
 					return ERR_OCL_API;
@@ -988,7 +990,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 )
+	if(miner_algo == cryptonight_ipbc || miner_algo == cryptonight_monero || miner_algo == cryptonight_aeon )
 	{
 		// Input
 		if ((ret = clSetKernelArg(ctx->Kernels[kernel_storage][1], 3, sizeof(cl_mem), &ctx->InputBuffer)) != CL_SUCCESS)

xmrstak/backend/amd/amd_gpu/opencl/cryptonight.cl

@@ -32,7 +32,8 @@ R"===(
  * The implemented function is modified because the last is in our case always a scalar.
  * We can ignore the bitwise AND operation.
  */
-inline uint2 amd_bitalign( const uint2 src0, const uint2 src1, const uint src2)
+
+inline uint2 __attribute__((overloadable)) amd_bitalign( const uint2 src0, const uint2 src1, const uint src2)
 {
 	uint2 result;
 	result.s0 =  (uint) (((((long)src0.s0) << 32) | (long)src1.s0) >> (src2));
@@ -59,7 +60,7 @@ inline uint2 amd_bitalign( const uint2 src0, const uint2 src1, const uint src2)
  *         dst.s0 = src0.s0 >> offset;
  *     similar operation applied to other components of the vectors
  */
-inline int amd_bfe(const uint src0, const uint offset, const uint width)
+inline uint __attribute__((overloadable)) amd_bfe(const uint src0, const uint offset, const uint width)
 {
 	/* casts are removed because we can implement everything as uint
 	 * int offset = src1;
@@ -553,8 +554,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
-#if(ALGO == 3 || ALGO == 5)
+// cryptonight_monero || cryptonight_aeon || cryptonight_ipbc
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 , __global ulong *input
 #endif
 )
@@ -574,7 +575,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 	}
 
 	barrier(CLK_LOCAL_MEM_FENCE);
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
     uint2 tweak1_2;
 #endif
 	uint4 b_x;
@@ -598,7 +599,7 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 		b[1] = states[3] ^ states[7];
 
 		b_x = ((uint4 *)b)[0];
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 		tweak1_2 = as_uint2(input[4]);
 		tweak1_2.s0 >>= 24;
 		tweak1_2.s0 |= tweak1_2.s1 << 8;
@@ -625,7 +626,7 @@ __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];
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 			uint table = 0x75310U;
 			uint index = ((b_x.s2 >> 26) & 12) | ((b_x.s2 >> 23) & 2);
 			b_x.s2 ^= ((table >> index) & 0x30U) << 24;
@@ -639,13 +640,18 @@ __kernel void JOIN(cn1,ALGO) (__global uint4 *Scratchpad, __global ulong *states
 			a[0] += mul_hi(c[0], as_ulong2(tmp).s0);
 
 
-#if(ALGO == 3 || ALGO == 5)
+#if(ALGO == 3 || ALGO == 5 || ALGO == 6)
 			((uint2 *)&(a[1]))[0] ^= tweak1_2;
 			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
 			((uint2 *)&(a[1]))[0] ^= tweak1_2;
 #else
 			Scratchpad[IDX((c[0] & MASK) >> 4)] = ((uint4 *)a)[0];
 #endif
+#if(ALGO == 6)
+			long prev = *((__global long*)(Scratchpad + (IDX((c[0] & MASK) >> 4))));
+			long cur = *((__global long*)(Scratchpad + (IDX((c[0] & MASK) >> 4))) + 1);
+			*((__global long*)(Scratchpad + (IDX((c[0] & MASK) >> 4))) + 1) = prev ^ cur;
+#endif
 
 			((uint4 *)a)[0] ^= tmp;
 			idx0 = a[0];

xmrstak/backend/amd/autoAdjust.hpp

@@ -83,7 +83,10 @@ class autoAdjust
 
 		constexpr size_t byteToMiB = 1024u * 1024u;
 
-		size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+		size_t hashMemSize = std::max(
+			cn_select_memory(::jconf::inst()->GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetMiningAlgoRoot())
+		);
 
 		std::string conf;
 		for(auto& ctx : devVec)

xmrstak/backend/amd/minethd.cpp

@@ -217,6 +217,7 @@ void minethd::work_main()
 		}
 
 		uint8_t new_version = oWork.getVersion();
+		if (::jconf::inst()->GetMiningAlgo() == cryptonight_ipbc) new_version = oWork.bWorkBlob[1];
 		if(new_version != version)
 		{
 			if(new_version >= ::jconf::inst()->GetMiningForkVersion())

xmrstak/backend/cpu/autoAdjust.hpp

@@ -36,7 +36,10 @@ class autoAdjust
 	bool printConfig()
 	{
 
-		const size_t hashMemSizeKB = cn_select_memory(::jconf::inst()->GetMiningAlgo()) / 1024u;
+		const size_t hashMemSizeKB = std::max(
+			cn_select_memory(::jconf::inst()->GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetMiningAlgoRoot())
+		) / 1024u;
 		const size_t halfHashMemSizeKB = hashMemSizeKB / 2u;
 
 		configEditor configTpl{};

xmrstak/backend/cpu/autoAdjustHwloc.hpp

@@ -28,7 +28,10 @@ class autoAdjust
 
 	autoAdjust()
 	{
-		hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+		hashMemSize = std::max(
+			cn_select_memory(::jconf::inst()->GetMiningAlgo()),
+			cn_select_memory(::jconf::inst()->GetMiningAlgoRoot())
+		);
 		halfHashMemSize = hashMemSize / 2u;
 	}
 

xmrstak/backend/cpu/crypto/cryptonight_aesni.h

@@ -444,7 +444,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) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32);
 		return;
@@ -453,7 +453,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)
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 	{
 		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);
@@ -482,7 +482,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)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		else
 			_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -506,10 +506,15 @@ 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)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			((uint64_t*)&l0[idx0 & MASK])[1] = ah0 ^ monero_const;
 		else
 			((uint64_t*)&l0[idx0 & MASK])[1] = ah0;
+
+		if (ALGO == cryptonight_ipbc) {
+			((uint64_t*)&l0[idx0 & MASK])[1] ^= ((uint64_t*)&l0[idx0 & MASK])[0];
+		}
+
 		ah0 ^= ch;
 
 		idx0 = al0;
@@ -544,7 +549,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) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 64);
 		return;
@@ -554,7 +559,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)
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 	{
 		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);
@@ -592,7 +597,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)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			cryptonight_monero_tweak((uint64_t*)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
 		else
 			_mm_store_si128((__m128i *)&l0[idx0 & MASK], _mm_xor_si128(bx0, cx));
@@ -610,7 +615,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)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			cryptonight_monero_tweak((uint64_t*)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
 		else
 			_mm_store_si128((__m128i *)&l1[idx1 & MASK], _mm_xor_si128(bx1, cx));
@@ -631,11 +636,15 @@ 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)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			((uint64_t*)&l0[idx0 & MASK])[1] = axh0 ^ monero_const_0;
 		else
 			((uint64_t*)&l0[idx0 & MASK])[1] = axh0;
 
+		if (ALGO == cryptonight_ipbc) {
+			((uint64_t*)&l0[idx0 & MASK])[1] ^= ((uint64_t*)&l0[idx0 & MASK])[0];
+		}
+
 		axh0 ^= ch;
 		axl0 ^= cl;
 		idx0 = axl0;
@@ -662,11 +671,15 @@ 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)
+		if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc)
 			((uint64_t*)&l1[idx1 & MASK])[1] = axh1 ^ monero_const_1;
 		else
 			((uint64_t*)&l1[idx1 & MASK])[1] = axh1;
 
+		if (ALGO == cryptonight_ipbc) {
+			((uint64_t*)&l1[idx1 & MASK])[1] ^= ((uint64_t*)&l1[idx1 & MASK])[0];
+		}
+
 		axh1 ^= ch;
 		axl1 ^= cl;
 		idx1 = axl1;
@@ -709,7 +722,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) \
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) \
 		cryptonight_monero_tweak((uint64_t*)ptr, b); \
 	else \
 		_mm_store_si128(ptr, b);\
@@ -724,10 +737,12 @@ 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) \
+	if(ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) \
 		_mm_store_si128(ptr, _mm_xor_si128(a, mc)); \
 	else \
 		_mm_store_si128(ptr, a);\
+    if (ALGO == cryptonight_ipbc) \
+        ((uint64_t*)&l[idx & MASK])[1] ^= ((uint64_t*)&l[idx & MASK])[0];\
 	a = _mm_xor_si128(a, b); \
 	idx = _mm_cvtsi128_si64(a);	\
 	if(ALGO == cryptonight_heavy) \
@@ -751,7 +766,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) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32 * 3);
 		return;
@@ -845,7 +860,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) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32 * 4);
 		return;
@@ -954,7 +969,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) && len < 43)
+	if((ALGO == cryptonight_monero || ALGO == cryptonight_aeon || ALGO == cryptonight_ipbc) && len < 43)
 	{
 		memset(output, 0, 32 * 5);
 		return;

xmrstak/backend/cpu/crypto/cryptonight_common.cpp

@@ -35,6 +35,7 @@ extern "C"
 #include "xmrstak/jconf.hpp"
 #include <stdio.h>
 #include <stdlib.h>
+#include <algorithm>
 
 #ifdef __GNUC__
 #include <mm_malloc.h>
@@ -202,7 +203,10 @@ size_t cryptonight_init(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
 
 cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, alloc_msg* msg)
 {
-	size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+	size_t hashMemSize = std::max(
+		cn_select_memory(::jconf::inst()->GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetMiningAlgoRoot())
+	);
 
 	cryptonight_ctx* ptr = (cryptonight_ctx*)_mm_malloc(sizeof(cryptonight_ctx), 4096);
 
@@ -278,7 +282,10 @@ cryptonight_ctx* cryptonight_alloc_ctx(size_t use_fast_mem, size_t use_mlock, al
 
 void cryptonight_free_ctx(cryptonight_ctx* ctx)
 {
-	size_t hashMemSize = cn_select_memory(::jconf::inst()->GetMiningAlgo());
+	size_t hashMemSize = std::max(
+		cn_select_memory(::jconf::inst()->GetMiningAlgo()),
+		cn_select_memory(::jconf::inst()->GetMiningAlgoRoot())
+	);
 
 	if(ctx->ctx_info[0] != 0)
 	{