#include <cstdint>

/*
     * For 512 <= d < 1024,
     *
     *   RECIPROCAL_TABLE_32[d - 512] = floor((2^24 - 2^14 + 2^9)/d)
     *
     * Total space at the moment is 512*2 = 1024 bytes.
     *
     * TODO: Investigate whether alternative storage layouts are better; examples:
     *
     *  - redundantly store each element in a uint32_t
     *  - pack two uint16_t values into each uint32_t
     *  - is __constant__ the right storage specifier? Maybe load into shared memory?
     *    Shared memory seems like an excellent choice (48k available per SM), though
     *    I'll need to be mindful of bank conflicts (perhaps circumvent by having
     *    many copies of the data in SM?).
     *  - perhaps reading an element from memory is slower than simply calculating
     *    floor((2^24 - 2^14 + 2^9)/d) in assembly?
     */
__device__ __constant__
    uint16_t
        RECIPROCAL_TABLE_32[0x200] =
            {
                0x7fe1, 0x7fa1, 0x7f61, 0x7f22, 0x7ee3, 0x7ea4, 0x7e65, 0x7e27,
                0x7de9, 0x7dab, 0x7d6d, 0x7d30, 0x7cf3, 0x7cb6, 0x7c79, 0x7c3d,
                0x7c00, 0x7bc4, 0x7b89, 0x7b4d, 0x7b12, 0x7ad7, 0x7a9c, 0x7a61,
                0x7a27, 0x79ec, 0x79b2, 0x7979, 0x793f, 0x7906, 0x78cc, 0x7894,
                0x785b, 0x7822, 0x77ea, 0x77b2, 0x777a, 0x7742, 0x770b, 0x76d3,
                0x769c, 0x7665, 0x762f, 0x75f8, 0x75c2, 0x758c, 0x7556, 0x7520,
                0x74ea, 0x74b5, 0x7480, 0x744b, 0x7416, 0x73e2, 0x73ad, 0x7379,
                0x7345, 0x7311, 0x72dd, 0x72aa, 0x7277, 0x7243, 0x7210, 0x71de,
                0x71ab, 0x7179, 0x7146, 0x7114, 0x70e2, 0x70b1, 0x707f, 0x704e,
                0x701c, 0x6feb, 0x6fba, 0x6f8a, 0x6f59, 0x6f29, 0x6ef9, 0x6ec8,
                0x6e99, 0x6e69, 0x6e39, 0x6e0a, 0x6ddb, 0x6dab, 0x6d7d, 0x6d4e,
                0x6d1f, 0x6cf1, 0x6cc2, 0x6c94, 0x6c66, 0x6c38, 0x6c0a, 0x6bdd,
                0x6bb0, 0x6b82, 0x6b55, 0x6b28, 0x6afb, 0x6acf, 0x6aa2, 0x6a76,
                0x6a49, 0x6a1d, 0x69f1, 0x69c6, 0x699a, 0x696e, 0x6943, 0x6918,
                0x68ed, 0x68c2, 0x6897, 0x686c, 0x6842, 0x6817, 0x67ed, 0x67c3,
                0x6799, 0x676f, 0x6745, 0x671b, 0x66f2, 0x66c8, 0x669f, 0x6676,
                0x664d, 0x6624, 0x65fc, 0x65d3, 0x65aa, 0x6582, 0x655a, 0x6532,
                0x650a, 0x64e2, 0x64ba, 0x6493, 0x646b, 0x6444, 0x641c, 0x63f5,
                0x63ce, 0x63a7, 0x6381, 0x635a, 0x6333, 0x630d, 0x62e7, 0x62c1,
                0x629a, 0x6275, 0x624f, 0x6229, 0x6203, 0x61de, 0x61b8, 0x6193,
                0x616e, 0x6149, 0x6124, 0x60ff, 0x60da, 0x60b6, 0x6091, 0x606d,
                0x6049, 0x6024, 0x6000, 0x5fdc, 0x5fb8, 0x5f95, 0x5f71, 0x5f4d,
                0x5f2a, 0x5f07, 0x5ee3, 0x5ec0, 0x5e9d, 0x5e7a, 0x5e57, 0x5e35,
                0x5e12, 0x5def, 0x5dcd, 0x5dab, 0x5d88, 0x5d66, 0x5d44, 0x5d22,
                0x5d00, 0x5cde, 0x5cbd, 0x5c9b, 0x5c7a, 0x5c58, 0x5c37, 0x5c16,
                0x5bf5, 0x5bd4, 0x5bb3, 0x5b92, 0x5b71, 0x5b51, 0x5b30, 0x5b10,
                0x5aef, 0x5acf, 0x5aaf, 0x5a8f, 0x5a6f, 0x5a4f, 0x5a2f, 0x5a0f,
                0x59ef, 0x59d0, 0x59b0, 0x5991, 0x5972, 0x5952, 0x5933, 0x5914,
                0x58f5, 0x58d6, 0x58b7, 0x5899, 0x587a, 0x585b, 0x583d, 0x581f,
                0x5800, 0x57e2, 0x57c4, 0x57a6, 0x5788, 0x576a, 0x574c, 0x572e,
                0x5711, 0x56f3, 0x56d5, 0x56b8, 0x569b, 0x567d, 0x5660, 0x5643,
                0x5626, 0x5609, 0x55ec, 0x55cf, 0x55b2, 0x5596, 0x5579, 0x555d,
                0x5540, 0x5524, 0x5507, 0x54eb, 0x54cf, 0x54b3, 0x5497, 0x547b,
                0x545f, 0x5443, 0x5428, 0x540c, 0x53f0, 0x53d5, 0x53b9, 0x539e,
                0x5383, 0x5368, 0x534c, 0x5331, 0x5316, 0x52fb, 0x52e0, 0x52c6,
                0x52ab, 0x5290, 0x5276, 0x525b, 0x5240, 0x5226, 0x520c, 0x51f1,
                0x51d7, 0x51bd, 0x51a3, 0x5189, 0x516f, 0x5155, 0x513b, 0x5121,
                0x5108, 0x50ee, 0x50d5, 0x50bb, 0x50a2, 0x5088, 0x506f, 0x5056,
                0x503c, 0x5023, 0x500a, 0x4ff1, 0x4fd8, 0x4fbf, 0x4fa6, 0x4f8e,
                0x4f75, 0x4f5c, 0x4f44, 0x4f2b, 0x4f13, 0x4efa, 0x4ee2, 0x4eca,
                0x4eb1, 0x4e99, 0x4e81, 0x4e69, 0x4e51, 0x4e39, 0x4e21, 0x4e09,
                0x4df1, 0x4dda, 0x4dc2, 0x4daa, 0x4d93, 0x4d7b, 0x4d64, 0x4d4d,
                0x4d35, 0x4d1e, 0x4d07, 0x4cf0, 0x4cd8, 0x4cc1, 0x4caa, 0x4c93,
                0x4c7d, 0x4c66, 0x4c4f, 0x4c38, 0x4c21, 0x4c0b, 0x4bf4, 0x4bde,
                0x4bc7, 0x4bb1, 0x4b9a, 0x4b84, 0x4b6e, 0x4b58, 0x4b41, 0x4b2b,
                0x4b15, 0x4aff, 0x4ae9, 0x4ad3, 0x4abd, 0x4aa8, 0x4a92, 0x4a7c,
                0x4a66, 0x4a51, 0x4a3b, 0x4a26, 0x4a10, 0x49fb, 0x49e5, 0x49d0,
                0x49bb, 0x49a6, 0x4990, 0x497b, 0x4966, 0x4951, 0x493c, 0x4927,
                0x4912, 0x48fe, 0x48e9, 0x48d4, 0x48bf, 0x48ab, 0x4896, 0x4881,
                0x486d, 0x4858, 0x4844, 0x482f, 0x481b, 0x4807, 0x47f3, 0x47de,
                0x47ca, 0x47b6, 0x47a2, 0x478e, 0x477a, 0x4766, 0x4752, 0x473e,
                0x472a, 0x4717, 0x4703, 0x46ef, 0x46db, 0x46c8, 0x46b4, 0x46a1,
                0x468d, 0x467a, 0x4666, 0x4653, 0x4640, 0x462c, 0x4619, 0x4606,
                0x45f3, 0x45e0, 0x45cd, 0x45ba, 0x45a7, 0x4594, 0x4581, 0x456e,
                0x455b, 0x4548, 0x4536, 0x4523, 0x4510, 0x44fe, 0x44eb, 0x44d8,
                0x44c6, 0x44b3, 0x44a1, 0x448f, 0x447c, 0x446a, 0x4458, 0x4445,
                0x4433, 0x4421, 0x440f, 0x43fd, 0x43eb, 0x43d9, 0x43c7, 0x43b5,
                0x43a3, 0x4391, 0x437f, 0x436d, 0x435c, 0x434a, 0x4338, 0x4327,
                0x4315, 0x4303, 0x42f2, 0x42e0, 0x42cf, 0x42bd, 0x42ac, 0x429b,
                0x4289, 0x4278, 0x4267, 0x4256, 0x4244, 0x4233, 0x4222, 0x4211,
                0x4200, 0x41ef, 0x41de, 0x41cd, 0x41bc, 0x41ab, 0x419a, 0x418a,
                0x4179, 0x4168, 0x4157, 0x4147, 0x4136, 0x4125, 0x4115, 0x4104,
                0x40f4, 0x40e3, 0x40d3, 0x40c2, 0x40b2, 0x40a2, 0x4091, 0x4081,
                0x4071, 0x4061, 0x4050, 0x4040, 0x4030, 0x4020, 0x4010, 0x4000};

