diff --git a/README.md b/README.md index 7b68299b9..e35c3fe62 100644 --- a/README.md +++ b/README.md @@ -276,8 +276,9 @@ void merge(uint32_t &a, uint32_t b, uint32_t r) { case 0: a = (a * 33) + b; break; case 1: a = (a ^ b) * 33; break; - case 2: a = ROTL32(a, ((r >> 16) % 32)) ^ b; break; - case 3: a = ROTR32(a, ((r >> 16) % 32)) ^ b; break; + // prevent rotate by 0 which is a NOP + case 2: a = ROTL32(a, ((r >> 16) % 31) + 1) ^ b; break; + case 3: a = ROTR32(a, ((r >> 16) % 31) + 1) ^ b; break; } } ``` @@ -354,8 +355,11 @@ void progPowLoop( if (i < PROGPOW_CNT_MATH) { // Random Math - int src1 = kiss99(prog_rnd) % PROGPOW_REGS; - int src2 = kiss99(prog_rnd) % PROGPOW_REGS; + // Generate 2 unique sources + int src_rnd = kiss99(prog_rnd) % (PROGPOW_REGS * (PROGPOW_REGS-1)); + int src1 = src_rnd % PROGPOW_REGS; // 0 <= src1 < PROGPOW_REGS + int src2 = src_rnd / PROGPOW_REGS; // 0 <= src2 < PROGPOW_REGS - 1 + if (src2 >= src1) ++src2; // src2 is now any reg other than src1 int sel1 = kiss99(prog_rnd); int dst = mix_seq_dst[(mix_seq_dst_cnt++)%PROGPOW_REGS]; int sel2 = kiss99(prog_rnd); @@ -396,7 +400,8 @@ A full run showing intermediate values can be seen in [result.log](test/result.l ## Change History -- 0.9.1 (current) - Shuffle what part of the DAG entry each lane accesses, suggested by [mbevand](https://github.com/ifdefelse/ProgPOW/pull/13) +- 0.9.2 (current) - Unique sources for math() and prevent rotation by 0 in merge(). Suggested by [SChernykh](https://github.com/ifdefelse/ProgPOW/issues/19) +- [0.9.1](https://github.com/ifdefelse/ProgPOW/blob/60bba1c3fdad6a54539fc3e9f05727547de9c58c/README.md) - Shuffle what part of the DAG entry each lane accesses. Suggested by [mbevand](https://github.com/ifdefelse/ProgPOW/pull/13) - [0.9.0](https://github.com/ifdefelse/ProgPOW/blob/a3f62349a1513f0393524683f9671cfe17cca895/README.md) - Unique cache address sources, re-tune parameters - [0.8.0](https://github.com/ifdefelse/ProgPOW/blob/620b4c7aafe60391f863372814d7517e94386379/README.md) - Original spec diff --git a/libprogpow/ProgPow.cpp b/libprogpow/ProgPow.cpp index 066e09cd3..d1ec3ec24 100644 --- a/libprogpow/ProgPow.cpp +++ b/libprogpow/ProgPow.cpp @@ -138,9 +138,9 @@ std::string ProgPow::getKern(uint64_t block_number, kernel_t kern) ret << "data_dag = g_dag[offset];\n"; ret << "// hack to prevent compiler from reordering LD and usage\n"; if (kern == KERNEL_CUDA) - ret << "if( hack_false ) __threadfence_block();\n"; + ret << "if (hack_false) __threadfence_block();\n"; else - ret << "if( hack_false ) barrier(CLK_LOCAL_MEM_FENCE);\n"; + ret << "if (hack_false) barrier(CLK_LOCAL_MEM_FENCE);\n"; for (int i = 0; (i < PROGPOW_CNT_CACHE) || (i < PROGPOW_CNT_MATH); i++) { @@ -159,15 +159,18 @@ std::string ProgPow::getKern(uint64_t block_number, kernel_t kern) if (i < PROGPOW_CNT_MATH) { // Random Math - // A tree combining random input registers together - // reduced to a single result - std::string src1 = mix_src(); - std::string src2 = mix_src(); + // Generate 2 unique sources + int src_rnd = rnd() % ((PROGPOW_REGS - 1) * PROGPOW_REGS); + int src1 = src_rnd % PROGPOW_REGS; // 0 <= src1 < PROGPOW_REGS + int src2 = src_rnd / PROGPOW_REGS; // 0 <= src2 < PROGPOW_REGS - 1 + if (src2 >= src1) ++src2; // src2 is now any reg other than src1 + std::string src1_str = "mix[" + std::to_string(src1) + "]"; + std::string src2_str = "mix[" + std::to_string(src2) + "]"; uint32_t r1 = rnd(); std::string dest = mix_dst(); uint32_t r2 = rnd(); ret << "// random math " << i << "\n"; - ret << math("data", src1, src2, r1); + ret << math("data", src1_str, src2_str, r1); ret << merge(dest, "data", r2); } } @@ -175,9 +178,9 @@ std::string ProgPow::getKern(uint64_t block_number, kernel_t kern) ret << "// consume global load data\n"; ret << "// hack to prevent compiler from reordering LD and usage\n"; if (kern == KERNEL_CUDA) - ret << "if( hack_false ) __threadfence_block();\n"; + ret << "if (hack_false) __threadfence_block();\n"; else - ret << "if( hack_false ) barrier(CLK_LOCAL_MEM_FENCE);\n"; + ret << "if (hack_false) barrier(CLK_LOCAL_MEM_FENCE);\n"; ret << merge("mix[0]", "data_dag.s[0]", rnd()); for (int i = 1; i < PROGPOW_DAG_LOADS; i++) { @@ -200,8 +203,8 @@ std::string ProgPow::merge(std::string a, std::string b, uint32_t r) { case 0: return a + " = (" + a + " * 33) + " + b + ";\n"; case 1: return a + " = (" + a + " ^ " + b + ") * 33;\n"; - case 2: return a + " = ROTL32(" + a + ", " + std::to_string((r >> 16) % 32) + ") ^ " + b + ";\n"; - case 3: return a + " = ROTR32(" + a + ", " + std::to_string((r >> 16) % 32) + ") ^ " + b + ";\n"; + case 2: return a + " = ROTL32(" + a + ", " + std::to_string(((r >> 16) % 31) + 1) + ") ^ " + b + ";\n"; + case 3: return a + " = ROTR32(" + a + ", " + std::to_string(((r >> 16) % 31) + 1) + ") ^ " + b + ";\n"; } return "#error\n"; } diff --git a/test/kernel.cu b/test/kernel.cu index aae003f67..7aa147310 100644 --- a/test/kernel.cu +++ b/test/kernel.cu @@ -1,133 +1,133 @@ // Inner loop for prog_seed 600 __device__ __forceinline__ void progPowLoop(const uint32_t loop, - uint32_t mix[PROGPOW_REGS], - const dag_t *g_dag, - const uint32_t c_dag[PROGPOW_CACHE_WORDS], - const bool hack_false) + uint32_t mix[PROGPOW_REGS], + const dag_t *g_dag, + const uint32_t c_dag[PROGPOW_CACHE_WORDS], + const bool hack_false) { -dag_t data_dag; -uint32_t offset, data; -const uint32_t lane_id = threadIdx.x & (PROGPOW_LANES-1); -// global load -offset = __shfl_sync(0xFFFFFFFF, mix[0], loop%PROGPOW_LANES, PROGPOW_LANES); -offset %= PROGPOW_DAG_ELEMENTS; -offset = offset * PROGPOW_LANES + (lane_id ^ loop) % PROGPOW_LANES; -data_dag = g_dag[offset]; -// hack to prevent compiler from reordering LD and usage -if( hack_false ) __threadfence_block(); -// cache load 0 -offset = mix[26] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[0] = (mix[0] ^ data) * 33; -// random math 0 -data = clz(mix[10]) + clz(mix[5]); -mix[4] = ROTR32(mix[4], 1) ^ data; -// cache load 1 -offset = mix[30] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[27] = (mix[27] * 33) + data; -// random math 1 -data = min(mix[24], mix[16]); -mix[26] = (mix[26] ^ data) * 33; -// cache load 2 -offset = mix[1] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[13] = ROTL32(mix[13], 4) ^ data; -// random math 2 -data = popcount(mix[3]) + popcount(mix[25]); -mix[15] = ROTR32(mix[15], 14) ^ data; -// cache load 3 -offset = mix[19] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[17] = (mix[17] ^ data) * 33; -// random math 3 -data = mix[13] * mix[23]; -mix[7] = (mix[7] * 33) + data; -// cache load 4 -offset = mix[11] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[14] = (mix[14] ^ data) * 33; -// random math 4 -data = mix[30] ^ mix[17]; -mix[8] = ROTR32(mix[8], 18) ^ data; -// cache load 5 -offset = mix[21] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[9] = (mix[9] * 33) + data; -// random math 5 -data = mix[0] | mix[20]; -mix[12] = (mix[12] * 33) + data; -// cache load 6 -offset = mix[15] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[3] = (mix[3] ^ data) * 33; -// random math 6 -data = mix[15] | mix[29]; -mix[10] = ROTR32(mix[10], 14) ^ data; -// cache load 7 -offset = mix[18] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[1] = (mix[1] ^ data) * 33; -// random math 7 -data = ROTL32(mix[30], mix[15]); -mix[11] = (mix[11] * 33) + data; -// cache load 8 -offset = mix[3] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[6] = (mix[6] * 33) + data; -// random math 8 -data = mix[11] + mix[9]; -mix[16] = ROTR32(mix[16], 4) ^ data; -// cache load 9 -offset = mix[17] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[28] = ROTL32(mix[28], 21) ^ data; -// random math 9 -data = mix[26] + mix[7]; -mix[31] = (mix[31] ^ data) * 33; -// cache load 10 -offset = mix[31] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[2] = ROTR32(mix[2], 21) ^ data; -// random math 10 -data = min(mix[11], mix[1]); -mix[19] = ROTR32(mix[19], 13) ^ data; -// cache load 11 -offset = mix[16] % PROGPOW_CACHE_WORDS; -data = c_dag[offset]; -mix[30] = ROTR32(mix[30], 7) ^ data; -// random math 11 -data = mix[10] & mix[14]; -mix[22] = (mix[22] * 33) + data; -// random math 12 -data = mix[28] ^ mix[30]; -mix[29] = ROTR32(mix[29], 28) ^ data; -// random math 13 -data = min(mix[17], mix[17]); -mix[5] = (mix[5] ^ data) * 33; -// random math 14 -data = mix[23] + mix[25]; -mix[24] = ROTL32(mix[24], 6) ^ data; -// random math 15 -data = mix[1] & mix[27]; -mix[18] = (mix[18] ^ data) * 33; -// random math 16 -data = ROTL32(mix[3], mix[16]); -mix[25] = (mix[25] * 33) + data; -// random math 17 -data = ROTL32(mix[0], mix[9]); -mix[23] = (mix[23] * 33) + data; -// random math 18 -data = mix[2] ^ mix[0]; -mix[21] = (mix[21] * 33) + data; -// random math 19 -data = mix[25] * mix[10]; -mix[20] = (mix[20] ^ data) * 33; -// consume global load data -// hack to prevent compiler from reordering LD and usage -if( hack_false ) __threadfence_block(); -mix[0] = (mix[0] ^ data_dag.s[0]) * 33; -mix[0] = (mix[0] * 33) + data_dag.s[1]; -mix[4] = ROTR32(mix[4], 1) ^ data_dag.s[2]; -mix[27] = ROTL32(mix[27], 31) ^ data_dag.s[3]; -} + dag_t data_dag; + uint32_t offset, data; + const uint32_t lane_id = threadIdx.x & (PROGPOW_LANES - 1); + // global load + offset = __shfl_sync(0xFFFFFFFF, mix[0], loop%PROGPOW_LANES, PROGPOW_LANES); + offset %= PROGPOW_DAG_ELEMENTS; + offset = offset * PROGPOW_LANES + (lane_id ^ loop) % PROGPOW_LANES; + data_dag = g_dag[offset]; + // hack to prevent compiler from reordering LD and usage + if (hack_false) __threadfence_block(); + // cache load 0 + offset = mix[26] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[0] = (mix[0] ^ data) * 33; + // random math 0 + data = mix[10] ^ mix[16]; + mix[4] = ROTL32(mix[4], 27) ^ data; + // cache load 1 + offset = mix[30] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[27] = ROTR32(mix[27], 7) ^ data; + // random math 1 + data = mix[24] & mix[14]; + mix[26] = (mix[26] * 33) + data; + // cache load 2 + offset = mix[1] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[13] = (mix[13] * 33) + data; + // random math 2 + data = mix[17] & mix[16]; + mix[15] = ROTR32(mix[15], 12) ^ data; + // cache load 3 + offset = mix[19] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[17] = (mix[17] ^ data) * 33; + // random math 3 + data = mul_hi(mix[31], mix[5]); + mix[7] = (mix[7] ^ data) * 33; + // cache load 4 + offset = mix[11] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[14] = (mix[14] ^ data) * 33; + // random math 4 + data = mix[23] * mix[19]; + mix[8] = (mix[8] * 33) + data; + // cache load 5 + offset = mix[21] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[9] = (mix[9] ^ data) * 33; + // random math 5 + data = clz(mix[30]) + clz(mix[15]); + mix[12] = ROTR32(mix[12], 16) ^ data; + // cache load 6 + offset = mix[15] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[3] = ROTR32(mix[3], 27) ^ data; + // random math 6 + data = clz(mix[12]) + clz(mix[5]); + mix[10] = (mix[10] * 33) + data; + // cache load 7 + offset = mix[18] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[1] = ROTR32(mix[1], 6) ^ data; + // random math 7 + data = min(mix[4], mix[25]); + mix[11] = ROTR32(mix[11], 27) ^ data; + // cache load 8 + offset = mix[3] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[6] = (mix[6] ^ data) * 33; + // random math 8 + data = mul_hi(mix[18], mix[16]); + mix[16] = (mix[16] ^ data) * 33; + // cache load 9 + offset = mix[17] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[28] = ROTL32(mix[28], 17) ^ data; + // random math 9 + data = ROTL32(mix[15], mix[23]); + mix[31] = (mix[31] * 33) + data; + // cache load 10 + offset = mix[31] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[2] = (mix[2] * 33) + data; + // random math 10 + data = mix[11] | mix[17]; + mix[19] = ROTL32(mix[19], 28) ^ data; + // cache load 11 + offset = mix[16] % PROGPOW_CACHE_WORDS; + data = c_dag[offset]; + mix[30] = ROTR32(mix[30], 18) ^ data; + // random math 11 + data = mix[22] * mix[7]; + mix[22] = ROTR32(mix[22], 30) ^ data; + // random math 12 + data = mix[27] & mix[16]; + mix[29] = ROTR32(mix[29], 25) ^ data; + // random math 13 + data = ROTL32(mix[11], mix[0]); + mix[5] = (mix[5] ^ data) * 33; + // random math 14 + data = ROTR32(mix[15], mix[25]); + mix[24] = ROTL32(mix[24], 13) ^ data; + // random math 15 + data = mix[14] & mix[26]; + mix[18] = (mix[18] * 33) + data; + // random math 16 + data = mix[28] * mix[16]; + mix[25] = (mix[25] ^ data) * 33; + // random math 17 + data = mix[11] * mix[0]; + mix[23] = (mix[23] ^ data) * 33; + // random math 18 + data = mix[2] + mix[24]; + mix[21] = ROTR32(mix[21], 20) ^ data; + // random math 19 + data = mix[25] + mix[4]; + mix[20] = ROTL32(mix[20], 22) ^ data; + // consume global load data + // hack to prevent compiler from reordering LD and usage + if (hack_false) __threadfence_block(); + mix[0] = (mix[0] ^ data_dag.s[0]) * 33; + mix[0] = ROTR32(mix[0], 21) ^ data_dag.s[1]; + mix[4] = (mix[4] * 33) + data_dag.s[2]; + mix[27] = (mix[27] ^ data_dag.s[3]) * 33; +} \ No newline at end of file diff --git a/test/result.log b/test/result.log index da93f0cfe..9c9695cd6 100644 --- a/test/result.log +++ b/test/result.log @@ -1,32 +1,32 @@ >ethminer.exe -U -M 30000 - m 23:42:31|main | ethminer version 0.15.0.dev0 - m 23:42:31|main | Build: windows / release +git. 575487f - cu 23:42:31|main | Using grid size 1024 , block size 512 + m 00:37:26|main | ethminer version 0.15.0.dev0 + m 00:37:26|main | Build: windows / release +git. 2c02a51 + cu 00:37:26|main | Using grid size 1024 , block size 512 Benchmarking on platform: CUDA Preparing DAG for block #30000 -Warming up... - i 23:42:31|cuda-0 | No work. - i 23:42:31|cuda-0 | Initialising miner 0 - cu 23:42:32|cuda-0 | Using device: GeForce GTX 1060 6GB (Compute 6.1) - cu 23:42:32|cuda-0 | Set Device to current - cu 23:42:32|cuda-0 | Resetting device - cu 23:42:32|cuda-0 | Compile log: - cu 23:42:32|cuda-0 | JIT info: + i Warming up...00:37:26|cuda-0 | No work. + + i 00:37:26|cuda-0 | Initialising miner 0 + cu 00:37:26|cuda-0 | Using device: GeForce GTX 1060 6GB (Compute 6.1) + cu 00:37:26|cuda-0 | Set Device to current + cu 00:37:26|cuda-0 | Resetting device + cu 00:37:27|cuda-0 | Allocating light with size: 16907456 + cu 00:37:27|cuda-0 | Generating mining buffers + cu 00:37:27|cuda-0 | Generating DAG for GPU # 0 with dagBytes: 1082130304 gridSize: 1024 + cu 00:37:30|cuda-0 | Finished DAG + cu 00:37:31|cuda-0 | Compile log: + cu 00:37:31|cuda-0 | JIT info: ptxas info : 202 bytes gmem, 96 bytes cmem[3] ptxas info : Compiling entry function '_Z14progpow_searchy8hash32_tyPK5dag_tPV14search_resultsb' for 'sm_61' ptxas info : Function properties for _Z14progpow_searchy8hash32_tyPK5dag_tPV14search_resultsb ptxas . 32 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads -ptxas info : Used 157 registers, 16384 bytes smem, 385 bytes cmem[0], 4 bytes cmem[2] +ptxas info : Used 158 registers, 16384 bytes smem, 385 bytes cmem[0], 4 bytes cmem[2] ptxas info : Function properties for _Z11keccak_f8008hash32_tyS_ ptxas . 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads - cu 23:42:32|cuda-0 | JIT err: + cu 00:37:31|cuda-0 | JIT err: - cu 23:42:32|cuda-0 | Mangled name: _Z14progpow_searchy8hash32_tyPK5dag_tPV14search_resultsb - cu 23:42:32|cuda-0 | done compiling - cu 23:42:32|cuda-0 | Allocating light with size: 16907456 - cu 23:42:32|cuda-0 | Generating mining buffers - cu 23:42:32|cuda-0 | Generating DAG for GPU # 0 with dagBytes: 1082130304 gridSize: 1024 - cu 23:42:36|cuda-0 | Finished DAG + cu 00:37:31|cuda-0 | Mangled name: _Z14progpow_searchy8hash32_tyPK5dag_tPV14search_resultsb + cu 00:37:31|cuda-0 | done compiling header ffeeddccbbaa9988776655443322110000112233445566778899aabbccddeeff nonce 123456789abcdef0 seed ee304846ddd0a47b @@ -55,84 +55,84 @@ fill_mix lane 13 mix[0]=4e46d05d ... mix[31]=3837dda3 fill_mix lane 14 mix[0]=1a9b1d40 ... mix[31]=08b4fc39 fill_mix lane 15 mix[0]=3344ce0f ... mix[31]=20201012 loop 0 dag entry 2043727 of 4227071 total -loop 1 dag entry 978957 of 4227071 total -loop 2 dag entry 2219106 of 4227071 total -loop 3 dag entry 2200566 of 4227071 total -loop 4 dag entry 3060295 of 4227071 total -loop 5 dag entry 1066534 of 4227071 total -loop 6 dag entry 1860140 of 4227071 total -loop 7 dag entry 3119727 of 4227071 total -loop 8 dag entry 1192860 of 4227071 total -loop 9 dag entry 148078 of 4227071 total -loop 10 dag entry 4019583 of 4227071 total -loop 11 dag entry 763952 of 4227071 total -loop 12 dag entry 938751 of 4227071 total -loop 13 dag entry 540327 of 4227071 total -loop 14 dag entry 3142750 of 4227071 total -loop 15 dag entry 780266 of 4227071 total -loop 16 dag entry 842914 of 4227071 total -loop 17 dag entry 302514 of 4227071 total -loop 18 dag entry 1341289 of 4227071 total -loop 19 dag entry 1042834 of 4227071 total -loop 20 dag entry 3313136 of 4227071 total -loop 21 dag entry 1552169 of 4227071 total -loop 22 dag entry 4108553 of 4227071 total -loop 23 dag entry 1435285 of 4227071 total -loop 24 dag entry 4069636 of 4227071 total -loop 25 dag entry 2928713 of 4227071 total -loop 26 dag entry 1770730 of 4227071 total -loop 27 dag entry 238606 of 4227071 total -loop 28 dag entry 3263664 of 4227071 total -loop 29 dag entry 1755218 of 4227071 total -loop 30 dag entry 3035329 of 4227071 total -loop 31 dag entry 652802 of 4227071 total -loop 32 dag entry 3283360 of 4227071 total -loop 33 dag entry 2327823 of 4227071 total -loop 34 dag entry 3339662 of 4227071 total -loop 35 dag entry 1453003 of 4227071 total -loop 36 dag entry 3602110 of 4227071 total -loop 37 dag entry 1821760 of 4227071 total -loop 38 dag entry 754663 of 4227071 total -loop 39 dag entry 184501 of 4227071 total -loop 40 dag entry 3600362 of 4227071 total -loop 41 dag entry 1816934 of 4227071 total -loop 42 dag entry 1381225 of 4227071 total -loop 43 dag entry 1304550 of 4227071 total -loop 44 dag entry 1353179 of 4227071 total -loop 45 dag entry 10429 of 4227071 total -loop 46 dag entry 894212 of 4227071 total -loop 47 dag entry 4224745 of 4227071 total -loop 48 dag entry 215407 of 4227071 total -loop 49 dag entry 506509 of 4227071 total -loop 50 dag entry 1702595 of 4227071 total -loop 51 dag entry 303727 of 4227071 total -loop 52 dag entry 41423 of 4227071 total -loop 53 dag entry 3805288 of 4227071 total -loop 54 dag entry 4093082 of 4227071 total -loop 55 dag entry 148616 of 4227071 total -loop 56 dag entry 3441343 of 4227071 total -loop 57 dag entry 8257 of 4227071 total -loop 58 dag entry 276369 of 4227071 total -loop 59 dag entry 3320925 of 4227071 total -loop 60 dag entry 1431023 of 4227071 total -loop 61 dag entry 2778561 of 4227071 total -loop 62 dag entry 452777 of 4227071 total -loop 63 dag entry 1275346 of 4227071 total -digest lane 0: f0768603 -digest lane 1: b64fd1eb -digest lane 2: 76cb337c -digest lane 3: 47e14410 -digest lane 4: c91decf1 -digest lane 5: 8b665796 -digest lane 6: 16ee0104 -digest lane 7: 9eb0d380 -digest lane 8: 5cd02f9e -digest lane 9: bad01c1e -digest lane 10: edfe0292 -digest lane 11: 3960e1e8 -digest lane 12: def7562f -digest lane 13: 7bf8fab5 -digest lane 14: 0dce42c8 -digest lane 15: fae46383 -digest: 44fa88669c864aa30ba7da46e557593289c4d1fb143a1c43813d512b14fb4636 -result (top 64 bits): b946ea7d74e3c619 +loop 1 dag entry 1878577 of 4227071 total +loop 2 dag entry 1972818 of 4227071 total +loop 3 dag entry 4192557 of 4227071 total +loop 4 dag entry 2908963 of 4227071 total +loop 5 dag entry 650106 of 4227071 total +loop 6 dag entry 3360110 of 4227071 total +loop 7 dag entry 2666972 of 4227071 total +loop 8 dag entry 3262571 of 4227071 total +loop 9 dag entry 1876031 of 4227071 total +loop 10 dag entry 1099946 of 4227071 total +loop 11 dag entry 1058639 of 4227071 total +loop 12 dag entry 4091582 of 4227071 total +loop 13 dag entry 2295331 of 4227071 total +loop 14 dag entry 2587683 of 4227071 total +loop 15 dag entry 950942 of 4227071 total +loop 16 dag entry 2427766 of 4227071 total +loop 17 dag entry 677253 of 4227071 total +loop 18 dag entry 3564299 of 4227071 total +loop 19 dag entry 2373221 of 4227071 total +loop 20 dag entry 2065878 of 4227071 total +loop 21 dag entry 2684534 of 4227071 total +loop 22 dag entry 3563556 of 4227071 total +loop 23 dag entry 909053 of 4227071 total +loop 24 dag entry 3867986 of 4227071 total +loop 25 dag entry 959685 of 4227071 total +loop 26 dag entry 2837635 of 4227071 total +loop 27 dag entry 3312470 of 4227071 total +loop 28 dag entry 3048893 of 4227071 total +loop 29 dag entry 3601694 of 4227071 total +loop 30 dag entry 3536836 of 4227071 total +loop 31 dag entry 533409 of 4227071 total +loop 32 dag entry 3736438 of 4227071 total +loop 33 dag entry 864961 of 4227071 total +loop 34 dag entry 188397 of 4227071 total +loop 35 dag entry 3814381 of 4227071 total +loop 36 dag entry 4108296 of 4227071 total +loop 37 dag entry 3950694 of 4227071 total +loop 38 dag entry 2069968 of 4227071 total +loop 39 dag entry 2745630 of 4227071 total +loop 40 dag entry 1008990 of 4227071 total +loop 41 dag entry 2675149 of 4227071 total +loop 42 dag entry 3352224 of 4227071 total +loop 43 dag entry 472995 of 4227071 total +loop 44 dag entry 3348733 of 4227071 total +loop 45 dag entry 2693084 of 4227071 total +loop 46 dag entry 3406906 of 4227071 total +loop 47 dag entry 2629222 of 4227071 total +loop 48 dag entry 911095 of 4227071 total +loop 49 dag entry 1358231 of 4227071 total +loop 50 dag entry 3574453 of 4227071 total +loop 51 dag entry 4109829 of 4227071 total +loop 52 dag entry 3503254 of 4227071 total +loop 53 dag entry 3897141 of 4227071 total +loop 54 dag entry 4030761 of 4227071 total +loop 55 dag entry 2855865 of 4227071 total +loop 56 dag entry 452603 of 4227071 total +loop 57 dag entry 1258735 of 4227071 total +loop 58 dag entry 2380407 of 4227071 total +loop 59 dag entry 3320172 of 4227071 total +loop 60 dag entry 3180940 of 4227071 total +loop 61 dag entry 1407113 of 4227071 total +loop 62 dag entry 1057525 of 4227071 total +loop 63 dag entry 574671 of 4227071 total +digest lane 0: 5883883e +digest lane 1: 2fb0fd2e +digest lane 2: eadb7563 +digest lane 3: 4a171075 +digest lane 4: ac2758f5 +digest lane 5: aa5b06cf +digest lane 6: 52156e93 +digest lane 7: 4f7a7fff +digest lane 8: fe91e36a +digest lane 9: 9964c8b6 +digest lane 10: 6a3d93e2 +digest lane 11: 3c6d641f +digest lane 12: e90da618 +digest lane 13: 80cd8ab9 +digest lane 14: ce72386f +digest lane 15: 95517d28 +digest: 11f19805c58ab46610ff9c719dcf0a5f18fa2f1605798eef770c47219274767d +result (top 64 bits): 5b7ccd472dbefdd9 \ No newline at end of file