Added RandomX JIT for AMD Navi GPUs

This commit is contained in:
SChernykh 2020-02-13 20:15:08 +01:00
parent 2f27d5d108
commit 7fa5e8706e
9 changed files with 2148 additions and 873 deletions

File diff suppressed because it is too large Load diff

View file

@ -43,6 +43,62 @@ along with RandomX OpenCL. If not, see <http://www.gnu.org/licenses/>.
#define RANDOMX_JUMP_BITS 8
#define RANDOMX_JUMP_OFFSET 8
#if GCN_VERSION >= 15
#define S_SETPC_B64_S12_13 0xbe80200cu
#define V_AND_B32_CALC_ADDRESS 0x3638000eu
#define GLOBAL_LOAD_DWORDX2_SCRATCHPAD_LOAD 0xdc348000u
#define S_WAITCNT_SCRATCHPAD_LOAD2 0xbf8c3f70u
#define V_READLANE_B32_SCRATCHPAD_LOAD2 0xd7600000u
#define S_MUL_HI_U32_IMUL_R 0x9a8f1010u
#define S_MUL_I32_IMUL 0x93000000u
#define S_MUL_HI_U32_IMUL_R_2 0x9a8fff10u
#define S_MUL_HI_U32_IMUL_M 0x9aa10e10u
#define S_MOV_B32_IMUL_RCP 0xbea003ffu
#define S_MUL_HI_U32_IMUL_RCP 0x9a8f2010u
#define S_XOR_B32_64 0x89000000u
#define S_MOV_B32_XOR_R 0xbebe03ffu
#define S_LSHR 0x90000000u
#define S_LSHL 0x8f000000u
#define S_OR 0x88000000u
#define S_AND 0x87000000u
#define S_BFE 0x94000000u
#define DS_SWIZZLE_B32_FSWAP_R 0xd8d48001u
#define V_ADD_F64 0xd564003cu
#define V_AND_B32 0x36000000u
#define GLOBAL_LOAD_DWORD_SCRATCHPAD_LOAD_FP 0xdc308000u
#define V_XOR_B32 0x3a000000u
#define V_MUL_F64 0xd5650044u
#else
#define S_SETPC_B64_S12_13 0xbe801d0cu
#define V_AND_B32_CALC_ADDRESS 0x2638000eu
#define GLOBAL_LOAD_DWORDX2_SCRATCHPAD_LOAD 0xdc548000u
#define S_WAITCNT_SCRATCHPAD_LOAD2 0xbf8c0f70u
#define V_READLANE_B32_SCRATCHPAD_LOAD2 0xd2890000u
#define S_MUL_HI_U32_IMUL_R 0x960f1010u
#define S_MUL_I32_IMUL 0x92000000u
#define S_MUL_HI_U32_IMUL_R_2 0x960fff10u
#define S_MUL_HI_U32_IMUL_M 0x96210e10u
#define S_MOV_B32_IMUL_RCP 0xbea000ffu
#define S_MUL_HI_U32_IMUL_RCP 0x960f2010u
#define S_XOR_B32_64 0x88000000u
#define S_MOV_B32_XOR_R 0xbebe00ffu
#define S_LSHR 0x8f000000u
#define S_LSHL 0x8e000000u
#define S_OR 0x87000000u
#define S_AND 0x86000000u
#define S_BFE 0x93000000u
#define DS_SWIZZLE_B32_FSWAP_R 0xd87a8001u
#define V_ADD_F64 0xd280003cu
#define V_AND_B32 0x26000000u
#define GLOBAL_LOAD_DWORD_SCRATCHPAD_LOAD_FP 0xdc508000u
#define V_XOR_B32 0x2a000000u
#define V_MUL_F64 0xd2810044u
#endif
__global uint* jit_scratchpad_calc_address(__global uint* p, uint src, uint imm32, uint mask_reg, uint batch_size)
{
// s_add_i32 s14, s(16 + src * 2), imm32
@ -50,7 +106,7 @@ __global uint* jit_scratchpad_calc_address(__global uint* p, uint src, uint imm3
*(p++) = imm32;
// v_and_b32 v28, s14, mask_reg
*(p++) = 0x2638000eu | (mask_reg << 9);
*(p++) = V_AND_B32_CALC_ADDRESS | (mask_reg << 9);
return p;
}
@ -70,7 +126,7 @@ __global uint* jit_scratchpad_load(__global uint* p, uint vgpr_index)
#if GCN_VERSION >= 14
// global_load_dwordx2 v[vgpr_index:vgpr_index+1], v28, s[0:1]
*(p++) = 0xdc548000u;
*(p++) = GLOBAL_LOAD_DWORDX2_SCRATCHPAD_LOAD;
*(p++) = 0x0000001cu | (vgpr_index << 24);
#else
*(p++) = 0x32543902u; // v_add_u32 v42, vcc, v2, v28
@ -87,14 +143,14 @@ __global uint* jit_scratchpad_load2(__global uint* p, uint vgpr_index, int vmcnt
{
// s_waitcnt vmcnt(N)
if (vmcnt >= 0)
*(p++) = 0xbf8c0f70u | (vmcnt & 15) | ((vmcnt >> 4) << 14);
*(p++) = S_WAITCNT_SCRATCHPAD_LOAD2 | (vmcnt & 15) | ((vmcnt >> 4) << 14);
// v_readlane_b32 s14, vgpr_index, 0
*(p++) = 0xd289000eu;
*(p++) = V_READLANE_B32_SCRATCHPAD_LOAD2 | 14;
*(p++) = 0x00010100u | vgpr_index;
// v_readlane_b32 s15, vgpr_index + 1, 0
*(p++) = 0xd289000fu;
*(p++) = V_READLANE_B32_SCRATCHPAD_LOAD2 | 15;
*(p++) = 0x00010100u | (vgpr_index + 1);
return p;
@ -107,9 +163,12 @@ __global uint* jit_scratchpad_calc_address_fp(__global uint* p, uint src, uint i
*(p++) = imm32;
// v_and_b32 v28, s14, mask_reg
*(p++) = 0x2638000eu | (mask_reg << 9);
*(p++) = V_AND_B32 | 0x38000eu | (mask_reg << 9);
#if GCN_VERSION >= 14
#if GCN_VERSION >= 15
// v_add_nc_u32 v28, v28, v44
*(p++) = 0x4a38591cu;
#elif GCN_VERSION == 14
// v_add_u32 v28, v28, v44
*(p++) = 0x6838591cu;
#else
@ -126,7 +185,7 @@ __global uint* jit_scratchpad_load_fp(__global uint* p, uint vgpr_index)
#if GCN_VERSION >= 14
// global_load_dword v(vgpr_index), v28, s[0:1]
*(p++) = 0xdc508000u;
*(p++) = GLOBAL_LOAD_DWORD_SCRATCHPAD_LOAD_FP;
*(p++) = 0x0000001cu | (vgpr_index << 24);
#else
*(p++) = 0x32543902u; // v_add_u32 v42, vcc, v2, v28
@ -143,7 +202,7 @@ __global uint* jit_scratchpad_load2_fp(__global uint* p, uint vgpr_index, int vm
{
// s_waitcnt vmcnt(N)
if (vmcnt >= 0)
*(p++) = 0xbf8c0f70u | (vmcnt & 15) | ((vmcnt >> 4) << 14);
*(p++) = S_WAITCNT_SCRATCHPAD_LOAD2 | (vmcnt & 15) | ((vmcnt >> 4) << 14);
// v_cvt_f64_i32 v[28:29], vgpr_index
*(p++) = 0x7e380900u | vgpr_index;
@ -164,7 +223,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (shift > 0) // p = 3/4
{
// s_lshl_b64 s[14:15], s[(16 + src * 2):(17 + src * 2)], shift
*(p++) = 0x8e8e8010u | (src << 1) | (shift << 8);
*(p++) = S_LSHL | 0x8e8010u | (src << 1) | (shift << 8);
// s_add_u32 s(16 + dst * 2), s(16 + dst * 2), s14
*(p++) = 0x80100e10u | (dst << 1) | (dst << 17);
@ -283,7 +342,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
{
#if GCN_VERSION >= 14
// s_mul_hi_u32 s15, s(16 + dst * 2), s(16 + src * 2)
*(p++) = 0x960f1010u | (dst << 1) | (src << 9);
*(p++) = S_MUL_HI_U32_IMUL_R | (dst << 1) | (src << 9);
#else
// v_mov_b32 v28, s(16 + dst * 2)
*(p++) = 0x7e380210u | (dst << 1);
@ -296,25 +355,25 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
#endif
// s_mul_i32 s14, s(16 + dst * 2), s(17 + src * 2)
*(p++) = 0x920e1110u | (dst << 1) | (src << 9);
*(p++) = S_MUL_I32_IMUL | 0x0e1110u | (dst << 1) | (src << 9);
// s_add_u32 s15, s15, s14
*(p++) = 0x800f0e0fu;
// s_mul_i32 s14, s(17 + dst * 2), s(16 + src * 2)
*(p++) = 0x920e1011u | (dst << 1) | (src << 9);
*(p++) = S_MUL_I32_IMUL | 0x0e1011u | (dst << 1) | (src << 9);
// s_add_u32 s(17 + dst * 2), s15, s14
*(p++) = 0x80110e0fu | (dst << 17);
// s_mul_i32 s(16 + dst * 2), s(16 + dst * 2), s(16 + src * 2)
*(p++) = 0x92101010u | (dst << 1) | (dst << 17) | (src << 9);
*(p++) = S_MUL_I32_IMUL | 0x101010u | (dst << 1) | (dst << 17) | (src << 9);
}
else // p = 1/8
{
#if GCN_VERSION >= 14
// s_mul_hi_u32 s15, s(16 + dst * 2), imm32
*(p++) = 0x960fff10u | (dst << 1);
*(p++) = S_MUL_HI_U32_IMUL_R_2 | (dst << 1);
*(p++) = inst.y;
#else
// v_mov_b32 v28, imm32
@ -335,14 +394,14 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
}
// s_mul_i32 s14, s(17 + dst * 2), imm32
*(p++) = 0x920eff11u | (dst << 1);
*(p++) = S_MUL_I32_IMUL | 0x0eff11u | (dst << 1);
*(p++) = inst.y;
// s_add_u32 s(17 + dst * 2), s15, s14
*(p++) = 0x80110e0fu | (dst << 17);
// s_mul_i32 s(16 + dst * 2), s(16 + dst * 2), imm32
*(p++) = 0x9210ff10u | (dst << 1) | (dst << 17);
*(p++) = S_MUL_I32_IMUL | 0x10ff10u | (dst << 1) | (dst << 17);
*(p++) = inst.y;
}
@ -369,7 +428,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
#if GCN_VERSION >= 14
// s_mul_hi_u32 s33, s(16 + dst * 2), s14
*(p++) = 0x96210e10u | (dst << 1);
*(p++) = S_MUL_HI_U32_IMUL_M | (dst << 1);
#else
// v_mov_b32 v28, s(16 + dst * 2)
*(p++) = 0x7e380210u | (dst << 1);
@ -382,19 +441,19 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
#endif
// s_mul_i32 s32, s(16 + dst * 2), s15
*(p++) = 0x92200f10u | (dst << 1);
*(p++) = S_MUL_I32_IMUL | 0x200f10u | (dst << 1);
// s_add_u32 s33, s33, s32
*(p++) = 0x80212021u;
// s_mul_i32 s32, s(17 + dst * 2), s14
*(p++) = 0x92200e11u | (dst << 1);
*(p++) = S_MUL_I32_IMUL | 0x200e11u | (dst << 1);
// s_add_u32 s(17 + dst * 2), s33, s32
*(p++) = 0x80112021u | (dst << 17);
// s_mul_i32 s(16 + dst * 2), s(16 + dst * 2), s14
*(p++) = 0x92100e10u | (dst << 1) | (dst << 17);
*(p++) = S_MUL_I32_IMUL | 0x100e10u | (dst << 1) | (dst << 17);
}
// (12*7/8 + 8*1/8 + 28) + 24 = 63.5 bytes on average
@ -404,10 +463,17 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_IMULH_R)
{
#if GCN_VERSION >= 15
*(p++) = 0xbe8e0410u | (dst << 1); // s_mov_b64 s[14:15], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbea60410u | (src << 1); // s_mov_b64 s[38:39], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbebc213au; // s_swappc_b64 s[60:61], s[58:59]
*(p++) = 0xbe90040eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#else
*(p++) = 0xbe8e0110u | (dst << 1); // s_mov_b64 s[14:15], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbea60110u | (src << 1); // s_mov_b64 s[38:39], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbebc1e3au; // s_swappc_b64 s[60:61], s[58:59]
*(p++) = 0xbe90010eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#endif
// 16 bytes
return p;
@ -430,9 +496,15 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
{
p = jit_scratchpad_load2(p, prefetch_vgpr_index ? -prefetch_vgpr_index : 28, prefetch_vgpr_index ? vmcnt : 0);
#if GCN_VERSION >= 15
*(p++) = 0xbea60410u | (dst << 1); // s_mov_b64 s[38:39], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbebc213au; // s_swappc_b64 s[60:61], s[58:59]
*(p++) = 0xbe90040eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#else
*(p++) = 0xbea60110u | (dst << 1); // s_mov_b64 s[38:39], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbebc1e3au; // s_swappc_b64 s[60:61], s[58:59]
*(p++) = 0xbe90010eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#endif
}
// (12*7/8 + 8*1/8 + 28) + 12 = 51.5 bytes on average
@ -442,10 +514,17 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_ISMULH_R)
{
#if GCN_VERSION >= 15
*(p++) = 0xbe8e0410u | (dst << 1); // s_mov_b64 s[14:15], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbea60410u | (src << 1); // s_mov_b64 s[38:39], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbebc2138u; // s_swappc_b64 s[60:61], s[56:57]
*(p++) = 0xbe90040eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#else
*(p++) = 0xbe8e0110u | (dst << 1); // s_mov_b64 s[14:15], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbea60110u | (src << 1); // s_mov_b64 s[38:39], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbebc1e38u; // s_swappc_b64 s[60:61], s[56:57]
*(p++) = 0xbe90010eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#endif
// 16 bytes
return p;
@ -468,9 +547,15 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
{
p = jit_scratchpad_load2(p, prefetch_vgpr_index ? -prefetch_vgpr_index : 28, prefetch_vgpr_index ? vmcnt : 0);
#if GCN_VERSION >= 15
*(p++) = 0xbea60410u | (dst << 1); // s_mov_b64 s[38:39], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbebc2138u; // s_swappc_b64 s[60:61], s[56:57]
*(p++) = 0xbe90040eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#else
*(p++) = 0xbea60110u | (dst << 1); // s_mov_b64 s[38:39], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbebc1e38u; // s_swappc_b64 s[60:61], s[56:57]
*(p++) = 0xbe90010eu | (dst << 17); // s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[14:15]
#endif
}
// (12*7/8 + 8*1/8 + 28) + 12 = 51.5 bytes on average
@ -484,10 +569,10 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
{
const uint2 rcp_value = as_uint2(imul_rcp_value(inst.y));
*(p++) = 0xbea000ffu; // s_mov_b32 s32, imm32
*(p++) = S_MOV_B32_IMUL_RCP; // s_mov_b32 s32, imm32
*(p++) = rcp_value.x;
#if GCN_VERSION >= 14
*(p++) = 0x960f2010u | (dst << 1); // s_mul_hi_u32 s15, s(16 + dst * 2), s32
*(p++) = S_MUL_HI_U32_IMUL_RCP | (dst << 1); // s_mul_hi_u32 s15, s(16 + dst * 2), s32
#else
// v_mov_b32 v28, s32
*(p++) = 0x7e380220u;
@ -498,12 +583,12 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
*(p++) = 0xd289000fu;
*(p++) = 0x0001011cu;
#endif
*(p++) = 0x920eff10u | (dst << 1); // s_mul_i32 s14, s(16 + dst * 2), imm32
*(p++) = S_MUL_I32_IMUL | 0x0eff10u | (dst << 1); // s_mul_i32 s14, s(16 + dst * 2), imm32
*(p++) = rcp_value.y;
*(p++) = 0x800f0e0fu; // s_add_u32 s15, s15, s14
*(p++) = 0x920e2011u | (dst << 1); // s_mul_i32 s14, s(17 + dst * 2), s32
*(p++) = S_MUL_I32_IMUL | 0x0e2011u | (dst << 1); // s_mul_i32 s14, s(17 + dst * 2), s32
*(p++) = 0x80110e0fu | (dst << 17); // s_add_u32 s(17 + dst * 2), s15, s14
*(p++) = 0x92102010u | (dst << 1) | (dst << 17);// s_mul_i32 s(16 + dst * 2), s(16 + dst * 2), s32
*(p++) = S_MUL_I32_IMUL | 0x102010u | (dst << 1) | (dst << 17);// s_mul_i32 s(16 + dst * 2), s(16 + dst * 2), s32
}
// 36 bytes
@ -526,23 +611,23 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (src != dst) // p = 7/8
{
// s_xor_b64 s[16 + dst * 2:17 + dst * 2], s[16 + dst * 2:17 + dst * 2], s[16 + src * 2:17 + src * 2]
*(p++) = 0x88901010u | (dst << 1) | (dst << 17) | (src << 9);
*(p++) = S_XOR_B32_64 | 0x901010u | (dst << 1) | (dst << 17) | (src << 9);
}
else // p = 1/8
{
if (as_int(inst.y) < 0) // p = 1/2
{
// s_mov_b32 s62, imm32
*(p++) = 0xbebe00ffu;
*(p++) = S_MOV_B32_XOR_R;
*(p++) = inst.y;
// s_xor_b64 s[16 + dst * 2:17 + dst * 2], s[16 + dst * 2:17 + dst * 2], s[62:63]
*(p++) = 0x88903e10u | (dst << 1) | (dst << 17);
*(p++) = S_XOR_B32_64 | 0x903e10u | (dst << 1) | (dst << 17);
}
else
{
// s_xor_b32 s(16 + dst * 2), s(16 + dst * 2), imm32
*(p++) = 0x8810ff10u | (dst << 1) | (dst << 17);
*(p++) = S_XOR_B32_64 | 0x10ff10u | (dst << 1) | (dst << 17);
*(p++) = inst.y;
}
}
@ -569,7 +654,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
p = jit_scratchpad_load2(p, prefetch_vgpr_index ? -prefetch_vgpr_index : 28, prefetch_vgpr_index ? vmcnt : 0);
// s_xor_b64 s[16 + dst * 2:17 + dst * 2], s[16 + dst * 2:17 + dst * 2], s[14:15]
*(p++) = 0x88900e10u | (dst << 1) | (dst << 17);
*(p++) = S_XOR_B32_64 | 0x900e10u | (dst << 1) | (dst << 17);
}
// (12*7/8 + 8*1/8 + 28) + 4 = 43.5 bytes on average
@ -584,24 +669,24 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_IROR_R)
{
// s_lshr_b64 s[32:33], s[16 + dst * 2:17 + dst * 2], s(16 + src * 2)
*(p++) = 0x8fa01010u | (dst << 1) | (src << 9);
*(p++) = S_LSHR | 0xa01010u | (dst << 1) | (src << 9);
// s_sub_u32 s15, 64, s(16 + src * 2)
*(p++) = 0x808f10c0u | (src << 9);
// s_lshl_b64 s[34:35], s[16 + dst * 2:17 + dst * 2], s15
*(p++) = 0x8ea20f10u | (dst << 1);
*(p++) = S_LSHL | 0xa20f10u | (dst << 1);
}
else
{
// s_lshl_b64 s[32:33], s[16 + dst * 2:17 + dst * 2], s(16 + src * 2)
*(p++) = 0x8ea01010u | (dst << 1) | (src << 9);
*(p++) = S_LSHL | 0xa01010u | (dst << 1) | (src << 9);
// s_sub_u32 s15, 64, s(16 + src * 2)
*(p++) = 0x808f10c0u | (src << 9);
// s_lshr_b64 s[34:35], s[16 + dst * 2:17 + dst * 2], s15
*(p++) = 0x8fa20f10u | (dst << 1);
*(p++) = S_LSHR | 0xa20f10u | (dst << 1);
}
}
else // p = 1/8
@ -609,14 +694,14 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
const uint shift = ((opcode < RANDOMX_FREQ_IROR_R) ? inst.y : -inst.y) & 63;
// s_lshr_b64 s[32:33], s[16 + dst * 2:17 + dst * 2], shift
*(p++) = 0x8fa08010u | (dst << 1) | (shift << 8);
*(p++) = S_LSHR | 0xa08010u | (dst << 1) | (shift << 8);
// s_lshl_b64 s[34:35], s[16 + dst * 2:17 + dst * 2], 64 - shift
*(p++) = 0x8ea28010u | (dst << 1) | ((64 - shift) << 8);
*(p++) = S_LSHL | 0xa28010u | (dst << 1) | ((64 - shift) << 8);
}
// s_or_b64 s[16 + dst * 2:17 + dst * 2], s[32:33], s[34:35]
*(p++) = 0x87902220u | (dst << 17);
*(p++) = S_OR | 0x902220u | (dst << 17);
// 12*7/8 + 8/8 + 4 = 15.5 bytes on average
return p;
@ -627,9 +712,15 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
{
if (src != dst)
{
#if GCN_VERSION >= 15
*(p++) = 0xbea00410u | (dst << 1); // s_mov_b64 s[32:33], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbe900410u | (src << 1) | (dst << 17);// s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbe900420u | (src << 17); // s_mov_b64 s[16 + src * 2:17 + Src * 2], s[32:33]
#else
*(p++) = 0xbea00110u | (dst << 1); // s_mov_b64 s[32:33], s[16 + dst * 2:17 + dst * 2]
*(p++) = 0xbe900110u | (src << 1) | (dst << 17);// s_mov_b64 s[16 + dst * 2:17 + dst * 2], s[16 + src * 2:17 + src * 2]
*(p++) = 0xbe900120u | (src << 17); // s_mov_b64 s[16 + src * 2:17 + Src * 2], s[32:33]
#endif
}
// 12*7/8 = 10.5 bytes on average
@ -640,11 +731,11 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_FSWAP_R)
{
// ds_swizzle_b32 v(60 + dst * 2), v(60 + dst * 2) offset:0x8001
*(p++) = 0xd87a8001u;
*(p++) = DS_SWIZZLE_B32_FSWAP_R;
*(p++) = 0x3c00003cu + (dst << 1) + (dst << 25);
// ds_swizzle_b32 v(61 + dst * 2), v(61 + dst * 2) offset:0x8001
*(p++) = 0xd87a8001u;
*(p++) = DS_SWIZZLE_B32_FSWAP_R;
*(p++) = 0x3d00003du + (dst << 1) + (dst << 25);
// s_waitcnt lgkmcnt(0)
@ -658,7 +749,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_FADD_R)
{
// v_add_f64 v[60 + dst * 2:61 + dst * 2], v[60 + dst * 2:61 + dst * 2], v[52 + src * 2:53 + src * 2]
*(p++) = 0xd280003cu + ((dst & 3) << 1);
*(p++) = V_ADD_F64 + ((dst & 3) << 1);
*(p++) = 0x0002693cu + ((dst & 3) << 1) + ((src & 3) << 10);
// 8 bytes
@ -679,7 +770,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
p = jit_scratchpad_load2_fp(p, prefetch_vgpr_index ? -prefetch_vgpr_index : 28, prefetch_vgpr_index ? vmcnt : 0);
// v_add_f64 v[60 + dst * 2:61 + dst * 2], v[60 + dst * 2:61 + dst * 2], v[28:29]
*(p++) = 0xd280003cu + ((dst & 3) << 1);
*(p++) = V_ADD_F64 + ((dst & 3) << 1);
*(p++) = 0x0002393cu + ((dst & 3) << 1);
}
@ -691,7 +782,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_FSUB_R)
{
// v_add_f64 v[60 + dst * 2:61 + dst * 2], v[60 + dst * 2:61 + dst * 2], -v[52 + src * 2:53 + src * 2]
*(p++) = 0xd280003cu + ((dst & 3) << 1);
*(p++) = V_ADD_F64 + ((dst & 3) << 1);
*(p++) = 0x4002693cu + ((dst & 3) << 1) + ((src & 3) << 10);
// 8 bytes
@ -712,7 +803,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
p = jit_scratchpad_load2_fp(p, prefetch_vgpr_index ? -prefetch_vgpr_index : 28, prefetch_vgpr_index ? vmcnt : 0);
// v_add_f64 v[60 + dst * 2:61 + dst * 2], v[60 + dst * 2:61 + dst * 2], -v[28:29]
*(p++) = 0xd280003cu + ((dst & 3) << 1);
*(p++) = V_ADD_F64 + ((dst & 3) << 1);
*(p++) = 0x4002393cu + ((dst & 3) << 1);
}
@ -724,7 +815,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_FSCAL_R)
{
// v_xor_b32 v(61 + dst * 2), v(61 + dst * 2), v51
*(p++) = 0x2a7a673du + ((dst & 3) << 1) + ((dst & 3) << 18);
*(p++) = (V_XOR_B32 | 0x7a673du) + ((dst & 3) << 1) + ((dst & 3) << 18);
// 4 bytes
return p;
@ -734,7 +825,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_FMUL_R)
{
// v_mul_f64 v[68 + dst * 2:69 + dst * 2], v[68 + dst * 2:69 + dst * 2], v[52 + src * 2:53 + src * 2]
*(p++) = 0xd2810044u + ((dst & 3) << 1);
*(p++) = V_MUL_F64 + ((dst & 3) << 1);
*(p++) = 0x00026944u + ((dst & 3) << 1) + ((src & 3) << 10);
// 8 bytes
@ -755,7 +846,11 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
p = jit_scratchpad_load2_fp(p, prefetch_vgpr_index ? -prefetch_vgpr_index : 28, prefetch_vgpr_index ? vmcnt : 0);
// s_swappc_b64 s[60:61], s[48 + dst * 2:49 + dst * 2]
#if GCN_VERSION >= 15
*(p++) = 0xbebc2130u + ((dst & 3) << 1);
#else
*(p++) = 0xbebc1e30u + ((dst & 3) << 1);
#endif
}
// 32 + 4 = 36 bytes
@ -766,7 +861,11 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
if (opcode < RANDOMX_FREQ_FSQRT_R)
{
// s_swappc_b64 s[60:61], s[40 + dst * 2:41 + dst * 2]
#if GCN_VERSION >= 15
*(p++) = 0xbebc2128u + ((dst & 3) << 1);
#else
*(p++) = 0xbebc1e28u + ((dst & 3) << 1);
#endif
// 4 bytes
return p;
@ -789,7 +888,7 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
const uint conditionMaskReg = 70 + (mod >> 4);
// s_and_b32 s14, s(16 + dst * 2), conditionMaskReg
*(p++) = 0x860e0010u | (dst << 1) | (conditionMaskReg << 8);
*(p++) = S_AND | 0x0e0010u | (dst << 1) | (conditionMaskReg << 8);
// s_cbranch_scc0 target
const int delta = ((last_branch_target - p) - 1);
@ -805,26 +904,30 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
const uint shift = inst.y & 63;
if (shift == 63)
{
*(p++) = 0x8e0e8110u | (src << 1); // s_lshl_b32 s14, s(16 + src * 2), 1
*(p++) = 0x8f0f9f11u | (src << 1); // s_lshr_b32 s15, s(17 + src * 2), 31
*(p++) = 0x870e0f0eu; // s_or_b32 s14, s14, s15
*(p++) = 0x860e830eu; // s_and_b32 s14, s14, 3
*(p++) = S_LSHL | 0x0e8110u | (src << 1); // s_lshl_b32 s14, s(16 + src * 2), 1
*(p++) = S_LSHR | 0x0f9f11u | (src << 1); // s_lshr_b32 s15, s(17 + src * 2), 31
*(p++) = S_OR | 0x0e0f0eu; // s_or_b32 s14, s14, s15
*(p++) = S_AND | 0x0e830eu; // s_and_b32 s14, s14, 3
}
else
{
// s_bfe_u64 s[14:15], s[16:17], (shift,width=2)
*(p++) = 0x938eff10u | (src << 1);
*(p++) = S_BFE | 0x8eff10u | (src << 1);
*(p++) = shift | (2 << 16);
}
// s_brev_b32 s14, s14
*(p++) = 0xbe8e080eu;
// s_lshr_b32 s66, s14, 30
*(p++) = 0x8f429e0eu;
// s_setreg_b32 hwreg(mode, 2, 2), s66
#if GCN_VERSION >= 15
*(p++) = 0xbe8e0b0eu;
*(p++) = 0x90429e0eu;
*(p++) = 0xb9c20881u;
#else
*(p++) = 0xbe8e080eu;
*(p++) = 0x8f429e0eu;
*(p++) = 0xb9420881u;
#endif
// 20 bytes
return p;
@ -843,6 +946,10 @@ __global uint* jit_emit_instruction(__global uint* p, __global uint* last_branch
// v28 = offset
#if GCN_VERSION >= 14
#if GCN_VERSION >= 15
// s_waitcnt vmcnt(0)
*(p++) = 0xbf8c3f70u;
#endif
// global_store_dwordx2 v28, v[vgpr_id:vgpr_id + 1], s[0:1]
*(p++) = 0xdc748000u;
*(p++) = 0x0000001cu | (vgpr_id << 8);
@ -1419,14 +1526,14 @@ __global uint* generate_jit_code(__global uint2* e, __global uint2* p0, __global
{
// Code size limit exceeded!!!
// Jump back to randomx_run kernel
*(p++) = 0xbe801d0cu; // s_setpc_b64 s[12:13]
*(p++) = S_SETPC_B64_S12_13; // s_setpc_b64 s[12:13]
return p;
}
} while (!done);
}
// Jump back to randomx_run kernel
*(p++) = 0xbe801d0cu; // s_setpc_b64 s[12:13]
*(p++) = S_SETPC_B64_S12_13; // s_setpc_b64 s[12:13]
return p;
}