/*
     * For 256 <= d < 512,
     *
     *   RECIPROCAL_TABLE_64[d - 256] = floor((2^19 - 3*2^9)/d)
     *
     * Total space ATM is 256*2 = 512 bytes. Entries range from 10 to 11
     * bits, so with some clever handling of hi bits, we could get three
     * entries per 32 bit word, reducing the size to about 256*11/8 = 352
     * bytes.
     *
     * TODO: Investigate whether alternative storage layouts are better;
     * see RECIPROCAL_TABLE_32 above for ideas.
     */
__device__ __constant__
    uint16_t
        RECIPROCAL_TABLE_64[0x100] =
            {
                0x7fd, 0x7f5, 0x7ed, 0x7e5, 0x7dd, 0x7d5, 0x7ce, 0x7c6,
                0x7bf, 0x7b7, 0x7b0, 0x7a8, 0x7a1, 0x79a, 0x792, 0x78b,
                0x784, 0x77d, 0x776, 0x76f, 0x768, 0x761, 0x75b, 0x754,
                0x74d, 0x747, 0x740, 0x739, 0x733, 0x72c, 0x726, 0x720,
                0x719, 0x713, 0x70d, 0x707, 0x700, 0x6fa, 0x6f4, 0x6ee,
                0x6e8, 0x6e2, 0x6dc, 0x6d6, 0x6d1, 0x6cb, 0x6c5, 0x6bf,
                0x6ba, 0x6b4, 0x6ae, 0x6a9, 0x6a3, 0x69e, 0x698, 0x693,
                0x68d, 0x688, 0x683, 0x67d, 0x678, 0x673, 0x66e, 0x669,
                0x664, 0x65e, 0x659, 0x654, 0x64f, 0x64a, 0x645, 0x640,
                0x63c, 0x637, 0x632, 0x62d, 0x628, 0x624, 0x61f, 0x61a,
                0x616, 0x611, 0x60c, 0x608, 0x603, 0x5ff, 0x5fa, 0x5f6,
                0x5f1, 0x5ed, 0x5e9, 0x5e4, 0x5e0, 0x5dc, 0x5d7, 0x5d3,
                0x5cf, 0x5cb, 0x5c6, 0x5c2, 0x5be, 0x5ba, 0x5b6, 0x5b2,
                0x5ae, 0x5aa, 0x5a6, 0x5a2, 0x59e, 0x59a, 0x596, 0x592,
                0x58e, 0x58a, 0x586, 0x583, 0x57f, 0x57b, 0x577, 0x574,
                0x570, 0x56c, 0x568, 0x565, 0x561, 0x55e, 0x55a, 0x556,
                0x553, 0x54f, 0x54c, 0x548, 0x545, 0x541, 0x53e, 0x53a,
                0x537, 0x534, 0x530, 0x52d, 0x52a, 0x526, 0x523, 0x520,
                0x51c, 0x519, 0x516, 0x513, 0x50f, 0x50c, 0x509, 0x506,
                0x503, 0x500, 0x4fc, 0x4f9, 0x4f6, 0x4f3, 0x4f0, 0x4ed,
                0x4ea, 0x4e7, 0x4e4, 0x4e1, 0x4de, 0x4db, 0x4d8, 0x4d5,
                0x4d2, 0x4cf, 0x4cc, 0x4ca, 0x4c7, 0x4c4, 0x4c1, 0x4be,
                0x4bb, 0x4b9, 0x4b6, 0x4b3, 0x4b0, 0x4ad, 0x4ab, 0x4a8,
                0x4a5, 0x4a3, 0x4a0, 0x49d, 0x49b, 0x498, 0x495, 0x493,
                0x490, 0x48d, 0x48b, 0x488, 0x486, 0x483, 0x481, 0x47e,
                0x47c, 0x479, 0x477, 0x474, 0x472, 0x46f, 0x46d, 0x46a,
                0x468, 0x465, 0x463, 0x461, 0x45e, 0x45c, 0x459, 0x457,
                0x455, 0x452, 0x450, 0x44e, 0x44b, 0x449, 0x447, 0x444,
                0x442, 0x440, 0x43e, 0x43b, 0x439, 0x437, 0x435, 0x432,
                0x430, 0x42e, 0x42c, 0x42a, 0x428, 0x425, 0x423, 0x421,
                0x41f, 0x41d, 0x41b, 0x419, 0x417, 0x414, 0x412, 0x410,
                0x40e, 0x40c, 0x40a, 0x408, 0x406, 0x404, 0x402, 0x400};

/*
     * This Table 2 from Koç, C. K., 1995, "Analysis of Sliding Window
     * Techniques for Exponentiation".
     *
     * The resolution of this table is higher than the one above because it's
     * used in the fixed exponent modexp code and can benefit from using the
     * precise bit length of the exponent, whereas the table above has to
     * accommodate multiple different exponents simultaneously.
     */
__constant__ int BYTES_TO_CLNW_WINDOW_SIZE_TABLE[] = {
    -1, // bits
    4,  //  128
    5,  //  256
    5,  //  384
    5,  //  512
    6,  //  640
    6,  //  768
    6,  //  896
    6,  // 1024
    6,  // 1152
    6,  // 1280
    6,  // 1408
    6,  // 1536
    6,  // 1664
    7,  // 1792
    7,  // 1920
    7,  // 2048
};
	.file	"tmpxft_0027a660_00000000-6_compile.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"RECIPROCAL_TABLE_32"
.LC1:
	.string	"RECIPROCAL_TABLE_64"
