mirror of
https://github.com/xmrig/xmrig.git
synced 2026-04-17 13:02:57 -04:00
Compare commits
29 Commits
f16a06eb67
...
dev
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
6bf43053f7 | ||
|
|
69b7e60d35 | ||
|
|
b2ca72480c | ||
|
|
92705f2dae | ||
|
|
4f58a7afff | ||
|
|
806cfc3f4d | ||
|
|
84352c71ca | ||
|
|
27d535d00f | ||
|
|
9d296c7f02 | ||
|
|
3b4e38ba18 | ||
|
|
8b33d2494b | ||
|
|
c534c669cb | ||
|
|
d7f7094c45 | ||
|
|
14dcd36296 | ||
|
|
a935641274 | ||
|
|
976a08efb4 | ||
|
|
17aa97e543 | ||
|
|
48b29fd68b | ||
|
|
6454a0abf2 | ||
|
|
8464d474d4 | ||
|
|
26ee1cd291 | ||
|
|
d9b39e8c32 | ||
|
|
05dd6dcc40 | ||
|
|
316a3673bc | ||
|
|
42c943c03f | ||
|
|
c2c8080783 | ||
|
|
d82d7f3f20 | ||
|
|
a189d84fcd | ||
|
|
cb6001945e |
14
CHANGELOG.md
14
CHANGELOG.md
@@ -1,3 +1,17 @@
|
||||
# v6.26.0
|
||||
- [#3769](https://github.com/xmrig/xmrig/pull/3769), [#3772](https://github.com/xmrig/xmrig/pull/3772), [#3774](https://github.com/xmrig/xmrig/pull/3774), [#3775](https://github.com/xmrig/xmrig/pull/3775), [#3776](https://github.com/xmrig/xmrig/pull/3776), [#3782](https://github.com/xmrig/xmrig/pull/3782), [#3783](https://github.com/xmrig/xmrig/pull/3783) **Added support for RandomX v2.**
|
||||
- [#3746](https://github.com/xmrig/xmrig/pull/3746) RISC-V: vectorized RandomX main loop.
|
||||
- [#3748](https://github.com/xmrig/xmrig/pull/3748) RISC-V: auto-detect and use vector code for all RandomX AES functions.
|
||||
- [#3749](https://github.com/xmrig/xmrig/pull/3749) RISC-V: detect and use hardware AES.
|
||||
- [#3750](https://github.com/xmrig/xmrig/pull/3750) RISC-V: use vector hardware AES instead of scalar.
|
||||
- [#3757](https://github.com/xmrig/xmrig/pull/3757) RISC-V: Fixed scratchpad prefetch, removed an unnecessary instruction.
|
||||
- [#3758](https://github.com/xmrig/xmrig/pull/3758) RandomX: added VAES-512 support for Zen5.
|
||||
- [#3759](https://github.com/xmrig/xmrig/pull/3759) RandomX: Optimized VAES code.
|
||||
- [#3762](https://github.com/xmrig/xmrig/pull/3762) Fixed keepalive timer logic.
|
||||
- [#3778](https://github.com/xmrig/xmrig/pull/3778) RandomX: ARM64 fixes.
|
||||
- [#3784](https://github.com/xmrig/xmrig/pull/3784) Fixed OpenCL address-space mismatch in `keccak_f800_round`.
|
||||
- [#3785](https://github.com/xmrig/xmrig/pull/3785) Don't reset nonce during donation rounds.
|
||||
|
||||
# v6.25.0
|
||||
- [#3680](https://github.com/xmrig/xmrig/pull/3680) Added `armv8l` to the list of 32-bit ARM targets.
|
||||
- [#3708](https://github.com/xmrig/xmrig/pull/3708) Minor Aarch64 JIT changes (better instruction selection, don't emit instructions that add 0, etc).
|
||||
|
||||
@@ -110,7 +110,7 @@ if (WITH_RANDOMX)
|
||||
endif()
|
||||
endif()
|
||||
|
||||
set_source_files_properties(src/crypto/randomx/jit_compiler_rv64_vector_static.S PROPERTIES COMPILE_FLAGS "-march=${RV64_VECTOR_FILE_ARCH}")
|
||||
set_source_files_properties(src/crypto/randomx/jit_compiler_rv64_vector_static.S PROPERTIES COMPILE_FLAGS "-march=${RV64_VECTOR_FILE_ARCH}_zvkned")
|
||||
set_source_files_properties(src/crypto/randomx/aes_hash_rv64_vector.cpp PROPERTIES COMPILE_FLAGS "-O3 -march=${RV64_VECTOR_FILE_ARCH}")
|
||||
set_source_files_properties(src/crypto/randomx/aes_hash_rv64_zvkned.cpp PROPERTIES COMPILE_FLAGS "-O3 -march=${RV64_VECTOR_FILE_ARCH}_zvkned")
|
||||
else()
|
||||
|
||||
@@ -363,8 +363,20 @@ void xmrig::CpuWorker<N>::start()
|
||||
}
|
||||
else
|
||||
# endif
|
||||
|
||||
if (value < job.target()) {
|
||||
JobResults::submit(job, current_job_nonces[i], m_hash + (i * 32), job.hasMinerSignature() ? miner_signature_saved : nullptr);
|
||||
uint8_t* extra_data = nullptr;
|
||||
|
||||
if (job.algorithm().family() == Algorithm::RANDOM_X) {
|
||||
if (RandomX_CurrentConfig.Tweak_V2_COMMITMENT) {
|
||||
extra_data = m_commitment;
|
||||
}
|
||||
else if (job.hasMinerSignature()) {
|
||||
extra_data = miner_signature_saved;
|
||||
}
|
||||
}
|
||||
|
||||
JobResults::submit(job, current_job_nonces[i], m_hash + (i * 32), extra_data);
|
||||
}
|
||||
}
|
||||
m_count += N;
|
||||
|
||||
@@ -74,7 +74,7 @@ void keccak_f800_round(uint32_t st[25], const int r)
|
||||
// Keccak - implemented as a variant of SHAKE
|
||||
// The width is 800, with a bitrate of 576, a capacity of 224, and no padding
|
||||
// Only need 64 bits of output for mining
|
||||
void keccak_f800(uint32_t* st)
|
||||
void keccak_f800(uint32_t st[25])
|
||||
{
|
||||
// Complete all 22 rounds as a separate impl to
|
||||
// evaluate only first 8 words is wasteful of registers
|
||||
|
||||
@@ -2,7 +2,7 @@
|
||||
|
||||
namespace xmrig {
|
||||
|
||||
static const char kawpow_cl[5944] = {
|
||||
static const char kawpow_cl[5947] = {
|
||||
0x23,0x69,0x66,0x64,0x65,0x66,0x20,0x63,0x6c,0x5f,0x63,0x6c,0x61,0x6e,0x67,0x5f,0x73,0x74,0x6f,0x72,0x61,0x67,0x65,0x5f,0x63,0x6c,0x61,0x73,0x73,0x5f,0x73,0x70,
|
||||
0x65,0x63,0x69,0x66,0x69,0x65,0x72,0x73,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x4f,0x50,0x45,0x4e,0x43,0x4c,0x20,0x45,0x58,0x54,0x45,0x4e,0x53,0x49,0x4f,
|
||||
0x4e,0x20,0x63,0x6c,0x5f,0x63,0x6c,0x61,0x6e,0x67,0x5f,0x73,0x74,0x6f,0x72,0x61,0x67,0x65,0x5f,0x63,0x6c,0x61,0x73,0x73,0x5f,0x73,0x70,0x65,0x63,0x69,0x66,0x69,
|
||||
@@ -77,118 +77,118 @@ static const char kawpow_cl[5944] = {
|
||||
0x29,0x0a,0x73,0x74,0x5b,0x6a,0x2b,0x69,0x5d,0x20,0x5e,0x3d,0x20,0x28,0x7e,0x62,0x63,0x5b,0x28,0x69,0x2b,0x31,0x29,0x20,0x25,0x20,0x35,0x5d,0x29,0x26,0x62,0x63,
|
||||
0x5b,0x28,0x69,0x2b,0x32,0x29,0x20,0x25,0x20,0x35,0x5d,0x3b,0x0a,0x7d,0x0a,0x73,0x74,0x5b,0x30,0x5d,0x20,0x5e,0x3d,0x20,0x6b,0x65,0x63,0x63,0x61,0x6b,0x66,0x5f,
|
||||
0x72,0x6e,0x64,0x63,0x5b,0x72,0x5d,0x3b,0x0a,0x7d,0x0a,0x76,0x6f,0x69,0x64,0x20,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,0x30,0x30,0x28,0x75,0x69,0x6e,0x74,
|
||||
0x33,0x32,0x5f,0x74,0x2a,0x20,0x73,0x74,0x29,0x0a,0x7b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x72,0x3d,0x30,0x3b,0x20,0x72,0x3c,0x32,0x32,0x3b,0x20,
|
||||
0x72,0x2b,0x2b,0x29,0x20,0x7b,0x0a,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,0x30,0x30,0x5f,0x72,0x6f,0x75,0x6e,0x64,0x28,0x73,0x74,0x2c,0x72,0x29,0x3b,0x0a,
|
||||
0x7d,0x0a,0x7d,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x66,0x6e,0x76,0x31,0x61,0x28,0x68,0x2c,0x20,0x64,0x29,0x20,0x28,0x68,0x20,0x3d,0x20,0x28,0x68,0x20,
|
||||
0x5e,0x20,0x64,0x29,0x20,0x2a,0x20,0x46,0x4e,0x56,0x5f,0x50,0x52,0x49,0x4d,0x45,0x29,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x73,0x74,0x72,0x75,0x63,0x74,
|
||||
0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x7a,0x2c,0x77,0x2c,0x6a,0x73,0x72,0x2c,0x6a,0x63,0x6f,0x6e,0x67,0x3b,0x0a,0x7d,0x20,0x6b,0x69,0x73,
|
||||
0x73,0x39,0x39,0x5f,0x74,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6b,0x69,0x73,0x73,0x39,0x39,0x28,0x6b,0x69,0x73,0x73,0x39,0x39,0x5f,0x74,0x2a,
|
||||
0x20,0x73,0x74,0x29,0x0a,0x7b,0x0a,0x73,0x74,0x2d,0x3e,0x7a,0x3d,0x33,0x36,0x39,0x36,0x39,0x2a,0x28,0x73,0x74,0x2d,0x3e,0x7a,0x26,0x36,0x35,0x35,0x33,0x35,0x29,
|
||||
0x2b,0x28,0x73,0x74,0x2d,0x3e,0x7a,0x3e,0x3e,0x31,0x36,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x77,0x3d,0x31,0x38,0x30,0x30,0x30,0x2a,0x28,0x73,0x74,0x2d,0x3e,0x77,
|
||||
0x26,0x36,0x35,0x35,0x33,0x35,0x29,0x2b,0x28,0x73,0x74,0x2d,0x3e,0x77,0x3e,0x3e,0x31,0x36,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x4d,0x57,
|
||||
0x43,0x3d,0x28,0x28,0x73,0x74,0x2d,0x3e,0x7a,0x3c,0x3c,0x31,0x36,0x29,0x2b,0x73,0x74,0x2d,0x3e,0x77,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x20,0x5e,
|
||||
0x3d,0x20,0x28,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x3c,0x3c,0x31,0x37,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x20,0x5e,0x3d,0x20,0x28,0x73,0x74,0x2d,
|
||||
0x3e,0x6a,0x73,0x72,0x3e,0x3e,0x31,0x33,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x20,0x5e,0x3d,0x20,0x28,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x3c,0x3c,
|
||||
0x35,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x63,0x6f,0x6e,0x67,0x3d,0x36,0x39,0x30,0x36,0x39,0x2a,0x73,0x74,0x2d,0x3e,0x6a,0x63,0x6f,0x6e,0x67,0x2b,0x31,0x32,
|
||||
0x33,0x34,0x35,0x36,0x37,0x3b,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x20,0x28,0x28,0x4d,0x57,0x43,0x5e,0x73,0x74,0x2d,0x3e,0x6a,0x63,0x6f,0x6e,0x67,0x29,0x2b,0x73,
|
||||
0x74,0x2d,0x3e,0x6a,0x73,0x72,0x29,0x3b,0x0a,0x7d,0x0a,0x76,0x6f,0x69,0x64,0x20,0x66,0x69,0x6c,0x6c,0x5f,0x6d,0x69,0x78,0x28,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,
|
||||
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x73,0x65,0x65,0x64,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x2c,0x75,
|
||||
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6d,0x69,0x78,0x29,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,
|
||||
0x68,0x3d,0x46,0x4e,0x56,0x5f,0x4f,0x46,0x46,0x53,0x45,0x54,0x5f,0x42,0x41,0x53,0x49,0x53,0x3b,0x0a,0x6b,0x69,0x73,0x73,0x39,0x39,0x5f,0x74,0x20,0x73,0x74,0x3b,
|
||||
0x0a,0x73,0x74,0x2e,0x7a,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x73,0x65,0x65,0x64,0x5b,0x30,0x5d,0x29,0x3b,0x0a,0x73,
|
||||
0x74,0x2e,0x77,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x73,0x65,0x65,0x64,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x73,0x74,0x2e,
|
||||
0x6a,0x73,0x72,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x29,0x3b,0x0a,0x73,0x74,0x2e,
|
||||
0x6a,0x63,0x6f,0x6e,0x67,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x29,0x3b,0x0a,0x23,
|
||||
0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x50,0x52,
|
||||
0x4f,0x47,0x50,0x4f,0x57,0x5f,0x52,0x45,0x47,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x6d,0x69,0x78,0x5b,0x69,0x5d,0x3d,0x6b,0x69,0x73,0x73,0x39,0x39,0x28,0x26,
|
||||
0x73,0x74,0x29,0x3b,0x0a,0x7d,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x73,0x74,0x72,0x75,0x63,0x74,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,
|
||||
0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x5d,0x3b,0x0a,0x7d,0x20,0x73,0x68,0x75,0x66,0x66,
|
||||
0x6c,0x65,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x73,0x74,0x72,0x75,0x63,0x74,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||
0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x5d,0x3b,0x0a,0x7d,0x20,
|
||||
0x68,0x61,0x73,0x68,0x33,0x32,0x5f,0x74,0x3b,0x0a,0x23,0x69,0x66,0x20,0x50,0x4c,0x41,0x54,0x46,0x4f,0x52,0x4d,0x20,0x21,0x3d,0x20,0x4f,0x50,0x45,0x4e,0x43,0x4c,
|
||||
0x5f,0x50,0x4c,0x41,0x54,0x46,0x4f,0x52,0x4d,0x5f,0x4e,0x56,0x49,0x44,0x49,0x41,0x20,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,
|
||||
0x28,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x28,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,
|
||||
0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,0x23,0x65,0x6e,0x64,0x69,0x66,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x70,0x72,0x6f,
|
||||
0x67,0x70,0x6f,0x77,0x5f,0x73,0x65,0x61,0x72,0x63,0x68,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x64,0x61,0x67,0x5f,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,
|
||||
0x2a,0x20,0x67,0x5f,0x64,0x61,0x67,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x2a,0x20,0x6a,0x6f,0x62,0x5f,0x62,0x6c,0x6f,0x62,0x2c,
|
||||
0x75,0x6c,0x6f,0x6e,0x67,0x20,0x74,0x61,0x72,0x67,0x65,0x74,0x2c,0x75,0x69,0x6e,0x74,0x20,0x68,0x61,0x63,0x6b,0x5f,0x66,0x61,0x6c,0x73,0x65,0x2c,0x76,0x6f,0x6c,
|
||||
0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x2a,0x20,0x72,0x65,0x73,0x75,0x6c,0x74,0x73,0x2c,0x76,0x6f,0x6c,
|
||||
0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x2a,0x20,0x73,0x74,0x6f,0x70,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,
|
||||
0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,
|
||||
0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,
|
||||
0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x73,0x74,0x6f,0x70,0x5b,0x30,0x5d,0x29,0x20,0x7b,0x0a,0x69,0x66,0x28,0x6c,0x69,0x64,0x3d,0x3d,0x30,0x29,0x20,0x7b,
|
||||
0x0a,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x69,0x6e,0x63,0x28,0x73,0x74,0x6f,0x70,0x2b,0x31,0x29,0x3b,0x0a,0x7d,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x3b,0x0a,0x7d,
|
||||
0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x73,0x68,0x75,0x66,0x66,0x6c,0x65,0x5f,0x74,0x20,0x73,0x68,0x61,0x72,0x65,0x5b,0x48,0x41,0x53,0x48,0x45,0x53,0x5f,
|
||||
0x50,0x45,0x52,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x63,0x5f,0x64,
|
||||
0x61,0x67,0x5b,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x41,0x43,0x48,0x45,0x5f,0x57,0x4f,0x52,0x44,0x53,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
|
||||
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x6c,0x69,0x64,0x26,0x28,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,
|
||||
0x45,0x53,0x2d,0x31,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x3d,0x6c,
|
||||
0x69,0x64,0x2f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||
0x77,0x6f,0x72,0x64,0x3d,0x6c,0x69,0x64,0x2a,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x3b,0x20,0x77,0x6f,0x72,0x64,
|
||||
0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x41,0x43,0x48,0x45,0x5f,0x57,0x4f,0x52,0x44,0x53,0x3b,0x20,0x77,0x6f,0x72,0x64,0x2b,0x3d,0x47,0x52,0x4f,0x55,
|
||||
0x50,0x5f,0x53,0x49,0x5a,0x45,0x2a,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x29,0x0a,0x7b,0x0a,0x64,0x61,0x67,0x5f,
|
||||
0x74,0x20,0x6c,0x6f,0x61,0x64,0x3d,0x67,0x5f,0x64,0x61,0x67,0x5b,0x77,0x6f,0x72,0x64,0x2f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x4c,0x4f,
|
||||
0x41,0x44,0x53,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,
|
||||
0x47,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x63,0x5f,0x64,0x61,0x67,0x5b,0x77,0x6f,0x72,0x64,0x2b,0x69,0x5d,0x3d,0x6c,0x6f,0x61,0x64,
|
||||
0x2e,0x73,0x5b,0x69,0x5d,0x3b,0x0a,0x7d,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x68,0x61,0x73,0x68,0x5f,0x73,0x65,0x65,0x64,0x5b,0x32,0x5d,0x3b,0x20,
|
||||
0x0a,0x68,0x61,0x73,0x68,0x33,0x32,0x5f,0x74,0x20,0x64,0x69,0x67,0x65,0x73,0x74,0x3b,0x20,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x74,
|
||||
0x65,0x32,0x5b,0x38,0x5d,0x3b,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x74,0x65,0x5b,0x32,0x35,0x5d,0x3b,0x20,0x0a,0x66,0x6f,
|
||||
0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x31,0x30,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,
|
||||
0x6a,0x6f,0x62,0x5f,0x62,0x6c,0x6f,0x62,0x5b,0x69,0x5d,0x3b,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x38,0x5d,0x3d,0x67,0x69,0x64,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,
|
||||
0x69,0x6e,0x74,0x20,0x69,0x3d,0x31,0x30,0x3b,0x20,0x69,0x3c,0x32,0x35,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x72,0x61,
|
||||
0x76,0x65,0x6e,0x63,0x6f,0x69,0x6e,0x5f,0x72,0x6e,0x64,0x63,0x5b,0x69,0x2d,0x31,0x30,0x5d,0x3b,0x0a,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,0x30,0x30,0x28,
|
||||
0x73,0x74,0x61,0x74,0x65,0x29,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,
|
||||
0x73,0x74,0x61,0x74,0x65,0x32,0x5b,0x69,0x5d,0x3d,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3b,0x0a,0x7d,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,
|
||||
0x72,0x6f,0x6c,0x6c,0x20,0x31,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x68,0x3d,0x30,0x3b,0x20,0x68,0x3c,0x50,0x52,0x4f,0x47,
|
||||
0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x20,0x68,0x2b,0x2b,0x29,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6d,0x69,0x78,0x5b,0x50,
|
||||
0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x52,0x45,0x47,0x53,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x3d,0x68,0x29,0x20,0x7b,0x0a,0x73,
|
||||
0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x30,0x5d,0x3d,0x73,0x74,0x61,0x74,0x65,0x32,
|
||||
0x5b,0x30,0x5d,0x3b,0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x31,0x5d,0x3d,
|
||||
0x73,0x74,0x61,0x74,0x65,0x32,0x5b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,
|
||||
0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x66,0x69,0x6c,0x6c,0x5f,0x6d,0x69,0x78,0x28,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,
|
||||
0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x2c,0x6d,0x69,0x78,0x29,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,
|
||||
0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x20,0x32,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x6f,0x6f,0x70,0x3d,0x30,0x3b,
|
||||
0x20,0x6c,0x6f,0x6f,0x70,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x4e,0x54,0x5f,0x44,0x41,0x47,0x3b,0x20,0x2b,0x2b,0x6c,0x6f,0x6f,0x70,0x29,0x0a,0x7b,
|
||||
0x0a,0x69,0x66,0x28,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x3d,0x28,0x6c,0x6f,0x6f,0x70,0x20,0x25,0x20,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,
|
||||
0x45,0x53,0x29,0x29,0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x30,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x3d,
|
||||
0x6d,0x69,0x78,0x5b,0x30,0x5d,0x3b,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,
|
||||
0x4e,0x43,0x45,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x73,0x68,0x61,0x72,0x65,0x5b,0x30,0x5d,0x2e,0x75,
|
||||
0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x3b,0x0a,0x6f,0x66,0x66,0x73,0x65,0x74,0x20,0x25,0x3d,0x20,0x50,0x52,0x4f,0x47,
|
||||
0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x45,0x4c,0x45,0x4d,0x45,0x4e,0x54,0x53,0x3b,0x0a,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x6f,0x66,0x66,0x73,0x65,0x74,0x2a,
|
||||
0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x2b,0x28,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x5e,0x6c,0x6f,0x6f,0x70,0x29,0x20,0x25,0x20,0x50,
|
||||
0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x0a,0x64,0x61,0x67,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x5f,0x64,0x61,0x67,0x3d,0x67,0x5f,0x64,
|
||||
0x61,0x67,0x5b,0x6f,0x66,0x66,0x73,0x65,0x74,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x68,0x61,0x63,0x6b,0x5f,0x66,0x61,0x6c,0x73,0x65,0x29,0x20,0x62,0x61,0x72,0x72,0x69,
|
||||
0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,
|
||||
0x74,0x20,0x64,0x61,0x74,0x61,0x3b,0x0a,0x58,0x4d,0x52,0x49,0x47,0x5f,0x49,0x4e,0x43,0x4c,0x55,0x44,0x45,0x5f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x52,0x41,
|
||||
0x4e,0x44,0x4f,0x4d,0x5f,0x4d,0x41,0x54,0x48,0x0a,0x69,0x66,0x28,0x68,0x61,0x63,0x6b,0x5f,0x66,0x61,0x6c,0x73,0x65,0x29,0x20,0x62,0x61,0x72,0x72,0x69,0x65,0x72,
|
||||
0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x58,0x4d,0x52,0x49,0x47,0x5f,0x49,0x4e,0x43,
|
||||
0x4c,0x55,0x44,0x45,0x5f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x54,0x41,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x0a,0x7d,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||
0x5f,0x74,0x20,0x6d,0x69,0x78,0x5f,0x68,0x61,0x73,0x68,0x3d,0x46,0x4e,0x56,0x5f,0x4f,0x46,0x46,0x53,0x45,0x54,0x5f,0x42,0x41,0x53,0x49,0x53,0x3b,0x0a,0x23,0x70,
|
||||
0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x50,0x52,0x4f,
|
||||
0x47,0x50,0x4f,0x57,0x5f,0x52,0x45,0x47,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x66,0x6e,0x76,0x31,0x61,0x28,0x6d,0x69,0x78,0x5f,0x68,0x61,0x73,0x68,0x2c,0x6d,
|
||||
0x69,0x78,0x5b,0x69,0x5d,0x29,0x3b,0x0a,0x68,0x61,0x73,0x68,0x33,0x32,0x5f,0x74,0x20,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,0x70,0x3b,0x0a,0x66,0x6f,
|
||||
0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,
|
||||
0x70,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x5d,0x3d,0x46,0x4e,0x56,0x5f,0x4f,0x46,0x46,0x53,0x45,0x54,0x5f,0x42,0x41,0x53,0x49,0x53,0x3b,0x0a,0x73,
|
||||
0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x5d,0x3d,
|
||||
0x6d,0x69,0x78,0x5f,0x68,0x61,0x73,0x68,0x3b,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,
|
||||
0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,
|
||||
0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x66,0x6e,0x76,0x31,0x61,
|
||||
0x28,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,0x70,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x20,0x25,0x20,0x38,0x5d,0x2c,0x73,0x68,0x61,0x72,
|
||||
0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x5d,0x29,0x3b,0x0a,0x69,0x66,0x28,0x68,0x3d,0x3d,0x6c,
|
||||
0x61,0x6e,0x65,0x5f,0x69,0x64,0x29,0x0a,0x64,0x69,0x67,0x65,0x73,0x74,0x3d,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,0x70,0x3b,0x0a,0x7d,0x0a,0x75,0x69,
|
||||
0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x72,0x65,0x73,0x75,0x6c,0x74,0x3b,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x74,0x65,0x5b,
|
||||
0x32,0x35,0x5d,0x3d,0x7b,0x30,0x78,0x30,0x7d,0x3b,0x20,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x69,
|
||||
0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x73,0x74,0x61,0x74,0x65,0x32,0x5b,0x69,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,
|
||||
0x20,0x69,0x3d,0x38,0x3b,0x20,0x69,0x3c,0x31,0x36,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x64,0x69,0x67,0x65,0x73,0x74,
|
||||
0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x2d,0x38,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x31,0x36,0x3b,0x20,0x69,0x3c,
|
||||
0x32,0x35,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x72,0x61,0x76,0x65,0x6e,0x63,0x6f,0x69,0x6e,0x5f,0x72,0x6e,0x64,0x63,
|
||||
0x5b,0x69,0x2d,0x31,0x36,0x5d,0x3b,0x0a,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,0x30,0x30,0x28,0x73,0x74,0x61,0x74,0x65,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,
|
||||
0x36,0x34,0x5f,0x74,0x20,0x72,0x65,0x73,0x3d,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x73,0x74,0x61,0x74,0x65,0x5b,0x31,0x5d,0x3c,0x3c,0x33,0x32,0x7c,
|
||||
0x73,0x74,0x61,0x74,0x65,0x5b,0x30,0x5d,0x3b,0x0a,0x72,0x65,0x73,0x75,0x6c,0x74,0x3d,0x61,0x73,0x5f,0x75,0x6c,0x6f,0x6e,0x67,0x28,0x61,0x73,0x5f,0x75,0x63,0x68,
|
||||
0x61,0x72,0x38,0x28,0x72,0x65,0x73,0x29,0x2e,0x73,0x37,0x36,0x35,0x34,0x33,0x32,0x31,0x30,0x29,0x3b,0x0a,0x7d,0x0a,0x69,0x66,0x28,0x72,0x65,0x73,0x75,0x6c,0x74,
|
||||
0x3c,0x3d,0x74,0x61,0x72,0x67,0x65,0x74,0x29,0x0a,0x7b,0x0a,0x2a,0x73,0x74,0x6f,0x70,0x3d,0x31,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x20,
|
||||
0x6b,0x3d,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x69,0x6e,0x63,0x28,0x72,0x65,0x73,0x75,0x6c,0x74,0x73,0x29,0x2b,0x31,0x3b,0x0a,0x69,0x66,0x28,0x6b,0x3c,0x3d,0x31,
|
||||
0x35,0x29,0x0a,0x72,0x65,0x73,0x75,0x6c,0x74,0x73,0x5b,0x6b,0x5d,0x3d,0x67,0x69,0x64,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x00
|
||||
0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x5b,0x32,0x35,0x5d,0x29,0x0a,0x7b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x72,0x3d,0x30,0x3b,0x20,0x72,0x3c,0x32,
|
||||
0x32,0x3b,0x20,0x72,0x2b,0x2b,0x29,0x20,0x7b,0x0a,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,0x30,0x30,0x5f,0x72,0x6f,0x75,0x6e,0x64,0x28,0x73,0x74,0x2c,0x72,
|
||||
0x29,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x66,0x6e,0x76,0x31,0x61,0x28,0x68,0x2c,0x20,0x64,0x29,0x20,0x28,0x68,0x20,0x3d,0x20,
|
||||
0x28,0x68,0x20,0x5e,0x20,0x64,0x29,0x20,0x2a,0x20,0x46,0x4e,0x56,0x5f,0x50,0x52,0x49,0x4d,0x45,0x29,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x73,0x74,0x72,
|
||||
0x75,0x63,0x74,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x7a,0x2c,0x77,0x2c,0x6a,0x73,0x72,0x2c,0x6a,0x63,0x6f,0x6e,0x67,0x3b,0x0a,0x7d,0x20,
|
||||
0x6b,0x69,0x73,0x73,0x39,0x39,0x5f,0x74,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6b,0x69,0x73,0x73,0x39,0x39,0x28,0x6b,0x69,0x73,0x73,0x39,0x39,
|
||||
0x5f,0x74,0x2a,0x20,0x73,0x74,0x29,0x0a,0x7b,0x0a,0x73,0x74,0x2d,0x3e,0x7a,0x3d,0x33,0x36,0x39,0x36,0x39,0x2a,0x28,0x73,0x74,0x2d,0x3e,0x7a,0x26,0x36,0x35,0x35,
|
||||
0x33,0x35,0x29,0x2b,0x28,0x73,0x74,0x2d,0x3e,0x7a,0x3e,0x3e,0x31,0x36,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x77,0x3d,0x31,0x38,0x30,0x30,0x30,0x2a,0x28,0x73,0x74,
|
||||
0x2d,0x3e,0x77,0x26,0x36,0x35,0x35,0x33,0x35,0x29,0x2b,0x28,0x73,0x74,0x2d,0x3e,0x77,0x3e,0x3e,0x31,0x36,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,
|
||||
0x20,0x4d,0x57,0x43,0x3d,0x28,0x28,0x73,0x74,0x2d,0x3e,0x7a,0x3c,0x3c,0x31,0x36,0x29,0x2b,0x73,0x74,0x2d,0x3e,0x77,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x73,
|
||||
0x72,0x20,0x5e,0x3d,0x20,0x28,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x3c,0x3c,0x31,0x37,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x20,0x5e,0x3d,0x20,0x28,
|
||||
0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x3e,0x3e,0x31,0x33,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x20,0x5e,0x3d,0x20,0x28,0x73,0x74,0x2d,0x3e,0x6a,0x73,
|
||||
0x72,0x3c,0x3c,0x35,0x29,0x3b,0x0a,0x73,0x74,0x2d,0x3e,0x6a,0x63,0x6f,0x6e,0x67,0x3d,0x36,0x39,0x30,0x36,0x39,0x2a,0x73,0x74,0x2d,0x3e,0x6a,0x63,0x6f,0x6e,0x67,
|
||||
0x2b,0x31,0x32,0x33,0x34,0x35,0x36,0x37,0x3b,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x20,0x28,0x28,0x4d,0x57,0x43,0x5e,0x73,0x74,0x2d,0x3e,0x6a,0x63,0x6f,0x6e,0x67,
|
||||
0x29,0x2b,0x73,0x74,0x2d,0x3e,0x6a,0x73,0x72,0x29,0x3b,0x0a,0x7d,0x0a,0x76,0x6f,0x69,0x64,0x20,0x66,0x69,0x6c,0x6c,0x5f,0x6d,0x69,0x78,0x28,0x6c,0x6f,0x63,0x61,
|
||||
0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x73,0x65,0x65,0x64,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x61,0x6e,0x65,0x5f,0x69,
|
||||
0x64,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6d,0x69,0x78,0x29,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x66,0x6e,0x76,0x5f,
|
||||
0x68,0x61,0x73,0x68,0x3d,0x46,0x4e,0x56,0x5f,0x4f,0x46,0x46,0x53,0x45,0x54,0x5f,0x42,0x41,0x53,0x49,0x53,0x3b,0x0a,0x6b,0x69,0x73,0x73,0x39,0x39,0x5f,0x74,0x20,
|
||||
0x73,0x74,0x3b,0x0a,0x73,0x74,0x2e,0x7a,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x73,0x65,0x65,0x64,0x5b,0x30,0x5d,0x29,
|
||||
0x3b,0x0a,0x73,0x74,0x2e,0x77,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x73,0x65,0x65,0x64,0x5b,0x31,0x5d,0x29,0x3b,0x0a,
|
||||
0x73,0x74,0x2e,0x6a,0x73,0x72,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x29,0x3b,0x0a,
|
||||
0x73,0x74,0x2e,0x6a,0x63,0x6f,0x6e,0x67,0x3d,0x66,0x6e,0x76,0x31,0x61,0x28,0x66,0x6e,0x76,0x5f,0x68,0x61,0x73,0x68,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x29,
|
||||
0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,
|
||||
0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x52,0x45,0x47,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x6d,0x69,0x78,0x5b,0x69,0x5d,0x3d,0x6b,0x69,0x73,0x73,0x39,
|
||||
0x39,0x28,0x26,0x73,0x74,0x29,0x3b,0x0a,0x7d,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x73,0x74,0x72,0x75,0x63,0x74,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,
|
||||
0x32,0x5f,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x5d,0x3b,0x0a,0x7d,0x20,0x73,0x68,
|
||||
0x75,0x66,0x66,0x6c,0x65,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x73,0x74,0x72,0x75,0x63,0x74,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||
0x5f,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x5d,0x3b,
|
||||
0x0a,0x7d,0x20,0x68,0x61,0x73,0x68,0x33,0x32,0x5f,0x74,0x3b,0x0a,0x23,0x69,0x66,0x20,0x50,0x4c,0x41,0x54,0x46,0x4f,0x52,0x4d,0x20,0x21,0x3d,0x20,0x4f,0x50,0x45,
|
||||
0x4e,0x43,0x4c,0x5f,0x50,0x4c,0x41,0x54,0x46,0x4f,0x52,0x4d,0x5f,0x4e,0x56,0x49,0x44,0x49,0x41,0x20,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,
|
||||
0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x28,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,
|
||||
0x49,0x5a,0x45,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,0x23,0x65,0x6e,0x64,0x69,0x66,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,
|
||||
0x70,0x72,0x6f,0x67,0x70,0x6f,0x77,0x5f,0x73,0x65,0x61,0x72,0x63,0x68,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x64,0x61,0x67,0x5f,0x74,0x20,0x63,0x6f,
|
||||
0x6e,0x73,0x74,0x2a,0x20,0x67,0x5f,0x64,0x61,0x67,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x2a,0x20,0x6a,0x6f,0x62,0x5f,0x62,0x6c,
|
||||
0x6f,0x62,0x2c,0x75,0x6c,0x6f,0x6e,0x67,0x20,0x74,0x61,0x72,0x67,0x65,0x74,0x2c,0x75,0x69,0x6e,0x74,0x20,0x68,0x61,0x63,0x6b,0x5f,0x66,0x61,0x6c,0x73,0x65,0x2c,
|
||||
0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x2a,0x20,0x72,0x65,0x73,0x75,0x6c,0x74,0x73,0x2c,
|
||||
0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x2a,0x20,0x73,0x74,0x6f,0x70,0x29,0x0a,0x7b,0x0a,
|
||||
0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,
|
||||
0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,
|
||||
0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x73,0x74,0x6f,0x70,0x5b,0x30,0x5d,0x29,0x20,0x7b,0x0a,0x69,0x66,0x28,0x6c,0x69,0x64,0x3d,0x3d,0x30,
|
||||
0x29,0x20,0x7b,0x0a,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x69,0x6e,0x63,0x28,0x73,0x74,0x6f,0x70,0x2b,0x31,0x29,0x3b,0x0a,0x7d,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,
|
||||
0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x73,0x68,0x75,0x66,0x66,0x6c,0x65,0x5f,0x74,0x20,0x73,0x68,0x61,0x72,0x65,0x5b,0x48,0x41,0x53,0x48,
|
||||
0x45,0x53,0x5f,0x50,0x45,0x52,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||
0x63,0x5f,0x64,0x61,0x67,0x5b,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x41,0x43,0x48,0x45,0x5f,0x57,0x4f,0x52,0x44,0x53,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,
|
||||
0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x6c,0x69,0x64,0x26,0x28,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,
|
||||
0x4c,0x41,0x4e,0x45,0x53,0x2d,0x31,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,
|
||||
0x64,0x3d,0x6c,0x69,0x64,0x2f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||
0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x3d,0x6c,0x69,0x64,0x2a,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x3b,0x20,0x77,
|
||||
0x6f,0x72,0x64,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x41,0x43,0x48,0x45,0x5f,0x57,0x4f,0x52,0x44,0x53,0x3b,0x20,0x77,0x6f,0x72,0x64,0x2b,0x3d,0x47,
|
||||
0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2a,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x29,0x0a,0x7b,0x0a,0x64,
|
||||
0x61,0x67,0x5f,0x74,0x20,0x6c,0x6f,0x61,0x64,0x3d,0x67,0x5f,0x64,0x61,0x67,0x5b,0x77,0x6f,0x72,0x64,0x2f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,
|
||||
0x5f,0x4c,0x4f,0x41,0x44,0x53,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,
|
||||
0x5f,0x44,0x41,0x47,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x63,0x5f,0x64,0x61,0x67,0x5b,0x77,0x6f,0x72,0x64,0x2b,0x69,0x5d,0x3d,0x6c,
|
||||
0x6f,0x61,0x64,0x2e,0x73,0x5b,0x69,0x5d,0x3b,0x0a,0x7d,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x68,0x61,0x73,0x68,0x5f,0x73,0x65,0x65,0x64,0x5b,0x32,
|
||||
0x5d,0x3b,0x20,0x0a,0x68,0x61,0x73,0x68,0x33,0x32,0x5f,0x74,0x20,0x64,0x69,0x67,0x65,0x73,0x74,0x3b,0x20,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,
|
||||
0x74,0x61,0x74,0x65,0x32,0x5b,0x38,0x5d,0x3b,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x74,0x65,0x5b,0x32,0x35,0x5d,0x3b,0x20,
|
||||
0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x31,0x30,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,
|
||||
0x69,0x5d,0x3d,0x6a,0x6f,0x62,0x5f,0x62,0x6c,0x6f,0x62,0x5b,0x69,0x5d,0x3b,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x38,0x5d,0x3d,0x67,0x69,0x64,0x3b,0x0a,0x66,0x6f,
|
||||
0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x31,0x30,0x3b,0x20,0x69,0x3c,0x32,0x35,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,
|
||||
0x3d,0x72,0x61,0x76,0x65,0x6e,0x63,0x6f,0x69,0x6e,0x5f,0x72,0x6e,0x64,0x63,0x5b,0x69,0x2d,0x31,0x30,0x5d,0x3b,0x0a,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,
|
||||
0x30,0x30,0x28,0x73,0x74,0x61,0x74,0x65,0x29,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x69,0x2b,
|
||||
0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x32,0x5b,0x69,0x5d,0x3d,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3b,0x0a,0x7d,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,
|
||||
0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x20,0x31,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x68,0x3d,0x30,0x3b,0x20,0x68,0x3c,0x50,
|
||||
0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x20,0x68,0x2b,0x2b,0x29,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6d,0x69,
|
||||
0x78,0x5b,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x52,0x45,0x47,0x53,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x3d,0x68,0x29,0x20,
|
||||
0x7b,0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x30,0x5d,0x3d,0x73,0x74,0x61,
|
||||
0x74,0x65,0x32,0x5b,0x30,0x5d,0x3b,0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,
|
||||
0x31,0x5d,0x3d,0x73,0x74,0x61,0x74,0x65,0x32,0x5b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,
|
||||
0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x66,0x69,0x6c,0x6c,0x5f,0x6d,0x69,0x78,0x28,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,
|
||||
0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x2c,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x2c,0x6d,0x69,0x78,0x29,0x3b,0x0a,0x23,0x70,0x72,
|
||||
0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x20,0x32,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x6f,0x6f,0x70,
|
||||
0x3d,0x30,0x3b,0x20,0x6c,0x6f,0x6f,0x70,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x43,0x4e,0x54,0x5f,0x44,0x41,0x47,0x3b,0x20,0x2b,0x2b,0x6c,0x6f,0x6f,0x70,
|
||||
0x29,0x0a,0x7b,0x0a,0x69,0x66,0x28,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x3d,0x3d,0x28,0x6c,0x6f,0x6f,0x70,0x20,0x25,0x20,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,
|
||||
0x4c,0x41,0x4e,0x45,0x53,0x29,0x29,0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x30,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,
|
||||
0x64,0x5d,0x3d,0x6d,0x69,0x78,0x5b,0x30,0x5d,0x3b,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,
|
||||
0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x73,0x68,0x61,0x72,0x65,0x5b,0x30,
|
||||
0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x3b,0x0a,0x6f,0x66,0x66,0x73,0x65,0x74,0x20,0x25,0x3d,0x20,0x50,
|
||||
0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x47,0x5f,0x45,0x4c,0x45,0x4d,0x45,0x4e,0x54,0x53,0x3b,0x0a,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x6f,0x66,0x66,0x73,
|
||||
0x65,0x74,0x2a,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x2b,0x28,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x5e,0x6c,0x6f,0x6f,0x70,0x29,0x20,
|
||||
0x25,0x20,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x0a,0x64,0x61,0x67,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x5f,0x64,0x61,0x67,0x3d,
|
||||
0x67,0x5f,0x64,0x61,0x67,0x5b,0x6f,0x66,0x66,0x73,0x65,0x74,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x68,0x61,0x63,0x6b,0x5f,0x66,0x61,0x6c,0x73,0x65,0x29,0x20,0x62,0x61,
|
||||
0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,
|
||||
0x33,0x32,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x3b,0x0a,0x58,0x4d,0x52,0x49,0x47,0x5f,0x49,0x4e,0x43,0x4c,0x55,0x44,0x45,0x5f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,
|
||||
0x5f,0x52,0x41,0x4e,0x44,0x4f,0x4d,0x5f,0x4d,0x41,0x54,0x48,0x0a,0x69,0x66,0x28,0x68,0x61,0x63,0x6b,0x5f,0x66,0x61,0x6c,0x73,0x65,0x29,0x20,0x62,0x61,0x72,0x72,
|
||||
0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x58,0x4d,0x52,0x49,0x47,0x5f,
|
||||
0x49,0x4e,0x43,0x4c,0x55,0x44,0x45,0x5f,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x44,0x41,0x54,0x41,0x5f,0x4c,0x4f,0x41,0x44,0x53,0x0a,0x7d,0x0a,0x75,0x69,0x6e,
|
||||
0x74,0x33,0x32,0x5f,0x74,0x20,0x6d,0x69,0x78,0x5f,0x68,0x61,0x73,0x68,0x3d,0x46,0x4e,0x56,0x5f,0x4f,0x46,0x46,0x53,0x45,0x54,0x5f,0x42,0x41,0x53,0x49,0x53,0x3b,
|
||||
0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,
|
||||
0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x52,0x45,0x47,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x66,0x6e,0x76,0x31,0x61,0x28,0x6d,0x69,0x78,0x5f,0x68,0x61,0x73,
|
||||
0x68,0x2c,0x6d,0x69,0x78,0x5b,0x69,0x5d,0x29,0x3b,0x0a,0x68,0x61,0x73,0x68,0x33,0x32,0x5f,0x74,0x20,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,0x70,0x3b,
|
||||
0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,
|
||||
0x74,0x65,0x6d,0x70,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x5d,0x3d,0x46,0x4e,0x56,0x5f,0x4f,0x46,0x46,0x53,0x45,0x54,0x5f,0x42,0x41,0x53,0x49,0x53,
|
||||
0x3b,0x0a,0x73,0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x6c,0x61,0x6e,0x65,0x5f,0x69,
|
||||
0x64,0x5d,0x3d,0x6d,0x69,0x78,0x5f,0x68,0x61,0x73,0x68,0x3b,0x0a,0x62,0x61,0x72,0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x4c,0x4f,0x43,0x41,0x4c,0x5f,0x4d,
|
||||
0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,
|
||||
0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x50,0x52,0x4f,0x47,0x50,0x4f,0x57,0x5f,0x4c,0x41,0x4e,0x45,0x53,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x66,0x6e,
|
||||
0x76,0x31,0x61,0x28,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,0x70,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x20,0x25,0x20,0x38,0x5d,0x2c,0x73,
|
||||
0x68,0x61,0x72,0x65,0x5b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x5d,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x5d,0x29,0x3b,0x0a,0x69,0x66,0x28,0x68,
|
||||
0x3d,0x3d,0x6c,0x61,0x6e,0x65,0x5f,0x69,0x64,0x29,0x0a,0x64,0x69,0x67,0x65,0x73,0x74,0x3d,0x64,0x69,0x67,0x65,0x73,0x74,0x5f,0x74,0x65,0x6d,0x70,0x3b,0x0a,0x7d,
|
||||
0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x72,0x65,0x73,0x75,0x6c,0x74,0x3b,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,
|
||||
0x74,0x65,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x30,0x78,0x30,0x7d,0x3b,0x20,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,
|
||||
0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x73,0x74,0x61,0x74,0x65,0x32,0x5b,0x69,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,
|
||||
0x69,0x6e,0x74,0x20,0x69,0x3d,0x38,0x3b,0x20,0x69,0x3c,0x31,0x36,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x64,0x69,0x67,
|
||||
0x65,0x73,0x74,0x2e,0x75,0x69,0x6e,0x74,0x33,0x32,0x73,0x5b,0x69,0x2d,0x38,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x31,0x36,0x3b,
|
||||
0x20,0x69,0x3c,0x32,0x35,0x3b,0x20,0x69,0x2b,0x2b,0x29,0x0a,0x73,0x74,0x61,0x74,0x65,0x5b,0x69,0x5d,0x3d,0x72,0x61,0x76,0x65,0x6e,0x63,0x6f,0x69,0x6e,0x5f,0x72,
|
||||
0x6e,0x64,0x63,0x5b,0x69,0x2d,0x31,0x36,0x5d,0x3b,0x0a,0x6b,0x65,0x63,0x63,0x61,0x6b,0x5f,0x66,0x38,0x30,0x30,0x28,0x73,0x74,0x61,0x74,0x65,0x29,0x3b,0x0a,0x75,
|
||||
0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x72,0x65,0x73,0x3d,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x73,0x74,0x61,0x74,0x65,0x5b,0x31,0x5d,0x3c,0x3c,
|
||||
0x33,0x32,0x7c,0x73,0x74,0x61,0x74,0x65,0x5b,0x30,0x5d,0x3b,0x0a,0x72,0x65,0x73,0x75,0x6c,0x74,0x3d,0x61,0x73,0x5f,0x75,0x6c,0x6f,0x6e,0x67,0x28,0x61,0x73,0x5f,
|
||||
0x75,0x63,0x68,0x61,0x72,0x38,0x28,0x72,0x65,0x73,0x29,0x2e,0x73,0x37,0x36,0x35,0x34,0x33,0x32,0x31,0x30,0x29,0x3b,0x0a,0x7d,0x0a,0x69,0x66,0x28,0x72,0x65,0x73,
|
||||
0x75,0x6c,0x74,0x3c,0x3d,0x74,0x61,0x72,0x67,0x65,0x74,0x29,0x0a,0x7b,0x0a,0x2a,0x73,0x74,0x6f,0x70,0x3d,0x31,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,
|
||||
0x6e,0x74,0x20,0x6b,0x3d,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x69,0x6e,0x63,0x28,0x72,0x65,0x73,0x75,0x6c,0x74,0x73,0x29,0x2b,0x31,0x3b,0x0a,0x69,0x66,0x28,0x6b,
|
||||
0x3c,0x3d,0x31,0x35,0x29,0x0a,0x72,0x65,0x73,0x75,0x6c,0x74,0x73,0x5b,0x6b,0x5d,0x3d,0x67,0x69,0x64,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x00
|
||||
};
|
||||
|
||||
} // namespace xmrig
|
||||
|
||||
@@ -81,7 +81,7 @@ xmrig::Client::Client(int id, const char *agent, IClientListener *listener) :
|
||||
BaseClient(id, listener),
|
||||
m_agent(agent),
|
||||
m_sendBuf(1024),
|
||||
m_tempBuf(256)
|
||||
m_tempBuf(320)
|
||||
{
|
||||
m_reader.setListener(this);
|
||||
m_key = m_storage.add(this);
|
||||
@@ -199,6 +199,7 @@ int64_t xmrig::Client::submit(const JobResult &result)
|
||||
char *nonce = m_tempBuf.data();
|
||||
char *data = m_tempBuf.data() + 16;
|
||||
char *signature = m_tempBuf.data() + 88;
|
||||
char *commitment = m_tempBuf.data() + 224;
|
||||
|
||||
Cvt::toHex(nonce, sizeof(uint32_t) * 2 + 1, reinterpret_cast<const uint8_t *>(&result.nonce), sizeof(uint32_t));
|
||||
Cvt::toHex(data, 65, result.result(), 32);
|
||||
@@ -206,6 +207,10 @@ int64_t xmrig::Client::submit(const JobResult &result)
|
||||
if (result.minerSignature()) {
|
||||
Cvt::toHex(signature, 129, result.minerSignature(), 64);
|
||||
}
|
||||
|
||||
if (result.commitment()) {
|
||||
Cvt::toHex(commitment, 65, result.commitment(), 32);
|
||||
}
|
||||
# endif
|
||||
|
||||
Document doc(kObjectType);
|
||||
@@ -227,6 +232,16 @@ int64_t xmrig::Client::submit(const JobResult &result)
|
||||
}
|
||||
# endif
|
||||
|
||||
# ifndef XMRIG_PROXY_PROJECT
|
||||
if (result.commitment()) {
|
||||
params.AddMember("commitment", StringRef(commitment), allocator);
|
||||
}
|
||||
# else
|
||||
if (result.commitment) {
|
||||
params.AddMember("commitment", StringRef(result.commitment), allocator);
|
||||
}
|
||||
# endif
|
||||
|
||||
if (has<EXT_ALGO>() && result.algorithm.isValid()) {
|
||||
params.AddMember("algo", StringRef(result.algorithm.name()), allocator);
|
||||
}
|
||||
|
||||
@@ -410,6 +410,7 @@ bool xmrig::DaemonClient::parseJob(const rapidjson::Value ¶ms, int *code)
|
||||
m_blocktemplate.offset(BlockTemplate::TX_EXTRA_NONCE_OFFSET) - k,
|
||||
m_blocktemplate.txExtraNonce().size(),
|
||||
m_blocktemplate.minerTxMerkleTreeBranch(),
|
||||
m_blocktemplate.minerTxMerkleTreePath(),
|
||||
m_blocktemplate.outputType() == 3
|
||||
);
|
||||
# endif
|
||||
|
||||
@@ -269,6 +269,7 @@ void xmrig::Job::copy(const Job &other)
|
||||
m_minerTxExtraNonceOffset = other.m_minerTxExtraNonceOffset;
|
||||
m_minerTxExtraNonceSize = other.m_minerTxExtraNonceSize;
|
||||
m_minerTxMerkleTreeBranch = other.m_minerTxMerkleTreeBranch;
|
||||
m_minerTxMerkleTreePath = other.m_minerTxMerkleTreePath;
|
||||
m_hasViewTag = other.m_hasViewTag;
|
||||
# else
|
||||
memcpy(m_ephPublicKey, other.m_ephPublicKey, sizeof(m_ephPublicKey));
|
||||
@@ -325,6 +326,7 @@ void xmrig::Job::move(Job &&other)
|
||||
m_minerTxExtraNonceOffset = other.m_minerTxExtraNonceOffset;
|
||||
m_minerTxExtraNonceSize = other.m_minerTxExtraNonceSize;
|
||||
m_minerTxMerkleTreeBranch = std::move(other.m_minerTxMerkleTreeBranch);
|
||||
m_minerTxMerkleTreePath = other.m_minerTxMerkleTreePath;
|
||||
m_hasViewTag = other.m_hasViewTag;
|
||||
# else
|
||||
memcpy(m_ephPublicKey, other.m_ephPublicKey, sizeof(m_ephPublicKey));
|
||||
@@ -349,7 +351,7 @@ void xmrig::Job::setSpendSecretKey(const uint8_t *key)
|
||||
}
|
||||
|
||||
|
||||
void xmrig::Job::setMinerTx(const uint8_t *begin, const uint8_t *end, size_t minerTxEphPubKeyOffset, size_t minerTxPubKeyOffset, size_t minerTxExtraNonceOffset, size_t minerTxExtraNonceSize, const Buffer &minerTxMerkleTreeBranch, bool hasViewTag)
|
||||
void xmrig::Job::setMinerTx(const uint8_t *begin, const uint8_t *end, size_t minerTxEphPubKeyOffset, size_t minerTxPubKeyOffset, size_t minerTxExtraNonceOffset, size_t minerTxExtraNonceSize, const Buffer &minerTxMerkleTreeBranch, uint32_t minerTxMerkleTreePath, bool hasViewTag)
|
||||
{
|
||||
m_minerTxPrefix.assign(begin, end);
|
||||
m_minerTxEphPubKeyOffset = minerTxEphPubKeyOffset;
|
||||
@@ -357,6 +359,7 @@ void xmrig::Job::setMinerTx(const uint8_t *begin, const uint8_t *end, size_t min
|
||||
m_minerTxExtraNonceOffset = minerTxExtraNonceOffset;
|
||||
m_minerTxExtraNonceSize = minerTxExtraNonceSize;
|
||||
m_minerTxMerkleTreeBranch = minerTxMerkleTreeBranch;
|
||||
m_minerTxMerkleTreePath = minerTxMerkleTreePath;
|
||||
m_hasViewTag = hasViewTag;
|
||||
}
|
||||
|
||||
@@ -401,7 +404,7 @@ void xmrig::Job::generateHashingBlob(String &blob) const
|
||||
{
|
||||
uint8_t root_hash[32];
|
||||
const uint8_t* p = m_minerTxPrefix.data();
|
||||
BlockTemplate::calculateRootHash(p, p + m_minerTxPrefix.size(), m_minerTxMerkleTreeBranch, root_hash);
|
||||
BlockTemplate::calculateRootHash(p, p + m_minerTxPrefix.size(), m_minerTxMerkleTreeBranch, m_minerTxMerkleTreePath, root_hash);
|
||||
|
||||
uint64_t root_hash_offset = nonceOffset() + nonceSize();
|
||||
|
||||
|
||||
@@ -121,7 +121,7 @@ public:
|
||||
inline bool hasViewTag() const { return m_hasViewTag; }
|
||||
|
||||
void setSpendSecretKey(const uint8_t* key);
|
||||
void setMinerTx(const uint8_t* begin, const uint8_t* end, size_t minerTxEphPubKeyOffset, size_t minerTxPubKeyOffset, size_t minerTxExtraNonceOffset, size_t minerTxExtraNonceSize, const Buffer& minerTxMerkleTreeBranch, bool hasViewTag);
|
||||
void setMinerTx(const uint8_t* begin, const uint8_t* end, size_t minerTxEphPubKeyOffset, size_t minerTxPubKeyOffset, size_t minerTxExtraNonceOffset, size_t minerTxExtraNonceSize, const Buffer& minerTxMerkleTreeBranch, uint32_t minerTxMerkleTreePath, bool hasViewTag);
|
||||
void setViewTagInMinerTx(uint8_t view_tag);
|
||||
void setExtraNonceInMinerTx(uint32_t extra_nonce);
|
||||
void generateSignatureData(String& signatureData, uint8_t& view_tag) const;
|
||||
@@ -179,6 +179,7 @@ private:
|
||||
size_t m_minerTxExtraNonceOffset = 0;
|
||||
size_t m_minerTxExtraNonceSize = 0;
|
||||
Buffer m_minerTxMerkleTreeBranch;
|
||||
uint32_t m_minerTxMerkleTreePath = 0;
|
||||
bool m_hasViewTag = false;
|
||||
# else
|
||||
// Miner signatures
|
||||
|
||||
@@ -48,69 +48,98 @@ void xmrig::BlockTemplate::calculateMinerTxHash(const uint8_t *prefix_begin, con
|
||||
}
|
||||
|
||||
|
||||
void xmrig::BlockTemplate::calculateRootHash(const uint8_t *prefix_begin, const uint8_t *prefix_end, const Buffer &miner_tx_merkle_tree_branch, uint8_t *root_hash)
|
||||
void xmrig::BlockTemplate::calculateRootHash(const uint8_t *prefix_begin, const uint8_t *prefix_end, const Buffer &miner_tx_merkle_tree_branch, uint32_t miner_tx_merkle_tree_path, uint8_t *root_hash)
|
||||
{
|
||||
calculateMinerTxHash(prefix_begin, prefix_end, root_hash);
|
||||
|
||||
for (size_t i = 0; i < miner_tx_merkle_tree_branch.size(); i += kHashSize) {
|
||||
const size_t depth = miner_tx_merkle_tree_branch.size() / kHashSize;
|
||||
|
||||
for (size_t d = 0; d < depth; ++d) {
|
||||
uint8_t h[kHashSize * 2];
|
||||
|
||||
memcpy(h, root_hash, kHashSize);
|
||||
memcpy(h + kHashSize, miner_tx_merkle_tree_branch.data() + i, kHashSize);
|
||||
const uint32_t t = (miner_tx_merkle_tree_path >> (depth - d - 1)) & 1;
|
||||
|
||||
memcpy(h + kHashSize * t, root_hash, kHashSize);
|
||||
memcpy(h + kHashSize * (t ^ 1), miner_tx_merkle_tree_branch.data() + d * kHashSize, kHashSize);
|
||||
|
||||
keccak(h, kHashSize * 2, root_hash, kHashSize);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void xmrig::BlockTemplate::calculateMerkleTreeHash()
|
||||
void xmrig::BlockTemplate::calculateMerkleTreeHash(uint32_t index)
|
||||
{
|
||||
m_minerTxMerkleTreeBranch.clear();
|
||||
m_minerTxMerkleTreePath = 0;
|
||||
|
||||
const uint64_t count = m_numHashes + 1;
|
||||
const size_t count = m_hashes.size() / kHashSize;
|
||||
const uint8_t *h = m_hashes.data();
|
||||
|
||||
if (count == 1) {
|
||||
if (count == 1) {
|
||||
memcpy(m_rootHash, h, kHashSize);
|
||||
}
|
||||
else if (count == 2) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), h + kHashSize, h + kHashSize * 2);
|
||||
}
|
||||
else if (count == 2) {
|
||||
keccak(h, kHashSize * 2, m_rootHash, kHashSize);
|
||||
}
|
||||
else {
|
||||
size_t i = 0;
|
||||
size_t j = 0;
|
||||
size_t cnt = 0;
|
||||
|
||||
for (i = 0, cnt = 1; cnt <= count; ++i, cnt <<= 1) {}
|
||||
m_minerTxMerkleTreeBranch.reserve(1);
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), h + kHashSize * (index ^ 1), h + kHashSize * ((index ^ 1) + 1));
|
||||
m_minerTxMerkleTreePath = static_cast<uint32_t>(index);
|
||||
}
|
||||
else {
|
||||
uint8_t h2[kHashSize];
|
||||
memcpy(h2, h + kHashSize * index, kHashSize);
|
||||
|
||||
cnt >>= 1;
|
||||
size_t cnt = 1, proof_max_size = 0;
|
||||
do {
|
||||
cnt <<= 1;
|
||||
++proof_max_size;
|
||||
} while (cnt <= count);
|
||||
cnt >>= 1;
|
||||
|
||||
m_minerTxMerkleTreeBranch.reserve(kHashSize * (i - 1));
|
||||
m_minerTxMerkleTreeBranch.reserve(proof_max_size);
|
||||
|
||||
Buffer ints(cnt * kHashSize);
|
||||
memcpy(ints.data(), h, (cnt * 2 - count) * kHashSize);
|
||||
|
||||
for (i = cnt * 2 - count, j = cnt * 2 - count; j < cnt; i += 2, ++j) {
|
||||
if (i == 0) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), h + kHashSize, h + kHashSize * 2);
|
||||
}
|
||||
keccak(h + i * kHashSize, kHashSize * 2, ints.data() + j * kHashSize, kHashSize);
|
||||
}
|
||||
const size_t k = cnt * 2 - count;
|
||||
memcpy(ints.data(), h, k * kHashSize);
|
||||
|
||||
while (cnt > 2) {
|
||||
cnt >>= 1;
|
||||
for (i = 0, j = 0; j < cnt; i += 2, ++j) {
|
||||
if (i == 0) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), ints.data() + kHashSize, ints.data() + kHashSize * 2);
|
||||
}
|
||||
keccak(ints.data() + i * kHashSize, kHashSize * 2, ints.data() + j * kHashSize, kHashSize);
|
||||
}
|
||||
}
|
||||
for (size_t i = k, j = k; j < cnt; i += 2, ++j) {
|
||||
keccak(h + i * kHashSize, kHashSize * 2, ints.data() + j * kHashSize, kHashSize);
|
||||
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), ints.data() + kHashSize, ints.data() + kHashSize * 2);
|
||||
keccak(ints.data(), kHashSize * 2, m_rootHash, kHashSize);
|
||||
}
|
||||
if (memcmp(h + i * kHashSize, h2, kHashSize) == 0) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), h + kHashSize * (i + 1), h + kHashSize * (i + 2));
|
||||
memcpy(h2, ints.data() + j * kHashSize, kHashSize);
|
||||
}
|
||||
else if (memcmp(h + (i + 1) * kHashSize, h2, kHashSize) == 0) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), h + kHashSize * i, h + kHashSize * (i + 1));
|
||||
memcpy(h2, ints.data() + j * kHashSize, kHashSize);
|
||||
m_minerTxMerkleTreePath = 1;
|
||||
}
|
||||
}
|
||||
|
||||
while (cnt >= 2) {
|
||||
cnt >>= 1;
|
||||
for (size_t i = 0, j = 0; j < cnt; i += 2, ++j) {
|
||||
uint8_t tmp[kHashSize];
|
||||
keccak(ints.data() + i * kHashSize, kHashSize * 2, tmp, kHashSize);
|
||||
|
||||
if (memcmp(ints.data() + i * kHashSize, h2, kHashSize) == 0) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), ints.data() + kHashSize * (i + 1), ints.data() + kHashSize * (i + 2));
|
||||
memcpy(h2, tmp, kHashSize);
|
||||
m_minerTxMerkleTreePath <<= 1;
|
||||
}
|
||||
else if (memcmp(ints.data() + (i + 1) * kHashSize, h2, kHashSize) == 0) {
|
||||
m_minerTxMerkleTreeBranch.insert(m_minerTxMerkleTreeBranch.end(), ints.data() + kHashSize * i, ints.data() + kHashSize * (i + 1));
|
||||
memcpy(h2, tmp, kHashSize);
|
||||
m_minerTxMerkleTreePath = (m_minerTxMerkleTreePath << 1) | 1;
|
||||
}
|
||||
|
||||
memcpy(ints.data() + j * kHashSize, tmp, kHashSize);
|
||||
}
|
||||
}
|
||||
|
||||
memcpy(m_rootHash, ints.data(), kHashSize);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -375,16 +404,42 @@ bool xmrig::BlockTemplate::parse(bool hashes)
|
||||
ar(m_numHashes);
|
||||
|
||||
if (hashes) {
|
||||
m_hashes.resize((m_numHashes + 1) * kHashSize);
|
||||
calculateMinerTxHash(blob(MINER_TX_PREFIX_OFFSET), blob(MINER_TX_PREFIX_END_OFFSET), m_hashes.data());
|
||||
// FCMP++ layout:
|
||||
//
|
||||
// index 0 fcmp_pp_n_tree_layers + 31 zero bytes
|
||||
// index 1 fcmp_pp_tree_root
|
||||
// index 2 coinbase transaction hash
|
||||
// index 3+ other transaction hashes
|
||||
//
|
||||
// pre-FCMP++ layout:
|
||||
//
|
||||
// index 0 coinbase transaction hash
|
||||
// index 1+ other transaction hashes
|
||||
//
|
||||
const uint32_t coinbase_tx_index = is_fcmp_pp ? 2 : 0;
|
||||
|
||||
m_hashes.clear();
|
||||
m_hashes.resize((coinbase_tx_index + m_numHashes + 1) * kHashSize);
|
||||
|
||||
uint8_t* data = m_hashes.data() + coinbase_tx_index * kHashSize;
|
||||
|
||||
calculateMinerTxHash(blob(MINER_TX_PREFIX_OFFSET), blob(MINER_TX_PREFIX_END_OFFSET), data);
|
||||
|
||||
for (uint64_t i = 1; i <= m_numHashes; ++i) {
|
||||
Span h;
|
||||
ar(h, kHashSize);
|
||||
memcpy(m_hashes.data() + i * kHashSize, h.data(), kHashSize);
|
||||
memcpy(data + i * kHashSize, h.data(), kHashSize);
|
||||
}
|
||||
|
||||
calculateMerkleTreeHash();
|
||||
if (is_fcmp_pp) {
|
||||
ar(m_FCMPTreeLayers);
|
||||
ar(m_FCMPTreeRoot);
|
||||
|
||||
m_hashes[0] = m_FCMPTreeLayers;
|
||||
memcpy(m_hashes.data() + kHashSize, m_FCMPTreeRoot, kHashSize);
|
||||
}
|
||||
|
||||
calculateMerkleTreeHash(coinbase_tx_index);
|
||||
}
|
||||
|
||||
return true;
|
||||
|
||||
@@ -93,6 +93,7 @@ public:
|
||||
inline uint64_t numHashes() const { return m_numHashes; }
|
||||
inline const Buffer &hashes() const { return m_hashes; }
|
||||
inline const Buffer &minerTxMerkleTreeBranch() const { return m_minerTxMerkleTreeBranch; }
|
||||
inline uint32_t minerTxMerkleTreePath() const { return m_minerTxMerkleTreePath; }
|
||||
inline const uint8_t *rootHash() const { return m_rootHash; }
|
||||
|
||||
inline Buffer generateHashingBlob() const
|
||||
@@ -104,13 +105,13 @@ public:
|
||||
}
|
||||
|
||||
static void calculateMinerTxHash(const uint8_t *prefix_begin, const uint8_t *prefix_end, uint8_t *hash);
|
||||
static void calculateRootHash(const uint8_t *prefix_begin, const uint8_t *prefix_end, const Buffer &miner_tx_merkle_tree_branch, uint8_t *root_hash);
|
||||
static void calculateRootHash(const uint8_t *prefix_begin, const uint8_t *prefix_end, const Buffer &miner_tx_merkle_tree_branch, uint32_t miner_tx_merkle_tree_path, uint8_t *root_hash);
|
||||
|
||||
bool parse(const Buffer &blocktemplate, const Coin &coin, bool hashes = kCalcHashes);
|
||||
bool parse(const char *blocktemplate, size_t size, const Coin &coin, bool hashes);
|
||||
bool parse(const rapidjson::Value &blocktemplate, const Coin &coin, bool hashes = kCalcHashes);
|
||||
bool parse(const String &blocktemplate, const Coin &coin, bool hashes = kCalcHashes);
|
||||
void calculateMerkleTreeHash();
|
||||
void calculateMerkleTreeHash(uint32_t index);
|
||||
void generateHashingBlob(Buffer &out) const;
|
||||
|
||||
private:
|
||||
@@ -147,9 +148,12 @@ private:
|
||||
uint64_t m_numHashes = 0;
|
||||
Buffer m_hashes;
|
||||
Buffer m_minerTxMerkleTreeBranch;
|
||||
uint32_t m_minerTxMerkleTreePath = 0;
|
||||
uint8_t m_rootHash[kHashSize]{};
|
||||
uint8_t m_carrotViewTag[3]{};
|
||||
uint8_t m_janusAnchor[16]{};
|
||||
uint8_t m_FCMPTreeLayers = 0;
|
||||
uint8_t m_FCMPTreeRoot[kHashSize]{};
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -59,6 +59,7 @@
|
||||
# include "crypto/rx/Profiler.h"
|
||||
# include "crypto/rx/Rx.h"
|
||||
# include "crypto/rx/RxConfig.h"
|
||||
# include "crypto/rx/RxAlgo.h"
|
||||
#endif
|
||||
|
||||
|
||||
@@ -556,11 +557,12 @@ void xmrig::Miner::setJob(const Job &job, bool donate)
|
||||
}
|
||||
|
||||
# ifdef XMRIG_ALGO_RANDOMX
|
||||
if (job.algorithm().family() == Algorithm::RANDOM_X && !Rx::isReady(job)) {
|
||||
if (job.algorithm().family() == Algorithm::RANDOM_X) {
|
||||
if (d_ptr->algorithm != job.algorithm()) {
|
||||
stop();
|
||||
RxAlgo::apply(job.algorithm());
|
||||
}
|
||||
else {
|
||||
else if (!Rx::isReady(job)) {
|
||||
Nonce::pause(true);
|
||||
Nonce::touch();
|
||||
}
|
||||
@@ -572,6 +574,7 @@ void xmrig::Miner::setJob(const Job &job, bool donate)
|
||||
mutex.lock();
|
||||
|
||||
const uint8_t index = donate ? 1 : 0;
|
||||
const bool same_job_index = d_ptr->job.index() == index;
|
||||
|
||||
d_ptr->reset = !(d_ptr->job.index() == 1 && index == 0 && d_ptr->userJobId == job.id());
|
||||
|
||||
@@ -591,7 +594,8 @@ void xmrig::Miner::setJob(const Job &job, bool donate)
|
||||
const bool ready = d_ptr->initRX();
|
||||
|
||||
// Always reset nonce on RandomX dataset change
|
||||
if (!ready) {
|
||||
// Except for switching to/from donation
|
||||
if (!ready && same_job_index) {
|
||||
d_ptr->reset = true;
|
||||
}
|
||||
# else
|
||||
|
||||
@@ -1,19 +1,8 @@
|
||||
/* XMRig
|
||||
* Copyright (c) 2018-2025 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2025 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
* Copyright (c) 2018-2026 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2026 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* This program is free software: you can redistribute it and/or modify
|
||||
* it under the terms of the GNU General Public License as published by
|
||||
* the Free Software Foundation, either version 3 of the License, or
|
||||
* (at your option) any later version.
|
||||
*
|
||||
* This program is distributed in the hope that it will be useful,
|
||||
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
* GNU General Public License for more details.
|
||||
*
|
||||
* You should have received a copy of the GNU General Public License
|
||||
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||
* SPDX-License-Identifier: GPL-3.0-or-later
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
@@ -31,7 +20,7 @@
|
||||
namespace xmrig {
|
||||
|
||||
|
||||
static const char short_options[] = "a:c:kBp:Px:r:R:s:t:T:o:u:O:v:l:Sx:46";
|
||||
static const char short_options[] = "a:c:kBp:x:r:R:s:t:T:o:u:O:v:l:S46";
|
||||
|
||||
|
||||
static const option options[] = {
|
||||
|
||||
@@ -75,7 +75,7 @@ void hashAes1Rx4(const void *input, size_t inputSize, void *hash)
|
||||
}
|
||||
|
||||
if (xmrig::Cpu::info()->hasRISCV_Vector()) {
|
||||
hashAes1Rx4_RVV<softAes>(input, inputSize, hash);
|
||||
hashAes1Rx4_RVV(input, inputSize, hash);
|
||||
return;
|
||||
}
|
||||
#endif
|
||||
@@ -156,7 +156,7 @@ void fillAes1Rx4(void *state, size_t outputSize, void *buffer)
|
||||
}
|
||||
|
||||
if (xmrig::Cpu::info()->hasRISCV_Vector()) {
|
||||
fillAes1Rx4_RVV<softAes>(state, outputSize, buffer);
|
||||
fillAes1Rx4_RVV(state, outputSize, buffer);
|
||||
return;
|
||||
}
|
||||
#endif
|
||||
@@ -213,7 +213,7 @@ void fillAes4Rx4(void *state, size_t outputSize, void *buffer)
|
||||
}
|
||||
|
||||
if (xmrig::Cpu::info()->hasRISCV_Vector()) {
|
||||
fillAes4Rx4_RVV<softAes>(state, outputSize, buffer);
|
||||
fillAes4Rx4_RVV(state, outputSize, buffer);
|
||||
return;
|
||||
}
|
||||
#endif
|
||||
@@ -297,7 +297,7 @@ void hashAndFillAes1Rx4(void *scratchpad, size_t scratchpadSize, void *hash, voi
|
||||
}
|
||||
|
||||
if (xmrig::Cpu::info()->hasRISCV_Vector()) {
|
||||
hashAndFillAes1Rx4_RVV<softAes, unroll>(scratchpad, scratchpadSize, hash, fill_state);
|
||||
hashAndFillAes1Rx4_RVV(scratchpad, scratchpadSize, hash, fill_state);
|
||||
return;
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -69,7 +69,16 @@ static constexpr uint32_t AES_HASH_1R_XKEY11[8] = { 0x61b263d1, 0x51f4e03c, 0xee
|
||||
static constexpr uint32_t AES_HASH_STRIDE_X2[8] = { 0, 4, 8, 12, 32, 36, 40, 44 };
|
||||
static constexpr uint32_t AES_HASH_STRIDE_X4[8] = { 0, 4, 8, 12, 64, 68, 72, 76 };
|
||||
|
||||
template<int softAes>
|
||||
#define lutEnc0 lutEnc[0]
|
||||
#define lutEnc1 lutEnc[1]
|
||||
#define lutEnc2 lutEnc[2]
|
||||
#define lutEnc3 lutEnc[3]
|
||||
|
||||
#define lutDec0 lutDec[0]
|
||||
#define lutDec1 lutDec[1]
|
||||
#define lutDec2 lutDec[2]
|
||||
#define lutDec3 lutDec[3]
|
||||
|
||||
void hashAes1Rx4_RVV(const void *input, size_t inputSize, void *hash) {
|
||||
const uint8_t* inptr = (const uint8_t*)input;
|
||||
const uint8_t* inputEnd = inptr + inputSize;
|
||||
@@ -113,10 +122,6 @@ void hashAes1Rx4_RVV(const void *input, size_t inputSize, void *hash) {
|
||||
__riscv_vsuxei32_v_u32m1((uint32_t*)hash + 4, stride, state13, 8);
|
||||
}
|
||||
|
||||
template void hashAes1Rx4_RVV<false>(const void *input, size_t inputSize, void *hash);
|
||||
template void hashAes1Rx4_RVV<true>(const void *input, size_t inputSize, void *hash);
|
||||
|
||||
template<int softAes>
|
||||
void fillAes1Rx4_RVV(void *state, size_t outputSize, void *buffer) {
|
||||
const uint8_t* outptr = (uint8_t*)buffer;
|
||||
const uint8_t* outputEnd = outptr + outputSize;
|
||||
@@ -153,10 +158,6 @@ void fillAes1Rx4_RVV(void *state, size_t outputSize, void *buffer) {
|
||||
__riscv_vsuxei32_v_u32m1((uint32_t*)state + 4, stride, state13, 8);
|
||||
}
|
||||
|
||||
template void fillAes1Rx4_RVV<false>(void *state, size_t outputSize, void *buffer);
|
||||
template void fillAes1Rx4_RVV<true>(void *state, size_t outputSize, void *buffer);
|
||||
|
||||
template<int softAes>
|
||||
void fillAes4Rx4_RVV(void *state, size_t outputSize, void *buffer) {
|
||||
const uint8_t* outptr = (uint8_t*)buffer;
|
||||
const uint8_t* outputEnd = outptr + outputSize;
|
||||
@@ -203,10 +204,6 @@ void fillAes4Rx4_RVV(void *state, size_t outputSize, void *buffer) {
|
||||
}
|
||||
}
|
||||
|
||||
template void fillAes4Rx4_RVV<false>(void *state, size_t outputSize, void *buffer);
|
||||
template void fillAes4Rx4_RVV<true>(void *state, size_t outputSize, void *buffer);
|
||||
|
||||
template<int softAes, int unroll>
|
||||
void hashAndFillAes1Rx4_RVV(void *scratchpad, size_t scratchpadSize, void *hash, void* fill_state) {
|
||||
uint8_t* scratchpadPtr = (uint8_t*)scratchpad;
|
||||
const uint8_t* scratchpadEnd = scratchpadPtr + scratchpadSize;
|
||||
@@ -244,54 +241,13 @@ void hashAndFillAes1Rx4_RVV(void *scratchpad, size_t scratchpadSize, void *hash,
|
||||
__riscv_vsuxei32_v_u32m1((uint32_t*)scratchpadPtr + k * 16 + 0, stride, fill_state02, 8); \
|
||||
__riscv_vsuxei32_v_u32m1((uint32_t*)scratchpadPtr + k * 16 + 4, stride, fill_state13, 8);
|
||||
|
||||
switch (softAes) {
|
||||
case 0:
|
||||
HASH_STATE(0);
|
||||
HASH_STATE(1);
|
||||
HASH_STATE(0);
|
||||
HASH_STATE(1);
|
||||
|
||||
FILL_STATE(0);
|
||||
FILL_STATE(1);
|
||||
FILL_STATE(0);
|
||||
FILL_STATE(1);
|
||||
|
||||
scratchpadPtr += 128;
|
||||
break;
|
||||
|
||||
default:
|
||||
switch (unroll) {
|
||||
case 4:
|
||||
HASH_STATE(0);
|
||||
FILL_STATE(0);
|
||||
|
||||
HASH_STATE(1);
|
||||
FILL_STATE(1);
|
||||
|
||||
HASH_STATE(2);
|
||||
FILL_STATE(2);
|
||||
|
||||
HASH_STATE(3);
|
||||
FILL_STATE(3);
|
||||
|
||||
scratchpadPtr += 64 * 4;
|
||||
break;
|
||||
|
||||
case 2:
|
||||
HASH_STATE(0);
|
||||
FILL_STATE(0);
|
||||
|
||||
HASH_STATE(1);
|
||||
FILL_STATE(1);
|
||||
|
||||
scratchpadPtr += 64 * 2;
|
||||
break;
|
||||
|
||||
default:
|
||||
HASH_STATE(0);
|
||||
FILL_STATE(0);
|
||||
|
||||
scratchpadPtr += 64;
|
||||
break;
|
||||
}
|
||||
break;
|
||||
}
|
||||
scratchpadPtr += 128;
|
||||
}
|
||||
|
||||
#undef HASH_STATE
|
||||
@@ -314,9 +270,3 @@ void hashAndFillAes1Rx4_RVV(void *scratchpad, size_t scratchpadSize, void *hash,
|
||||
__riscv_vsuxei32_v_u32m1((uint32_t*)hash + 0, stride, hash_state02, 8);
|
||||
__riscv_vsuxei32_v_u32m1((uint32_t*)hash + 4, stride, hash_state13, 8);
|
||||
}
|
||||
|
||||
template void hashAndFillAes1Rx4_RVV<0,2>(void* scratchpad, size_t scratchpadSize, void* hash, void* fill_state);
|
||||
template void hashAndFillAes1Rx4_RVV<1,1>(void* scratchpad, size_t scratchpadSize, void* hash, void* fill_state);
|
||||
template void hashAndFillAes1Rx4_RVV<2,1>(void* scratchpad, size_t scratchpadSize, void* hash, void* fill_state);
|
||||
template void hashAndFillAes1Rx4_RVV<2,2>(void* scratchpad, size_t scratchpadSize, void* hash, void* fill_state);
|
||||
template void hashAndFillAes1Rx4_RVV<2,4>(void* scratchpad, size_t scratchpadSize, void* hash, void* fill_state);
|
||||
|
||||
@@ -29,14 +29,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
|
||||
#pragma once
|
||||
|
||||
template<int softAes>
|
||||
void hashAes1Rx4_RVV(const void *input, size_t inputSize, void *hash);
|
||||
|
||||
template<int softAes>
|
||||
void fillAes1Rx4_RVV(void *state, size_t outputSize, void *buffer);
|
||||
|
||||
template<int softAes>
|
||||
void fillAes4Rx4_RVV(void *state, size_t outputSize, void *buffer);
|
||||
|
||||
template<int softAes, int unroll>
|
||||
void hashAndFillAes1Rx4_RVV(void *scratchpad, size_t scratchpadSize, void *hash, void* fill_state);
|
||||
|
||||
@@ -13,12 +13,6 @@
|
||||
mov rbp, qword ptr [rsi] ;# "mx", "ma"
|
||||
mov rdi, qword ptr [rsi+8] ;# uint8_t* dataset
|
||||
|
||||
;# dataset prefetch for the first iteration of the main loop
|
||||
mov rax, rbp
|
||||
shr rax, 32
|
||||
and eax, RANDOMX_DATASET_BASE_MASK
|
||||
prefetchnta byte ptr [rdi+rax]
|
||||
|
||||
mov rsi, rdx ;# uint8_t* scratchpad
|
||||
|
||||
mov rax, rbp
|
||||
|
||||
@@ -25,12 +25,6 @@
|
||||
mov rbp, qword ptr [rdx] ;# "mx", "ma"
|
||||
mov rdi, qword ptr [rdx+8] ;# uint8_t* dataset
|
||||
|
||||
;# dataset prefetch for the first iteration of the main loop
|
||||
mov rax, rbp
|
||||
shr rax, 32
|
||||
and eax, RANDOMX_DATASET_BASE_MASK
|
||||
prefetchnta byte ptr [rdi+rax]
|
||||
|
||||
mov rsi, r8 ;# uint8_t* scratchpad
|
||||
mov rbx, r9 ;# loop counter
|
||||
|
||||
|
||||
@@ -420,7 +420,7 @@ inline void* rx_aligned_alloc(size_t size, size_t align) {
|
||||
# define rx_aligned_free(a) free(a)
|
||||
#endif
|
||||
|
||||
inline void rx_prefetch_nta(void* ptr) {
|
||||
inline void rx_prefetch_nta(const void* ptr) {
|
||||
asm volatile ("prfm pldl1strm, [%0]\n" : : "r" (ptr));
|
||||
}
|
||||
|
||||
@@ -577,8 +577,13 @@ inline void* rx_aligned_alloc(size_t size, size_t align) {
|
||||
# define rx_aligned_free(a) free(a)
|
||||
#endif
|
||||
|
||||
#if defined(__GNUC__) && (!defined(__clang__) || __has_builtin(__builtin_prefetch))
|
||||
#define rx_prefetch_nta(x) __builtin_prefetch((x), 0, 0)
|
||||
#define rx_prefetch_t0(x) __builtin_prefetch((x), 0, 3)
|
||||
#else
|
||||
#define rx_prefetch_nta(x)
|
||||
#define rx_prefetch_t0(x)
|
||||
#endif
|
||||
|
||||
FORCE_INLINE rx_vec_f128 rx_load_vec_f128(const double* pd) {
|
||||
rx_vec_f128 x;
|
||||
|
||||
@@ -34,6 +34,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
#include "crypto/randomx/reciprocal.h"
|
||||
#include "crypto/randomx/superscalar.hpp"
|
||||
#include "crypto/randomx/virtual_memory.hpp"
|
||||
#include "crypto/randomx/soft_aes.h"
|
||||
|
||||
static bool hugePagesJIT = false;
|
||||
static int optimizedDatasetInit = -1;
|
||||
@@ -114,7 +115,7 @@ JitCompilerA64::~JitCompilerA64()
|
||||
freePagedMemory(code, allocatedSize);
|
||||
}
|
||||
|
||||
void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& config, uint32_t)
|
||||
void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& config, uint32_t flags)
|
||||
{
|
||||
if (!allocatedSize) {
|
||||
allocate(CodeSize);
|
||||
@@ -125,6 +126,8 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
|
||||
}
|
||||
#endif
|
||||
|
||||
vm_flags = flags;
|
||||
|
||||
uint32_t codePos = MainLoopBegin + 4;
|
||||
|
||||
uint32_t mask = ((RandomX_CurrentConfig.Log2_ScratchpadL3 - 7) << 10);
|
||||
@@ -156,9 +159,9 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
|
||||
emit32(ARMV8A::B | (offset / 4), code, codePos);
|
||||
|
||||
mask = ((RandomX_CurrentConfig.Log2_DatasetBaseSize - 7) << 10);
|
||||
// and w20, w9, CacheLineAlignMask
|
||||
// and w20, w20, CacheLineAlignMask
|
||||
codePos = (((uint8_t*)randomx_program_aarch64_cacheline_align_mask1) - ((uint8_t*)randomx_program_aarch64));
|
||||
emit32(0x121A0000 | 20 | (9 << 5) | mask, code, codePos);
|
||||
emit32(0x121A0000 | 20 | (20 << 5) | mask, code, codePos);
|
||||
|
||||
// and w10, w10, CacheLineAlignMask
|
||||
codePos = (((uint8_t*)randomx_program_aarch64_cacheline_align_mask2) - ((uint8_t*)randomx_program_aarch64));
|
||||
@@ -169,6 +172,43 @@ void JitCompilerA64::generateProgram(Program& program, ProgramConfiguration& con
|
||||
codePos = ((uint8_t*)randomx_program_aarch64_update_spMix1) - ((uint8_t*)randomx_program_aarch64);
|
||||
emit32(ARMV8A::EOR | 10 | (IntRegMap[config.readReg0] << 5) | (IntRegMap[config.readReg1] << 16), code, codePos);
|
||||
|
||||
codePos = ((uint8_t*)randomx_program_aarch64_v2_FE_mix) - ((uint8_t*)randomx_program_aarch64);
|
||||
|
||||
// Enable RandomX v2 AES tweak
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES) {
|
||||
if (flags & RANDOMX_FLAG_HARD_AES) {
|
||||
// Disable the jump to RandomX v1 FE mix code by writing "movi v28.4s, 0" instruction
|
||||
emit32(0x4F00041C, code, codePos);
|
||||
}
|
||||
else {
|
||||
// Jump to RandomX v2 FE mix soft AES code by writing "b randomx_program_aarch64_v2_FE_mix_soft_aes" instruction
|
||||
uint32_t offset = (uint8_t*)randomx_program_aarch64_v2_FE_mix_soft_aes - (uint8_t*)randomx_program_aarch64_v2_FE_mix;
|
||||
emit32(ARMV8A::B | (offset / 4), code, codePos);
|
||||
|
||||
offset = (uint8_t*)randomx_program_aarch64_aes_lut_pointers - (uint8_t*)randomx_program_aarch64;
|
||||
|
||||
*(uint64_t*)(code + offset + 0) = (uint64_t) &lutEnc[0][0];
|
||||
*(uint64_t*)(code + offset + 8) = (uint64_t) &lutDec[0][0];
|
||||
}
|
||||
}
|
||||
else {
|
||||
// Restore the jump to RandomX v1 FE mix code
|
||||
const uint32_t offset = (uint8_t*)randomx_program_aarch64_v1_FE_mix - (uint8_t*)randomx_program_aarch64_v2_FE_mix;
|
||||
emit32(ARMV8A::B | (offset / 4), code, codePos);
|
||||
}
|
||||
|
||||
// Apply v2 prefetch tweak
|
||||
if (RandomX_CurrentConfig.Tweak_V2_PREFETCH) {
|
||||
uint32_t dst = (((uint8_t*)randomx_program_aarch64_vm_instructions_end) - ((uint8_t*)randomx_program_aarch64));
|
||||
uint32_t src = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_v2) - ((uint8_t*)randomx_program_aarch64));
|
||||
memcpy(code + dst, code + src, 16);
|
||||
}
|
||||
else {
|
||||
uint32_t dst = (((uint8_t*)randomx_program_aarch64_vm_instructions_end) - ((uint8_t*)randomx_program_aarch64));
|
||||
uint32_t src = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_v1) - ((uint8_t*)randomx_program_aarch64));
|
||||
memcpy(code + dst, code + src, 16);
|
||||
}
|
||||
|
||||
# ifndef XMRIG_OS_APPLE
|
||||
xmrig::VirtualMemory::flushInstructionCache(reinterpret_cast<char*>(code + MainLoopBegin), codePos - MainLoopBegin);
|
||||
# endif
|
||||
@@ -210,19 +250,56 @@ void JitCompilerA64::generateProgramLight(Program& program, ProgramConfiguration
|
||||
// eor w20, config.readReg2, config.readReg3
|
||||
emit32(ARMV8A::EOR32 | 20 | (IntRegMap[config.readReg2] << 5) | (IntRegMap[config.readReg3] << 16), code, codePos);
|
||||
|
||||
// Apply v2 prefetch tweak
|
||||
if (RandomX_CurrentConfig.Tweak_V2_PREFETCH) {
|
||||
uint32_t dst = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_light_tweak) - ((uint8_t*)randomx_program_aarch64));
|
||||
uint32_t src = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_light_v2) - ((uint8_t*)randomx_program_aarch64));
|
||||
memcpy(code + dst, code + src, 8);
|
||||
}
|
||||
else {
|
||||
uint32_t dst = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_light_tweak) - ((uint8_t*)randomx_program_aarch64));
|
||||
uint32_t src = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_light_v1) - ((uint8_t*)randomx_program_aarch64));
|
||||
memcpy(code + dst, code + src, 8);
|
||||
}
|
||||
|
||||
// Jump back to the main loop
|
||||
const uint32_t offset = (((uint8_t*)randomx_program_aarch64_vm_instructions_end_light) - ((uint8_t*)randomx_program_aarch64)) - codePos;
|
||||
emit32(ARMV8A::B | (offset / 4), code, codePos);
|
||||
|
||||
// and w2, w9, CacheLineAlignMask
|
||||
// and w2, w2, CacheLineAlignMask
|
||||
codePos = (((uint8_t*)randomx_program_aarch64_light_cacheline_align_mask) - ((uint8_t*)randomx_program_aarch64));
|
||||
emit32(0x121A0000 | 2 | (9 << 5) | ((RandomX_CurrentConfig.Log2_DatasetBaseSize - 7) << 10), code, codePos);
|
||||
emit32(0x121A0000 | 2 | (2 << 5) | ((RandomX_CurrentConfig.Log2_DatasetBaseSize - 7) << 10), code, codePos);
|
||||
|
||||
// Update spMix1
|
||||
// eor x10, config.readReg0, config.readReg1
|
||||
codePos = ((uint8_t*)randomx_program_aarch64_update_spMix1) - ((uint8_t*)randomx_program_aarch64);
|
||||
emit32(ARMV8A::EOR | 10 | (IntRegMap[config.readReg0] << 5) | (IntRegMap[config.readReg1] << 16), code, codePos);
|
||||
|
||||
codePos = ((uint8_t*)randomx_program_aarch64_v2_FE_mix) - ((uint8_t*)randomx_program_aarch64);
|
||||
|
||||
// Enable RandomX v2 AES tweak
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES) {
|
||||
if (vm_flags & RANDOMX_FLAG_HARD_AES) {
|
||||
// Disable the jump to RandomX v1 FE mix code by writing "movi v28.4s, 0" instruction
|
||||
emit32(0x4F00041C, code, codePos);
|
||||
}
|
||||
else {
|
||||
// Jump to RandomX v2 FE mix soft AES code by writing "b randomx_program_aarch64_v2_FE_mix_soft_aes" instruction
|
||||
uint32_t offset = (uint8_t*)randomx_program_aarch64_v2_FE_mix_soft_aes - (uint8_t*)randomx_program_aarch64_v2_FE_mix;
|
||||
emit32(ARMV8A::B | (offset / 4), code, codePos);
|
||||
|
||||
offset = (uint8_t*)randomx_program_aarch64_aes_lut_pointers - (uint8_t*)randomx_program_aarch64;
|
||||
|
||||
*(uint64_t*)(code + offset + 0) = (uint64_t) &lutEnc[0][0];
|
||||
*(uint64_t*)(code + offset + 8) = (uint64_t) &lutDec[0][0];
|
||||
}
|
||||
}
|
||||
else {
|
||||
// Restore the jump to RandomX v1 FE mix code
|
||||
const uint32_t offset = (uint8_t*)randomx_program_aarch64_v1_FE_mix - (uint8_t*)randomx_program_aarch64_v2_FE_mix;
|
||||
emit32(ARMV8A::B | (offset / 4), code, codePos);
|
||||
}
|
||||
|
||||
// Apply dataset offset
|
||||
codePos = ((uint8_t*)randomx_program_aarch64_light_dataset_offset) - ((uint8_t*)randomx_program_aarch64);
|
||||
|
||||
@@ -1035,20 +1112,20 @@ void JitCompilerA64::h_CFROUND(Instruction& instr, uint32_t& codePos)
|
||||
constexpr uint32_t tmp_reg = 20;
|
||||
constexpr uint32_t fpcr_tmp_reg = 8;
|
||||
|
||||
if (instr.getImm32() & 63)
|
||||
{
|
||||
// ror tmp_reg, src, imm
|
||||
emit32(ARMV8A::ROR_IMM | tmp_reg | (src << 5) | ((instr.getImm32() & 63) << 10) | (src << 16), code, k);
|
||||
// ror tmp_reg, src, imm
|
||||
emit32(ARMV8A::ROR_IMM | tmp_reg | (src << 5) | ((instr.getImm32() & 63) << 10) | (src << 16), code, k);
|
||||
|
||||
// bfi fpcr_tmp_reg, tmp_reg, 40, 2
|
||||
emit32(0xB3580400 | fpcr_tmp_reg | (tmp_reg << 5), code, k);
|
||||
}
|
||||
else // no rotation
|
||||
{
|
||||
// bfi fpcr_tmp_reg, src, 40, 2
|
||||
emit32(0xB3580400 | fpcr_tmp_reg | (src << 5), code, k);
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
// tst tmp_reg, 60
|
||||
emit32(0xF27E0E9F, code, k);
|
||||
|
||||
// bne next
|
||||
emit32(0x54000081, code, k);
|
||||
}
|
||||
|
||||
// bfi fpcr_tmp_reg, tmp_reg, 40, 2
|
||||
emit32(0xB3580400 | fpcr_tmp_reg | (tmp_reg << 5), code, k);
|
||||
|
||||
// rbit tmp_reg, fpcr_tmp_reg
|
||||
emit32(0xDAC00000 | tmp_reg | (fpcr_tmp_reg << 5), code, k);
|
||||
|
||||
|
||||
@@ -83,6 +83,7 @@ namespace randomx {
|
||||
uint32_t literalPos;
|
||||
uint32_t num32bitLiterals = 0;
|
||||
size_t allocatedSize = 0;
|
||||
uint32_t vm_flags = 0;
|
||||
|
||||
void allocate(size_t size);
|
||||
|
||||
|
||||
@@ -31,7 +31,7 @@
|
||||
#define DECL(x) x
|
||||
#endif
|
||||
|
||||
.arch armv8-a
|
||||
.arch armv8-a+crypto
|
||||
.text
|
||||
.global DECL(randomx_program_aarch64)
|
||||
.global DECL(randomx_program_aarch64_main_loop)
|
||||
@@ -41,9 +41,18 @@
|
||||
.global DECL(randomx_program_aarch64_cacheline_align_mask1)
|
||||
.global DECL(randomx_program_aarch64_cacheline_align_mask2)
|
||||
.global DECL(randomx_program_aarch64_update_spMix1)
|
||||
.global DECL(randomx_program_aarch64_v2_FE_mix)
|
||||
.global DECL(randomx_program_aarch64_v1_FE_mix)
|
||||
.global DECL(randomx_program_aarch64_v2_FE_mix_soft_aes)
|
||||
.global DECL(randomx_program_aarch64_aes_lut_pointers)
|
||||
.global DECL(randomx_program_aarch64_vm_instructions_end_light)
|
||||
.global DECL(randomx_program_aarch64_vm_instructions_end_light_tweak)
|
||||
.global DECL(randomx_program_aarch64_light_cacheline_align_mask)
|
||||
.global DECL(randomx_program_aarch64_light_dataset_offset)
|
||||
.global DECL(randomx_program_aarch64_vm_instructions_end_v1)
|
||||
.global DECL(randomx_program_aarch64_vm_instructions_end_v2)
|
||||
.global DECL(randomx_program_aarch64_vm_instructions_end_light_v1)
|
||||
.global DECL(randomx_program_aarch64_vm_instructions_end_light_v2)
|
||||
.global DECL(randomx_init_dataset_aarch64)
|
||||
.global DECL(randomx_init_dataset_aarch64_end)
|
||||
.global DECL(randomx_calc_dataset_item_aarch64)
|
||||
@@ -242,8 +251,8 @@ DECL(randomx_program_aarch64_main_loop):
|
||||
# Execute VM instructions
|
||||
DECL(randomx_program_aarch64_vm_instructions):
|
||||
|
||||
# 16 KB buffer for generated instructions
|
||||
.fill 4096,4,0
|
||||
# 24 KB buffer for generated instructions
|
||||
.fill 6144,4,0
|
||||
|
||||
literal_x0: .fill 1,8,0
|
||||
literal_x11: .fill 1,8,0
|
||||
@@ -285,17 +294,19 @@ DECL(randomx_program_aarch64_vm_instructions_end):
|
||||
eor x9, x9, x20
|
||||
|
||||
# Calculate dataset pointer for dataset prefetch
|
||||
mov w20, w9
|
||||
|
||||
# mx <-> ma
|
||||
ror x9, x9, 32
|
||||
|
||||
DECL(randomx_program_aarch64_cacheline_align_mask1):
|
||||
# Actual mask will be inserted by JIT compiler
|
||||
and x20, x9, 1
|
||||
and x20, x20, 1
|
||||
add x20, x20, x1
|
||||
|
||||
# Prefetch dataset data
|
||||
prfm pldl2strm, [x20]
|
||||
|
||||
# mx <-> ma
|
||||
ror x9, x9, 32
|
||||
|
||||
DECL(randomx_program_aarch64_cacheline_align_mask2):
|
||||
# Actual mask will be inserted by JIT compiler
|
||||
and x10, x10, 1
|
||||
@@ -326,12 +337,93 @@ DECL(randomx_program_aarch64_update_spMix1):
|
||||
stp x12, x13, [x17, 32]
|
||||
stp x14, x15, [x17, 48]
|
||||
|
||||
# xor group F and group E registers
|
||||
# RandomX v2 AES tweak (mix group F and group E registers using AES)
|
||||
DECL(randomx_program_aarch64_v2_FE_mix):
|
||||
|
||||
# Jump to v1 FE mix code if we're running RandomX v1
|
||||
# JIT compiler will write a "movi v28.4s, 0" (set v28 to all 0) here if we're running RandomX v2
|
||||
# Or, JIT compiler will write a "b randomx_program_aarch64_v2_FE_mix_soft_aes" if we're running RandomX v2 with soft AES
|
||||
b DECL(randomx_program_aarch64_v1_FE_mix)
|
||||
|
||||
# f0 = aesenc(f0, e0), f1 = aesdec(f1, e0), f2 = aesenc(f2, e0), f3 = aesdec(f3, e0)
|
||||
|
||||
aese v16.16b, v28.16b
|
||||
aesd v17.16b, v28.16b
|
||||
aese v18.16b, v28.16b
|
||||
aesd v19.16b, v28.16b
|
||||
|
||||
aesmc v16.16b, v16.16b
|
||||
aesimc v17.16b, v17.16b
|
||||
aesmc v18.16b, v18.16b
|
||||
aesimc v19.16b, v19.16b
|
||||
|
||||
eor v16.16b, v16.16b, v20.16b
|
||||
eor v17.16b, v17.16b, v20.16b
|
||||
eor v18.16b, v18.16b, v20.16b
|
||||
eor v19.16b, v19.16b, v20.16b
|
||||
|
||||
# f0 = aesenc(f0, e1), f1 = aesdec(f1, e1), f2 = aesenc(f2, e1), f3 = aesdec(f3, e1)
|
||||
|
||||
aese v16.16b, v28.16b
|
||||
aesd v17.16b, v28.16b
|
||||
aese v18.16b, v28.16b
|
||||
aesd v19.16b, v28.16b
|
||||
|
||||
aesmc v16.16b, v16.16b
|
||||
aesimc v17.16b, v17.16b
|
||||
aesmc v18.16b, v18.16b
|
||||
aesimc v19.16b, v19.16b
|
||||
|
||||
eor v16.16b, v16.16b, v21.16b
|
||||
eor v17.16b, v17.16b, v21.16b
|
||||
eor v18.16b, v18.16b, v21.16b
|
||||
eor v19.16b, v19.16b, v21.16b
|
||||
|
||||
# f0 = aesenc(f0, e2), f1 = aesdec(f1, e2), f2 = aesenc(f2, e2), f3 = aesdec(f3, e2)
|
||||
|
||||
aese v16.16b, v28.16b
|
||||
aesd v17.16b, v28.16b
|
||||
aese v18.16b, v28.16b
|
||||
aesd v19.16b, v28.16b
|
||||
|
||||
aesmc v16.16b, v16.16b
|
||||
aesimc v17.16b, v17.16b
|
||||
aesmc v18.16b, v18.16b
|
||||
aesimc v19.16b, v19.16b
|
||||
|
||||
eor v16.16b, v16.16b, v22.16b
|
||||
eor v17.16b, v17.16b, v22.16b
|
||||
eor v18.16b, v18.16b, v22.16b
|
||||
eor v19.16b, v19.16b, v22.16b
|
||||
|
||||
# f0 = aesenc(f0, e3), f1 = aesdec(f1, e3), f2 = aesenc(f2, e3), f3 = aesdec(f3, e3)
|
||||
|
||||
aese v16.16b, v28.16b
|
||||
aesd v17.16b, v28.16b
|
||||
aese v18.16b, v28.16b
|
||||
aesd v19.16b, v28.16b
|
||||
|
||||
aesmc v16.16b, v16.16b
|
||||
aesimc v17.16b, v17.16b
|
||||
aesmc v18.16b, v18.16b
|
||||
aesimc v19.16b, v19.16b
|
||||
|
||||
eor v16.16b, v16.16b, v23.16b
|
||||
eor v17.16b, v17.16b, v23.16b
|
||||
eor v18.16b, v18.16b, v23.16b
|
||||
eor v19.16b, v19.16b, v23.16b
|
||||
|
||||
# Skip v1 FE mix code because we already did v2 FE mix
|
||||
b randomx_program_aarch64_FE_store
|
||||
|
||||
DECL(randomx_program_aarch64_v1_FE_mix):
|
||||
eor v16.16b, v16.16b, v20.16b
|
||||
eor v17.16b, v17.16b, v21.16b
|
||||
eor v18.16b, v18.16b, v22.16b
|
||||
eor v19.16b, v19.16b, v23.16b
|
||||
|
||||
randomx_program_aarch64_FE_store:
|
||||
|
||||
# Store FP registers to scratchpad (spAddr0)
|
||||
stp q16, q17, [x16, 0]
|
||||
stp q18, q19, [x16, 32]
|
||||
@@ -376,6 +468,13 @@ DECL(randomx_program_aarch64_vm_instructions_end_light):
|
||||
stp x0, x1, [sp, 64]
|
||||
stp x2, x30, [sp, 80]
|
||||
|
||||
lsr x2, x9, 32
|
||||
|
||||
DECL(randomx_program_aarch64_light_cacheline_align_mask):
|
||||
# Actual mask will be inserted by JIT compiler
|
||||
and w2, w2, 1
|
||||
|
||||
DECL(randomx_program_aarch64_vm_instructions_end_light_tweak):
|
||||
# mx ^= r[readReg2] ^ r[readReg3];
|
||||
eor x9, x9, x20
|
||||
|
||||
@@ -388,10 +487,6 @@ DECL(randomx_program_aarch64_vm_instructions_end_light):
|
||||
# x1 -> pointer to output
|
||||
mov x1, sp
|
||||
|
||||
DECL(randomx_program_aarch64_light_cacheline_align_mask):
|
||||
# Actual mask will be inserted by JIT compiler
|
||||
and w2, w9, 1
|
||||
|
||||
# x2 -> item number
|
||||
lsr x2, x2, 6
|
||||
|
||||
@@ -409,6 +504,237 @@ DECL(randomx_program_aarch64_light_dataset_offset):
|
||||
|
||||
b DECL(randomx_program_aarch64_xor_with_dataset_line)
|
||||
|
||||
DECL(randomx_program_aarch64_vm_instructions_end_v1):
|
||||
lsr x10, x9, 32
|
||||
eor x9, x9, x20
|
||||
mov w20, w9
|
||||
ror x9, x9, 32
|
||||
|
||||
DECL(randomx_program_aarch64_vm_instructions_end_v2):
|
||||
lsr x10, x9, 32
|
||||
ror x9, x9, 32
|
||||
eor x9, x9, x20
|
||||
mov w20, w9
|
||||
|
||||
DECL(randomx_program_aarch64_vm_instructions_end_light_v1):
|
||||
eor x9, x9, x20
|
||||
ror x9, x9, 32
|
||||
|
||||
DECL(randomx_program_aarch64_vm_instructions_end_light_v2):
|
||||
ror x9, x9, 32
|
||||
eor x9, x9, x20
|
||||
|
||||
DECL(randomx_program_aarch64_v2_FE_mix_soft_aes):
|
||||
sub sp, sp, 176
|
||||
|
||||
stp x0, x1, [sp]
|
||||
stp x2, x3, [sp, 16]
|
||||
stp x4, x5, [sp, 32]
|
||||
stp x6, x7, [sp, 48]
|
||||
stp x8, x9, [sp, 64]
|
||||
stp x10, x11, [sp, 80]
|
||||
stp x12, x13, [sp, 96]
|
||||
stp x14, x15, [sp, 112]
|
||||
stp x16, x30, [sp, 128]
|
||||
stp q0, q1, [sp, 144]
|
||||
|
||||
adr x19, DECL(randomx_program_aarch64_aes_lut_pointers)
|
||||
ldp x19, x20, [x19]
|
||||
|
||||
# f0 = aesenc(f0, e0), f0 = aesenc(f0, e1), f0 = aesenc(f0, e2), f0 = aesenc(f0, e3)
|
||||
mov v0.16b, v16.16b
|
||||
mov v1.16b, v20.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v1.16b, v21.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v1.16b, v22.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v1.16b, v23.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v16.16b, v0.16b
|
||||
|
||||
# f1 = aesdec(f1, e0), f1 = aesdec(f1, e1), f1 = aesdec(f1, e2), f1 = aesdec(f1, e3)
|
||||
mov v0.16b, v17.16b
|
||||
mov v1.16b, v20.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v1.16b, v21.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v1.16b, v22.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v1.16b, v23.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v17.16b, v0.16b
|
||||
|
||||
# f2 = aesenc(f2, e0), f2 = aesenc(f2, e1), f2 = aesenc(f2, e2), f2 = aesenc(f2, e3)
|
||||
mov v0.16b, v18.16b
|
||||
mov v1.16b, v20.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v1.16b, v21.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v1.16b, v22.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v1.16b, v23.16b
|
||||
bl randomx_soft_aesenc
|
||||
mov v18.16b, v0.16b
|
||||
|
||||
# f3 = aesdec(f3, e0), f3 = aesdec(f3, e1), f3 = aesdec(f3, e2), f3 = aesdec(f3, e3)
|
||||
mov v0.16b, v19.16b
|
||||
mov v1.16b, v20.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v1.16b, v21.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v1.16b, v22.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v1.16b, v23.16b
|
||||
bl randomx_soft_aesdec
|
||||
mov v19.16b, v0.16b
|
||||
|
||||
ldp x0, x1, [sp]
|
||||
ldp x2, x3, [sp, 16]
|
||||
ldp x4, x5, [sp, 32]
|
||||
ldp x6, x7, [sp, 48]
|
||||
ldp x8, x9, [sp, 64]
|
||||
ldp x10, x11, [sp, 80]
|
||||
ldp x12, x13, [sp, 96]
|
||||
ldp x14, x15, [sp, 112]
|
||||
ldp x16, x30, [sp, 128]
|
||||
ldp q0, q1, [sp, 144]
|
||||
|
||||
add sp, sp, 176
|
||||
|
||||
b randomx_program_aarch64_FE_store
|
||||
|
||||
|
||||
randomx_soft_aesenc:
|
||||
umov w4, v0.b[5]
|
||||
umov w1, v0.b[10]
|
||||
umov w12, v0.b[15]
|
||||
umov w9, v0.b[9]
|
||||
umov w2, v0.b[14]
|
||||
umov w11, v0.b[3]
|
||||
umov w5, v0.b[0]
|
||||
umov w16, v0.b[4]
|
||||
add x4, x4, 256
|
||||
add x1, x1, 512
|
||||
add x12, x12, 768
|
||||
umov w3, v0.b[13]
|
||||
umov w8, v0.b[2]
|
||||
umov w7, v0.b[7]
|
||||
add x9, x9, 256
|
||||
add x2, x2, 512
|
||||
add x11, x11, 768
|
||||
ldr w10, [x19, x4, lsl 2]
|
||||
ldr w15, [x19, x5, lsl 2]
|
||||
umov w13, v0.b[8]
|
||||
ldr w14, [x19, x12, lsl 2]
|
||||
umov w6, v0.b[1]
|
||||
ldr w1, [x19, x1, lsl 2]
|
||||
eor w10, w10, w15
|
||||
ldr w2, [x19, x2, lsl 2]
|
||||
umov w5, v0.b[6]
|
||||
ldr w9, [x19, x9, lsl 2]
|
||||
umov w4, v0.b[11]
|
||||
ldr w12, [x19, x16, lsl 2]
|
||||
eor w1, w1, w14
|
||||
ldr w11, [x19, x11, lsl 2]
|
||||
eor w1, w1, w10
|
||||
add x8, x8, 512
|
||||
add x3, x3, 256
|
||||
add x7, x7, 768
|
||||
eor w9, w9, w12
|
||||
fmov s28, w1
|
||||
eor w1, w2, w11
|
||||
umov w10, v0.b[12]
|
||||
eor w1, w1, w9
|
||||
ldr w3, [x19, x3, lsl 2]
|
||||
add x6, x6, 256
|
||||
ldr w9, [x19, x13, lsl 2]
|
||||
ins v28.s[1], w1
|
||||
ldr w2, [x19, x8, lsl 2]
|
||||
add x5, x5, 512
|
||||
ldr w7, [x19, x7, lsl 2]
|
||||
add x4, x4, 768
|
||||
eor w1, w3, w9
|
||||
ldr w3, [x19, x6, lsl 2]
|
||||
eor w2, w2, w7
|
||||
ldr w6, [x19, x10, lsl 2]
|
||||
eor w2, w2, w1
|
||||
ldr w1, [x19, x5, lsl 2]
|
||||
ldr w0, [x19, x4, lsl 2]
|
||||
eor w3, w3, w6
|
||||
ins v28.s[2], w2
|
||||
eor w0, w1, w0
|
||||
eor w0, w0, w3
|
||||
ins v28.s[3], w0
|
||||
eor v0.16b, v1.16b, v28.16b
|
||||
ret
|
||||
|
||||
randomx_soft_aesdec:
|
||||
umov w1, v0.b[10]
|
||||
umov w3, v0.b[7]
|
||||
umov w12, v0.b[13]
|
||||
umov w2, v0.b[14]
|
||||
umov w9, v0.b[11]
|
||||
umov w11, v0.b[1]
|
||||
umov w4, v0.b[0]
|
||||
umov w16, v0.b[4]
|
||||
add x3, x3, 768
|
||||
add x1, x1, 512
|
||||
add x12, x12, 256
|
||||
umov w8, v0.b[5]
|
||||
umov w6, v0.b[2]
|
||||
umov w7, v0.b[15]
|
||||
add x9, x9, 768
|
||||
add x2, x2, 512
|
||||
add x11, x11, 256
|
||||
ldr w15, [x20, x3, lsl 2]
|
||||
ldr w10, [x20, x4, lsl 2]
|
||||
umov w13, v0.b[8]
|
||||
ldr w14, [x20, x12, lsl 2]
|
||||
umov w5, v0.b[9]
|
||||
ldr w1, [x20, x1, lsl 2]
|
||||
umov w3, v0.b[6]
|
||||
ldr w12, [x20, x9, lsl 2]
|
||||
umov w4, v0.b[3]
|
||||
ldr w9, [x20, x16, lsl 2]
|
||||
eor w1, w1, w15
|
||||
ldr w2, [x20, x2, lsl 2]
|
||||
eor w10, w10, w14
|
||||
ldr w11, [x20, x11, lsl 2]
|
||||
eor w1, w1, w10
|
||||
add x8, x8, 256
|
||||
add x6, x6, 512
|
||||
add x7, x7, 768
|
||||
eor w2, w2, w12
|
||||
fmov s28, w1
|
||||
eor w1, w9, w11
|
||||
eor w1, w2, w1
|
||||
umov w9, v0.b[12]
|
||||
ldr w2, [x20, x13, lsl 2]
|
||||
add x5, x5, 256
|
||||
ldr w8, [x20, x8, lsl 2]
|
||||
ins v28.s[1], w1
|
||||
ldr w6, [x20, x6, lsl 2]
|
||||
add x3, x3, 512
|
||||
ldr w7, [x20, x7, lsl 2]
|
||||
add x4, x4, 768
|
||||
eor w2, w2, w8
|
||||
ldr w1, [x20, x9, lsl 2]
|
||||
eor w6, w6, w7
|
||||
ldr w3, [x20, x3, lsl 2]
|
||||
eor w2, w2, w6
|
||||
ldr w4, [x20, x4, lsl 2]
|
||||
ldr w5, [x20, x5, lsl 2]
|
||||
ins v28.s[2], w2
|
||||
eor w0, w1, w5
|
||||
eor w1, w3, w4
|
||||
eor w0, w0, w1
|
||||
ins v28.s[3], w0
|
||||
eor v0.16b, v1.16b, v28.16b
|
||||
ret
|
||||
|
||||
DECL(randomx_program_aarch64_aes_lut_pointers):
|
||||
.fill 2, 8, 0
|
||||
|
||||
|
||||
# Input parameters
|
||||
|
||||
@@ -38,9 +38,18 @@ extern "C" {
|
||||
void randomx_program_aarch64_cacheline_align_mask1();
|
||||
void randomx_program_aarch64_cacheline_align_mask2();
|
||||
void randomx_program_aarch64_update_spMix1();
|
||||
void randomx_program_aarch64_v2_FE_mix();
|
||||
void randomx_program_aarch64_v1_FE_mix();
|
||||
void randomx_program_aarch64_v2_FE_mix_soft_aes();
|
||||
void randomx_program_aarch64_aes_lut_pointers();
|
||||
void randomx_program_aarch64_vm_instructions_end_light();
|
||||
void randomx_program_aarch64_vm_instructions_end_light_tweak();
|
||||
void randomx_program_aarch64_light_cacheline_align_mask();
|
||||
void randomx_program_aarch64_light_dataset_offset();
|
||||
void randomx_program_aarch64_vm_instructions_end_v1();
|
||||
void randomx_program_aarch64_vm_instructions_end_v2();
|
||||
void randomx_program_aarch64_vm_instructions_end_light_v1();
|
||||
void randomx_program_aarch64_vm_instructions_end_light_v2();
|
||||
void randomx_init_dataset_aarch64();
|
||||
void randomx_init_dataset_aarch64_end();
|
||||
void randomx_calc_dataset_item_aarch64();
|
||||
|
||||
@@ -39,6 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
#include "crypto/randomx/program.hpp"
|
||||
#include "crypto/randomx/reciprocal.h"
|
||||
#include "crypto/randomx/virtual_memory.hpp"
|
||||
#include "crypto/randomx/soft_aes.h"
|
||||
#include "crypto/common/VirtualMemory.h"
|
||||
|
||||
|
||||
@@ -253,10 +254,12 @@ namespace randomx {
|
||||
static const uint8_t* codePrologue = (uint8_t*)&randomx_riscv64_prologue;
|
||||
static const uint8_t* codeLoopBegin = (uint8_t*)&randomx_riscv64_loop_begin;
|
||||
static const uint8_t* codeDataRead = (uint8_t*)&randomx_riscv64_data_read;
|
||||
static const uint8_t* codeDataRead2 = (uint8_t*)&randomx_riscv64_data_read_v2_tweak;
|
||||
static const uint8_t* codeDataReadLight = (uint8_t*)&randomx_riscv64_data_read_light;
|
||||
static const uint8_t* codeDataReadLight1 = (uint8_t*)&randomx_riscv64_data_read_light_v1;
|
||||
static const uint8_t* codeDataReadLight2 = (uint8_t*)&randomx_riscv64_data_read_light_v2;
|
||||
static const uint8_t* codeFixLoopCall = (uint8_t*)&randomx_riscv64_fix_loop_call;
|
||||
static const uint8_t* codeSpadStore = (uint8_t*)&randomx_riscv64_spad_store;
|
||||
static const uint8_t* codeSpadStoreHardAes = (uint8_t*)&randomx_riscv64_spad_store_hardaes;
|
||||
static const uint8_t* codeSpadStoreSoftAes = (uint8_t*)&randomx_riscv64_spad_store_softaes;
|
||||
static const uint8_t* codeLoopEnd = (uint8_t*)&randomx_riscv64_loop_end;
|
||||
static const uint8_t* codeFixContinueLoop = (uint8_t*)&randomx_riscv64_fix_continue_loop;
|
||||
@@ -272,9 +275,13 @@ namespace randomx {
|
||||
static const int32_t sizeDataInit = codePrologue - codeDataInit;
|
||||
static const int32_t sizePrologue = codeLoopBegin - codePrologue;
|
||||
static const int32_t sizeLoopBegin = codeDataRead - codeLoopBegin;
|
||||
static const int32_t sizeDataRead = codeDataReadLight - codeDataRead;
|
||||
static const int32_t sizeDataReadLight = codeSpadStore - codeDataReadLight;
|
||||
static const int32_t sizeSpadStore = codeSpadStoreHardAes - codeSpadStore;
|
||||
static const int32_t sizeDataRead = codeDataRead2 - codeDataRead;
|
||||
static const int32_t sizeDataRead2 = codeDataReadLight - codeDataRead2;
|
||||
static const int32_t sizeDataReadLight = codeDataReadLight1 - codeDataReadLight;
|
||||
static const int32_t sizeDataReadLight1 = codeDataReadLight2 - codeDataReadLight1;
|
||||
static const int32_t sizeDataReadLight2 = codeFixLoopCall - codeDataReadLight2;
|
||||
static const int32_t sizeFixLoopCall = codeSpadStore - codeFixLoopCall;
|
||||
static const int32_t sizeSpadStore = codeSpadStoreSoftAes - codeSpadStore;
|
||||
static const int32_t sizeSpadStoreSoftAes = codeLoopEnd - codeSpadStoreSoftAes;
|
||||
static const int32_t sizeLoopEnd = codeEpilogue - codeLoopEnd;
|
||||
static const int32_t sizeEpilogue = codeSoftAes - codeEpilogue;
|
||||
@@ -284,7 +291,6 @@ namespace randomx {
|
||||
static const int32_t sizeSshPrefetch = codeSshEnd - codeSshPrefetch;
|
||||
|
||||
static const int32_t offsetFixDataCall = codeFixDataCall - codeDataInit;
|
||||
static const int32_t offsetFixLoopCall = codeFixLoopCall - codeDataReadLight;
|
||||
static const int32_t offsetFixContinueLoop = codeFixContinueLoop - codeLoopEnd;
|
||||
|
||||
static const int32_t LoopTopPos = LiteralPoolSize + sizeDataInit + sizePrologue;
|
||||
@@ -479,8 +485,15 @@ namespace randomx {
|
||||
static void emitProgramPrefix(CompilerState& state, Program& prog, ProgramConfiguration& pcfg) {
|
||||
state.codePos = RandomXCodePos;
|
||||
state.rcpCount = 0;
|
||||
|
||||
state.emitAt(LiteralPoolOffset + sizeLiterals, pcfg.eMask[0]);
|
||||
state.emitAt(LiteralPoolOffset + sizeLiterals + 8, pcfg.eMask[1]);
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES) {
|
||||
state.emitAt(LiteralPoolOffset + sizeLiterals + 16, (uint64_t) &lutEnc[2][0]);
|
||||
state.emitAt(LiteralPoolOffset + sizeLiterals + 24, (uint64_t) &lutDec[2][0]);
|
||||
}
|
||||
|
||||
for (unsigned i = 0; i < RegistersCount; ++i) {
|
||||
state.registerUsage[i] = -1;
|
||||
}
|
||||
@@ -493,7 +506,13 @@ namespace randomx {
|
||||
}
|
||||
|
||||
static void emitProgramSuffix(CompilerState& state, ProgramConfiguration& pcfg) {
|
||||
state.emit(codeSpadStore, sizeSpadStore);
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES) {
|
||||
state.emit(codeSpadStoreSoftAes, sizeSpadStoreSoftAes);
|
||||
}
|
||||
else {
|
||||
state.emit(codeSpadStore, sizeSpadStore);
|
||||
}
|
||||
|
||||
int32_t fixPos = state.codePos;
|
||||
state.emit(codeLoopEnd, sizeLoopEnd);
|
||||
//xor x26, x{readReg0}, x{readReg1}
|
||||
@@ -502,6 +521,10 @@ namespace randomx {
|
||||
//j LoopTop
|
||||
emitJump(state, 0, fixPos, LoopTopPos);
|
||||
state.emit(codeEpilogue, sizeEpilogue);
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES) {
|
||||
state.emit(codeSoftAes, sizeSoftAes);
|
||||
}
|
||||
}
|
||||
|
||||
static void generateSuperscalarCode(CodeBuffer& buf, Instruction isn, bool lastLiteral) {
|
||||
@@ -669,6 +692,9 @@ namespace randomx {
|
||||
state.emit(codeDataRead, sizeDataRead);
|
||||
//xor x8, x{readReg2}, x{readReg3}
|
||||
state.emitAt(fixPos, rvi(rv64::XOR, Tmp1Reg, regR(pcfg.readReg2), regR(pcfg.readReg3)));
|
||||
int32_t fixPos2 = state.codePos;
|
||||
state.emit(codeDataRead2, sizeDataRead2);
|
||||
state.emitAt(fixPos2, (uint16_t)(RandomX_CurrentConfig.Tweak_V2_PREFETCH ? 0x1402 : 0x0001));
|
||||
emitProgramSuffix(state, pcfg);
|
||||
clearCache(state);
|
||||
}
|
||||
@@ -691,7 +717,14 @@ namespace randomx {
|
||||
state.emitAt(fixPos + 4, rv64::LUI | (uimm << 12) | rvrd(Tmp2Reg));
|
||||
//addi x9, x9, {limm}
|
||||
state.emitAt(fixPos + 8, rvi(rv64::ADDI, Tmp2Reg, Tmp2Reg, limm));
|
||||
fixPos += offsetFixLoopCall;
|
||||
if (RandomX_CurrentConfig.Tweak_V2_PREFETCH) {
|
||||
state.emit(codeDataReadLight2, sizeDataReadLight2);
|
||||
}
|
||||
else {
|
||||
state.emit(codeDataReadLight1, sizeDataReadLight1);
|
||||
}
|
||||
fixPos = state.codePos;
|
||||
state.emit(codeFixLoopCall, sizeFixLoopCall);
|
||||
//jal x1, SuperscalarHash
|
||||
emitJump(state, ReturnReg, fixPos, SuperScalarHashOffset);
|
||||
emitProgramSuffix(state, pcfg);
|
||||
@@ -1175,10 +1208,22 @@ namespace randomx {
|
||||
//c.or x8, x9
|
||||
state.emit(rvc(rv64::C_OR, Tmp1Reg + OffsetXC, Tmp2Reg + OffsetXC));
|
||||
#endif
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
//andi x9, x8, 240
|
||||
state.emit(rvi(rv64::ANDI, Tmp2Reg, Tmp1Reg, 240));
|
||||
//c.bnez x9, +12
|
||||
state.emit(uint16_t(0xE491));
|
||||
}
|
||||
//c.andi x8, 12
|
||||
state.emit(rvc(rv64::C_ANDI, Tmp1Reg + OffsetXC, 12));
|
||||
}
|
||||
else {
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
//andi x9, x{src}, 240
|
||||
state.emit(rvi(rv64::ANDI, Tmp2Reg, regR(isn.src), 240));
|
||||
//c.bnez x9, +14
|
||||
state.emit(uint16_t(0xE499));
|
||||
}
|
||||
//and x8, x{src}, 12
|
||||
state.emit(rvi(rv64::ANDI, Tmp1Reg, regR(isn.src), 12));
|
||||
}
|
||||
|
||||
@@ -40,10 +40,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
.global DECL(randomx_riscv64_prologue)
|
||||
.global DECL(randomx_riscv64_loop_begin)
|
||||
.global DECL(randomx_riscv64_data_read)
|
||||
.global DECL(randomx_riscv64_data_read_v2_tweak)
|
||||
.global DECL(randomx_riscv64_data_read_light)
|
||||
.global DECL(randomx_riscv64_data_read_light_v1)
|
||||
.global DECL(randomx_riscv64_data_read_light_v2)
|
||||
.global DECL(randomx_riscv64_fix_loop_call)
|
||||
.global DECL(randomx_riscv64_spad_store)
|
||||
.global DECL(randomx_riscv64_spad_store_hardaes)
|
||||
.global DECL(randomx_riscv64_spad_store_softaes)
|
||||
.global DECL(randomx_riscv64_loop_end)
|
||||
.global DECL(randomx_riscv64_fix_continue_loop)
|
||||
@@ -408,7 +410,9 @@ DECL(randomx_riscv64_data_read):
|
||||
slli x8, x8, 32
|
||||
srli x8, x8, 32
|
||||
#endif
|
||||
/* update "mx" */
|
||||
DECL(randomx_riscv64_data_read_v2_tweak):
|
||||
slli x8, x8, 32 /* JIT compiler will replace it with "nop" for RandomX v1 */
|
||||
/* update "mp" */
|
||||
xor x25, x25, x8
|
||||
/* read dataset and update registers */
|
||||
ld x8, 0(x7)
|
||||
@@ -456,13 +460,22 @@ DECL(randomx_riscv64_data_read_light):
|
||||
slli x25, x25, 32
|
||||
or x25, x25, x31
|
||||
#endif
|
||||
DECL(randomx_riscv64_data_read_light_v1):
|
||||
slli x8, x8, 32
|
||||
/* update "mx" */
|
||||
/* update "mp" */
|
||||
xor x25, x25, x8
|
||||
/* the next dataset item */
|
||||
and x7, x25, x1
|
||||
srli x7, x7, 6
|
||||
add x7, x7, x9
|
||||
DECL(randomx_riscv64_data_read_light_v2):
|
||||
/* the next dataset item */
|
||||
and x7, x25, x1
|
||||
srli x7, x7, 6
|
||||
add x7, x7, x9
|
||||
and x8, x8, x1
|
||||
/* update "mp" */
|
||||
xor x25, x25, x8
|
||||
DECL(randomx_riscv64_fix_loop_call):
|
||||
jal superscalar_hash /* JIT compiler will adjust the offset */
|
||||
xor x16, x16, x8
|
||||
@@ -536,9 +549,6 @@ DECL(randomx_riscv64_spad_store):
|
||||
sd x30, 56(x26)
|
||||
fmv.d.x f7, x30
|
||||
|
||||
DECL(randomx_riscv64_spad_store_hardaes):
|
||||
nop /* not implemented */
|
||||
|
||||
DECL(randomx_riscv64_spad_store_softaes):
|
||||
/* store integer registers */
|
||||
sd x16, 0(x27)
|
||||
|
||||
@@ -36,10 +36,12 @@ extern "C" {
|
||||
void randomx_riscv64_prologue();
|
||||
void randomx_riscv64_loop_begin();
|
||||
void randomx_riscv64_data_read();
|
||||
void randomx_riscv64_data_read_v2_tweak();
|
||||
void randomx_riscv64_data_read_light();
|
||||
void randomx_riscv64_data_read_light_v1();
|
||||
void randomx_riscv64_data_read_light_v2();
|
||||
void randomx_riscv64_fix_loop_call();
|
||||
void randomx_riscv64_spad_store();
|
||||
void randomx_riscv64_spad_store_hardaes();
|
||||
void randomx_riscv64_spad_store_softaes();
|
||||
void randomx_riscv64_loop_end();
|
||||
void randomx_riscv64_fix_continue_loop();
|
||||
|
||||
@@ -34,12 +34,16 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
#include "crypto/randomx/reciprocal.h"
|
||||
#include "crypto/randomx/superscalar.hpp"
|
||||
#include "crypto/randomx/program.hpp"
|
||||
#include "crypto/randomx/soft_aes.h"
|
||||
#include "backend/cpu/Cpu.h"
|
||||
|
||||
namespace randomx {
|
||||
|
||||
#define ADDR(x) ((uint8_t*) &(x))
|
||||
#define DIST(x, y) (ADDR(y) - ADDR(x))
|
||||
|
||||
#define JUMP(offset) (0x6F | (((offset) & 0x7FE) << 20) | (((offset) & 0x800) << 9) | ((offset) & 0xFF000))
|
||||
|
||||
void* generateDatasetInitVectorRV64(uint8_t* buf, SuperscalarProgram* programs, size_t num_programs)
|
||||
{
|
||||
uint8_t* p = buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_sshash_generated_instructions);
|
||||
@@ -205,8 +209,7 @@ void* generateDatasetInitVectorRV64(uint8_t* buf, SuperscalarProgram* programs,
|
||||
|
||||
// Emit "J randomx_riscv64_vector_sshash_generated_instructions_end" instruction
|
||||
const uint8_t* e = buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_sshash_generated_instructions_end);
|
||||
const uint32_t k = e - p;
|
||||
const uint32_t j = 0x6F | ((k & 0x7FE) << 20) | ((k & 0x800) << 9) | (k & 0xFF000);
|
||||
const uint32_t j = JUMP(e - p);
|
||||
memcpy(p, &j, 4);
|
||||
|
||||
char* result = (char*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_sshash_dataset_init));
|
||||
@@ -323,6 +326,26 @@ void* generateProgramVectorRV64(uint8_t* buf, Program& prog, ProgramConfiguratio
|
||||
params[3] = RandomX_CurrentConfig.DatasetBaseSize - 64;
|
||||
params[4] = (1 << RandomX_ConfigurationBase::JumpBits) - 1;
|
||||
|
||||
const bool hasAES = xmrig::Cpu::info()->hasAES();
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES && !hasAES) {
|
||||
params[5] = (uint64_t) &lutEnc[2][0];
|
||||
params[6] = (uint64_t) &lutDec[2][0];
|
||||
params[7] = (uint64_t) lutEncIndex;
|
||||
params[8] = (uint64_t) lutDecIndex;
|
||||
|
||||
uint32_t* p1 = (uint32_t*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_v2_soft_aes_init));
|
||||
|
||||
// Restore vsetivli zero, 4, e32, m1, ta, ma
|
||||
*p1 = 0xCD027057;
|
||||
}
|
||||
else {
|
||||
uint32_t* p1 = (uint32_t*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_v2_soft_aes_init));
|
||||
|
||||
// Emit "J randomx_riscv64_vector_program_main_loop" instruction
|
||||
*p1 = JUMP(DIST(randomx_riscv64_vector_program_v2_soft_aes_init, randomx_riscv64_vector_program_main_loop));
|
||||
}
|
||||
|
||||
uint64_t* imul_rcp_literals = (uint64_t*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_imul_rcp_literals));
|
||||
uint64_t* cur_literal = imul_rcp_literals;
|
||||
|
||||
@@ -338,6 +361,16 @@ void* generateProgramVectorRV64(uint8_t* buf, Program& prog, ProgramConfiguratio
|
||||
*mx_xor = mx_xor_value;
|
||||
*mx_xor_light = mx_xor_value;
|
||||
|
||||
// "slli x5, x5, 32" for RandomX v2, "nop" for RandomX v1
|
||||
const uint16_t mp_reg_value = RandomX_CurrentConfig.Tweak_V2_PREFETCH ? 0x1282 : 0x0001;
|
||||
|
||||
memcpy(((uint8_t*)mx_xor) + 8, &mp_reg_value, sizeof(mp_reg_value));
|
||||
memcpy(((uint8_t*)mx_xor_light) + 8, &mp_reg_value, sizeof(mp_reg_value));
|
||||
|
||||
// "srli x5, x14, 32" for RandomX v2, "srli x5, x14, 0" for RandomX v1
|
||||
const uint32_t mp_reg_value2 = RandomX_CurrentConfig.Tweak_V2_PREFETCH ? 0x02075293 : 0x00075293;
|
||||
memcpy(((uint8_t*)mx_xor) + 14, &mp_reg_value2, sizeof(mp_reg_value2));
|
||||
|
||||
if (entryDataInitScalar) {
|
||||
void* light_mode_data = buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_main_loop_light_mode_data);
|
||||
|
||||
@@ -760,10 +793,24 @@ void* generateProgramVectorRV64(uint8_t* buf, Program& prog, ProgramConfiguratio
|
||||
emit32(0x0062E2B3);
|
||||
#endif // __riscv_zbb
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
// andi x6, x5, 120
|
||||
emit32(0x0782F313);
|
||||
// bnez x6, +24
|
||||
emit32(0x00031C63);
|
||||
}
|
||||
|
||||
// andi x5, x5, 6
|
||||
emit32(0x0062F293);
|
||||
}
|
||||
else {
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
// andi x6, x20 + src, 120
|
||||
emit32(0x078A7313 + (src << 15));
|
||||
// bnez x6, +24
|
||||
emit32(0x00031C63);
|
||||
}
|
||||
|
||||
// andi x5, x20 + src, 6
|
||||
emit32(0x006A7293 + (src << 15));
|
||||
}
|
||||
@@ -813,6 +860,9 @@ void* generateProgramVectorRV64(uint8_t* buf, Program& prog, ProgramConfiguratio
|
||||
}
|
||||
break;
|
||||
|
||||
case InstructionType::NOP:
|
||||
break;
|
||||
|
||||
default:
|
||||
UNREACHABLE;
|
||||
}
|
||||
@@ -829,8 +879,26 @@ void* generateProgramVectorRV64(uint8_t* buf, Program& prog, ProgramConfiguratio
|
||||
e = buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_main_loop_instructions_end);
|
||||
}
|
||||
|
||||
const uint32_t k = e - p;
|
||||
emit32(0x6F | ((k & 0x7FE) << 20) | ((k & 0x800) << 9) | (k & 0xFF000));
|
||||
emit32(JUMP(e - p));
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_AES) {
|
||||
uint32_t* p1 = (uint32_t*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_main_loop_fe_mix));
|
||||
|
||||
if (hasAES) {
|
||||
// Restore vsetivli zero, 4, e32, m1, ta, ma
|
||||
*p1 = 0xCD027057;
|
||||
}
|
||||
else {
|
||||
// Emit "J randomx_riscv64_vector_program_main_loop_fe_mix_v2_soft_aes" instruction
|
||||
*p1 = JUMP(DIST(randomx_riscv64_vector_program_main_loop_fe_mix, randomx_riscv64_vector_program_main_loop_fe_mix_v2_soft_aes));
|
||||
}
|
||||
}
|
||||
else {
|
||||
uint32_t* p1 = (uint32_t*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_main_loop_fe_mix));
|
||||
|
||||
// Emit "J randomx_riscv64_vector_program_main_loop_fe_mix_v1" instruction
|
||||
*p1 = JUMP(DIST(randomx_riscv64_vector_program_main_loop_fe_mix, randomx_riscv64_vector_program_main_loop_fe_mix_v1));
|
||||
}
|
||||
|
||||
#ifdef __GNUC__
|
||||
char* p1 = (char*)(buf + DIST(randomx_riscv64_vector_code_begin, randomx_riscv64_vector_program_params));
|
||||
|
||||
@@ -66,16 +66,22 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
.global DECL(randomx_riscv64_vector_program_params)
|
||||
.global DECL(randomx_riscv64_vector_program_imul_rcp_literals)
|
||||
.global DECL(randomx_riscv64_vector_program_begin)
|
||||
.global DECL(randomx_riscv64_vector_program_v2_soft_aes_init)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_instructions)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_instructions_end)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_mx_xor)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_spaddr_xor)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_fe_mix)
|
||||
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_light_mode_data)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_instructions_end_light_mode)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_mx_xor_light_mode)
|
||||
.global DECL(randomx_riscv64_vector_program_scratchpad_prefetch)
|
||||
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_fe_mix_v1)
|
||||
.global DECL(randomx_riscv64_vector_program_main_loop_fe_mix_v2_soft_aes)
|
||||
|
||||
.global DECL(randomx_riscv64_vector_program_end)
|
||||
|
||||
.global DECL(randomx_riscv64_vector_code_end)
|
||||
@@ -375,12 +381,21 @@ v12 = E 'and' mask = 0x00ffffffffffffff'00ffffffffffffff
|
||||
v13 = E 'or' mask = 0x3*00000000******'3*00000000******
|
||||
v14 = scale mask = 0x80f0000000000000'80f0000000000000
|
||||
|
||||
v15 = unused
|
||||
v15 = all zeroes
|
||||
v16 = temporary
|
||||
v17 = unused
|
||||
v18 = temporary
|
||||
|
||||
v19-v31 = unused
|
||||
v19 = unused
|
||||
v20 = randomx_aes_lut_enc_index[0]
|
||||
v21 = randomx_aes_lut_enc_index[1]
|
||||
v22 = randomx_aes_lut_enc_index[2]
|
||||
v23 = randomx_aes_lut_enc_index[3]
|
||||
v24 = randomx_aes_lut_dec_index[0]
|
||||
v25 = randomx_aes_lut_dec_index[1]
|
||||
v26 = randomx_aes_lut_dec_index[2]
|
||||
v27 = randomx_aes_lut_dec_index[3]
|
||||
v28-v31 = temporary in aesenc_soft/aesdec_soft
|
||||
*/
|
||||
|
||||
.balign 8
|
||||
@@ -390,6 +405,11 @@ DECL(randomx_riscv64_vector_program_params):
|
||||
// JIT compiler will adjust these values for different RandomX variants
|
||||
randomx_masks: .dword 16376, 262136, 2097144, 2147483584, 255
|
||||
|
||||
randomx_aes_lut_enc_ptr: .dword 0
|
||||
randomx_aes_lut_dec_ptr: .dword 0
|
||||
randomx_aes_lut_enc_index_ptr: .dword 0
|
||||
randomx_aes_lut_dec_index_ptr: .dword 0
|
||||
|
||||
DECL(randomx_riscv64_vector_program_imul_rcp_literals):
|
||||
|
||||
imul_rcp_literals: .fill RANDOMX_PROGRAM_MAX_SIZE, 8, 0
|
||||
@@ -507,7 +527,44 @@ DECL(randomx_riscv64_vector_program_begin):
|
||||
fld f30, 192(x18)
|
||||
fld f31, 200(x18)
|
||||
|
||||
randomx_riscv64_vector_program_main_loop:
|
||||
// Set v15 to zero
|
||||
vxor.vv v15, v15, v15
|
||||
|
||||
DECL(randomx_riscv64_vector_program_v2_soft_aes_init):
|
||||
// JIT compiler will place a jump to the main loop here if needed
|
||||
|
||||
// Load randomx_aes_lut_enc_index/randomx_aes_lut_dec_index
|
||||
vsetivli zero, 4, e32, m1, ta, ma
|
||||
|
||||
lla x5, randomx_aes_lut_enc_index_ptr
|
||||
ld x5, (x5)
|
||||
vle32.v v20, (x5)
|
||||
|
||||
addi x6, x5, 32
|
||||
vle32.v v21, (x6)
|
||||
|
||||
addi x6, x5, 64
|
||||
vle32.v v22, (x6)
|
||||
|
||||
addi x6, x5, 96
|
||||
vle32.v v23, (x6)
|
||||
|
||||
lla x5, randomx_aes_lut_dec_index_ptr
|
||||
ld x5, (x5)
|
||||
vle32.v v24, (x5)
|
||||
|
||||
addi x6, x5, 32
|
||||
vle32.v v25, (x6)
|
||||
|
||||
addi x6, x5, 64
|
||||
vle32.v v26, (x6)
|
||||
|
||||
addi x6, x5, 96
|
||||
vle32.v v27, (x6)
|
||||
|
||||
vsetivli zero, 2, e64, m1, ta, ma
|
||||
|
||||
DECL(randomx_riscv64_vector_program_main_loop):
|
||||
and x5, x15, x9 // x5 = spAddr0 & 64-byte aligned L3 mask
|
||||
add x5, x5, x12 // x5 = &scratchpad[spAddr0 & 64-byte aligned L3 mask]
|
||||
|
||||
@@ -609,11 +666,13 @@ DECL(randomx_riscv64_vector_program_main_loop_instructions_end):
|
||||
|
||||
DECL(randomx_riscv64_vector_program_main_loop_mx_xor):
|
||||
xor x5, x24, x26 // x5 = readReg2 ^ readReg3 (JIT compiler will substitute the actual registers)
|
||||
|
||||
and x5, x5, x19 // x5 = (readReg2 ^ readReg3) & dataset mask
|
||||
xor x14, x14, x5 // mx ^= (readReg2 ^ readReg3) & dataset mask
|
||||
slli x5, x5, 32 // JIT compiler will replace it with "nop" for v1
|
||||
xor x14, x14, x5 // mp ^= (readReg2 ^ readReg3) & dataset mask
|
||||
|
||||
add x5, x14, x11 // x5 = &dataset[mx & dataset mask]
|
||||
srli x5, x14, 32 // JIT compiler will replace it with "srli x5, x14, 0" for v1
|
||||
and x5, x5, x19 // x5 = mp & dataset mask
|
||||
add x5, x5, x11 // x5 = &dataset[mp & dataset mask]
|
||||
|
||||
#ifdef __riscv_zicbop
|
||||
prefetch.r (x5)
|
||||
@@ -689,11 +748,47 @@ DECL(randomx_riscv64_vector_program_main_loop_spaddr_xor):
|
||||
xor x15, x20, x22 // spAddr0-spAddr1 = readReg0 ^ readReg1 (JIT compiler will substitute the actual registers)
|
||||
|
||||
// store registers f0-f3 to the scratchpad (f0-f3 are first combined with e0-e3)
|
||||
vxor.vv v0, v0, v4
|
||||
|
||||
// v2 FE mix code is the main code path
|
||||
// JIT compiler will place a jump to v1 or v2 soft AES code here if needed
|
||||
DECL(randomx_riscv64_vector_program_main_loop_fe_mix):
|
||||
vsetivli zero, 4, e32, m1, ta, ma
|
||||
|
||||
// f0 = aesenc(f0, e0), f1 = aesdec(f1, e0), f2 = aesenc(f2, e0), f3 = aesdec(f3, e0)
|
||||
vaesem.vv v0, v4
|
||||
vaesdm.vv v1, v15
|
||||
vaesem.vv v2, v4
|
||||
vaesdm.vv v3, v15
|
||||
vxor.vv v1, v1, v4
|
||||
vxor.vv v3, v3, v4
|
||||
|
||||
// f0 = aesenc(f0, e1), f1 = aesdec(f1, e1), f2 = aesenc(f2, e1), f3 = aesdec(f3, e1)
|
||||
vaesem.vv v0, v5
|
||||
vaesdm.vv v1, v15
|
||||
vaesem.vv v2, v5
|
||||
vaesdm.vv v3, v15
|
||||
vxor.vv v1, v1, v5
|
||||
vxor.vv v2, v2, v6
|
||||
vxor.vv v3, v3, v5
|
||||
|
||||
// f0 = aesenc(f0, e2), f1 = aesdec(f1, e2), f2 = aesenc(f2, e2), f3 = aesdec(f3, e2)
|
||||
vaesem.vv v0, v6
|
||||
vaesdm.vv v1, v15
|
||||
vaesem.vv v2, v6
|
||||
vaesdm.vv v3, v15
|
||||
vxor.vv v1, v1, v6
|
||||
vxor.vv v3, v3, v6
|
||||
|
||||
// f0 = aesenc(f0, e3), f1 = aesdec(f1, e3), f2 = aesenc(f2, e3), f3 = aesdec(f3, e3)
|
||||
vaesem.vv v0, v7
|
||||
vaesdm.vv v1, v15
|
||||
vaesem.vv v2, v7
|
||||
vaesdm.vv v3, v15
|
||||
vxor.vv v1, v1, v7
|
||||
vxor.vv v3, v3, v7
|
||||
|
||||
vsetivli zero, 2, e64, m1, ta, ma
|
||||
|
||||
randomx_riscv64_vector_program_main_loop_fe_store:
|
||||
vse64.v v0, (x5)
|
||||
|
||||
addi x6, x5, 16
|
||||
@@ -782,6 +877,7 @@ DECL(randomx_riscv64_vector_program_main_loop_instructions_end_light_mode):
|
||||
DECL(randomx_riscv64_vector_program_main_loop_mx_xor_light_mode):
|
||||
xor x5, x24, x26 // x5 = readReg2 ^ readReg3 (JIT compiler will substitute the actual registers)
|
||||
and x5, x5, x19 // x5 = (readReg2 ^ readReg3) & dataset mask
|
||||
slli x5, x5, 32 // JIT compiler will replace it with "nop" for v1
|
||||
xor x14, x14, x5 // mx ^= (readReg2 ^ readReg3) & dataset mask
|
||||
|
||||
// Save all registers modified when calling dataset_init_scalar_func_ptr
|
||||
@@ -864,6 +960,131 @@ DECL(randomx_riscv64_vector_program_main_loop_mx_xor_light_mode):
|
||||
|
||||
j randomx_riscv64_vector_program_scratchpad_prefetch
|
||||
|
||||
DECL(randomx_riscv64_vector_program_main_loop_fe_mix_v1):
|
||||
vxor.vv v0, v0, v4
|
||||
vxor.vv v1, v1, v5
|
||||
vxor.vv v2, v2, v6
|
||||
vxor.vv v3, v3, v7
|
||||
j randomx_riscv64_vector_program_main_loop_fe_store
|
||||
|
||||
/*
|
||||
aesenc middle round
|
||||
|
||||
x5 = pointer to aesenc LUT
|
||||
v16 = input and return value
|
||||
*/
|
||||
.macro aesenc_soft input, key
|
||||
vsetivli zero, 16, e8, m1, ta, ma
|
||||
|
||||
vrgather.vv v28, \input, v20
|
||||
vrgather.vv v29, \input, v21
|
||||
vrgather.vv v30, \input, v22
|
||||
vrgather.vv v31, \input, v23
|
||||
|
||||
vsetivli zero, 4, e32, m1, ta, ma
|
||||
|
||||
vsll.vi v28, v28, 2
|
||||
vsll.vi v29, v29, 2
|
||||
vsll.vi v30, v30, 2
|
||||
vsll.vi v31, v31, 2
|
||||
|
||||
addi x6, x5, -2048
|
||||
vluxei32.v v28, (x6), v28
|
||||
|
||||
addi x6, x5, -1024
|
||||
vluxei32.v v29, (x6), v29
|
||||
|
||||
vluxei32.v v30, (x5), v30
|
||||
|
||||
addi x6, x5, 1024
|
||||
vluxei32.v v31, (x6), v31
|
||||
|
||||
vxor.vv v28, v28, v29
|
||||
vxor.vv v30, v30, v31
|
||||
vxor.vv \input, v28, v30
|
||||
vxor.vv \input, \input, \key
|
||||
.endm
|
||||
|
||||
/*
|
||||
aesdec middle round
|
||||
|
||||
x5 = pointer to aesdec LUT
|
||||
v16 = input and return value
|
||||
*/
|
||||
.macro aesdec_soft input, key
|
||||
vsetivli zero, 16, e8, m1, ta, ma
|
||||
|
||||
vrgather.vv v28, \input, v24
|
||||
vrgather.vv v29, \input, v25
|
||||
vrgather.vv v30, \input, v26
|
||||
vrgather.vv v31, \input, v27
|
||||
|
||||
vsetivli zero, 4, e32, m1, ta, ma
|
||||
|
||||
vsll.vi v28, v28, 2
|
||||
vsll.vi v29, v29, 2
|
||||
vsll.vi v30, v30, 2
|
||||
vsll.vi v31, v31, 2
|
||||
|
||||
addi x6, x5, -2048
|
||||
vluxei32.v v28, (x6), v28
|
||||
|
||||
addi x6, x5, -1024
|
||||
vluxei32.v v29, (x6), v29
|
||||
|
||||
vluxei32.v v30, (x5), v30
|
||||
|
||||
addi x6, x5, 1024
|
||||
vluxei32.v v31, (x6), v31
|
||||
|
||||
vxor.vv v28, v28, v29
|
||||
vxor.vv v30, v30, v31
|
||||
vxor.vv \input, v28, v30
|
||||
vxor.vv \input, \input, \key
|
||||
.endm
|
||||
|
||||
DECL(randomx_riscv64_vector_program_main_loop_fe_mix_v2_soft_aes):
|
||||
// save x5
|
||||
vmv.s.x v16, x5
|
||||
|
||||
lla x5, randomx_aes_lut_enc_ptr
|
||||
ld x5, (x5)
|
||||
|
||||
// f0 = aesenc(f0, e0), f0 = aesenc(f0, e1), f0 = aesenc(f0, e2), f0 = aesenc(f0, e3)
|
||||
aesenc_soft v0, v4
|
||||
aesenc_soft v0, v5
|
||||
aesenc_soft v0, v6
|
||||
aesenc_soft v0, v7
|
||||
|
||||
// f2 = aesenc(f2, e0), f2 = aesenc(f2, e1), f2 = aesenc(f2, e2), f2 = aesenc(f2, e3)
|
||||
aesenc_soft v2, v4
|
||||
aesenc_soft v2, v5
|
||||
aesenc_soft v2, v6
|
||||
aesenc_soft v2, v7
|
||||
|
||||
lla x5, randomx_aes_lut_dec_ptr
|
||||
ld x5, (x5)
|
||||
|
||||
// f1 = aesdec(f1, e0), f1 = aesdec(f1, e1), f1 = aesdec(f1, e2), f1 = aesdec(f1, e3)
|
||||
aesdec_soft v1, v4
|
||||
aesdec_soft v1, v5
|
||||
aesdec_soft v1, v6
|
||||
aesdec_soft v1, v7
|
||||
|
||||
// f3 = aesdec(f3, e0), f3 = aesdec(f3, e1), f3 = aesdec(f3, e2), f3 = aesdec(f3, e3)
|
||||
aesdec_soft v3, v4
|
||||
aesdec_soft v3, v5
|
||||
aesdec_soft v3, v6
|
||||
aesdec_soft v3, v7
|
||||
|
||||
// Set vector registers back to 2x64 bit
|
||||
vsetivli zero, 2, e64, m1, ta, ma
|
||||
|
||||
// restore x5
|
||||
vmv.x.s x5, v16
|
||||
|
||||
j randomx_riscv64_vector_program_main_loop_fe_store
|
||||
|
||||
DECL(randomx_riscv64_vector_program_end):
|
||||
|
||||
DECL(randomx_riscv64_vector_code_end):
|
||||
|
||||
@@ -57,16 +57,22 @@ void randomx_riscv64_vector_sshash_end();
|
||||
void randomx_riscv64_vector_program_params();
|
||||
void randomx_riscv64_vector_program_imul_rcp_literals();
|
||||
void randomx_riscv64_vector_program_begin();
|
||||
void randomx_riscv64_vector_program_v2_soft_aes_init();
|
||||
void randomx_riscv64_vector_program_main_loop();
|
||||
void randomx_riscv64_vector_program_main_loop_instructions();
|
||||
void randomx_riscv64_vector_program_main_loop_instructions_end();
|
||||
void randomx_riscv64_vector_program_main_loop_mx_xor();
|
||||
void randomx_riscv64_vector_program_main_loop_spaddr_xor();
|
||||
void randomx_riscv64_vector_program_main_loop_fe_mix();
|
||||
void randomx_riscv64_vector_program_main_loop_light_mode_data();
|
||||
void randomx_riscv64_vector_program_main_loop_instructions_end_light_mode();
|
||||
void randomx_riscv64_vector_program_main_loop_mx_xor_light_mode();
|
||||
void randomx_riscv64_vector_program_end();
|
||||
void randomx_riscv64_vector_program_scratchpad_prefetch();
|
||||
|
||||
void randomx_riscv64_vector_program_main_loop_fe_mix_v1();
|
||||
void randomx_riscv64_vector_program_main_loop_fe_mix_v2_soft_aes();
|
||||
|
||||
void randomx_riscv64_vector_code_end();
|
||||
|
||||
#if defined(__cplusplus)
|
||||
|
||||
@@ -1379,10 +1379,10 @@ namespace randomx {
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
if (BranchesWithin32B) {
|
||||
const uint32_t branch_begin = static_cast<uint32_t>(pos + 2) & 31;
|
||||
const uint32_t branch_begin = pos & 31;
|
||||
|
||||
// If the jump crosses or touches 32-byte boundary, align it
|
||||
if (branch_begin >= 30) {
|
||||
if (branch_begin >= 28) {
|
||||
const uint32_t alignment_size = 32 - branch_begin;
|
||||
emit(NOPX[alignment_size - 1], alignment_size, code, pos);
|
||||
}
|
||||
@@ -1438,10 +1438,10 @@ namespace randomx {
|
||||
|
||||
if (RandomX_CurrentConfig.Tweak_V2_CFROUND) {
|
||||
if (BranchesWithin32B) {
|
||||
const uint32_t branch_begin = static_cast<uint32_t>(pos + 2) & 31;
|
||||
const uint32_t branch_begin = pos & 31;
|
||||
|
||||
// If the jump crosses or touches 32-byte boundary, align it
|
||||
if (branch_begin >= 30) {
|
||||
if (branch_begin >= 28) {
|
||||
const uint32_t alignment_size = 32 - branch_begin;
|
||||
emit(NOPX[alignment_size - 1], alignment_size, code, pos);
|
||||
}
|
||||
|
||||
@@ -61,6 +61,17 @@ namespace randomx {
|
||||
# if defined(XMRIG_ARM) || defined(XMRIG_RISCV)
|
||||
memcpy(reg.f, config.eMask, sizeof(config.eMask));
|
||||
# endif
|
||||
|
||||
const uint8_t* p = mem.memory;
|
||||
|
||||
// dataset prefetch for the first iteration of the main loop
|
||||
rx_prefetch_nta(p + (mem.ma & (RandomX_ConfigurationBase::DatasetBaseSize - 64)));
|
||||
|
||||
// dataset prefetch for the second iteration of the main loop (RandomX v2)
|
||||
if (RandomX_CurrentConfig.Tweak_V2_PREFETCH) {
|
||||
rx_prefetch_nta(p + (mem.mx & (RandomX_ConfigurationBase::DatasetBaseSize - 64)));
|
||||
}
|
||||
|
||||
compiler.getProgramFunc()(reg, mem, scratchpad, RandomX_CurrentConfig.ProgramIterations);
|
||||
}
|
||||
|
||||
|
||||
@@ -17,6 +17,7 @@
|
||||
*/
|
||||
|
||||
#include "crypto/rx/RxConfig.h"
|
||||
#include "crypto/randomx/randomx.h"
|
||||
#include "3rdparty/rapidjson/document.h"
|
||||
#include "backend/cpu/Cpu.h"
|
||||
#include "base/io/json/Json.h"
|
||||
@@ -25,6 +26,7 @@
|
||||
#include <array>
|
||||
#include <algorithm>
|
||||
#include <cmath>
|
||||
#include <uv.h>
|
||||
|
||||
|
||||
#ifdef _MSC_VER
|
||||
@@ -183,11 +185,20 @@ rapidjson::Value xmrig::RxConfig::toJSON(rapidjson::Document &doc) const
|
||||
#ifdef XMRIG_FEATURE_HWLOC
|
||||
std::vector<uint32_t> xmrig::RxConfig::nodeset() const
|
||||
{
|
||||
auto info = Cpu::info();
|
||||
|
||||
constexpr uint64_t dataset_mem = RandomX_ConfigurationBase::DatasetBaseSize + RandomX_ConfigurationBase::DatasetExtraSize;
|
||||
constexpr uint64_t cache_mem = RandomX_ConfigurationBase::ArgonMemory * 1024;
|
||||
const uint64_t threads_mem = info->threads() << 21;
|
||||
|
||||
const uint64_t freem_mem = uv_get_free_memory();
|
||||
|
||||
if (!m_nodeset.empty()) {
|
||||
return m_nodeset;
|
||||
return (freem_mem > m_nodeset.size() * dataset_mem + cache_mem + threads_mem) ? m_nodeset : std::vector<uint32_t>();
|
||||
}
|
||||
|
||||
return (m_numa && Cpu::info()->nodes() > 1) ? Cpu::info()->nodeset() : std::vector<uint32_t>();
|
||||
const uint64_t n = info->nodes();
|
||||
return (m_numa && (n > 1) && (freem_mem > n * dataset_mem + cache_mem + threads_mem)) ? Cpu::info()->nodeset() : std::vector<uint32_t>();
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
@@ -47,8 +47,8 @@ public:
|
||||
inline RxSeed(const Algorithm &algorithm, const Buffer &seed) : m_algorithm(algorithm), m_data(seed) {}
|
||||
inline RxSeed(const Job &job) : m_algorithm(job.algorithm()), m_data(job.seed()) {}
|
||||
|
||||
inline bool isEqual(const Job &job) const { return m_algorithm == job.algorithm() && m_data == job.seed(); }
|
||||
inline bool isEqual(const RxSeed &other) const { return m_algorithm == other.m_algorithm && m_data == other.m_data; }
|
||||
inline bool isEqual(const Job &job) const { return isEqualSeedAlgo(job.algorithm()) && m_data == job.seed(); }
|
||||
inline bool isEqual(const RxSeed &other) const { return isEqualSeedAlgo(other.m_algorithm) && m_data == other.m_data; }
|
||||
inline const Algorithm &algorithm() const { return m_algorithm; }
|
||||
inline const Buffer &data() const { return m_data; }
|
||||
|
||||
@@ -60,6 +60,12 @@ public:
|
||||
private:
|
||||
Algorithm m_algorithm;
|
||||
Buffer m_data;
|
||||
|
||||
inline bool isEqualSeedAlgo(Algorithm other) const {
|
||||
return (m_algorithm == other) ||
|
||||
((m_algorithm == Algorithm::RX_0) && (other == Algorithm::RX_V2)) ||
|
||||
((m_algorithm == Algorithm::RX_V2) && (other == Algorithm::RX_0));
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -34,6 +34,8 @@
|
||||
#include "base/tools/String.h"
|
||||
#include "base/net/stratum/Job.h"
|
||||
|
||||
#include "crypto/randomx/randomx.h"
|
||||
|
||||
|
||||
namespace xmrig {
|
||||
|
||||
@@ -43,7 +45,7 @@ class JobResult
|
||||
public:
|
||||
JobResult() = delete;
|
||||
|
||||
inline JobResult(const Job &job, uint64_t nonce, const uint8_t *result, const uint8_t* header_hash = nullptr, const uint8_t *mix_hash = nullptr, const uint8_t* miner_signature = nullptr) :
|
||||
inline JobResult(const Job &job, uint64_t nonce, const uint8_t *result, const uint8_t* header_hash = nullptr, const uint8_t *mix_hash = nullptr, const uint8_t* extra_data = nullptr) :
|
||||
algorithm(job.algorithm()),
|
||||
index(job.index()),
|
||||
clientId(job.clientId()),
|
||||
@@ -62,9 +64,15 @@ public:
|
||||
memcpy(m_mixHash, mix_hash, sizeof(m_mixHash));
|
||||
}
|
||||
|
||||
if (miner_signature) {
|
||||
m_hasMinerSignature = true;
|
||||
memcpy(m_minerSignature, miner_signature, sizeof(m_minerSignature));
|
||||
if (extra_data) {
|
||||
if (algorithm == Algorithm::RX_V2) {
|
||||
m_hasCommitment = true;
|
||||
memcpy(m_extraData, extra_data, RANDOMX_HASH_SIZE);
|
||||
}
|
||||
else if (algorithm == Algorithm::RX_WOW) {
|
||||
m_hasMinerSignature = true;
|
||||
memcpy(m_extraData, extra_data, RANDOMX_HASH_SIZE * 2);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -85,7 +93,8 @@ public:
|
||||
inline const uint8_t *headerHash() const { return m_headerHash; }
|
||||
inline const uint8_t *mixHash() const { return m_mixHash; }
|
||||
|
||||
inline const uint8_t *minerSignature() const { return m_hasMinerSignature ? m_minerSignature : nullptr; }
|
||||
inline const uint8_t *minerSignature() const { return m_hasMinerSignature ? m_extraData : nullptr; }
|
||||
inline const uint8_t *commitment() const { return m_hasCommitment ? m_extraData : nullptr; }
|
||||
|
||||
const Algorithm algorithm;
|
||||
const uint8_t index;
|
||||
@@ -100,8 +109,10 @@ private:
|
||||
uint8_t m_headerHash[32] = { 0 };
|
||||
uint8_t m_mixHash[32] = { 0 };
|
||||
|
||||
uint8_t m_minerSignature[64] = { 0 };
|
||||
uint8_t m_extraData[RANDOMX_HASH_SIZE * 2] = { 0 };
|
||||
|
||||
bool m_hasMinerSignature = false;
|
||||
bool m_hasCommitment = false;
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -339,9 +339,9 @@ void xmrig::JobResults::submit(const Job &job, uint32_t nonce, const uint8_t *re
|
||||
}
|
||||
|
||||
|
||||
void xmrig::JobResults::submit(const Job& job, uint32_t nonce, const uint8_t* result, const uint8_t* miner_signature)
|
||||
void xmrig::JobResults::submit(const Job& job, uint32_t nonce, const uint8_t* result, const uint8_t* extra_data)
|
||||
{
|
||||
submit(JobResult(job, nonce, result, nullptr, nullptr, miner_signature));
|
||||
submit(JobResult(job, nonce, result, nullptr, nullptr, extra_data));
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -46,7 +46,7 @@ public:
|
||||
static void setListener(IJobResultListener *listener, bool hwAES);
|
||||
static void stop();
|
||||
static void submit(const Job &job, uint32_t nonce, const uint8_t *result);
|
||||
static void submit(const Job& job, uint32_t nonce, const uint8_t* result, const uint8_t* miner_signature);
|
||||
static void submit(const Job& job, uint32_t nonce, const uint8_t* result, const uint8_t* extra_data);
|
||||
static void submit(const JobResult &result);
|
||||
|
||||
# if defined(XMRIG_FEATURE_OPENCL) || defined(XMRIG_FEATURE_CUDA)
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
/* XMRig
|
||||
* Copyright (c) 2018-2025 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2025 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
* Copyright (c) 2018-2026 SChernykh <https://github.com/SChernykh>
|
||||
* Copyright (c) 2016-2026 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||
*
|
||||
* SPDX-License-Identifier: GPL-3.0-or-later
|
||||
*/
|
||||
@@ -11,14 +11,14 @@
|
||||
#define APP_ID "xmrig"
|
||||
#define APP_NAME "XMRig"
|
||||
#define APP_DESC "XMRig miner"
|
||||
#define APP_VERSION "6.25.1-dev"
|
||||
#define APP_VERSION "6.26.1-dev"
|
||||
#define APP_DOMAIN "xmrig.com"
|
||||
#define APP_SITE "www.xmrig.com"
|
||||
#define APP_COPYRIGHT "Copyright (C) 2016-2025 xmrig.com"
|
||||
#define APP_COPYRIGHT "Copyright (C) 2016-2026 xmrig.com"
|
||||
#define APP_KIND "miner"
|
||||
|
||||
#define APP_VER_MAJOR 6
|
||||
#define APP_VER_MINOR 25
|
||||
#define APP_VER_MINOR 26
|
||||
#define APP_VER_PATCH 1
|
||||
|
||||
#ifdef _MSC_VER
|
||||
|
||||
Reference in New Issue
Block a user