View file

@ -0,0 +1,742 @@
/*
Copyright (c) 2019-2020 SChernykh
This file is part of RandomX OpenCL.
RandomX OpenCL 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.
RandomX OpenCL 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 RandomX OpenCL. If not, see <http://www.gnu.org/licenses/>.
*/
.rocm
.gpu GFX1010
.arch_minor 1
.arch_stepping 0
.eflags 53
.llvm10binfmt
.metadatav3
.md_version 1, 0
.globaldata
.fill 64, 1, 0
.kernel randomx_run
.config
.dims x
.sgprsnum 96
.vgprsnum 128
.shared_vgprs 0
.dx10clamp
.ieeemode
.floatmode 0xf0
.priority 0
.exceptions 0
.userdatanum 6
# https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-compute-pgm-rsrc1-gfx6-gfx10-table
# https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-compute-pgm-rsrc2-gfx6-gfx10-table
# https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-compute-pgm-rsrc3-gfx10-table
.pgmrsrc1 0x40af0105
.pgmrsrc2 0x0000008c
.pgmrsrc3 0x00000000
.group_segment_fixed_size 256
.private_segment_fixed_size 0
.kernel_code_entry_offset 0x10c0
.use_private_segment_buffer
.use_kernarg_segment_ptr
.use_wave32
.config
.md_symname "randomx_run.kd"
.md_language "OpenCL C", 1, 2
.reqd_work_group_size 32, 1, 1
.md_kernarg_segment_size 104
.md_kernarg_segment_align 8
.md_group_segment_fixed_size 256
.md_private_segment_fixed_size 0
.md_wavefront_size 32
.md_sgprsnum 96
.md_vgprsnum 128
.spilledsgprs 0
.spilledvgprs 0
.max_flat_work_group_size 32
.arg dataset, "uchar*", 8, 0, globalbuf, u8, global, default const
.arg scratchpad, "uchar*", 8, 8, globalbuf, u8, global, default
.arg registers, "ulong*", 8, 16, globalbuf, u64, global, default
.arg rounding_modes, "uint*", 8, 24, globalbuf, u32, global, default
.arg programs, "uint*", 8, 32, globalbuf, u32, global, default
.arg batch_size, "uint", 4, 40, value, u32
.arg rx_parameters, "uint", 4, 44, value, u32
.arg , "", 8, 48, gox, i64
.arg , "", 8, 56, goy, i64
.arg , "", 8, 64, goz, i64
.arg , "", 8, 72, none, i8
.arg , "", 8, 80, none, i8
.arg , "", 8, 88, none, i8
.arg , "", 8, 96, multigridsyncarg, i8
.text
randomx_run:
# clear all caches
s_dcache_wb
s_waitcnt vmcnt(0) & lgkmcnt(0)
s_waitcnt_vscnt null, 0x0
s_icache_inv
s_branch begin
# pgmrsrc2 = 0x0000008c, bits 1:5 = 6, so first 6 SGPRs (s0-s7) contain user data
# s6 contains group id
# v0 contains local id
begin:
# s[0:1] - pointer to registers
# s[2:3] - pointer to rounding modes
s_load_dwordx4 s[0:3], s[4:5], 0x10
# s[8:9] - group_id*group_size
s_mov_b32 s9, 0
s_lshl_b32 s8, s6, 5
# v0 - local id (sub)
# v39 - R[sub]
v_lshlrev_b32 v39, 3, v0
s_mov_b32 s12, s7
# vcc_lo = "if (sub < 8)"
v_cmp_gt_u32 vcc_lo, 8, v0
s_waitcnt lgkmcnt(0)
# load rounding mode
s_lshl_b32 s16, s6, 2
s_add_u32 s64, s2, s16
s_addc_u32 s65, s3, 0
v_mov_b32 v1, 0
global_load_dword v1, v1, s[64:65]
s_waitcnt vmcnt(0)
v_readlane_b32 s66, v1, 0
s_setreg_b32 hwreg(mode, 2, 2), s66
s_mov_b32 s67, 0
# ((__local ulong*) R)[sub] = ((__global ulong*) registers)[sub];
s_lshl_b64 s[2:3], s[8:9], 3
s_mov_b32 s32, s12
s_add_u32 s0, s0, s2
s_addc_u32 s1, s1, s3
v_add_co_u32 v1, s0, s0, v39
v_add_co_ci_u32 v2, s0, s1, 0, s0
global_load_dwordx2 v[4:5], v[1:2], off
s_waitcnt vmcnt(0)
ds_write_b64 v39, v[4:5]
s_waitcnt vmcnt(0) & lgkmcnt(0)
s_waitcnt_vscnt null, 0x0
# "if (sub >= 8) return"
s_and_saveexec_b32 s0, vcc_lo
s_cbranch_execz program_end
# s[8:9] - pointer to dataset
# s[10:11] - pointer to scratchpads
# s[0:1] - pointer to programs
s_load_dwordx4 s[8:11], s[4:5], 0x0
s_load_dwordx2 s[0:1], s[4:5], 0x20
# rx_parameters
s_load_dword s20, s[4:5], 0x2c
v_mov_b32 v5, 0
v_mov_b32 v10, 0
s_waitcnt_vscnt null, 0x0
ds_read_b64 v[8:9], v39
v_cmp_gt_u32 vcc_lo, 4, v0
v_lshlrev_b32 v0, 3, v0
ds_read2_b64 v[25:28], v5 offset0:16 offset1:17
ds_read_b32 v11, v5 offset:152
ds_read_b64 v[35:36], v5 offset:168
ds_read2_b64 v[20:23], v5 offset0:18 offset1:20
v_cndmask_b32 v4, 0xffffff, -1, vcc_lo
v_add_nc_u32 v5, v39, v0
s_waitcnt lgkmcnt(0)
v_mov_b32 v13, s11
v_mov_b32 v7, s1
v_mov_b32 v6, s0
# Scratchpad L1 size
s_bfe_u32 s21, s20, 0x050000
s_lshl_b32 s21, 1, s21
# Scratchpad L2 size
s_bfe_u32 s22, s20, 0x050005
s_lshl_b32 s22, 1, s22
# Scratchpad L3 size
s_bfe_u32 s0, s20, 0x05000A
s_lshl_b32 s23, 1, s0
# program iterations
s_bfe_u32 s24, s20, 0x04000F
s_lshl_b32 s24, 1, s24
v_mov_b32 v12, s10
v_mad_u64_u32 v[6:7], s2, 10048, s6, v[6:7]
# s[4:5] - pointer to current program
v_readlane_b32 s4, v6, 0
v_readlane_b32 s5, v7, 0
s_lshl_b32 s2, 1, s0
v_add_co_u32 v14, s0, s8, v11
v_cndmask_b32 v34, v36, 0, vcc_lo
v_cndmask_b32 v24, v23, 0, vcc_lo
v_cndmask_b32 v3, v22, 0, vcc_lo
s_add_i32 s3, s2, 64
v_add_co_ci_u32 v29, s0, s9, v10, s0
v_cndmask_b32 v35, v35, 0, vcc_lo
v_add_co_u32 v22, vcc_lo, v14, v0
# v[12:13] - pointer to current scratchpad
v_mad_u64_u32 v[12:13], s2, s3, s6, v[12:13]
v_mov_b32 v10, v26
v_mov_b32 v11, v25
v_lshlrev_b32 v36, 3, v27
v_lshlrev_b32 v37, 3, v28
v_lshlrev_b32 v20, 3, v20
v_lshlrev_b32 v21, 3, v21
v_add_co_ci_u32 v23, vcc_lo, 0, v29, vcc_lo
# rename registers
# v6 - R[sub]
v_mov_b32 v6, v39
# loop counter
s_sub_u32 s2, s24, 1
# used in IXOR_R instruction
s_mov_b32 s63, -1
# used in CBRANCH instruction
s_mov_b32 s70, (0xFF << 8)
s_mov_b32 s71, (0xFF << 9)
s_mov_b32 s72, (0xFF << 10)
s_mov_b32 s73, (0xFF << 11)
s_mov_b32 s74, (0xFF << 12)
s_mov_b32 s75, (0xFF << 13)
s_mov_b32 s76, (0xFF << 14)
s_mov_b32 s77, (0xFF << 15)
s_mov_b32 s78, (0xFF << 16)
s_mov_b32 s79, (0xFF << 17)
s_mov_b32 s80, (0xFF << 18)
s_mov_b32 s81, (0xFF << 19)
s_mov_b32 s82, (0xFF << 20)
s_mov_b32 s83, (0xFF << 21)
s_mov_b32 s84, (0xFF << 22)
s_mov_b32 s85, (0xFF << 23)
# ScratchpadL3Mask64
s_sub_u32 s86, s23, 64
# Scratchpad masks for scratchpads
v_sub_nc_u32 v38, s21, 8
v_sub_nc_u32 v39, s22, 8
v_sub_nc_u32 v50, s23, 8
# mask for FSCAL_R
v_mov_b32 v51, 0x80F00000
# load scratchpad base address
v_readlane_b32 s0, v12, 0
v_readlane_b32 s1, v13, 0
# v41, v44 = 0
v_mov_b32 v41, 0
v_mov_b32 v44, 0
# v41 = 0 on lane 0, set it to 8 on lane 1
# v44 = 0 on lane 0, set it to 4 on lane 1
s_mov_b64 exec, 2
v_mov_b32 v41, 8
v_mov_b32 v44, 4
# load group A registers
# Read low 8 bytes into lane 0 and high 8 bytes into lane 1
s_mov_b64 exec, 3
ds_read2_b64 v[52:55], v41 offset0:24 offset1:26
ds_read2_b64 v[56:59], v41 offset0:28 offset1:30
# xmantissaMask
v_mov_b32 v77, (1 << 24) - 1
# xexponentMask
ds_read_b64 v[78:79], v41 offset:160
# Restore execution mask
s_mov_b64 exec, 255
# sign mask (used in FSQRT_R)
v_mov_b32 v82, 0x80000000
# used in FSQRT_R to check for "positive normal value" (v_cmpx_class_f64)
s_mov_b32 s68, 256
s_mov_b32 s69, 0
# High 32 bits of "1.0" constant (used in FDIV_M)
v_mov_b32 v83, (1023 << 20)
# Used to multiply FP64 values by 0.5
v_mov_b32 v84, (1 << 20)
s_getpc_b64 s[14:15]
cur_addr:
# get addresses of FSQRT_R subroutines
s_add_u32 s40, s14, fsqrt_r_sub0 - cur_addr
s_addc_u32 s41, s15, 0
s_add_u32 s42, s14, fsqrt_r_sub1 - cur_addr
s_addc_u32 s43, s15, 0
s_add_u32 s44, s14, fsqrt_r_sub2 - cur_addr
s_addc_u32 s45, s15, 0
s_add_u32 s46, s14, fsqrt_r_sub3 - cur_addr
s_addc_u32 s47, s15, 0
# get addresses of FDIV_M subroutines
s_add_u32 s48, s14, fdiv_m_sub0 - cur_addr
s_addc_u32 s49, s15, 0
s_add_u32 s50, s14, fdiv_m_sub1 - cur_addr
s_addc_u32 s51, s15, 0
s_add_u32 s52, s14, fdiv_m_sub2 - cur_addr
s_addc_u32 s53, s15, 0
s_add_u32 s54, s14, fdiv_m_sub3 - cur_addr
s_addc_u32 s55, s15, 0
# get address for ISMULH_R subroutine
s_add_u32 s56, s14, ismulh_r_sub - cur_addr
s_addc_u32 s57, s15, 0
# get address for IMULH_R subroutine
s_add_u32 s58, s14, imulh_r_sub - cur_addr
s_addc_u32 s59, s15, 0
/*
used: v0-v6, v8-v37
not used: v7
*/
main_loop:
s_waitcnt_vscnt null, 0x0
# v[27:28] = R[readReg0]
# v[29:30] = R[readReg1]
ds_read_b64 v[27:28], v37
ds_read_b64 v[29:30], v36
s_waitcnt lgkmcnt(0)
# R[readReg0] ^ R[readReg0] (high 32 bits)
v_xor_b32 v28, v28, v30
# spAddr1
v_xor_b32 v25, v28, v25
v_and_b32 v25, s86, v25
v_add_nc_u32 v25, v25, v0
v_add_co_u32 v16, vcc_lo, s0, v25
# R[readReg0] ^ R[readReg0] (low 32 bits)
v_xor_b32 v25, v27, v29
v_mov_b32 v29, v11
v_add_co_ci_u32 v17, vcc_lo, 0, s1, vcc_lo
v_xor_b32 v25, v25, v26
# load from spAddr1
global_load_dwordx2 v[27:28], v[16:17], off
# spAddr0
v_and_b32 v25, s86, v25
v_add_nc_u32 v25, v25, v0
v_add_co_u32 v31, vcc_lo, s0, v25
v_add_co_ci_u32 v32, vcc_lo, 0, s1, vcc_lo
v_add_co_u32 v29, vcc_lo, v22, v29
# load from spAddr0
global_load_dwordx2 v[25:26], v[31:32], off
v_add_co_ci_u32 v30, vcc_lo, 0, v23, vcc_lo
v_mov_b32 v33, v11
s_and_b32 vcc_lo, exec_lo, vcc_lo
s_waitcnt vmcnt(1)
v_cvt_f64_i32 v[14:15], v28
v_cvt_f64_i32 v[12:13], v27
v_or_b32 v14, v14, v35
s_waitcnt vmcnt(0)
# R[sub] ^= *p0;
v_xor_b32 v8, v25, v8
v_xor_b32 v9, v26, v9
v_and_b32 v26, v4, v15
v_and_b32 v19, v4, v13
v_or_b32 v15, v26, v34
v_or_b32 v18, v12, v3
v_mov_b32 v26, 0
v_or_b32 v19, v19, v24
v_mov_b32 v25, v26
ds_write2_b64 v5, v[18:19], v[14:15] offset0:8 offset1:9
# load from dataset
global_load_dwordx2 v[18:19], v[29:30], off
# load group F,E registers
# Read low 8 bytes into lane 0 and high 8 bytes into lane 1
s_mov_b64 exec, 3
s_waitcnt lgkmcnt(0)
ds_read2_b64 v[60:63], v41 offset0:8 offset1:10
ds_read2_b64 v[64:67], v41 offset0:12 offset1:14
ds_read2_b64 v[68:71], v41 offset0:16 offset1:18
ds_read2_b64 v[72:75], v41 offset0:20 offset1:22
# load VM integer registers
v_readlane_b32 s16, v8, 0
v_readlane_b32 s17, v9, 0
v_readlane_b32 s18, v8, 1
v_readlane_b32 s19, v9, 1
v_readlane_b32 s20, v8, 2
v_readlane_b32 s21, v9, 2
v_readlane_b32 s22, v8, 3
v_readlane_b32 s23, v9, 3
v_readlane_b32 s24, v8, 4
v_readlane_b32 s25, v9, 4
v_readlane_b32 s26, v8, 5
v_readlane_b32 s27, v9, 5
v_readlane_b32 s28, v8, 6
v_readlane_b32 s29, v9, 6
v_readlane_b32 s30, v8, 7
v_readlane_b32 s31, v9, 7
s_waitcnt lgkmcnt(0)
# Use only first 2 lanes for the program
s_mov_b64 exec, 3
# call JIT code
s_swappc_b64 s[12:13], s[4:5]
# Write out group F,E registers
# Write low 8 bytes from lane 0 and high 8 bytes from lane 1
ds_write2_b64 v41, v[60:61], v[62:63] offset0:8 offset1:10
ds_write2_b64 v41, v[64:65], v[66:67] offset0:12 offset1:14
ds_write2_b64 v41, v[68:69], v[70:71] offset0:16 offset1:18
ds_write2_b64 v41, v[72:73], v[74:75] offset0:20 offset1:22
# store VM integer registers
v_writelane_b32 v8, s16, 0
v_writelane_b32 v9, s17, 0
v_writelane_b32 v8, s18, 1
v_writelane_b32 v9, s19, 1
v_writelane_b32 v8, s20, 2
v_writelane_b32 v9, s21, 2
v_writelane_b32 v8, s22, 3
v_writelane_b32 v9, s23, 3
v_writelane_b32 v8, s24, 4
v_writelane_b32 v9, s25, 4
v_writelane_b32 v8, s26, 5
v_writelane_b32 v9, s27, 5
v_writelane_b32 v8, s28, 6
v_writelane_b32 v9, s29, 6
v_writelane_b32 v8, s30, 7
v_writelane_b32 v9, s31, 7
# Turn back on 8 execution lanes
s_mov_b64 exec, 255
# Write out VM integer registers
ds_write_b64 v6, v[8:9]
s_waitcnt lgkmcnt(0)
# R[readReg2], R[readReg3]
ds_read_b32 v11, v21
ds_read_b32 v27, v20
s_waitcnt lgkmcnt(0)
# mx ^= R[readReg2] ^ R[readReg3];
v_xor_b32 v11, v11, v27
v_xor_b32 v10, v10, v11
# v[27:28] = R[sub]
# v[29:30] = F[sub]
ds_read2_b64 v[27:30], v6 offset1:8
# mx &= CacheLineAlignMask;
v_and_b32 v11, 0x7fffffc0, v10
v_mov_b32 v10, v33
s_waitcnt lgkmcnt(0)
# const ulong next_r = R[sub] ^ data;
s_waitcnt lgkmcnt(0)
v_xor_b32 v8, v27, v18
v_xor_b32 v9, v28, v19
# *p1 = next_r;
global_store_dwordx2 v[16:17], v[8:9], off
# v[27:28] = E[sub]
ds_read_b64 v[27:28], v6 offset:128
# R[sub] = next_r;
ds_write_b64 v6, v[8:9]
s_waitcnt lgkmcnt(1)
# *p0 = as_ulong(F[sub]) ^ as_ulong(E[sub]);
v_xor_b32 v29, v27, v29
v_xor_b32 v30, v28, v30
global_store_dwordx2 v[31:32], v[29:30], off
s_sub_u32 s2, s2, 1
s_cbranch_scc0 main_loop
main_loop_end:
global_store_dwordx2 v[1:2], v[8:9], off
global_store_dwordx2 v[1:2], v[29:30], off inst_offset:64
global_store_dwordx2 v[1:2], v[27:28], off inst_offset:128
# store rounding mode
v_mov_b32 v0, 0
v_mov_b32 v1, s66
global_store_dword v0, v1, s[64:65]
program_end:
s_endpgm
fsqrt_r_sub0:
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rsq_f64 v[28:29], v[68:69]
# Improve initial approximation (can be skipped)
#v_mul_f64 v[42:43], v[28:29], v[68:69]
#v_mul_f64 v[48:49], v[28:29], -0.5
#v_fma_f64 v[48:49], v[48:49], v[42:43], 0.5
#v_fma_f64 v[28:29], v[28:29], v[48:49], v[28:29]
v_mul_f64 v[42:43], v[28:29], v[68:69]
v_mov_b32 v48, v28
v_sub_nc_u32 v49, v29, v84
v_mov_b32 v46, v28
v_xor_b32 v47, v49, v82
v_fma_f64 v[46:47], v[46:47], v[42:43], 0.5
v_fma_f64 v[42:43], v[42:43], v[46:47], v[42:43]
v_fma_f64 v[48:49], v[48:49], v[46:47], v[48:49]
v_fma_f64 v[46:47], -v[42:43], v[42:43], v[68:69]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[46:47], v[48:49], v[42:43]
v_cmpx_class_f64 v[68:69], s[68:69]
v_mov_b32 v68, v42
v_mov_b32 v69, v43
s_mov_b64 exec, 3
s_setpc_b64 s[60:61]
fsqrt_r_sub1:
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rsq_f64 v[28:29], v[70:71]
# Improve initial approximation (can be skipped)
#v_mul_f64 v[42:43], v[28:29], v[70:71]
#v_mul_f64 v[48:49], v[28:29], -0.5
#v_fma_f64 v[48:49], v[48:49], v[42:43], 0.5
#v_fma_f64 v[28:29], v[28:29], v[48:49], v[28:29]
v_mul_f64 v[42:43], v[28:29], v[70:71]
v_mov_b32 v48, v28
v_sub_nc_u32 v49, v29, v84
v_mov_b32 v46, v28
v_xor_b32 v47, v49, v82
v_fma_f64 v[46:47], v[46:47], v[42:43], 0.5
v_fma_f64 v[42:43], v[42:43], v[46:47], v[42:43]
v_fma_f64 v[48:49], v[48:49], v[46:47], v[48:49]
v_fma_f64 v[46:47], -v[42:43], v[42:43], v[70:71]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[46:47], v[48:49], v[42:43]
v_cmpx_class_f64 v[70:71], s[68:69]
v_mov_b32 v70, v42
v_mov_b32 v71, v43
s_mov_b64 exec, 3
s_setpc_b64 s[60:61]
fsqrt_r_sub2:
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rsq_f64 v[28:29], v[72:73]
# Improve initial approximation (can be skipped)
#v_mul_f64 v[42:43], v[28:29], v[72:73]
#v_mul_f64 v[48:49], v[28:29], -0.5
#v_fma_f64 v[48:49], v[48:49], v[42:43], 0.5
#v_fma_f64 v[28:29], v[28:29], v[48:49], v[28:29]
v_mul_f64 v[42:43], v[28:29], v[72:73]
v_mov_b32 v48, v28
v_sub_nc_u32 v49, v29, v84
v_mov_b32 v46, v28
v_xor_b32 v47, v49, v82
v_fma_f64 v[46:47], v[46:47], v[42:43], 0.5
v_fma_f64 v[42:43], v[42:43], v[46:47], v[42:43]
v_fma_f64 v[48:49], v[48:49], v[46:47], v[48:49]
v_fma_f64 v[46:47], -v[42:43], v[42:43], v[72:73]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[46:47], v[48:49], v[42:43]
v_cmpx_class_f64 v[72:73], s[68:69]
v_mov_b32 v72, v42
v_mov_b32 v73, v43
s_mov_b64 exec, 3
s_setpc_b64 s[60:61]
fsqrt_r_sub3:
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rsq_f64 v[28:29], v[74:75]
# Improve initial approximation (can be skipped)
#v_mul_f64 v[42:43], v[28:29], v[74:75]
#v_mul_f64 v[48:49], v[28:29], -0.5
#v_fma_f64 v[48:49], v[48:49], v[42:43], 0.5
#v_fma_f64 v[28:29], v[28:29], v[48:49], v[28:29]
v_mul_f64 v[42:43], v[28:29], v[74:75]
v_mov_b32 v48, v28
v_sub_nc_u32 v49, v29, v84
v_mov_b32 v46, v28
v_xor_b32 v47, v49, v82
v_fma_f64 v[46:47], v[46:47], v[42:43], 0.5
v_fma_f64 v[42:43], v[42:43], v[46:47], v[42:43]
v_fma_f64 v[48:49], v[48:49], v[46:47], v[48:49]
v_fma_f64 v[46:47], -v[42:43], v[42:43], v[74:75]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[46:47], v[48:49], v[42:43]
v_cmpx_class_f64 v[74:75], s[68:69]
v_mov_b32 v74, v42
v_mov_b32 v75, v43
s_mov_b64 exec, 3
s_setpc_b64 s[60:61]
fdiv_m_sub0:
v_or_b32 v28, v28, v78
v_and_or_b32 v29, v29, v77, v79
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rcp_f64 v[48:49], v[28:29]
v_fma_f64 v[80:81], -v[28:29], v[48:49], 1.0
v_fma_f64 v[48:49], v[48:49], v[80:81], v[48:49]
v_mul_f64 v[80:81], v[68:69], v[48:49]
v_fma_f64 v[42:43], -v[28:29], v[80:81], v[68:69]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[42:43], v[48:49], v[80:81]
v_div_fixup_f64 v[80:81], v[42:43], v[28:29], v[68:69]
v_cmpx_eq_f64 v[68:69], v[28:29]
v_mov_b32 v80, 0
v_mov_b32 v81, v83
s_mov_b64 exec, 3
v_mov_b32 v68, v80
v_mov_b32 v69, v81
s_setpc_b64 s[60:61]
fdiv_m_sub1:
v_or_b32 v28, v28, v78
v_and_or_b32 v29, v29, v77, v79
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rcp_f64 v[48:49], v[28:29]
v_fma_f64 v[80:81], -v[28:29], v[48:49], 1.0
v_fma_f64 v[48:49], v[48:49], v[80:81], v[48:49]
v_mul_f64 v[80:81], v[70:71], v[48:49]
v_fma_f64 v[42:43], -v[28:29], v[80:81], v[70:71]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[42:43], v[48:49], v[80:81]
v_div_fixup_f64 v[80:81], v[42:43], v[28:29], v[70:71]
v_cmpx_eq_f64 v[70:71], v[28:29]
v_mov_b32 v80, 0
v_mov_b32 v81, v83
s_mov_b64 exec, 3
v_mov_b32 v70, v80
v_mov_b32 v71, v81
s_setpc_b64 s[60:61]
fdiv_m_sub2:
v_or_b32 v28, v28, v78
v_and_or_b32 v29, v29, v77, v79
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rcp_f64 v[48:49], v[28:29]
v_fma_f64 v[80:81], -v[28:29], v[48:49], 1.0
v_fma_f64 v[48:49], v[48:49], v[80:81], v[48:49]
v_mul_f64 v[80:81], v[72:73], v[48:49]
v_fma_f64 v[42:43], -v[28:29], v[80:81], v[72:73]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[42:43], v[48:49], v[80:81]
v_div_fixup_f64 v[80:81], v[42:43], v[28:29], v[72:73]
v_cmpx_eq_f64 v[72:73], v[28:29]
v_mov_b32 v80, 0
v_mov_b32 v81, v83
s_mov_b64 exec, 3
v_mov_b32 v72, v80
v_mov_b32 v73, v81
s_setpc_b64 s[60:61]
fdiv_m_sub3:
v_or_b32 v28, v28, v78
v_and_or_b32 v29, v29, v77, v79
s_setreg_b32 hwreg(mode, 2, 2), s67
v_rcp_f64 v[48:49], v[28:29]
v_fma_f64 v[80:81], -v[28:29], v[48:49], 1.0
v_fma_f64 v[48:49], v[48:49], v[80:81], v[48:49]
v_mul_f64 v[80:81], v[74:75], v[48:49]
v_fma_f64 v[42:43], -v[28:29], v[80:81], v[74:75]
s_setreg_b32 hwreg(mode, 2, 2), s66
v_fma_f64 v[42:43], v[42:43], v[48:49], v[80:81]
v_div_fixup_f64 v[80:81], v[42:43], v[28:29], v[74:75]
v_cmpx_eq_f64 v[74:75], v[28:29]
v_mov_b32 v80, 0
v_mov_b32 v81, v83
s_mov_b64 exec, 3
v_mov_b32 v74, v80
v_mov_b32 v75, v81
s_setpc_b64 s[60:61]
ismulh_r_sub:
s_mov_b64 exec, 1
v_mov_b32 v45, s14
v_mul_hi_u32 v40, s38, v45
v_mov_b32 v47, s15
v_mad_u64_u32 v[42:43], s32, s38, v47, v[40:41]
v_mov_b32 v40, v42
v_mad_u64_u32 v[45:46], s32, s39, v45, v[40:41]
v_mad_u64_u32 v[42:43], s32, s39, v47, v[43:44]
v_add_co_u32 v42, vcc_lo, v42, v46
v_add_co_ci_u32 v43, vcc_lo, 0, v43, vcc_lo
v_readlane_b32 s32, v42, 0
v_readlane_b32 s33, v43, 0
s_cmp_lt_i32 s15, 0
s_cselect_b64 s[34:35], s[38:39], 0
s_sub_u32 s32, s32, s34
s_subb_u32 s33, s33, s35
s_cmp_lt_i32 s39, 0
s_cselect_b64 s[34:35], s[14:15], 0
s_sub_u32 s14, s32, s34
s_subb_u32 s15, s33, s35
s_mov_b64 exec, 3
s_setpc_b64 s[60:61]
imulh_r_sub:
s_mov_b64 exec, 1
v_mov_b32 v45, s38
v_mul_hi_u32 v40, s14, v45
v_mov_b32 v47, s39
v_mad_u64_u32 v[42:43], s32, s14, v47, v[40:41]
v_mov_b32 v40, v42
v_mad_u64_u32 v[45:46], s32, s15, v45, v[40:41]
v_mad_u64_u32 v[42:43], s32, s15, v47, v[43:44]
v_add_co_u32 v42, vcc_lo, v42, v46
v_add_co_ci_u32 v43, vcc_lo, 0, v43, vcc_lo
v_readlane_b32 s14, v42, 0
v_readlane_b32 s15, v43, 0
s_mov_b64 exec, 3
s_setpc_b64 s[60:61]