.LC2:
	.string	"BYTES_TO_CLNW_WINDOW_SIZE_TABLE"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	movl	$1024, %r9d
	xorl	%r8d, %r8d
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	movq	%rdx, %rcx
	leaq	_ZL19RECIPROCAL_TABLE_32(%rip), %rsi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$512, %r9d
	leaq	_ZL19RECIPROCAL_TABLE_64(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$68, %r9d
	movq	%rdx, %rcx
	leaq	_ZL31BYTES_TO_CLNW_WINDOW_SIZE_TABLE(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.local	_ZL31BYTES_TO_CLNW_WINDOW_SIZE_TABLE
	.comm	_ZL31BYTES_TO_CLNW_WINDOW_SIZE_TABLE,68,32
	.local	_ZL19RECIPROCAL_TABLE_64
	.comm	_ZL19RECIPROCAL_TABLE_64,512,32
	.local	_ZL19RECIPROCAL_TABLE_32
	.comm	_ZL19RECIPROCAL_TABLE_32,1024,32
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <cstdint>

/*
     * For 512 <= d < 1024,
     *
     *   RECIPROCAL_TABLE_32[d - 512] = floor((2^24 - 2^14 + 2^9)/d)
     *
     * Total space at the moment is 512*2 = 1024 bytes.
     *
     * TODO: Investigate whether alternative storage layouts are better; examples:
     *
     *  - redundantly store each element in a uint32_t
     *  - pack two uint16_t values into each uint32_t
     *  - is __constant__ the right storage specifier? Maybe load into shared memory?
     *    Shared memory seems like an excellent choice (48k available per SM), though
     *    I'll need to be mindful of bank conflicts (perhaps circumvent by having
     *    many copies of the data in SM?).
     *  - perhaps reading an element from memory is slower than simply calculating
     *    floor((2^24 - 2^14 + 2^9)/d) in assembly?
     */
__device__ __constant__
    uint16_t
        RECIPROCAL_TABLE_32[0x200] =
            {
                0x7fe1, 0x7fa1, 0x7f61, 0x7f22, 0x7ee3, 0x7ea4, 0x7e65, 0x7e27,
                0x7de9, 0x7dab, 0x7d6d, 0x7d30, 0x7cf3, 0x7cb6, 0x7c79, 0x7c3d,
                0x7c00, 0x7bc4, 0x7b89, 0x7b4d, 0x7b12, 0x7ad7, 0x7a9c, 0x7a61,
                0x7a27, 0x79ec, 0x79b2, 0x7979, 0x793f, 0x7906, 0x78cc, 0x7894,
                0x785b, 0x7822, 0x77ea, 0x77b2, 0x777a, 0x7742, 0x770b, 0x76d3,
                0x769c, 0x7665, 0x762f, 0x75f8, 0x75c2, 0x758c, 0x7556, 0x7520,
                0x74ea, 0x74b5, 0x7480, 0x744b, 0x7416, 0x73e2, 0x73ad, 0x7379,
                0x7345, 0x7311, 0x72dd, 0x72aa, 0x7277, 0x7243, 0x7210, 0x71de,
                0x71ab, 0x7179, 0x7146, 0x7114, 0x70e2, 0x70b1, 0x707f, 0x704e,
                0x701c, 0x6feb, 0x6fba, 0x6f8a, 0x6f59, 0x6f29, 0x6ef9, 0x6ec8,
                0x6e99, 0x6e69, 0x6e39, 0x6e0a, 0x6ddb, 0x6dab, 0x6d7d, 0x6d4e,
                0x6d1f, 0x6cf1, 0x6cc2, 0x6c94, 0x6c66, 0x6c38, 0x6c0a, 0x6bdd,
                0x6bb0, 0x6b82, 0x6b55, 0x6b28, 0x6afb, 0x6acf, 0x6aa2, 0x6a76,
                0x6a49, 0x6a1d, 0x69f1, 0x69c6, 0x699a, 0x696e, 0x6943, 0x6918,
                0x68ed, 0x68c2, 0x6897, 0x686c, 0x6842, 0x6817, 0x67ed, 0x67c3,
                0x6799, 0x676f, 0x6745, 0x671b, 0x66f2, 0x66c8, 0x669f, 0x6676,
                0x664d, 0x6624, 0x65fc, 0x65d3, 0x65aa, 0x6582, 0x655a, 0x6532,
                0x650a, 0x64e2, 0x64ba, 0x6493, 0x646b, 0x6444, 0x641c, 0x63f5,
                0x63ce, 0x63a7, 0x6381, 0x635a, 0x6333, 0x630d, 0x62e7, 0x62c1,
                0x629a, 0x6275, 0x624f, 0x6229, 0x6203, 0x61de, 0x61b8, 0x6193,
                0x616e, 0x6149, 0x6124, 0x60ff, 0x60da, 0x60b6, 0x6091, 0x606d,
                0x6049, 0x6024, 0x6000, 0x5fdc, 0x5fb8, 0x5f95, 0x5f71, 0x5f4d,
                0x5f2a, 0x5f07, 0x5ee3, 0x5ec0, 0x5e9d, 0x5e7a, 0x5e57, 0x5e35,
                0x5e12, 0x5def, 0x5dcd, 0x5dab, 0x5d88, 0x5d66, 0x5d44, 0x5d22,
                0x5d00, 0x5cde, 0x5cbd, 0x5c9b, 0x5c7a, 0x5c58, 0x5c37, 0x5c16,
                0x5bf5, 0x5bd4, 0x5bb3, 0x5b92, 0x5b71, 0x5b51, 0x5b30, 0x5b10,
                0x5aef, 0x5acf, 0x5aaf, 0x5a8f, 0x5a6f, 0x5a4f, 0x5a2f, 0x5a0f,
                0x59ef, 0x59d0, 0x59b0, 0x5991, 0x5972, 0x5952, 0x5933, 0x5914,
                0x58f5, 0x58d6, 0x58b7, 0x5899, 0x587a, 0x585b, 0x583d, 0x581f,
                0x5800, 0x57e2, 0x57c4, 0x57a6, 0x5788, 0x576a, 0x574c, 0x572e,
                0x5711, 0x56f3, 0x56d5, 0x56b8, 0x569b, 0x567d, 0x5660, 0x5643,
                0x5626, 0x5609, 0x55ec, 0x55cf, 0x55b2, 0x5596, 0x5579, 0x555d,
                0x5540, 0x5524, 0x5507, 0x54eb, 0x54cf, 0x54b3, 0x5497, 0x547b,
                0x545f, 0x5443, 0x5428, 0x540c, 0x53f0, 0x53d5, 0x53b9, 0x539e,
                0x5383, 0x5368, 0x534c, 0x5331, 0x5316, 0x52fb, 0x52e0, 0x52c6,
                0x52ab, 0x5290, 0x5276, 0x525b, 0x5240, 0x5226, 0x520c, 0x51f1,
                0x51d7, 0x51bd, 0x51a3, 0x5189, 0x516f, 0x5155, 0x513b, 0x5121,
                0x5108, 0x50ee, 0x50d5, 0x50bb, 0x50a2, 0x5088, 0x506f, 0x5056,
                0x503c, 0x5023, 0x500a, 0x4ff1, 0x4fd8, 0x4fbf, 0x4fa6, 0x4f8e,
                0x4f75, 0x4f5c, 0x4f44, 0x4f2b, 0x4f13, 0x4efa, 0x4ee2, 0x4eca,
                0x4eb1, 0x4e99, 0x4e81, 0x4e69, 0x4e51, 0x4e39, 0x4e21, 0x4e09,
                0x4df1, 0x4dda, 0x4dc2, 0x4daa, 0x4d93, 0x4d7b, 0x4d64, 0x4d4d,
                0x4d35, 0x4d1e, 0x4d07, 0x4cf0, 0x4cd8, 0x4cc1, 0x4caa, 0x4c93,
                0x4c7d, 0x4c66, 0x4c4f, 0x4c38, 0x4c21, 0x4c0b, 0x4bf4, 0x4bde,
                0x4bc7, 0x4bb1, 0x4b9a, 0x4b84, 0x4b6e, 0x4b58, 0x4b41, 0x4b2b,
                0x4b15, 0x4aff, 0x4ae9, 0x4ad3, 0x4abd, 0x4aa8, 0x4a92, 0x4a7c,
                0x4a66, 0x4a51, 0x4a3b, 0x4a26, 0x4a10, 0x49fb, 0x49e5, 0x49d0,
                0x49bb, 0x49a6, 0x4990, 0x497b, 0x4966, 0x4951, 0x493c, 0x4927,
                0x4912, 0x48fe, 0x48e9, 0x48d4, 0x48bf, 0x48ab, 0x4896, 0x4881,
                0x486d, 0x4858, 0x4844, 0x482f, 0x481b, 0x4807, 0x47f3, 0x47de,
                0x47ca, 0x47b6, 0x47a2, 0x478e, 0x477a, 0x4766, 0x4752, 0x473e,
                0x472a, 0x4717, 0x4703, 0x46ef, 0x46db, 0x46c8, 0x46b4, 0x46a1,
                0x468d, 0x467a, 0x4666, 0x4653, 0x4640, 0x462c, 0x4619, 0x4606,
                0x45f3, 0x45e0, 0x45cd, 0x45ba, 0x45a7, 0x4594, 0x4581, 0x456e,
                0x455b, 0x4548, 0x4536, 0x4523, 0x4510, 0x44fe, 0x44eb, 0x44d8,
                0x44c6, 0x44b3, 0x44a1, 0x448f, 0x447c, 0x446a, 0x4458, 0x4445,
                0x4433, 0x4421, 0x440f, 0x43fd, 0x43eb, 0x43d9, 0x43c7, 0x43b5,
                0x43a3, 0x4391, 0x437f, 0x436d, 0x435c, 0x434a, 0x4338, 0x4327,
                0x4315, 0x4303, 0x42f2, 0x42e0, 0x42cf, 0x42bd, 0x42ac, 0x429b,
                0x4289, 0x4278, 0x4267, 0x4256, 0x4244, 0x4233, 0x4222, 0x4211,
                0x4200, 0x41ef, 0x41de, 0x41cd, 0x41bc, 0x41ab, 0x419a, 0x418a,
                0x4179, 0x4168, 0x4157, 0x4147, 0x4136, 0x4125, 0x4115, 0x4104,
                0x40f4, 0x40e3, 0x40d3, 0x40c2, 0x40b2, 0x40a2, 0x4091, 0x4081,
                0x4071, 0x4061, 0x4050, 0x4040, 0x4030, 0x4020, 0x4010, 0x4000};

/*
     * For 256 <= d < 512,
     *
     *   RECIPROCAL_TABLE_64[d - 256] = floor((2^19 - 3*2^9)/d)
     *
     * Total space ATM is 256*2 = 512 bytes. Entries range from 10 to 11
     * bits, so with some clever handling of hi bits, we could get three
     * entries per 32 bit word, reducing the size to about 256*11/8 = 352
     * bytes.
     *
     * TODO: Investigate whether alternative storage layouts are better;
     * see RECIPROCAL_TABLE_32 above for ideas.
     */
__device__ __constant__
    uint16_t
        RECIPROCAL_TABLE_64[0x100] =
            {
                0x7fd, 0x7f5, 0x7ed, 0x7e5, 0x7dd, 0x7d5, 0x7ce, 0x7c6,
                0x7bf, 0x7b7, 0x7b0, 0x7a8, 0x7a1, 0x79a, 0x792, 0x78b,
                0x784, 0x77d, 0x776, 0x76f, 0x768, 0x761, 0x75b, 0x754,
                0x74d, 0x747, 0x740, 0x739, 0x733, 0x72c, 0x726, 0x720,
                0x719, 0x713, 0x70d, 0x707, 0x700, 0x6fa, 0x6f4, 0x6ee,
                0x6e8, 0x6e2, 0x6dc, 0x6d6, 0x6d1, 0x6cb, 0x6c5, 0x6bf,
                0x6ba, 0x6b4, 0x6ae, 0x6a9, 0x6a3, 0x69e, 0x698, 0x693,
                0x68d, 0x688, 0x683, 0x67d, 0x678, 0x673, 0x66e, 0x669,
                0x664, 0x65e, 0x659, 0x654, 0x64f, 0x64a, 0x645, 0x640,
                0x63c, 0x637, 0x632, 0x62d, 0x628, 0x624, 0x61f, 0x61a,
                0x616, 0x611, 0x60c, 0x608, 0x603, 0x5ff, 0x5fa, 0x5f6,
                0x5f1, 0x5ed, 0x5e9, 0x5e4, 0x5e0, 0x5dc, 0x5d7, 0x5d3,
                0x5cf, 0x5cb, 0x5c6, 0x5c2, 0x5be, 0x5ba, 0x5b6, 0x5b2,
                0x5ae, 0x5aa, 0x5a6, 0x5a2, 0x59e, 0x59a, 0x596, 0x592,
                0x58e, 0x58a, 0x586, 0x583, 0x57f, 0x57b, 0x577, 0x574,
                0x570, 0x56c, 0x568, 0x565, 0x561, 0x55e, 0x55a, 0x556,
                0x553, 0x54f, 0x54c, 0x548, 0x545, 0x541, 0x53e, 0x53a,
                0x537, 0x534, 0x530, 0x52d, 0x52a, 0x526, 0x523, 0x520,
                0x51c, 0x519, 0x516, 0x513, 0x50f, 0x50c, 0x509, 0x506,
                0x503, 0x500, 0x4fc, 0x4f9, 0x4f6, 0x4f3, 0x4f0, 0x4ed,
                0x4ea, 0x4e7, 0x4e4, 0x4e1, 0x4de, 0x4db, 0x4d8, 0x4d5,
                0x4d2, 0x4cf, 0x4cc, 0x4ca, 0x4c7, 0x4c4, 0x4c1, 0x4be,
                0x4bb, 0x4b9, 0x4b6, 0x4b3, 0x4b0, 0x4ad, 0x4ab, 0x4a8,
                0x4a5, 0x4a3, 0x4a0, 0x49d, 0x49b, 0x498, 0x495, 0x493,
                0x490, 0x48d, 0x48b, 0x488, 0x486, 0x483, 0x481, 0x47e,
                0x47c, 0x479, 0x477, 0x474, 0x472, 0x46f, 0x46d, 0x46a,
                0x468, 0x465, 0x463, 0x461, 0x45e, 0x45c, 0x459, 0x457,
                0x455, 0x452, 0x450, 0x44e, 0x44b, 0x449, 0x447, 0x444,
                0x442, 0x440, 0x43e, 0x43b, 0x439, 0x437, 0x435, 0x432,
                0x430, 0x42e, 0x42c, 0x42a, 0x428, 0x425, 0x423, 0x421,
                0x41f, 0x41d, 0x41b, 0x419, 0x417, 0x414, 0x412, 0x410,
                0x40e, 0x40c, 0x40a, 0x408, 0x406, 0x404, 0x402, 0x400};

/*
     * This Table 2 from Koç, C. K., 1995, "Analysis of Sliding Window
     * Techniques for Exponentiation".
     *
     * The resolution of this table is higher than the one above because it's
     * used in the fixed exponent modexp code and can benefit from using the
     * precise bit length of the exponent, whereas the table above has to
     * accommodate multiple different exponents simultaneously.
     */
__constant__ int BYTES_TO_CLNW_WINDOW_SIZE_TABLE[] = {
    -1, // bits
    4,  //  128
    5,  //  256
    5,  //  384
    5,  //  512
    6,  //  640
    6,  //  768
    6,  //  896
    6,  // 1024
    6,  // 1152
    6,  // 1280
    6,  // 1408
    6,  // 1536
    6,  // 1664
    7,  // 1792
    7,  // 1920
    7,  // 2048
};
	.text
	.file	"compile.hip"
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB0_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB0_2:
	movl	$RECIPROCAL_TABLE_32, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$1024, %r9d                     # imm = 0x400
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$RECIPROCAL_TABLE_64, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$512, %r9d                      # imm = 0x200
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$BYTES_TO_CLNW_WINDOW_SIZE_TABLE, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$68, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end0:
	.size	__hip_module_ctor, .Lfunc_end0-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB1_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB1_2:
	retq
.Lfunc_end1:
	.size	__hip_module_dtor, .Lfunc_end1-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	RECIPROCAL_TABLE_32,@object     # @RECIPROCAL_TABLE_32
	.local	RECIPROCAL_TABLE_32
	.comm	RECIPROCAL_TABLE_32,1024,16
	.type	RECIPROCAL_TABLE_64,@object     # @RECIPROCAL_TABLE_64
	.local	RECIPROCAL_TABLE_64
	.comm	RECIPROCAL_TABLE_64,512,16
	.type	BYTES_TO_CLNW_WINDOW_SIZE_TABLE,@object # @BYTES_TO_CLNW_WINDOW_SIZE_TABLE
	.local	BYTES_TO_CLNW_WINDOW_SIZE_TABLE
	.comm	BYTES_TO_CLNW_WINDOW_SIZE_TABLE,68,16
	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"RECIPROCAL_TABLE_32"
	.size	.L__unnamed_1, 20

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"RECIPROCAL_TABLE_64"
	.size	.L__unnamed_2, 20

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"BYTES_TO_CLNW_WINDOW_SIZE_TABLE"
	.size	.L__unnamed_3, 32

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym RECIPROCAL_TABLE_32
	.addrsig_sym RECIPROCAL_TABLE_64
	.addrsig_sym BYTES_TO_CLNW_WINDOW_SIZE_TABLE
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	RECIPROCAL_TABLE_32     ; @RECIPROCAL_TABLE_32
	.type	RECIPROCAL_TABLE_32,@object
	.data
	.globl	RECIPROCAL_TABLE_32
	.p2align	4, 0x0
RECIPROCAL_TABLE_32:
	.short	32737                           ; 0x7fe1
	.short	32673                           ; 0x7fa1
	.short	32609                           ; 0x7f61
	.short	32546                           ; 0x7f22
	.short	32483                           ; 0x7ee3
	.short	32420                           ; 0x7ea4
	.short	32357                           ; 0x7e65
	.short	32295                           ; 0x7e27
	.short	32233                           ; 0x7de9
	.short	32171                           ; 0x7dab
	.short	32109                           ; 0x7d6d
	.short	32048                           ; 0x7d30
	.short	31987                           ; 0x7cf3
	.short	31926                           ; 0x7cb6
	.short	31865                           ; 0x7c79
	.short	31805                           ; 0x7c3d
	.short	31744                           ; 0x7c00
	.short	31684                           ; 0x7bc4
	.short	31625                           ; 0x7b89
	.short	31565                           ; 0x7b4d
	.short	31506                           ; 0x7b12
	.short	31447                           ; 0x7ad7
	.short	31388                           ; 0x7a9c
	.short	31329                           ; 0x7a61
	.short	31271                           ; 0x7a27
	.short	31212                           ; 0x79ec
	.short	31154                           ; 0x79b2
	.short	31097                           ; 0x7979
	.short	31039                           ; 0x793f
	.short	30982                           ; 0x7906
	.short	30924                           ; 0x78cc
	.short	30868                           ; 0x7894
	.short	30811                           ; 0x785b
	.short	30754                           ; 0x7822
	.short	30698                           ; 0x77ea
	.short	30642                           ; 0x77b2
	.short	30586                           ; 0x777a
	.short	30530                           ; 0x7742
	.short	30475                           ; 0x770b
	.short	30419                           ; 0x76d3
	.short	30364                           ; 0x769c
	.short	30309                           ; 0x7665
	.short	30255                           ; 0x762f
	.short	30200                           ; 0x75f8
	.short	30146                           ; 0x75c2
	.short	30092                           ; 0x758c
	.short	30038                           ; 0x7556
	.short	29984                           ; 0x7520
	.short	29930                           ; 0x74ea
	.short	29877                           ; 0x74b5
	.short	29824                           ; 0x7480
	.short	29771                           ; 0x744b
	.short	29718                           ; 0x7416
	.short	29666                           ; 0x73e2
	.short	29613                           ; 0x73ad
	.short	29561                           ; 0x7379
	.short	29509                           ; 0x7345
	.short	29457                           ; 0x7311
	.short	29405                           ; 0x72dd
	.short	29354                           ; 0x72aa
	.short	29303                           ; 0x7277
	.short	29251                           ; 0x7243
	.short	29200                           ; 0x7210
	.short	29150                           ; 0x71de
	.short	29099                           ; 0x71ab
	.short	29049                           ; 0x7179
	.short	28998                           ; 0x7146
	.short	28948                           ; 0x7114
	.short	28898                           ; 0x70e2
	.short	28849                           ; 0x70b1
	.short	28799                           ; 0x707f
	.short	28750                           ; 0x704e
	.short	28700                           ; 0x701c
	.short	28651                           ; 0x6feb
	.short	28602                           ; 0x6fba
	.short	28554                           ; 0x6f8a
	.short	28505                           ; 0x6f59
	.short	28457                           ; 0x6f29
	.short	28409                           ; 0x6ef9
	.short	28360                           ; 0x6ec8
	.short	28313                           ; 0x6e99
	.short	28265                           ; 0x6e69
	.short	28217                           ; 0x6e39
	.short	28170                           ; 0x6e0a
	.short	28123                           ; 0x6ddb
	.short	28075                           ; 0x6dab
	.short	28029                           ; 0x6d7d
	.short	27982                           ; 0x6d4e
	.short	27935                           ; 0x6d1f
	.short	27889                           ; 0x6cf1
	.short	27842                           ; 0x6cc2
	.short	27796                           ; 0x6c94
	.short	27750                           ; 0x6c66
	.short	27704                           ; 0x6c38
	.short	27658                           ; 0x6c0a
	.short	27613                           ; 0x6bdd
	.short	27568                           ; 0x6bb0
	.short	27522                           ; 0x6b82
	.short	27477                           ; 0x6b55
	.short	27432                           ; 0x6b28
	.short	27387                           ; 0x6afb
	.short	27343                           ; 0x6acf
	.short	27298                           ; 0x6aa2
	.short	27254                           ; 0x6a76
	.short	27209                           ; 0x6a49
	.short	27165                           ; 0x6a1d
	.short	27121                           ; 0x69f1
	.short	27078                           ; 0x69c6
	.short	27034                           ; 0x699a
	.short	26990                           ; 0x696e
	.short	26947                           ; 0x6943
	.short	26904                           ; 0x6918
	.short	26861                           ; 0x68ed
	.short	26818                           ; 0x68c2
	.short	26775                           ; 0x6897
	.short	26732                           ; 0x686c
	.short	26690                           ; 0x6842
	.short	26647                           ; 0x6817
	.short	26605                           ; 0x67ed
	.short	26563                           ; 0x67c3
	.short	26521                           ; 0x6799
	.short	26479                           ; 0x676f
	.short	26437                           ; 0x6745
	.short	26395                           ; 0x671b
	.short	26354                           ; 0x66f2
	.short	26312                           ; 0x66c8
	.short	26271                           ; 0x669f
	.short	26230                           ; 0x6676
	.short	26189                           ; 0x664d
	.short	26148                           ; 0x6624
	.short	26108                           ; 0x65fc
	.short	26067                           ; 0x65d3
	.short	26026                           ; 0x65aa
	.short	25986                           ; 0x6582
	.short	25946                           ; 0x655a
	.short	25906                           ; 0x6532
	.short	25866                           ; 0x650a
	.short	25826                           ; 0x64e2
	.short	25786                           ; 0x64ba
	.short	25747                           ; 0x6493
	.short	25707                           ; 0x646b
	.short	25668                           ; 0x6444
	.short	25628                           ; 0x641c
	.short	25589                           ; 0x63f5
	.short	25550                           ; 0x63ce
	.short	25511                           ; 0x63a7
	.short	25473                           ; 0x6381
	.short	25434                           ; 0x635a
	.short	25395                           ; 0x6333
	.short	25357                           ; 0x630d
	.short	25319                           ; 0x62e7
	.short	25281                           ; 0x62c1
	.short	25242                           ; 0x629a
	.short	25205                           ; 0x6275
	.short	25167                           ; 0x624f
	.short	25129                           ; 0x6229
	.short	25091                           ; 0x6203
	.short	25054                           ; 0x61de
	.short	25016                           ; 0x61b8
	.short	24979                           ; 0x6193
	.short	24942                           ; 0x616e
	.short	24905                           ; 0x6149
	.short	24868                           ; 0x6124
	.short	24831                           ; 0x60ff
	.short	24794                           ; 0x60da
	.short	24758                           ; 0x60b6
	.short	24721                           ; 0x6091
	.short	24685                           ; 0x606d
	.short	24649                           ; 0x6049
	.short	24612                           ; 0x6024
	.short	24576                           ; 0x6000
	.short	24540                           ; 0x5fdc
	.short	24504                           ; 0x5fb8
	.short	24469                           ; 0x5f95
	.short	24433                           ; 0x5f71
	.short	24397                           ; 0x5f4d
	.short	24362                           ; 0x5f2a
	.short	24327                           ; 0x5f07
	.short	24291                           ; 0x5ee3
	.short	24256                           ; 0x5ec0
	.short	24221                           ; 0x5e9d
	.short	24186                           ; 0x5e7a
	.short	24151                           ; 0x5e57
	.short	24117                           ; 0x5e35
	.short	24082                           ; 0x5e12
	.short	24047                           ; 0x5def
	.short	24013                           ; 0x5dcd
	.short	23979                           ; 0x5dab
	.short	23944                           ; 0x5d88
	.short	23910                           ; 0x5d66
	.short	23876                           ; 0x5d44
	.short	23842                           ; 0x5d22
	.short	23808                           ; 0x5d00
	.short	23774                           ; 0x5cde
	.short	23741                           ; 0x5cbd
	.short	23707                           ; 0x5c9b
	.short	23674                           ; 0x5c7a
	.short	23640                           ; 0x5c58
	.short	23607                           ; 0x5c37
	.short	23574                           ; 0x5c16
	.short	23541                           ; 0x5bf5
	.short	23508                           ; 0x5bd4
	.short	23475                           ; 0x5bb3
	.short	23442                           ; 0x5b92
	.short	23409                           ; 0x5b71
	.short	23377                           ; 0x5b51
	.short	23344                           ; 0x5b30
	.short	23312                           ; 0x5b10
	.short	23279                           ; 0x5aef
	.short	23247                           ; 0x5acf
	.short	23215                           ; 0x5aaf
	.short	23183                           ; 0x5a8f
	.short	23151                           ; 0x5a6f
	.short	23119                           ; 0x5a4f
	.short	23087                           ; 0x5a2f
	.short	23055                           ; 0x5a0f
	.short	23023                           ; 0x59ef
	.short	22992                           ; 0x59d0
	.short	22960                           ; 0x59b0
	.short	22929                           ; 0x5991
	.short	22898                           ; 0x5972
	.short	22866                           ; 0x5952
	.short	22835                           ; 0x5933
	.short	22804                           ; 0x5914
	.short	22773                           ; 0x58f5
	.short	22742                           ; 0x58d6
	.short	22711                           ; 0x58b7
	.short	22681                           ; 0x5899
	.short	22650                           ; 0x587a
	.short	22619                           ; 0x585b
	.short	22589                           ; 0x583d
	.short	22559                           ; 0x581f
	.short	22528                           ; 0x5800
	.short	22498                           ; 0x57e2
	.short	22468                           ; 0x57c4
	.short	22438                           ; 0x57a6
	.short	22408                           ; 0x5788
	.short	22378                           ; 0x576a
	.short	22348                           ; 0x574c
	.short	22318                           ; 0x572e
	.short	22289                           ; 0x5711
	.short	22259                           ; 0x56f3
	.short	22229                           ; 0x56d5
	.short	22200                           ; 0x56b8
	.short	22171                           ; 0x569b
	.short	22141                           ; 0x567d
	.short	22112                           ; 0x5660
	.short	22083                           ; 0x5643
	.short	22054                           ; 0x5626
	.short	22025                           ; 0x5609
	.short	21996                           ; 0x55ec
	.short	21967                           ; 0x55cf
	.short	21938                           ; 0x55b2
	.short	21910                           ; 0x5596
	.short	21881                           ; 0x5579
	.short	21853                           ; 0x555d
	.short	21824                           ; 0x5540
	.short	21796                           ; 0x5524
	.short	21767                           ; 0x5507
	.short	21739                           ; 0x54eb
	.short	21711                           ; 0x54cf
	.short	21683                           ; 0x54b3
	.short	21655                           ; 0x5497
	.short	21627                           ; 0x547b
	.short	21599                           ; 0x545f
	.short	21571                           ; 0x5443
	.short	21544                           ; 0x5428
	.short	21516                           ; 0x540c
	.short	21488                           ; 0x53f0
	.short	21461                           ; 0x53d5
	.short	21433                           ; 0x53b9
	.short	21406                           ; 0x539e
	.short	21379                           ; 0x5383
	.short	21352                           ; 0x5368
	.short	21324                           ; 0x534c
	.short	21297                           ; 0x5331
	.short	21270                           ; 0x5316
	.short	21243                           ; 0x52fb
	.short	21216                           ; 0x52e0
	.short	21190                           ; 0x52c6
	.short	21163                           ; 0x52ab
	.short	21136                           ; 0x5290
	.short	21110                           ; 0x5276
	.short	21083                           ; 0x525b
	.short	21056                           ; 0x5240
	.short	21030                           ; 0x5226
	.short	21004                           ; 0x520c
	.short	20977                           ; 0x51f1
	.short	20951                           ; 0x51d7
	.short	20925                           ; 0x51bd
	.short	20899                           ; 0x51a3
	.short	20873                           ; 0x5189
	.short	20847                           ; 0x516f
	.short	20821                           ; 0x5155
	.short	20795                           ; 0x513b
	.short	20769                           ; 0x5121
	.short	20744                           ; 0x5108
	.short	20718                           ; 0x50ee
	.short	20693                           ; 0x50d5
	.short	20667                           ; 0x50bb
	.short	20642                           ; 0x50a2
	.short	20616                           ; 0x5088
	.short	20591                           ; 0x506f
	.short	20566                           ; 0x5056
	.short	20540                           ; 0x503c
	.short	20515                           ; 0x5023
	.short	20490                           ; 0x500a
	.short	20465                           ; 0x4ff1
	.short	20440                           ; 0x4fd8
	.short	20415                           ; 0x4fbf
	.short	20390                           ; 0x4fa6
	.short	20366                           ; 0x4f8e
	.short	20341                           ; 0x4f75
	.short	20316                           ; 0x4f5c
	.short	20292                           ; 0x4f44
	.short	20267                           ; 0x4f2b
	.short	20243                           ; 0x4f13
	.short	20218                           ; 0x4efa
	.short	20194                           ; 0x4ee2
	.short	20170                           ; 0x4eca
	.short	20145                           ; 0x4eb1
	.short	20121                           ; 0x4e99
	.short	20097                           ; 0x4e81
	.short	20073                           ; 0x4e69
	.short	20049                           ; 0x4e51
	.short	20025                           ; 0x4e39
	.short	20001                           ; 0x4e21
	.short	19977                           ; 0x4e09
	.short	19953                           ; 0x4df1
	.short	19930                           ; 0x4dda
	.short	19906                           ; 0x4dc2
	.short	19882                           ; 0x4daa
	.short	19859                           ; 0x4d93
	.short	19835                           ; 0x4d7b
	.short	19812                           ; 0x4d64
	.short	19789                           ; 0x4d4d
	.short	19765                           ; 0x4d35
	.short	19742                           ; 0x4d1e
	.short	19719                           ; 0x4d07
	.short	19696                           ; 0x4cf0
	.short	19672                           ; 0x4cd8
	.short	19649                           ; 0x4cc1
	.short	19626                           ; 0x4caa
	.short	19603                           ; 0x4c93
	.short	19581                           ; 0x4c7d
	.short	19558                           ; 0x4c66
	.short	19535                           ; 0x4c4f
	.short	19512                           ; 0x4c38
	.short	19489                           ; 0x4c21
	.short	19467                           ; 0x4c0b
	.short	19444                           ; 0x4bf4
	.short	19422                           ; 0x4bde
	.short	19399                           ; 0x4bc7
	.short	19377                           ; 0x4bb1
	.short	19354                           ; 0x4b9a
	.short	19332                           ; 0x4b84
	.short	19310                           ; 0x4b6e
	.short	19288                           ; 0x4b58
	.short	19265                           ; 0x4b41
	.short	19243                           ; 0x4b2b
	.short	19221                           ; 0x4b15
	.short	19199                           ; 0x4aff
	.short	19177                           ; 0x4ae9
	.short	19155                           ; 0x4ad3
	.short	19133                           ; 0x4abd
	.short	19112                           ; 0x4aa8
	.short	19090                           ; 0x4a92
	.short	19068                           ; 0x4a7c
	.short	19046                           ; 0x4a66
	.short	19025                           ; 0x4a51
	.short	19003                           ; 0x4a3b
	.short	18982                           ; 0x4a26
	.short	18960                           ; 0x4a10
	.short	18939                           ; 0x49fb
	.short	18917                           ; 0x49e5
	.short	18896                           ; 0x49d0
	.short	18875                           ; 0x49bb
	.short	18854                           ; 0x49a6
	.short	18832                           ; 0x4990
	.short	18811                           ; 0x497b
	.short	18790                           ; 0x4966
	.short	18769                           ; 0x4951
	.short	18748                           ; 0x493c
	.short	18727                           ; 0x4927
	.short	18706                           ; 0x4912
	.short	18686                           ; 0x48fe
	.short	18665                           ; 0x48e9
	.short	18644                           ; 0x48d4
	.short	18623                           ; 0x48bf
	.short	18603                           ; 0x48ab
	.short	18582                           ; 0x4896
	.short	18561                           ; 0x4881
	.short	18541                           ; 0x486d
	.short	18520                           ; 0x4858
	.short	18500                           ; 0x4844
	.short	18479                           ; 0x482f
	.short	18459                           ; 0x481b
	.short	18439                           ; 0x4807
	.short	18419                           ; 0x47f3
	.short	18398                           ; 0x47de
	.short	18378                           ; 0x47ca
	.short	18358                           ; 0x47b6
	.short	18338                           ; 0x47a2
	.short	18318                           ; 0x478e
	.short	18298                           ; 0x477a
	.short	18278                           ; 0x4766
	.short	18258                           ; 0x4752
	.short	18238                           ; 0x473e
	.short	18218                           ; 0x472a
	.short	18199                           ; 0x4717
	.short	18179                           ; 0x4703
	.short	18159                           ; 0x46ef
	.short	18139                           ; 0x46db
	.short	18120                           ; 0x46c8
	.short	18100                           ; 0x46b4
	.short	18081                           ; 0x46a1
	.short	18061                           ; 0x468d
	.short	18042                           ; 0x467a
	.short	18022                           ; 0x4666
	.short	18003                           ; 0x4653
	.short	17984                           ; 0x4640
	.short	17964                           ; 0x462c
	.short	17945                           ; 0x4619
	.short	17926                           ; 0x4606
	.short	17907                           ; 0x45f3
	.short	17888                           ; 0x45e0
	.short	17869                           ; 0x45cd
	.short	17850                           ; 0x45ba
	.short	17831                           ; 0x45a7
	.short	17812                           ; 0x4594
	.short	17793                           ; 0x4581
	.short	17774                           ; 0x456e
	.short	17755                           ; 0x455b
	.short	17736                           ; 0x4548
	.short	17718                           ; 0x4536
	.short	17699                           ; 0x4523
	.short	17680                           ; 0x4510
	.short	17662                           ; 0x44fe
	.short	17643                           ; 0x44eb
	.short	17624                           ; 0x44d8
	.short	17606                           ; 0x44c6
	.short	17587                           ; 0x44b3
	.short	17569                           ; 0x44a1
	.short	17551                           ; 0x448f
	.short	17532                           ; 0x447c
	.short	17514                           ; 0x446a
	.short	17496                           ; 0x4458
	.short	17477                           ; 0x4445
	.short	17459                           ; 0x4433
	.short	17441                           ; 0x4421
	.short	17423                           ; 0x440f
	.short	17405                           ; 0x43fd
	.short	17387                           ; 0x43eb
	.short	17369                           ; 0x43d9
	.short	17351                           ; 0x43c7
	.short	17333                           ; 0x43b5
	.short	17315                           ; 0x43a3
	.short	17297                           ; 0x4391
	.short	17279                           ; 0x437f
	.short	17261                           ; 0x436d
	.short	17244                           ; 0x435c
	.short	17226                           ; 0x434a
	.short	17208                           ; 0x4338
	.short	17191                           ; 0x4327
	.short	17173                           ; 0x4315
	.short	17155                           ; 0x4303
	.short	17138                           ; 0x42f2
	.short	17120                           ; 0x42e0
	.short	17103                           ; 0x42cf
	.short	17085                           ; 0x42bd
	.short	17068                           ; 0x42ac
	.short	17051                           ; 0x429b
	.short	17033                           ; 0x4289
	.short	17016                           ; 0x4278
	.short	16999                           ; 0x4267
	.short	16982                           ; 0x4256
	.short	16964                           ; 0x4244
	.short	16947                           ; 0x4233
	.short	16930                           ; 0x4222
	.short	16913                           ; 0x4211
	.short	16896                           ; 0x4200
	.short	16879                           ; 0x41ef
	.short	16862                           ; 0x41de
	.short	16845                           ; 0x41cd
	.short	16828                           ; 0x41bc
	.short	16811                           ; 0x41ab
	.short	16794                           ; 0x419a
	.short	16778                           ; 0x418a
	.short	16761                           ; 0x4179
	.short	16744                           ; 0x4168
	.short	16727                           ; 0x4157
	.short	16711                           ; 0x4147
	.short	16694                           ; 0x4136
	.short	16677                           ; 0x4125
	.short	16661                           ; 0x4115
	.short	16644                           ; 0x4104
	.short	16628                           ; 0x40f4
	.short	16611                           ; 0x40e3
	.short	16595                           ; 0x40d3
	.short	16578                           ; 0x40c2
	.short	16562                           ; 0x40b2
	.short	16546                           ; 0x40a2
	.short	16529                           ; 0x4091
	.short	16513                           ; 0x4081
	.short	16497                           ; 0x4071
	.short	16481                           ; 0x4061
	.short	16464                           ; 0x4050
	.short	16448                           ; 0x4040
	.short	16432                           ; 0x4030
	.short	16416                           ; 0x4020
	.short	16400                           ; 0x4010
	.short	16384                           ; 0x4000
	.size	RECIPROCAL_TABLE_32, 1024

	.protected	RECIPROCAL_TABLE_64     ; @RECIPROCAL_TABLE_64
	.type	RECIPROCAL_TABLE_64,@object
	.globl	RECIPROCAL_TABLE_64
	.p2align	4, 0x0
RECIPROCAL_TABLE_64:
	.short	2045                            ; 0x7fd
	.short	2037                            ; 0x7f5
	.short	2029                            ; 0x7ed
	.short	2021                            ; 0x7e5
	.short	2013                            ; 0x7dd
	.short	2005                            ; 0x7d5
	.short	1998                            ; 0x7ce
	.short	1990                            ; 0x7c6
	.short	1983                            ; 0x7bf
	.short	1975                            ; 0x7b7
	.short	1968                            ; 0x7b0
	.short	1960                            ; 0x7a8
	.short	1953                            ; 0x7a1
	.short	1946                            ; 0x79a
	.short	1938                            ; 0x792
	.short	1931                            ; 0x78b
	.short	1924                            ; 0x784
	.short	1917                            ; 0x77d
	.short	1910                            ; 0x776
	.short	1903                            ; 0x76f
	.short	1896                            ; 0x768
	.short	1889                            ; 0x761
	.short	1883                            ; 0x75b
	.short	1876                            ; 0x754
	.short	1869                            ; 0x74d
	.short	1863                            ; 0x747
	.short	1856                            ; 0x740
	.short	1849                            ; 0x739
	.short	1843                            ; 0x733
	.short	1836                            ; 0x72c
	.short	1830                            ; 0x726
	.short	1824                            ; 0x720
	.short	1817                            ; 0x719
	.short	1811                            ; 0x713
	.short	1805                            ; 0x70d
	.short	1799                            ; 0x707
	.short	1792                            ; 0x700
	.short	1786                            ; 0x6fa
	.short	1780                            ; 0x6f4
	.short	1774                            ; 0x6ee
	.short	1768                            ; 0x6e8
	.short	1762                            ; 0x6e2
	.short	1756                            ; 0x6dc
	.short	1750                            ; 0x6d6
	.short	1745                            ; 0x6d1
	.short	1739                            ; 0x6cb
	.short	1733                            ; 0x6c5
	.short	1727                            ; 0x6bf
	.short	1722                            ; 0x6ba
	.short	1716                            ; 0x6b4
	.short	1710                            ; 0x6ae
	.short	1705                            ; 0x6a9
	.short	1699                            ; 0x6a3
	.short	1694                            ; 0x69e
	.short	1688                            ; 0x698
	.short	1683                            ; 0x693
	.short	1677                            ; 0x68d
	.short	1672                            ; 0x688
	.short	1667                            ; 0x683
	.short	1661                            ; 0x67d
	.short	1656                            ; 0x678
	.short	1651                            ; 0x673
	.short	1646                            ; 0x66e
	.short	1641                            ; 0x669
	.short	1636                            ; 0x664
	.short	1630                            ; 0x65e
	.short	1625                            ; 0x659
	.short	1620                            ; 0x654
	.short	1615                            ; 0x64f
	.short	1610                            ; 0x64a
	.short	1605                            ; 0x645
	.short	1600                            ; 0x640
	.short	1596                            ; 0x63c
	.short	1591                            ; 0x637
	.short	1586                            ; 0x632
	.short	1581                            ; 0x62d
	.short	1576                            ; 0x628
	.short	1572                            ; 0x624
	.short	1567                            ; 0x61f
	.short	1562                            ; 0x61a
	.short	1558                            ; 0x616
	.short	1553                            ; 0x611
	.short	1548                            ; 0x60c
	.short	1544                            ; 0x608
	.short	1539                            ; 0x603
	.short	1535                            ; 0x5ff
	.short	1530                            ; 0x5fa
	.short	1526                            ; 0x5f6
	.short	1521                            ; 0x5f1
	.short	1517                            ; 0x5ed
	.short	1513                            ; 0x5e9
	.short	1508                            ; 0x5e4
	.short	1504                            ; 0x5e0
	.short	1500                            ; 0x5dc
	.short	1495                            ; 0x5d7
	.short	1491                            ; 0x5d3
	.short	1487                            ; 0x5cf
	.short	1483                            ; 0x5cb
	.short	1478                            ; 0x5c6
	.short	1474                            ; 0x5c2
	.short	1470                            ; 0x5be
	.short	1466                            ; 0x5ba
	.short	1462                            ; 0x5b6
	.short	1458                            ; 0x5b2
	.short	1454                            ; 0x5ae
	.short	1450                            ; 0x5aa
	.short	1446                            ; 0x5a6
	.short	1442                            ; 0x5a2
	.short	1438                            ; 0x59e
	.short	1434                            ; 0x59a
	.short	1430                            ; 0x596
	.short	1426                            ; 0x592
	.short	1422                            ; 0x58e
	.short	1418                            ; 0x58a
	.short	1414                            ; 0x586
	.short	1411                            ; 0x583
	.short	1407                            ; 0x57f
	.short	1403                            ; 0x57b
	.short	1399                            ; 0x577
	.short	1396                            ; 0x574
	.short	1392                            ; 0x570
	.short	1388                            ; 0x56c
	.short	1384                            ; 0x568
	.short	1381                            ; 0x565
	.short	1377                            ; 0x561
	.short	1374                            ; 0x55e
	.short	1370                            ; 0x55a
	.short	1366                            ; 0x556
	.short	1363                            ; 0x553
	.short	1359                            ; 0x54f
	.short	1356                            ; 0x54c
	.short	1352                            ; 0x548
	.short	1349                            ; 0x545
	.short	1345                            ; 0x541
	.short	1342                            ; 0x53e
	.short	1338                            ; 0x53a
	.short	1335                            ; 0x537
	.short	1332                            ; 0x534
	.short	1328                            ; 0x530
	.short	1325                            ; 0x52d
	.short	1322                            ; 0x52a
	.short	1318                            ; 0x526
	.short	1315                            ; 0x523
	.short	1312                            ; 0x520
	.short	1308                            ; 0x51c
	.short	1305                            ; 0x519
	.short	1302                            ; 0x516
	.short	1299                            ; 0x513
	.short	1295                            ; 0x50f
	.short	1292                            ; 0x50c
	.short	1289                            ; 0x509
	.short	1286                            ; 0x506
	.short	1283                            ; 0x503
	.short	1280                            ; 0x500
	.short	1276                            ; 0x4fc
	.short	1273                            ; 0x4f9
	.short	1270                            ; 0x4f6
	.short	1267                            ; 0x4f3
	.short	1264                            ; 0x4f0
	.short	1261                            ; 0x4ed
	.short	1258                            ; 0x4ea
	.short	1255                            ; 0x4e7
	.short	1252                            ; 0x4e4
	.short	1249                            ; 0x4e1
	.short	1246                            ; 0x4de
	.short	1243                            ; 0x4db
	.short	1240                            ; 0x4d8
	.short	1237                            ; 0x4d5
	.short	1234                            ; 0x4d2
	.short	1231                            ; 0x4cf
	.short	1228                            ; 0x4cc
	.short	1226                            ; 0x4ca
	.short	1223                            ; 0x4c7
	.short	1220                            ; 0x4c4
	.short	1217                            ; 0x4c1
	.short	1214                            ; 0x4be
	.short	1211                            ; 0x4bb
	.short	1209                            ; 0x4b9
	.short	1206                            ; 0x4b6
	.short	1203                            ; 0x4b3
	.short	1200                            ; 0x4b0
	.short	1197                            ; 0x4ad
	.short	1195                            ; 0x4ab
	.short	1192                            ; 0x4a8
	.short	1189                            ; 0x4a5
	.short	1187                            ; 0x4a3
	.short	1184                            ; 0x4a0
	.short	1181                            ; 0x49d
	.short	1179                            ; 0x49b
	.short	1176                            ; 0x498
	.short	1173                            ; 0x495
	.short	1171                            ; 0x493
	.short	1168                            ; 0x490
	.short	1165                            ; 0x48d
	.short	1163                            ; 0x48b
	.short	1160                            ; 0x488
	.short	1158                            ; 0x486
	.short	1155                            ; 0x483
	.short	1153                            ; 0x481
	.short	1150                            ; 0x47e
	.short	1148                            ; 0x47c
	.short	1145                            ; 0x479
	.short	1143                            ; 0x477
	.short	1140                            ; 0x474
	.short	1138                            ; 0x472
	.short	1135                            ; 0x46f
	.short	1133                            ; 0x46d
	.short	1130                            ; 0x46a
	.short	1128                            ; 0x468
	.short	1125                            ; 0x465
	.short	1123                            ; 0x463
	.short	1121                            ; 0x461
	.short	1118                            ; 0x45e
	.short	1116                            ; 0x45c
	.short	1113                            ; 0x459
	.short	1111                            ; 0x457
	.short	1109                            ; 0x455
	.short	1106                            ; 0x452
	.short	1104                            ; 0x450
	.short	1102                            ; 0x44e
	.short	1099                            ; 0x44b
	.short	1097                            ; 0x449
	.short	1095                            ; 0x447
	.short	1092                            ; 0x444
	.short	1090                            ; 0x442
	.short	1088                            ; 0x440
	.short	1086                            ; 0x43e
	.short	1083                            ; 0x43b
	.short	1081                            ; 0x439
	.short	1079                            ; 0x437
	.short	1077                            ; 0x435
	.short	1074                            ; 0x432
	.short	1072                            ; 0x430
	.short	1070                            ; 0x42e
	.short	1068                            ; 0x42c
	.short	1066                            ; 0x42a
	.short	1064                            ; 0x428
	.short	1061                            ; 0x425
	.short	1059                            ; 0x423
	.short	1057                            ; 0x421
	.short	1055                            ; 0x41f
	.short	1053                            ; 0x41d
	.short	1051                            ; 0x41b
	.short	1049                            ; 0x419
	.short	1047                            ; 0x417
	.short	1044                            ; 0x414
	.short	1042                            ; 0x412
	.short	1040                            ; 0x410
	.short	1038                            ; 0x40e
	.short	1036                            ; 0x40c
	.short	1034                            ; 0x40a
	.short	1032                            ; 0x408
	.short	1030                            ; 0x406
	.short	1028                            ; 0x404
	.short	1026                            ; 0x402
	.short	1024                            ; 0x400
	.size	RECIPROCAL_TABLE_64, 512

	.protected	BYTES_TO_CLNW_WINDOW_SIZE_TABLE ; @BYTES_TO_CLNW_WINDOW_SIZE_TABLE
	.type	BYTES_TO_CLNW_WINDOW_SIZE_TABLE,@object
	.globl	BYTES_TO_CLNW_WINDOW_SIZE_TABLE
	.p2align	4, 0x0
BYTES_TO_CLNW_WINDOW_SIZE_TABLE:
	.long	4294967295                      ; 0xffffffff
	.long	4                               ; 0x4
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	5                               ; 0x5
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	6                               ; 0x6
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.long	7                               ; 0x7
	.size	BYTES_TO_CLNW_WINDOW_SIZE_TABLE, 68

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