View file

@ -0,0 +1,303 @@
/*
This file was auto-generated from randomx_run_gfx1010.asm:
clrxasm randomx_run_gfx1010.asm -o randomx_run_gfx1010.bin
bin2h -c randomx_run_gfx1010_bin < randomx_run_gfx1010.bin > randomx_run_gfx1010.h
clrxasm can be downloaded here: https://github.com/CLRX/CLRX-mirror/releases
bin2h can be downloaded here: http://www.deadnode.org/sw/bin2h/
*/
static unsigned char randomx_run_gfx1010_bin[]={
0x7f,0x45,0x4c,0x46,0x02,0x01,0x01,0x40,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x03,0x00,0xe0,0x00,0x01,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x40
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x60,0x21,0x00,0x00,0x00,0x00,0x00,0x00,0x35,0x00,0x00,0x00,0x40,0x00,0x38,0x00,0x08,0x00,0x40,0x00,0x0c,0x00,0x0a,0x00,0x06
,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x40,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x40,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x40,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0xc0
,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0xc0,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x09,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x09,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x05,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x28,0x0a,0x00,0x00,0x00,0x00,0x00,0x00,0x28,0x0a,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x06,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x60,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x60,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x02
,0x00,0x00,0x00,0x06,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x60
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x60,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x52,0xe5,0x74,0x64,0x04,0x00,0x00,0x00,0x00
,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x60,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x60
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x51,0xe5,0x74,0x64,0x06,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x34,0x06,0x00,0x00,0x00,0x00,0x00,0x00,0x34,0x06,0x00,0x00,0x00,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x07
,0x00,0x00,0x00,0x20,0x06,0x00,0x00,0x20,0x00,0x00,0x00,0x41,0x4d,0x44,0x47,0x50,0x55,0x00,0x00,0x82,0xae,0x61,0x6d,0x64,0x68,0x73,0x61,0x2e,0x6b,0x65,0x72,0x6e
,0x65,0x6c,0x73,0x91,0xde,0x00,0x10,0xa5,0x2e,0x61,0x72,0x67,0x73,0x9e,0x88,0xae,0x2e,0x61,0x64,0x64,0x72,0x65,0x73,0x73,0x5f,0x73,0x70,0x61,0x63,0x65,0xa6,0x67
,0x6c,0x6f,0x62,0x61,0x6c,0xa9,0x2e,0x69,0x73,0x5f,0x63,0x6f,0x6e,0x73,0x74,0xc3,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xa7,0x64,0x61,0x74,0x61,0x73,0x65,0x74,0xa7,0x2e
,0x6f,0x66,0x66,0x73,0x65,0x74,0x00,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xaa,0x2e,0x74,0x79,0x70,0x65,0x5f,0x6e,0x61,0x6d,0x65,0xa6,0x75,0x63,0x68,0x61,0x72,0x2a
,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xad,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x62,0x75,0x66,0x66,0x65,0x72,0xab,0x2e,0x76,0x61,0x6c,0x75
,0x65,0x5f,0x74,0x79,0x70,0x65,0xa2,0x75,0x38,0x87,0xae,0x2e,0x61,0x64,0x64,0x72,0x65,0x73,0x73,0x5f,0x73,0x70,0x61,0x63,0x65,0xa6,0x67,0x6c,0x6f,0x62,0x61,0x6c
,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xaa,0x73,0x63,0x72,0x61,0x74,0x63,0x68,0x70,0x61,0x64,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x08,0xa5,0x2e,0x73,0x69,0x7a,0x65
,0x08,0xaa,0x2e,0x74,0x79,0x70,0x65,0x5f,0x6e,0x61,0x6d,0x65,0xa6,0x75,0x63,0x68,0x61,0x72,0x2a,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xad
,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x62,0x75,0x66,0x66,0x65,0x72,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa2,0x75,0x38,0x87,0xae,0x2e,0x61
,0x64,0x64,0x72,0x65,0x73,0x73,0x5f,0x73,0x70,0x61,0x63,0x65,0xa6,0x67,0x6c,0x6f,0x62,0x61,0x6c,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xa9,0x72,0x65,0x67,0x69,0x73,0x74
,0x65,0x72,0x73,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x10,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xaa,0x2e,0x74,0x79,0x70,0x65,0x5f,0x6e,0x61,0x6d,0x65,0xa6,0x75
,0x6c,0x6f,0x6e,0x67,0x2a,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xad,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x62,0x75,0x66,0x66,0x65,0x72,0xab
,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa3,0x75,0x36,0x34,0x87,0xae,0x2e,0x61,0x64,0x64,0x72,0x65,0x73,0x73,0x5f,0x73,0x70,0x61,0x63,0x65,0xa6
,0x67,0x6c,0x6f,0x62,0x61,0x6c,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xae,0x72,0x6f,0x75,0x6e,0x64,0x69,0x6e,0x67,0x5f,0x6d,0x6f,0x64,0x65,0x73,0xa7,0x2e,0x6f,0x66,0x66
,0x73,0x65,0x74,0x18,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xaa,0x2e,0x74,0x79,0x70,0x65,0x5f,0x6e,0x61,0x6d,0x65,0xa5,0x75,0x69,0x6e,0x74,0x2a,0xab,0x2e,0x76,0x61
,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xad,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x62,0x75,0x66,0x66,0x65,0x72,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79
,0x70,0x65,0xa3,0x75,0x33,0x32,0x87,0xae,0x2e,0x61,0x64,0x64,0x72,0x65,0x73,0x73,0x5f,0x73,0x70,0x61,0x63,0x65,0xa6,0x67,0x6c,0x6f,0x62,0x61,0x6c,0xa5,0x2e,0x6e
,0x61,0x6d,0x65,0xa8,0x70,0x72,0x6f,0x67,0x72,0x61,0x6d,0x73,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x20,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xaa,0x2e,0x74,0x79
,0x70,0x65,0x5f,0x6e,0x61,0x6d,0x65,0xa5,0x75,0x69,0x6e,0x74,0x2a,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xad,0x67,0x6c,0x6f,0x62,0x61,0x6c
,0x5f,0x62,0x75,0x66,0x66,0x65,0x72,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa3,0x75,0x33,0x32,0x86,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xaa,0x62
,0x61,0x74,0x63,0x68,0x5f,0x73,0x69,0x7a,0x65,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x28,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x04,0xaa,0x2e,0x74,0x79,0x70,0x65,0x5f
,0x6e,0x61,0x6d,0x65,0xa4,0x75,0x69,0x6e,0x74,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xa8,0x62,0x79,0x5f,0x76,0x61,0x6c,0x75,0x65,0xab,0x2e
,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa3,0x75,0x33,0x32,0x86,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xad,0x72,0x78,0x5f,0x70,0x61,0x72,0x61,0x6d,0x65,0x74
,0x65,0x72,0x73,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x2c,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x04,0xaa,0x2e,0x74,0x79,0x70,0x65,0x5f,0x6e,0x61,0x6d,0x65,0xa4,0x75
,0x69,0x6e,0x74,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xa8,0x62,0x79,0x5f,0x76,0x61,0x6c,0x75,0x65,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f
,0x74,0x79,0x70,0x65,0xa3,0x75,0x33,0x32,0x84,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x30,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65
,0x5f,0x6b,0x69,0x6e,0x64,0xb6,0x68,0x69,0x64,0x64,0x65,0x6e,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x5f,0x78,0xab,0x2e,0x76,0x61
,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa3,0x69,0x36,0x34,0x84,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x38,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xab,0x2e,0x76
,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xb6,0x68,0x69,0x64,0x64,0x65,0x6e,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x5f,0x79
,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa3,0x69,0x36,0x34,0x84,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x40,0xa5,0x2e,0x73,0x69,0x7a,0x65
,0x08,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xb6,0x68,0x69,0x64,0x64,0x65,0x6e,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x6f,0x66,0x66,0x73
,0x65,0x74,0x5f,0x7a,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa3,0x69,0x36,0x34,0x84,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x48,0xa5,0x2e
,0x73,0x69,0x7a,0x65,0x08,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xab,0x68,0x69,0x64,0x64,0x65,0x6e,0x5f,0x6e,0x6f,0x6e,0x65,0xab,0x2e,0x76
,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa2,0x69,0x38,0x84,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x50,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xab,0x2e,0x76
,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xab,0x68,0x69,0x64,0x64,0x65,0x6e,0x5f,0x6e,0x6f,0x6e,0x65,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70
,0x65,0xa2,0x69,0x38,0x84,0xa7,0x2e,0x6f,0x66,0x66,0x73,0x65,0x74,0x58,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e
,0x64,0xab,0x68,0x69,0x64,0x64,0x65,0x6e,0x5f,0x6e,0x6f,0x6e,0x65,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa2,0x69,0x38,0x84,0xa7,0x2e,0x6f
,0x66,0x66,0x73,0x65,0x74,0x60,0xa5,0x2e,0x73,0x69,0x7a,0x65,0x08,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x6b,0x69,0x6e,0x64,0xb9,0x68,0x69,0x64,0x64,0x65,0x6e
,0x5f,0x6d,0x75,0x6c,0x74,0x69,0x67,0x72,0x69,0x64,0x5f,0x73,0x79,0x6e,0x63,0x5f,0x61,0x72,0x67,0xab,0x2e,0x76,0x61,0x6c,0x75,0x65,0x5f,0x74,0x79,0x70,0x65,0xa2
,0x69,0x38,0xb9,0x2e,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x65,0x67,0x6d,0x65,0x6e,0x74,0x5f,0x66,0x69,0x78,0x65,0x64,0x5f,0x73,0x69,0x7a,0x65,0xcd,0x01,0x00,0xb6
,0x2e,0x6b,0x65,0x72,0x6e,0x61,0x72,0x67,0x5f,0x73,0x65,0x67,0x6d,0x65,0x6e,0x74,0x5f,0x61,0x6c,0x69,0x67,0x6e,0x08,0xb5,0x2e,0x6b,0x65,0x72,0x6e,0x61,0x72,0x67
,0x5f,0x73,0x65,0x67,0x6d,0x65,0x6e,0x74,0x5f,0x73,0x69,0x7a,0x65,0x68,0xa9,0x2e,0x6c,0x61,0x6e,0x67,0x75,0x61,0x67,0x65,0xa8,0x4f,0x70,0x65,0x6e,0x43,0x4c,0x20
,0x43,0xb1,0x2e,0x6c,0x61,0x6e,0x67,0x75,0x61,0x67,0x65,0x5f,0x76,0x65,0x72,0x73,0x69,0x6f,0x6e,0x92,0x01,0x02,0xb8,0x2e,0x6d,0x61,0x78,0x5f,0x66,0x6c,0x61,0x74
,0x5f,0x77,0x6f,0x72,0x6b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x20,0xa5,0x2e,0x6e,0x61,0x6d,0x65,0xab,0x72,0x61,0x6e,0x64,0x6f,0x6d,0x78,0x5f,0x72
,0x75,0x6e,0xbb,0x2e,0x70,0x72,0x69,0x76,0x61,0x74,0x65,0x5f,0x73,0x65,0x67,0x6d,0x65,0x6e,0x74,0x5f,0x66,0x69,0x78,0x65,0x64,0x5f,0x73,0x69,0x7a,0x65,0x00,0xb4
,0x2e,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x93,0x20,0x01,0x01,0xab,0x2e,0x73,0x67,0x70,0x72,0x5f,0x63
,0x6f,0x75,0x6e,0x74,0x60,0xb1,0x2e,0x73,0x67,0x70,0x72,0x5f,0x73,0x70,0x69,0x6c,0x6c,0x5f,0x63,0x6f,0x75,0x6e,0x74,0x00,0xa7,0x2e,0x73,0x79,0x6d,0x62,0x6f,0x6c
,0xae,0x72,0x61,0x6e,0x64,0x6f,0x6d,0x78,0x5f,0x72,0x75,0x6e,0x2e,0x6b,0x64,0xab,0x2e,0x76,0x67,0x70,0x72,0x5f,0x63,0x6f,0x75,0x6e,0x74,0xcc,0x80,0xb1,0x2e,0x76
,0x67,0x70,0x72,0x5f,0x73,0x70,0x69,0x6c,0x6c,0x5f,0x63,0x6f,0x75,0x6e,0x74,0x00,0xaf,0x2e,0x77,0x61,0x76,0x65,0x66,0x72,0x6f,0x6e,0x74,0x5f,0x73,0x69,0x7a,0x65
,0x20,0xae,0x61,0x6d,0x64,0x68,0x73,0x61,0x2e,0x76,0x65,0x72,0x73,0x69,0x6f,0x6e,0x92,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x12,0x03,0x06,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x0d,0x00,0x00,0x00,0x11,0x03,0x05,0x00,0xc0,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01
,0x00,0x00,0x00,0x03,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x72,0x61,0x6e,0x64,0x6f,0x6d,0x78,0x5f
,0x72,0x75,0x6e,0x00,0x72,0x61,0x6e,0x64,0x6f,0x6d,0x78,0x5f,0x72,0x75,0x6e,0x2e,0x6b,0x64,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x40,0x07,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0xcf,0x03,0xaf,0x40,0x8c,0x00,0x00,0x00,0x09,0x04,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x84,0xf4,0x00,0x00,0x00,0x00,0x70,0x00,0x8c,0xbf,0x00,0x00,0xfd,0xbb,0x00,0x00,0x93,0xbf,0x00,0x00,0x82,0xbf,0x02,0x00,0x08,0xf4,0x10,0x00,0x00,0xfa,0x80
,0x03,0x89,0xbe,0x06,0x85,0x08,0x8f,0x83,0x00,0x4e,0x34,0x07,0x03,0x8c,0xbe,0x88,0x00,0x88,0x7d,0x7f,0xc0,0x8c,0xbf,0x06,0x82,0x10,0x8f,0x02,0x10,0x40,0x80,0x03
,0x80,0x41,0x82,0x80,0x02,0x02,0x7e,0x00,0x80,0x30,0xdc,0x01,0x00,0x40,0x01,0x70,0x3f,0x8c,0xbf,0x42,0x00,0x60,0xd7,0x01,0x01,0x01,0x00,0x81,0x08,0xc2,0xb9,0x80
,0x03,0xc3,0xbe,0x08,0x83,0x82,0x8f,0x0c,0x03,0xa0,0xbe,0x00,0x02,0x00,0x80,0x01,0x03,0x01,0x82,0x01,0x00,0x0f,0xd7,0x00,0x4e,0x02,0x00,0x02,0x00,0x28,0xd5,0x01
,0x00,0x01,0x00,0x00,0x80,0x34,0xdc,0x01,0x00,0x7d,0x04,0x70,0x3f,0x8c,0xbf,0x00,0x00,0x34,0xd9,0x27,0x04,0x00,0x00,0x70,0x00,0x8c,0xbf,0x00,0x00,0xfd,0xbb,0x6a
,0x3c,0x80,0xbe,0x65,0x01,0x88,0xbf,0x02,0x02,0x08,0xf4,0x00,0x00,0x00,0xfa,0x02,0x00,0x04,0xf4,0x20,0x00,0x00,0xfa,0x02,0x05,0x00,0xf4,0x2c,0x00,0x00,0xfa,0x80
,0x02,0x0a,0x7e,0x80,0x02,0x14,0x7e,0x00,0x00,0xfd,0xbb,0x00,0x00,0xd8,0xd9,0x27,0x00,0x00,0x08,0x84,0x00,0x88,0x7d,0x83,0x00,0x00,0x34,0x10,0x11,0xdc,0xd9,0x05
,0x00,0x00,0x19,0x98,0x00,0xd8,0xd8,0x05,0x00,0x00,0x0b,0xa8,0x00,0xd8,0xd9,0x05,0x00,0x00,0x23,0x12,0x14,0xdc,0xd9,0x05,0x00,0x00,0x14,0x04,0x00,0x01,0xd5,0xff
,0x82,0xa9,0x01,0xff,0xff,0xff,0x00,0x27,0x01,0x0a,0x4a,0x7f,0xc0,0x8c,0xbf,0x0b,0x02,0x1a,0x7e,0x01,0x02,0x0e,0x7e,0x00,0x02,0x0c,0x7e,0x14,0xff,0x95,0x93,0x00
,0x00,0x05,0x00,0x81,0x15,0x15,0x8f,0x14,0xff,0x96,0x93,0x05,0x00,0x05,0x00,0x81,0x16,0x16,0x8f,0x14,0xff,0x80,0x93,0x0a,0x00,0x05,0x00,0x81,0x00,0x17,0x8f,0x14
,0xff,0x98,0x93,0x0f,0x00,0x04,0x00,0x81,0x18,0x18,0x8f,0x0a,0x02,0x18,0x7e,0x06,0x02,0x76,0xd5,0xff,0x0c,0x18,0x04,0x40,0x27,0x00,0x00,0x04,0x00,0x60,0xd7,0x06
,0x01,0x01,0x00,0x05,0x00,0x60,0xd7,0x07,0x01,0x01,0x00,0x81,0x00,0x02,0x8f,0x0e,0x00,0x0f,0xd7,0x08,0x16,0x02,0x00,0x22,0x00,0x01,0xd5,0x24,0x01,0xa9,0x01,0x18
,0x00,0x01,0xd5,0x17,0x01,0xa9,0x01,0x03,0x00,0x01,0xd5,0x16,0x01,0xa9,0x01,0x02,0xc0,0x03,0x81,0x1d,0x00,0x28,0xd5,0x09,0x14,0x02,0x00,0x23,0x00,0x01,0xd5,0x23
,0x01,0xa9,0x01,0x16,0x6a,0x0f,0xd7,0x0e,0x01,0x02,0x00,0x0c,0x02,0x76,0xd5,0x03,0x0c,0x30,0x04,0x1a,0x03,0x14,0x7e,0x19,0x03,0x16,0x7e,0x83,0x36,0x48,0x34,0x83
,0x38,0x4a,0x34,0x83,0x28,0x28,0x34,0x83,0x2a,0x2a,0x34,0x80,0x3a,0x2e,0x50,0x27,0x03,0x0c,0x7e,0x18,0x81,0x82,0x80,0xc1,0x03,0xbf,0xbe,0xff,0x03,0xc6,0xbe,0x00
,0xff,0x00,0x00,0xff,0x03,0xc7,0xbe,0x00,0xfe,0x01,0x00,0xff,0x03,0xc8,0xbe,0x00,0xfc,0x03,0x00,0xff,0x03,0xc9,0xbe,0x00,0xf8,0x07,0x00,0xff,0x03,0xca,0xbe,0x00
,0xf0,0x0f,0x00,0xff,0x03,0xcb,0xbe,0x00,0xe0,0x1f,0x00,0xff,0x03,0xcc,0xbe,0x00,0xc0,0x3f,0x00,0xff,0x03,0xcd,0xbe,0x00,0x80,0x7f,0x00,0xff,0x03,0xce,0xbe,0x00
,0x00,0xff,0x00,0xff,0x03,0xcf,0xbe,0x00,0x00,0xfe,0x01,0xff,0x03,0xd0,0xbe,0x00,0x00,0xfc,0x03,0xff,0x03,0xd1,0xbe,0x00,0x00,0xf8,0x07,0xff,0x03,0xd2,0xbe,0x00
,0x00,0xf0,0x0f,0xff,0x03,0xd3,0xbe,0x00,0x00,0xe0,0x1f,0xff,0x03,0xd4,0xbe,0x00,0x00,0xc0,0x3f,0xff,0x03,0xd5,0xbe,0x00,0x00,0x80,0x7f,0x17,0xc0,0xd6,0x80,0x26
,0x00,0x26,0xd5,0x15,0x10,0x01,0x00,0x27,0x00,0x26,0xd5,0x16,0x10,0x01,0x00,0x32,0x00,0x26,0xd5,0x17,0x10,0x01,0x00,0xff,0x02,0x66,0x7e,0x00,0x00,0xf0,0x80,0x00
,0x00,0x60,0xd7,0x0c,0x01,0x01,0x00,0x01,0x00,0x60,0xd7,0x0d,0x01,0x01,0x00,0x80,0x02,0x52,0x7e,0x80,0x02,0x58,0x7e,0x82,0x04,0xfe,0xbe,0x88,0x02,0x52,0x7e,0x84
,0x02,0x58,0x7e,0x83,0x04,0xfe,0xbe,0x18,0x1a,0xdc,0xd9,0x29,0x00,0x00,0x34,0x1c,0x1e,0xdc,0xd9,0x29,0x00,0x00,0x38,0xff,0x02,0x9a,0x7e,0xff,0xff,0xff,0x00,0xa0
,0x00,0xd8,0xd9,0x29,0x00,0x00,0x4e,0xff,0x04,0xfe,0xbe,0xff,0x00,0x00,0x00,0xff,0x02,0xa4,0x7e,0x00,0x00,0x00,0x80,0xff,0x03,0xc4,0xbe,0x00,0x01,0x00,0x00,0x80
,0x03,0xc5,0xbe,0xff,0x02,0xa6,0x7e,0x00,0x00,0xf0,0x3f,0xff,0x02,0xa8,0x7e,0x00,0x00,0x10,0x00,0x00,0x1f,0x8e,0xbe,0x0e,0xff,0x28,0x80,0x48,0x03,0x00,0x00,0x0f
,0x80,0x29,0x82,0x0e,0xff,0x2a,0x80,0xac,0x03,0x00,0x00,0x0f,0x80,0x2b,0x82,0x0e,0xff,0x2c,0x80,0x10,0x04,0x00,0x00,0x0f,0x80,0x2d,0x82,0x0e,0xff,0x2e,0x80,0x74
,0x04,0x00,0x00,0x0f,0x80,0x2f,0x82,0x0e,0xff,0x30,0x80,0xd8,0x04,0x00,0x00,0x0f,0x80,0x31,0x82,0x0e,0xff,0x32,0x80,0x3c,0x05,0x00,0x00,0x0f,0x80,0x33,0x82,0x0e
,0xff,0x34,0x80,0xa0,0x05,0x00,0x00,0x0f,0x80,0x35,0x82,0x0e,0xff,0x36,0x80,0x04,0x06,0x00,0x00,0x0f,0x80,0x37,0x82,0x0e,0xff,0x38,0x80,0x68,0x06,0x00,0x00,0x0f
,0x80,0x39,0x82,0x0e,0xff,0x3a,0x80,0xdc,0x06,0x00,0x00,0x0f,0x80,0x3b,0x82,0x00,0x00,0xfd,0xbb,0x00,0x00,0xd8,0xd9,0x25,0x00,0x00,0x1b,0x00,0x00,0xd8,0xd9,0x24
,0x00,0x00,0x1d,0x7f,0xc0,0x8c,0xbf,0x1c,0x3d,0x38,0x3a,0x1c,0x33,0x32,0x3a,0x56,0x32,0x32,0x36,0x19,0x01,0x32,0x4a,0x10,0x6a,0x0f,0xd7,0x00,0x32,0x02,0x00,0x1b
,0x3b,0x32,0x3a,0x0b,0x03,0x3a,0x7e,0x11,0x6a,0x28,0xd5,0x80,0x02,0xa8,0x01,0x19,0x35,0x32,0x3a,0x00,0x80,0x34,0xdc,0x10,0x00,0x7d,0x1b,0x56,0x32,0x32,0x36,0x19
,0x01,0x32,0x4a,0x1f,0x6a,0x0f,0xd7,0x00,0x32,0x02,0x00,0x20,0x6a,0x28,0xd5,0x80,0x02,0xa8,0x01,0x1d,0x6a,0x0f,0xd7,0x16,0x3b,0x02,0x00,0x00,0x80,0x34,0xdc,0x1f
,0x00,0x7d,0x19,0x80,0x2e,0x3c,0x50,0x0b,0x03,0x42,0x7e,0x7e,0x6a,0x6a,0x87,0x71,0x3f,0x8c,0xbf,0x1c,0x09,0x1c,0x7e,0x1b,0x09,0x18,0x7e,0x0e,0x47,0x1c,0x38,0x70
,0x3f,0x8c,0xbf,0x19,0x11,0x10,0x3a,0x1a,0x13,0x12,0x3a,0x04,0x1f,0x34,0x36,0x04,0x1b,0x26,0x36,0x1a,0x45,0x1e,0x38,0x0c,0x07,0x24,0x38,0x80,0x02,0x34,0x7e,0x13
,0x31,0x26,0x38,0x1a,0x03,0x32,0x7e,0x08,0x09,0x38,0xd9,0x05,0x12,0x0e,0x00,0x00,0x80,0x34,0xdc,0x1d,0x00,0x7d,0x12,0x83,0x04,0xfe,0xbe,0x7f,0xc0,0x8c,0xbf,0x08
,0x0a,0xdc,0xd9,0x29,0x00,0x00,0x3c,0x0c,0x0e,0xdc,0xd9,0x29,0x00,0x00,0x40,0x10,0x12,0xdc,0xd9,0x29,0x00,0x00,0x44,0x14,0x16,0xdc,0xd9,0x29,0x00,0x00,0x48,0x10
,0x00,0x60,0xd7,0x08,0x01,0x01,0x00,0x11,0x00,0x60,0xd7,0x09,0x01,0x01,0x00,0x12,0x00,0x60,0xd7,0x08,0x03,0x01,0x00,0x13,0x00,0x60,0xd7,0x09,0x03,0x01,0x00,0x14
,0x00,0x60,0xd7,0x08,0x05,0x01,0x00,0x15,0x00,0x60,0xd7,0x09,0x05,0x01,0x00,0x16,0x00,0x60,0xd7,0x08,0x07,0x01,0x00,0x17,0x00,0x60,0xd7,0x09,0x07,0x01,0x00,0x18
,0x00,0x60,0xd7,0x08,0x09,0x01,0x00,0x19,0x00,0x60,0xd7,0x09,0x09,0x01,0x00,0x1a,0x00,0x60,0xd7,0x08,0x0b,0x01,0x00,0x1b,0x00,0x60,0xd7,0x09,0x0b,0x01,0x00,0x1c
,0x00,0x60,0xd7,0x08,0x0d,0x01,0x00,0x1d,0x00,0x60,0xd7,0x09,0x0d,0x01,0x00,0x1e,0x00,0x60,0xd7,0x08,0x0f,0x01,0x00,0x1f,0x00,0x60,0xd7,0x09,0x0f,0x01,0x00,0x7f
,0xc0,0x8c,0xbf,0x83,0x04,0xfe,0xbe,0x04,0x21,0x8c,0xbe,0x08,0x0a,0x38,0xd9,0x29,0x3c,0x3e,0x00,0x0c,0x0e,0x38,0xd9,0x29,0x40,0x42,0x00,0x10,0x12,0x38,0xd9,0x29
,0x44,0x46,0x00,0x14,0x16,0x38,0xd9,0x29,0x48,0x4a,0x00,0x08,0x00,0x61,0xd7,0x10,0x00,0x01,0x00,0x09,0x00,0x61,0xd7,0x11,0x00,0x01,0x00,0x08,0x00,0x61,0xd7,0x12
,0x02,0x01,0x00,0x09,0x00,0x61,0xd7,0x13,0x02,0x01,0x00,0x08,0x00,0x61,0xd7,0x14,0x04,0x01,0x00,0x09,0x00,0x61,0xd7,0x15,0x04,0x01,0x00,0x08,0x00,0x61,0xd7,0x16
,0x06,0x01,0x00,0x09,0x00,0x61,0xd7,0x17,0x06,0x01,0x00,0x08,0x00,0x61,0xd7,0x18,0x08,0x01,0x00,0x09,0x00,0x61,0xd7,0x19,0x08,0x01,0x00,0x08,0x00,0x61,0xd7,0x1a
,0x0a,0x01,0x00,0x09,0x00,0x61,0xd7,0x1b,0x0a,0x01,0x00,0x08,0x00,0x61,0xd7,0x1c,0x0c,0x01,0x00,0x09,0x00,0x61,0xd7,0x1d,0x0c,0x01,0x00,0x08,0x00,0x61,0xd7,0x1e
,0x0e,0x01,0x00,0x09,0x00,0x61,0xd7,0x1f,0x0e,0x01,0x00,0xff,0x04,0xfe,0xbe,0xff,0x00,0x00,0x00,0x00,0x00,0x34,0xd9,0x06,0x08,0x00,0x00,0x7f,0xc0,0x8c,0xbf,0x00
,0x00,0xd8,0xd8,0x15,0x00,0x00,0x0b,0x00,0x00,0xd8,0xd8,0x14,0x00,0x00,0x1b,0x7f,0xc0,0x8c,0xbf,0x0b,0x37,0x16,0x3a,0x0a,0x17,0x14,0x3a,0x00,0x08,0xdc,0xd9,0x06
,0x00,0x00,0x1b,0xff,0x14,0x16,0x36,0xc0,0xff,0xff,0x7f,0x21,0x03,0x14,0x7e,0x7f,0xc0,0x8c,0xbf,0x7f,0xc0,0x8c,0xbf,0x1b,0x25,0x10,0x3a,0x1c,0x27,0x12,0x3a,0x00
,0x80,0x74,0xdc,0x10,0x08,0x7d,0x00,0x80,0x00,0xd8,0xd9,0x06,0x00,0x00,0x1b,0x00,0x00,0x34,0xd9,0x06,0x08,0x00,0x00,0x7f,0xc1,0x8c,0xbf,0x1b,0x3b,0x3a,0x3a,0x1c
,0x3d,0x3c,0x3a,0x00,0x80,0x74,0xdc,0x1f,0x1d,0x7d,0x00,0x02,0x81,0x82,0x80,0x57,0xff,0x84,0xbf,0x00,0x80,0x74,0xdc,0x01,0x08,0x7d,0x00,0x40,0x80,0x74,0xdc,0x01
,0x1d,0x7d,0x00,0x80,0x80,0x74,0xdc,0x01,0x1b,0x7d,0x00,0x80,0x02,0x00,0x7e,0x42,0x02,0x02,0x7e,0x00,0x80,0x70,0xdc,0x00,0x01,0x40,0x00,0x00,0x00,0x81,0xbf,0x81
,0x08,0xc3,0xb9,0x44,0x63,0x38,0x7e,0x2a,0x00,0x65,0xd5,0x1c,0x89,0x02,0x00,0x1c,0x03,0x60,0x7e,0x1d,0xa9,0x62,0x4c,0x1c,0x03,0x5c,0x7e,0x31,0xa5,0x5e,0x3a,0x2e
,0x00,0x4c,0xd5,0x2e,0x55,0xc2,0x03,0x2a,0x00,0x4c,0xd5,0x2a,0x5d,0xaa,0x04,0x30,0x00,0x4c,0xd5,0x30,0x5d,0xc2,0x04,0x2e,0x00,0x4c,0xd5,0x2a,0x55,0x12,0x25,0x81
,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2e,0x61,0xaa,0x04,0x00,0x00,0xb8,0xd4,0x44,0x89,0x00,0x00,0x2a,0x03,0x88,0x7e,0x2b,0x03,0x8a,0x7e,0x83,0x04,0xfe,0xbe,0x3c
,0x20,0x80,0xbe,0x81,0x08,0xc3,0xb9,0x46,0x63,0x38,0x7e,0x2a,0x00,0x65,0xd5,0x1c,0x8d,0x02,0x00,0x1c,0x03,0x60,0x7e,0x1d,0xa9,0x62,0x4c,0x1c,0x03,0x5c,0x7e,0x31
,0xa5,0x5e,0x3a,0x2e,0x00,0x4c,0xd5,0x2e,0x55,0xc2,0x03,0x2a,0x00,0x4c,0xd5,0x2a,0x5d,0xaa,0x04,0x30,0x00,0x4c,0xd5,0x30,0x5d,0xc2,0x04,0x2e,0x00,0x4c,0xd5,0x2a
,0x55,0x1a,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2e,0x61,0xaa,0x04,0x00,0x00,0xb8,0xd4,0x46,0x89,0x00,0x00,0x2a,0x03,0x8c,0x7e,0x2b,0x03,0x8e,0x7e,0x83
,0x04,0xfe,0xbe,0x3c,0x20,0x80,0xbe,0x81,0x08,0xc3,0xb9,0x48,0x63,0x38,0x7e,0x2a,0x00,0x65,0xd5,0x1c,0x91,0x02,0x00,0x1c,0x03,0x60,0x7e,0x1d,0xa9,0x62,0x4c,0x1c
,0x03,0x5c,0x7e,0x31,0xa5,0x5e,0x3a,0x2e,0x00,0x4c,0xd5,0x2e,0x55,0xc2,0x03,0x2a,0x00,0x4c,0xd5,0x2a,0x5d,0xaa,0x04,0x30,0x00,0x4c,0xd5,0x30,0x5d,0xc2,0x04,0x2e
,0x00,0x4c,0xd5,0x2a,0x55,0x22,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2e,0x61,0xaa,0x04,0x00,0x00,0xb8,0xd4,0x48,0x89,0x00,0x00,0x2a,0x03,0x90,0x7e,0x2b
,0x03,0x92,0x7e,0x83,0x04,0xfe,0xbe,0x3c,0x20,0x80,0xbe,0x81,0x08,0xc3,0xb9,0x4a,0x63,0x38,0x7e,0x2a,0x00,0x65,0xd5,0x1c,0x95,0x02,0x00,0x1c,0x03,0x60,0x7e,0x1d
,0xa9,0x62,0x4c,0x1c,0x03,0x5c,0x7e,0x31,0xa5,0x5e,0x3a,0x2e,0x00,0x4c,0xd5,0x2e,0x55,0xc2,0x03,0x2a,0x00,0x4c,0xd5,0x2a,0x5d,0xaa,0x04,0x30,0x00,0x4c,0xd5,0x30
,0x5d,0xc2,0x04,0x2e,0x00,0x4c,0xd5,0x2a,0x55,0x2a,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2e,0x61,0xaa,0x04,0x00,0x00,0xb8,0xd4,0x4a,0x89,0x00,0x00,0x2a
,0x03,0x94,0x7e,0x2b,0x03,0x96,0x7e,0x83,0x04,0xfe,0xbe,0x3c,0x20,0x80,0xbe,0x1c,0x9d,0x38,0x38,0x1d,0x00,0x71,0xd7,0x1d,0x9b,0x3e,0x05,0x81,0x08,0xc3,0xb9,0x1c
,0x5f,0x60,0x7e,0x50,0x00,0x4c,0xd5,0x1c,0x61,0xca,0x23,0x30,0x00,0x4c,0xd5,0x30,0xa1,0xc2,0x04,0x50,0x00,0x65,0xd5,0x44,0x61,0x02,0x00,0x2a,0x00,0x4c,0xd5,0x1c
,0xa1,0x12,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2a,0x61,0x42,0x05,0x50,0x00,0x60,0xd5,0x2a,0x39,0x12,0x05,0x44,0x39,0x64,0x7c,0x80,0x02,0xa0,0x7e,0x53
,0x03,0xa2,0x7e,0x83,0x04,0xfe,0xbe,0x50,0x03,0x88,0x7e,0x51,0x03,0x8a,0x7e,0x3c,0x20,0x80,0xbe,0x1c,0x9d,0x38,0x38,0x1d,0x00,0x71,0xd7,0x1d,0x9b,0x3e,0x05,0x81
,0x08,0xc3,0xb9,0x1c,0x5f,0x60,0x7e,0x50,0x00,0x4c,0xd5,0x1c,0x61,0xca,0x23,0x30,0x00,0x4c,0xd5,0x30,0xa1,0xc2,0x04,0x50,0x00,0x65,0xd5,0x46,0x61,0x02,0x00,0x2a
,0x00,0x4c,0xd5,0x1c,0xa1,0x1a,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2a,0x61,0x42,0x05,0x50,0x00,0x60,0xd5,0x2a,0x39,0x1a,0x05,0x46,0x39,0x64,0x7c,0x80
,0x02,0xa0,0x7e,0x53,0x03,0xa2,0x7e,0x83,0x04,0xfe,0xbe,0x50,0x03,0x8c,0x7e,0x51,0x03,0x8e,0x7e,0x3c,0x20,0x80,0xbe,0x1c,0x9d,0x38,0x38,0x1d,0x00,0x71,0xd7,0x1d
,0x9b,0x3e,0x05,0x81,0x08,0xc3,0xb9,0x1c,0x5f,0x60,0x7e,0x50,0x00,0x4c,0xd5,0x1c,0x61,0xca,0x23,0x30,0x00,0x4c,0xd5,0x30,0xa1,0xc2,0x04,0x50,0x00,0x65,0xd5,0x48
,0x61,0x02,0x00,0x2a,0x00,0x4c,0xd5,0x1c,0xa1,0x22,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2a,0x61,0x42,0x05,0x50,0x00,0x60,0xd5,0x2a,0x39,0x22,0x05,0x48
,0x39,0x64,0x7c,0x80,0x02,0xa0,0x7e,0x53,0x03,0xa2,0x7e,0x83,0x04,0xfe,0xbe,0x50,0x03,0x90,0x7e,0x51,0x03,0x92,0x7e,0x3c,0x20,0x80,0xbe,0x1c,0x9d,0x38,0x38,0x1d
,0x00,0x71,0xd7,0x1d,0x9b,0x3e,0x05,0x81,0x08,0xc3,0xb9,0x1c,0x5f,0x60,0x7e,0x50,0x00,0x4c,0xd5,0x1c,0x61,0xca,0x23,0x30,0x00,0x4c,0xd5,0x30,0xa1,0xc2,0x04,0x50
,0x00,0x65,0xd5,0x4a,0x61,0x02,0x00,0x2a,0x00,0x4c,0xd5,0x1c,0xa1,0x2a,0x25,0x81,0x08,0xc2,0xb9,0x2a,0x00,0x4c,0xd5,0x2a,0x61,0x42,0x05,0x50,0x00,0x60,0xd5,0x2a
,0x39,0x2a,0x05,0x4a,0x39,0x64,0x7c,0x80,0x02,0xa0,0x7e,0x53,0x03,0xa2,0x7e,0x83,0x04,0xfe,0xbe,0x50,0x03,0x94,0x7e,0x51,0x03,0x96,0x7e,0x3c,0x20,0x80,0xbe,0x81
,0x04,0xfe,0xbe,0x0e,0x02,0x5a,0x7e,0x28,0x00,0x6a,0xd5,0x26,0x5a,0x02,0x00,0x0f,0x02,0x5e,0x7e,0x2a,0x20,0x76,0xd5,0x26,0x5e,0xa2,0x04,0x2a,0x03,0x50,0x7e,0x2d
,0x20,0x76,0xd5,0x27,0x5a,0xa2,0x04,0x2a,0x20,0x76,0xd5,0x27,0x5e,0xae,0x04,0x2a,0x6a,0x0f,0xd7,0x2a,0x5d,0x02,0x00,0x80,0x56,0x56,0x50,0x20,0x00,0x60,0xd7,0x2a
,0x01,0x01,0x00,0x21,0x00,0x60,0xd7,0x2b,0x01,0x01,0x00,0x0f,0x80,0x04,0xbf,0x26,0x80,0xa2,0x85,0x20,0x22,0xa0,0x80,0x21,0x23,0xa1,0x82,0x27,0x80,0x04,0xbf,0x0e
,0x80,0xa2,0x85,0x20,0x22,0x8e,0x80,0x21,0x23,0x8f,0x82,0x83,0x04,0xfe,0xbe,0x3c,0x20,0x80,0xbe,0x81,0x04,0xfe,0xbe,0x26,0x02,0x5a,0x7e,0x28,0x00,0x6a,0xd5,0x0e
,0x5a,0x02,0x00,0x27,0x02,0x5e,0x7e,0x2a,0x20,0x76,0xd5,0x0e,0x5e,0xa2,0x04,0x2a,0x03,0x50,0x7e,0x2d,0x20,0x76,0xd5,0x0f,0x5a,0xa2,0x04,0x2a,0x20,0x76,0xd5,0x0f
,0x5e,0xae,0x04,0x2a,0x6a,0x0f,0xd7,0x2a,0x5d,0x02,0x00,0x80,0x56,0x56,0x50,0x0e,0x00,0x60,0xd7,0x2a,0x01,0x01,0x00,0x0f,0x00,0x60,0xd7,0x2b,0x01,0x01,0x00,0x83
,0x04,0xfe,0xbe,0x3c,0x20,0x80,0xbe,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x06
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x38,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x0b,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x18,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x05
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x98,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x0a,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x1c,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x04
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x80,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x43
,0x4c,0x52,0x58,0x20,0x52,0x4f,0x43,0x6d,0x42,0x69,0x6e,0x47,0x65,0x6e,0x65,0x72,0x61,0x74,0x6f,0x72,0x20,0x30,0x2e,0x31,0x2e,0x39,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x02,0x07,0x00,0x00
,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x0a,0x00,0x00,0x00,0x12,0x03,0x06,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x16,0x00,0x00,0x00,0x11,0x03,0x05,0x00,0xc0,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x2e,0x6e,0x6f,0x74,0x65,0x00,0x2e,0x64,0x79,0x6e,0x73,0x79,0x6d,0x00,0x2e,0x68,0x61,0x73,0x68,0x00,0x2e,0x64,0x79,0x6e,0x73,0x74,0x72,0x00,0x2e,0x72,0x6f,0x64
,0x61,0x74,0x61,0x00,0x2e,0x74,0x65,0x78,0x74,0x00,0x2e,0x64,0x79,0x6e,0x61,0x6d,0x69,0x63,0x00,0x2e,0x63,0x6f,0x6d,0x6d,0x65,0x6e,0x74,0x00,0x2e,0x73,0x79,0x6d
,0x74,0x61,0x62,0x00,0x2e,0x73,0x68,0x73,0x74,0x72,0x74,0x61,0x62,0x00,0x2e,0x73,0x74,0x72,0x74,0x61,0x62,0x00,0x00,0x5f,0x44,0x59,0x4e,0x41,0x4d,0x49,0x43,0x00
,0x72,0x61,0x6e,0x64,0x6f,0x6d,0x78,0x5f,0x72,0x75,0x6e,0x00,0x72,0x61,0x6e,0x64,0x6f,0x6d,0x78,0x5f,0x72,0x75,0x6e,0x2e,0x6b,0x64,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01
,0x00,0x00,0x00,0x07,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x34
,0x06,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x07
,0x00,0x00,0x00,0x0b,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x38,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x38,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x48
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x18,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x0f
,0x00,0x00,0x00,0x05,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x80,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x80,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x18
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x15
,0x00,0x00,0x00,0x03,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x98,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x98,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x1c
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x1d
,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0xc0,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0xc0,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x40
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x40,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x25
,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x06,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x28
,0x0a,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x2b
,0x00,0x00,0x00,0x06,0x00,0x00,0x00,0x03,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x60
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x04,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x10,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x34
,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x30,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x60,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x1b
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x3d
,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x80,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x60
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x0b,0x00,0x00,0x00,0x02,0x00,0x00,0x00,0x08,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x18,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x45
,0x00,0x00,0x00,0x03,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0xe0,0x20,0x00,0x00,0x00,0x00,0x00,0x00,0x57
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x4f
,0x00,0x00,0x00,0x03,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x37,0x21,0x00,0x00,0x00,0x00,0x00,0x00,0x25
,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x01,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00};
const int randomx_run_gfx1010_bin_size=9312;

View file

@ -63,6 +63,7 @@ bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorith
case OclDevice::Navi_10:
case OclDevice::Navi_12:
case OclDevice::Navi_14:
gcnAsm = true;
isNavi = true;
break;

View file

@ -30,12 +30,10 @@
#include "crypto/rx/RxAlgo.h"
void xmrig::RxRunKernel::enqueue(cl_command_queue queue, size_t threads)
void xmrig::RxRunKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
{
const size_t gthreads = threads * 64;
static const size_t lthreads = 64;
enqueueNDRange(queue, 1, nullptr, &gthreads, &lthreads);
const size_t gthreads = threads * workgroup_size;
enqueueNDRange(queue, 1, nullptr, &gthreads, &workgroup_size);
}

View file

@ -40,7 +40,7 @@ class RxRunKernel : public OclKernel
public:
inline RxRunKernel(cl_program program) : OclKernel(program, "randomx_run") {}
void enqueue(cl_command_queue queue, size_t threads);
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
void setArgs(cl_mem dataset, cl_mem scratchpads, cl_mem registers, cl_mem rounding, cl_mem programs, uint32_t batch_size, const Algorithm &algorithm);
};

View file

@ -56,6 +56,10 @@ xmrig::OclRxBaseRunner::OclRxBaseRunner(size_t index, const OclLaunchData &data)
m_gcn_version = 14;
}
if (data.device.type() == OclDevice::Navi_10 || data.device.type() == OclDevice::Navi_12 || data.device.type() == OclDevice::Navi_14) {
m_gcn_version = 15;
}
m_options += " -DALGO=" + std::to_string(RxAlgo::id(m_algorithm));
m_options += " -DWORKERS_PER_HASH=" + std::to_string(m_worksize);
m_options += " -DGCN_VERSION=" + std::to_string(m_gcn_version);

View file

@ -26,6 +26,7 @@
#include "backend/opencl/cl/rx/randomx_run_gfx803.h"
#include "backend/opencl/cl/rx/randomx_run_gfx900.h"
#include "backend/opencl/cl/rx/randomx_run_gfx1010.h"
#include "backend/opencl/kernels/rx/Blake2bHashRegistersKernel.h"
#include "backend/opencl/kernels/rx/HashAesKernel.h"
#include "backend/opencl/kernels/rx/RxJitKernel.h"
@ -84,7 +85,7 @@ void xmrig::OclRxJitRunner::execute(uint32_t iteration)
OclLib::finish(m_queue);
m_randomx_run->enqueue(m_queue, m_intensity);
m_randomx_run->enqueue(m_queue, m_intensity, (m_gcn_version == 15) ? 32 : 64);
}
@ -120,8 +121,23 @@ bool xmrig::OclRxJitRunner::loadAsmProgram()
elf_header_flags = *reinterpret_cast<uint32_t*>((binary_data.data() + elf_header_flags_offset));
}
const size_t len = (m_gcn_version == 14) ? randomx_run_gfx900_bin_size : randomx_run_gfx803_bin_size;
unsigned char *binary = (m_gcn_version == 14) ? randomx_run_gfx900_bin : randomx_run_gfx803_bin;
size_t len;
unsigned char *binary;
switch (m_gcn_version) {
case 14:
len = randomx_run_gfx900_bin_size;
binary = randomx_run_gfx900_bin;
break;
case 15:
len = randomx_run_gfx1010_bin_size;
binary = randomx_run_gfx1010_bin;
break;
default:
len = randomx_run_gfx803_bin_size;
binary = randomx_run_gfx803_bin;
break;
}
// Set correct internal device ID in the pre-compiled binary
if (elf_header_flags) {