1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
|
code for sm_100a
.target sm_100a
Function : _Z10div_kernelPKfS0_Pfi
.headerflags @"EF_CUDA_ACCELERATORS EF_CUDA_SM100 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM100)"
//////////
// 线程绑定
//////////
/*0000*/ LDC R1, c[0x0][0x37c] ; /* 0x0000df00ff017b82 */ // 从常量内存段加载参数列表指针，取栈帧/参数基址
/* 0x000fe20000000800 */
/*0010*/ S2R R3, SR_TID.X ; /* 0x0000000000037919 */ // 读入 threadIdx.x，准备计算线程线性索引
/* 0x000e2e0000002100 */
/*0020*/ S2UR UR4, SR_CTAID.X ; /* 0x00000000000479c3 */ // 将 blockIdx.x 读入 uniform 寄存器，用于后续 IMAD
/* 0x000e220000002500 */
/*0030*/ LDCU UR5, c[0x0][0x398] ; /* 0x00007300ff0577ac */ // 读参数 n到 uniform 寄存器 UR5
/* 0x000e6e0008000800 */
/*0040*/ LDC R2, c[0x0][0x360] ; /* 0x0000d800ff027b82 */ // 读 blockDim.x
/* 0x000e240000000800 */
/*0050*/ IMAD R2, R2, UR4, R3 ; /* 0x0000000402027c24 */ // 计算 idx = blockIdx.x * blockDim.x + threadIdx.x
/* 0x001fca000f8e0203 */
/*0060*/ ISETP.GE.AND P0, PT, R2, UR5, PT ; /* 0x0000000502007c0c */ // 判断 idx >= n，置位 P0 用于边界检查
/* 0x002fda000bf06270 */
/*0070*/ @P0 EXIT ; /* 0x000000000000094d */ // 若越界则提前退出线程
/* 0x000fea0003800000 */
/////////////////
// div 除法指令展开
/////////////////
/*0080*/ LDC.64 R6, c[0x0][0x388] ; /* 0x0000e200ff067b82 */ // 读分母基地址（device pointer）低/高 64 位
/* 0x000e220000000a00 */
/*0090*/ LDCU.64 UR4, c[0x0][0x358] ; /* 0x00006b00ff0477ac */ // 读取全局内存描述符（全局数据句柄）到 uniform 寄存器
/* 0x000e6e0008000a00 */
/*00a0*/ LDC.64 R4, c[0x0][0x380] ; /* 0x0000e000ff047b82 */ // 读分子基地址（device pointer）
/* 0x000ea20000000a00 */
/*00b0*/ IMAD.WIDE R6, R2, 0x4, R6 ; /* 0x0000000402067825 */ // 计算 denominator[idx] 的字节地址：base + idx*4
/* 0x001fca00078e0206 */
/*00c0*/ LDG.E.CONSTANT R3, desc[UR4][R6.64] ; /* 0x0000000406037981 */ // dnum: 通过描述符从全局内存读取分母值（缓存策略 constant）
/* 0x002ee2000c1e9900 */
/*00d0*/ IMAD.WIDE R4, R2, 0x4, R4 ; /* 0x0000000402047825 */ // 乘累加R4 = R2 * 0x4 + R4，得到 numerator[idx] 的地址
/* 0x004fca00078e0204 */
/*00e0*/ LDG.E.CONSTANT R0, desc[UR4][R4.64] ; /* 0x0000000404007981 */ // num: 读取分子值
/* 0x000ea2000c1e9900 */
// 【收敛同步点 B0】 用 SFU 单元求解倒数 + 1 次牛顿迭代
/*00f0*/ BSSY.RECONVERGENT B0, 0x1c0 ; /* 0x000000c000007945 */
/* 0x000fe20003800200 */
// 倒数初值 R8: 1/dnum
/*0100*/ MUFU.RCP R8, R3 ; /* 0x0000000300087308 */ // 调用 SFU单元， 计算分母 dnum 的倒数近似
/* 0x008e300000001000 */
/*0110*/ FCHK P0, R0, R3 ; /* 0x0000000300007302 */ // 检查除法特殊情况（如 0、NaN、Inf），更新谓词 P0 （关联 /*0170*/）
/* 0x004e620000000000 */
// 牛顿迭代
/*0120*/ FFMA R9, -R3, R8, 1 ; /* 0x3f80000003097423 */ // 牛顿迭代第一步：r = 1 - d * r0
/* 0x001fc80000000108 */
/*0130*/ FFMA R9, R8, R9, R8 ; /* 0x0000000908097223 */ // 更新倒数 r = r0 + r0 * r
/* 0x000fc80000000008 */
/*0140*/ FFMA R8, R0, R9, RZ ; /* 0x0000000900087223 */ // 将分子乘改进后的倒数，得到初始商 R8 = RZ + R0 * R9
/* 0x000fc800000000ff */ // RZ 表示加 0
/*0150*/ FFMA R10, -R3, R8, R0 ; /* 0x00000008030a7223 */ // 计算残差 = num - dnum * r0
/* 0x000fc80000000100 */
// 牛顿迭代结束，得到 num/dnum：R9
/*0160*/ FFMA R9, R9, R10, R8 ; /* 0x0000000a09097223 */ // 最终修正后的商输出到 R9
/* 0x000fe20000000008 */
// 若无异常 （P0!=0）, 方法1正常结束，将商返回（jmp 1b0）；否则继续，修正异常
/*0170*/ @!P0 BRA 0x1b0 ; /* 0x00000000000c8947 */
/* 0x002fec0003800000 */
//////////////////////////////////
// 快速路径--转入异常修复（慢速路径入口）
//////////////////////////////////
/*0180*/ MOV R4, 0x1a0 ; /* 0x000001a000047802 */ // 为调用特殊处理子程序准备固定参数
/* 0x000fce0000000f00 */
/*0190*/ CALL.REL.NOINC 0x200 ; /* 0x0000000000187944 */ // 调用 IEEE 754 特殊情况处理例程
/* 0x000fea0003c00000 */
/*01a0*/ IMAD.MOV.U32 R9, RZ, RZ, R0 ; /* 0x000000ffff097224 */ // 将特殊例程返回的结果放入 R9
/* 0x000fce00078e0000 */
///////////////////////////////////////
// 无异常收尾入口，计算结束，退出线程
///////////////////////////////////////
/*01b0*/ BSYNC.RECONVERGENT B0 ; /* 0x0000000000007941 */
/* 0x000fea0003800200 */
/*01c0*/ LDC.64 R4, c[0x0][0x390] ; /* 0x0000e400ff047b82 */ // 读取output输出数组基址
/* 0x000e240000000a00 */
/*01d0*/ IMAD.WIDE R2, R2, 0x4, R4 ; /* 0x0000000402027825 */ // 乘累加R2 = R4 + R2 * 0x4，得到 num[idx] 的地址
/* 0x001fca00078e0204 */
/*01e0*/ STG.E desc[UR4][R2.64], R9 ; /* 0x0000000902007986 */ // 将最终商写回全局内存
/* 0x000fe2000c101904 */
/*01f0*/ EXIT ; /* 0x000000000000794d */ // 内核主路径返回，线程结束
/* 0x000fea0003800000 */
//////////////////////////////////////////////////////////// 慢速路径 ////////////////////////////////////////////////////////////
//////////////////////
// IEEE 754 异常处理入口
//////////////////////
/*0200*/ SHF.R.U32.HI R6, RZ, 0x17, R3.reuse ; /* 0x00000017ff067819 */ // 提取分母指数 R6: 提取高位（移位 0x17 位）
/* 0x100fe20000011603 */
// 【收敛同步点 B1】 特殊处理流程的同步点
/*0210*/ BSSY.RECONVERGENT B1, 0x860 ; /* 0x0000064000017945 */
/* 0x000fe20003800200 */
//////////////////////
// NaN/Inf/零 的组合判断
//////////////////////
/*0220*/ SHF.R.U32.HI R5, RZ, 0x17, R0.reuse ; /* 0x00000017ff057819 */ // 提取分子指数 R5: 提取高位（移位 0x17 位）
/* 0x100fe20000011600 */
/*0230*/ IMAD.MOV.U32 R7, RZ, RZ, R0 ; /* 0x000000ffff077224 */ // 复制分子到 R7，用于后续处理
/* 0x000fe200078e0000 */
/*0240*/ LOP3.LUT R6, R6, 0xff, RZ, 0xc0, !PT ; /* 0x000000ff06067812 */ // 保留命中的指数字节，掩掉符号位
/* 0x000fe200078ec0ff */
/*0250*/ IMAD.MOV.U32 R8, RZ, RZ, R3 ; /* 0x000000ffff087224 */ // 复制分母到 R8
/* 0x000fe200078e0003 */
/*0260*/ LOP3.LUT R5, R5, 0xff, RZ, 0xc0, !PT ; /* 0x000000ff05057812 */ // 同样提取分子指数到 R5
/* 0x000fc600078ec0ff */
/*0270*/ VIADD R11, R6, 0xffffffff ; /* 0xffffffff060b7836 */ // 分母指数 R6 - 1，用于检测 INF/NaN
/* 0x000fe40000000000 */
/*0280*/ VIADD R10, R5, 0xffffffff ; /* 0xffffffff050a7836 */ // 分子指数 R5 -1
/* 0x000fc60000000000 */
/*0290*/ ISETP.GT.U32.AND P0, PT, R11, 0xfd, PT ; /* 0x000000fd0b00780c */ // 更新谓词PT：判断分母指数是否 > 253（表示 INF/NaN）
/* 0x000fc80003f04070 */
/*02a0*/ ISETP.GT.U32.OR P0, PT, R10, 0xfd, P0 ; /* 0x000000fd0a00780c */ // 更新谓词P0：将分子指数检查合并
/* 0x000fda0000704470 */
/*02b0*/ @!P0 IMAD.MOV.U32 R9, RZ, RZ, RZ ; /* 0x000000ffff098224 */ // 若无 INF/NaN，初始化结果寄存器为 0，后续回退
/* 0x000fe200078e00ff */
// 若无 INF/NaN，跳过复杂分支，跳到归一修正流程
/*02c0*/ @!P0 BRA 0x440 ; /* 0x00000000005c8947 */
/* 0x000fec0003800000 */
// 检查分子无穷大
/*02d0*/ FSETP.GTU.FTZ.AND P0, PT, |R0|, +INF , PT ; /* 0x7f8000000000780b */ // 检查分子是否为 NaN/Inf，更新 P0
/* 0x000fe40003f1c200 */ //
// 检查分母无穷大
/*02e0*/ FSETP.GTU.FTZ.AND P1, PT, |R3|, +INF , PT ; /* 0x7f8000000300780b */ // 检查分母是否为 NaN/Inf
/* 0x000fc80003f3c200 */ //
/*02f0*/ PLOP3.LUT P0, PT, P0, P1, PT, 0xf8, 0x8f ; /* 0x00000000008f781c */ // 组合 NaN/Inf 标志，识别任何异常
/* 0x000fda0000703f70 */ // 逻辑树配置
// 若检测到 NaN/Inf，跳到0x840处理
/*0300*/ @P0 BRA 0x840 ; /* 0x00000004004c0947 */
/* 0x000fea0003800000 */
/*0310*/ LOP3.LUT P0, RZ, R8, 0x7fffffff, R7, 0xc8, !PT ; /* 0x7fffffff08ff7812 */ // 检测分子/分母符号与零的组合（0 与 Inf）
/* 0x000fda000780c807 */
// 如果是 zero/Inf 组合，跳 0x820；否则继续
/*0320*/ @!P0 BRA 0x820 ; /* 0x00000004003c8947 */
/* 0x000fea0003800000 */
/*0330*/ FSETP.NEU.FTZ.AND P2, PT, |R0|.reuse, +INF , PT ; /* 0x7f8000000000780b */ // 判断分子是否非 Inf
/* 0x040fe40003f5d200 */
/*0340*/ FSETP.NEU.FTZ.AND P1, PT, |R3|, +INF , PT ; /* 0x7f8000000300780b */ // 判断分母是否非 Inf
/* 0x000fe40003f3d200 */
/*0350*/ FSETP.NEU.FTZ.AND P0, PT, |R0|, +INF , PT ; /* 0x7f8000000000780b */ // 再次检测分子，用于组合逻辑
/* 0x000fd60003f1d200 */
// 双条件分支： 如果【分母是 Inf】 且【分子不是 Inf】，跳 0x820
/*0360*/ @!P1 BRA !P2, 0x820 ; /* 0x00000004002c9947 */
/* 0x000fea0005000000 */
/*0370*/ LOP3.LUT P2, RZ, R7, 0x7fffffff, RZ, 0xc0, !PT ; /* 0x7fffffff07ff7812 */ // 检查分子是否为 0（取绝对值判断）
/* 0x000fc8000784c0ff */
/*0380*/ PLOP3.LUT P1, PT, P1, P2, PT, 0x2f, 0xf2 ; /* 0x0000000000f2781c */ // LUT 组合定义: 合成“分母 Inf 且分子 0”的条件
/* 0x000fda0000f24577 */
// 如果分子分母同时0/0 或 Inf/Inf， 跳生成 NaN 的分支（0x800）
/*0390*/ @P1 BRA 0x800 ; /* 0x0000000400181947 */
/* 0x000fea0003800000 */
/*03a0*/ LOP3.LUT P1, RZ, R8, 0x7fffffff, RZ, 0xc0, !PT ; /* 0x7fffffff08ff7812 */ // 检查分母是否为 0
/* 0x000fc8000782c0ff */
/*03b0*/ PLOP3.LUT P0, PT, P0, P1, PT, 0x2f, 0xf2 ; /* 0x0000000000f2781c */ // 组合“【分子非inf】且【分母为 0】”条件
/* 0x000fda0000702577 */
// 分母为 0，跳到生成带符号的无穷结果 （0x7d0）
/*03c0*/ @P0 BRA 0x7d0 ; /* 0x0000000400000947 */
/* 0x000fea0003800000 */
/*03d0*/ ISETP.GE.AND P0, PT, R10, RZ, PT ; /* 0x000000ff0a00720c */ // 判断分子是否为非次正规（指数>=0）
/* 0x000fe40003f06270 */
/*03e0*/ ISETP.GE.AND P1, PT, R11, RZ, PT ; /* 0x000000ff0b00720c */ // 判断分母是否为非次正规
/* 0x000fd60003f26270 */
/*03f0*/ @P0 IMAD.MOV.U32 R9, RZ, RZ, RZ ; /* 0x000000ffff090224 */ // 如果分子已正规化，则重置 R9 为 0 准备后续组合
/* 0x000fe400078e00ff */
/*0400*/ @!P0 FFMA R7, R0, 1.84467440737095516160e+19, RZ ; /* 0x5f80000000078823 */ // 若分子为次正规，乘以 2^64 进行归一化
/* 0x000fe200000000ff */ // 指明使用 FMA 归一化
/*0410*/ @!P0 IMAD.MOV.U32 R9, RZ, RZ, -0x40 ; /* 0xffffffc0ff098424 */ // 累计指数调整量（-64）
/* 0x000fe400078e00ff */
/*0420*/ @!P1 FFMA R8, R3, 1.84467440737095516160e+19, RZ ; /* 0x5f80000003089823 */ // 若分母为次正规，放大归一化
/* 0x000fe400000000ff */
/*0430*/ @!P1 VIADD R9, R9, 0x40 ; /* 0x0000004009099836 */ // 分母归一化时抵消前面的 -64，确保指数差正确
/* 0x000fce0000000000 */
// IEEE 754 归一化分支
/*0440*/ LEA R3, R6, 0xc0800000, 0x17 ; /* 0xc080000006037811 */ // 生成分母尾数（补 1）并左移，将指数/符号打包，合成 IEEE 754 形态
/* 0x000fe200078eb8ff */
/*0450*/ VIADD R5, R5, 0xffffff81 ; /* 0xffffff8105057836 */ // 分子指数减去偏置 127
/* 0x000fe20000000000 */
//【收敛同步点 B2】IEEE 754 舍入流程
/*0460*/ BSSY.RECONVERGENT B2, 0x7c0 ; /* 0x0000035000027945 */ // 设置另一个 reconvergence，用于舍入流程
/* 0x000fe60003800200 */
/*0470*/ IMAD.IADD R3, R8, 0x1, -R3 ; /* 0x0000000108037824 */ // 计算归一化后的分母尾数（补偿符号）
/* 0x000fe200078e0a03 */
/*0480*/ IADD3 R6, PT, PT, R5.reuse, 0x7f, -R6 ; /* 0x0000007f05067810 */ // 组合指数差：numExp - denExp + 偏置
/* 0x040fe20007ffe806 */
/*0490*/ IMAD R0, R5, -0x800000, R7 ; /* 0xff80000005007824 */ // 还原分子尾数的隐含 1，并结合符号
/* 0x000fe400078e0207 */
/*04a0*/ MUFU.RCP R8, R3 ; /* 0x0000000300087308 */ // 用SFU单元，对归一化后的分母尾数取倒数近似
/* 0x000e220000001000 */
/*04b0*/ FADD.FTZ R11, -R3, -RZ ; /* 0x800000ff030b7221 */ // 取反得到 remainder 初值（-denormMant）
/* 0x000fe20000010100 */ // FTZ 表示 flush-to-zero
/*04c0*/ IMAD.IADD R6, R6, 0x1, R9 ; /* 0x0000000106067824 */ // 在指数差上叠加归一化修正项 R9
/* 0x000fc600078e0209 */
// 牛顿迭代
/*04d0*/ FFMA R13, R8, R11, 1 ; /* 0x3f800000080d7423 */ // 牛顿迭代第一步
/* 0x001fc8000000000b */
/*04e0*/ FFMA R10, R8, R13, R8 ; /* 0x0000000d080a7223 */ // 迭代继续，得到更精确的倒数
/* 0x000fc80000000008 */
/*04f0*/ FFMA R7, R0, R10, RZ ; /* 0x0000000a00077223 */ // 计算商尾数（未舍入）
/* 0x000fc800000000ff */
/*0500*/ FFMA R8, R11, R7, R0 ; /* 0x000000070b087223 */ // 计算残差用于舍入
/* 0x000fc80000000000 */
/*0510*/ FFMA R7, R10, R8, R7 ; /* 0x000000080a077223 */ // 二次修正商尾数
/* 0x000fc80000000007 */
/*0520*/ FFMA R8, R11, R7, R0 ; /* 0x000000070b087223 */ // 再次计算残差，准备粘滞位
/* 0x000fc80000000000 */
/*0530*/ FFMA R0, R10, R8, R7 ; /* 0x000000080a007223 */ // 得到最终未打包的商尾数 R0
/* 0x000fca0000000007 */
/*0540*/ SHF.R.U32.HI R3, RZ, 0x17, R0 ; /* 0x00000017ff037819 */ // 提取商尾数的高位指数部分
/* 0x000fc80000011600 */
/*0550*/ LOP3.LUT R3, R3, 0xff, RZ, 0xc0, !PT ; /* 0x000000ff03037812 */ // 获取指数，准备重新组合
/* 0x000fca00078ec0ff */
/*0560*/ IMAD.IADD R9, R3, 0x1, R6 ; /* 0x0000000103097824 */ // 合并指数差，得到最终指数
/* 0x000fc800078e0206 */
/*0570*/ VIADD R3, R9, 0xffffffff ; /* 0xffffffff09037836 */ // 指数-1，用于溢出判定
/* 0x000fca0000000000 */
/*0580*/ ISETP.GE.U32.AND P0, PT, R3, 0xfe, PT ; /* 0x000000fe0300780c */ // 检测指数是否 >= 254（溢出）
/* 0x000fda0003f06070 */
// 指数范围正常，则跳至指数打包 0x7a0
/*0590*/ @!P0 BRA 0x7a0 ; /* 0x0000000000808947 */
/* 0x000fea0003800000 */
/*05a0*/ ISETP.GT.AND P0, PT, R9, 0xfe, PT ; /* 0x000000fe0900780c */ // 判断是否真正溢出
/* 0x000fda0003f04270 */
// 真溢出，则跳到生成 Inf (0x770)
/*05b0*/ @P0 BRA 0x770 ; /* 0x00000000006c0947 */
/* 0x000fea0003800000 */
/*05c0*/ ISETP.GE.AND P0, PT, R9, 0x1, PT ; /* 0x000000010900780c */ // 检查是否仍在正规数范围内
/* 0x000fda0003f06270 */
// 正规数，跳到尾数打包 0x7b0
/*05d0*/ @P0 BRA 0x7b0 ; /* 0x0000000000740947 */
/* 0x000fea0003800000 */
/*05e0*/ ISETP.GE.AND P0, PT, R9, -0x18, PT ; /* 0xffffffe80900780c */ // 判断是否可以通过移位表示为次正规数
/* 0x000fe40003f06270 */
/*05f0*/ LOP3.LUT R0, R0, 0x80000000, RZ, 0xc0, !PT ; /* 0x8000000000007812 */ // 提取商的符号位
/* 0x000fd600078ec0ff */
// 若无法表示次正规（即下溢到 0），跳到尾数打包 0x7b0
/*0600*/ @!P0 BRA 0x7b0 ; /* 0x0000000000688947 */
/* 0x000fea0003800000 */
/*0610*/ FFMA.RZ R3, R10.reuse, R8.reuse, R7.reuse ; /* 0x000000080a037223 */ // 计算舍入残差，判断粘滞位
/* 0x1c0fe2000000c007 */
/*0620*/ FFMA.RM R6, R10.reuse, R8.reuse, R7.reuse ; /* 0x000000080a067223 */ // 取向下舍入结果
/* 0x1c0fe20000004007 */
/*0630*/ ISETP.NE.AND P2, PT, R9.reuse, RZ, PT ; /* 0x000000ff0900720c */ // 判断分子指数是否非零，用于判断是否正规数
/* 0x040fe40003f45270 */
/*0640*/ ISETP.NE.AND P1, PT, R9, RZ, PT ; /* 0x000000ff0900720c */ // 分母
/* 0x000fe40003f25270 */
/*0650*/ LOP3.LUT R5, R3, 0x7fffff, RZ, 0xc0, !PT ; /* 0x007fffff03057812 */ // 取出尾数低 23 位
/* 0x000fe200078ec0ff */
/*0660*/ FFMA.RP R3, R10, R8, R7 ; /* 0x000000080a037223 */ // 向上舍入版本
/* 0x000fe20000008007 */
/*0670*/ VIADD R8, R9, 0x20 ; /* 0x0000002009087836 */ // 计算次正规右移量（24 + 指数）
/* 0x000fe40000000000 */
/*0680*/ LOP3.LUT R5, R5, 0x800000, RZ, 0xfc, !PT ; /* 0x0080000005057812 */ // 补回隐含的最高位 1
/* 0x000fe200078efcff */
/*0690*/ IMAD.MOV R7, RZ, RZ, -R9 ; /* 0x000000ffff077224 */ // 计算负指数，用于移位
/* 0x000fe200078e0a09 */
/*06a0*/ FSETP.NEU.FTZ.AND P0, PT, R3, R6, PT ; /* 0x000000060300720b */ // 比较不同舍入结果，判断是否出现 tie
/* 0x000fc40003f1d000 */
/*06b0*/ SHF.L.U32 R8, R5, R8, RZ ; /* 0x0000000805087219 */ // 将尾数左移（构造次正规数位布）
/* 0x000fe400000006ff */
/*06c0*/ SEL R6, R7, RZ, P2 ; /* 0x000000ff07067207 */ // 若仍为正规数，则保持移位量，否则置 0
/* 0x000fe40001000000 */
/*06d0*/ ISETP.NE.AND P1, PT, R8, RZ, P1 ; /* 0x000000ff0800720c */ // 判断移位后是否仍有非零位（粘滞位）
/* 0x000fe40000f25270 */
/*06e0*/ SHF.R.U32.HI R6, RZ, R6, R5 ; /* 0x00000006ff067219 */ // 根据移位量构造右移掩码
/* 0x000fe40000011605 */
/*06f0*/ PLOP3.LUT P0, PT, P0, P1, PT, 0xf8, 0x8f ; /* 0x00000000008f781c */ // 综合 tie 和粘滞位条件，决定是否进位舍入
/* 0x000fe40000703f70 *
/*0700*/ SHF.R.U32.HI R8, RZ, 0x1, R6 ; /* 0x00000001ff087819 */ // 生成舍入增量（右移一位）
/* 0x000fc40000011606 */
/*0710*/ SEL R3, RZ, 0x1, !P0 ; /* 0x00000001ff037807 */ // 若需要进位，则选择 1，否则 0
/* 0x000fc80004000000 */
/*0720*/ LOP3.LUT R3, R3, 0x1, R8, 0xf8, !PT ; /* 0x0000000103037812 */ // 将舍入增量合并进尾数
/* 0x000fc800078ef808 */
/*0730*/ LOP3.LUT R3, R3, R6, RZ, 0xc0, !PT ; /* 0x0000000603037212 */ // 清理多余位，仅保留有效部分
/* 0x000fca00078ec0ff */
/*0740*/ IMAD.IADD R3, R8, 0x1, R3 ; /* 0x0000000108037824 */ // 将进位加到尾数中
/* 0x000fca00078e0203 */
/*0750*/ LOP3.LUT R0, R3, R0, RZ, 0xfc, !PT ; /* 0x0000000003007212 */ // 将符号位与尾数合并
/* 0x000fe200078efcff */
// 跳到尾数打包 0x7b0
/*0760*/ BRA 0x7b0 ; /* 0x0000000000107947 */
/* 0x000fec0003800000 */
////////////
// 溢出到 inf
////////////
/*0770*/ LOP3.LUT R0, R0, 0x80000000, RZ, 0xc0, !PT ; /* 0x8000000000007812 */ // 溢出路径：保留符号位
/* 0x000fc800078ec0ff */
/*0780*/ LOP3.LUT R0, R0, 0x7f800000, RZ, 0xfc, !PT ; /* 0x7f80000000007812 */ // 将指数字段设为全 1，得到 ±Inf
/* 0x000fe200078efcff */
// 跳到尾数打包 0x7b0
/*0790*/ BRA 0x7b0 ; /* 0x0000000000047947 */
/* 0x000fec0003800000 */
/////////
// 指数打包
/////////
/*07a0*/ IMAD R0, R6, 0x800000, R0 ; /* 0x0080000006007824 */ // 正规数路径：将指数偏置并与尾数组合
/* 0x000fce00078e0200 */
/*07b0*/ BSYNC.RECONVERGENT B2 ; /* 0x0000000000027941 */ // 在B2舍入流程结束处同步
/* 0x000fea0003800200 */
// 跳回主流程的收尾部分
/*07c0*/ BRA 0x850 ; /* 0x0000000000207947 */
/* 0x000fea0003800000 */
///////////////////
// 生成带符号的无穷结果
///////////////////
/*07d0*/ LOP3.LUT R0, R8, 0x80000000, R7, 0x48, !PT ; /* 0x8000000008007812 */ // 分母为 0 时：根据输入符号构造 ±Inf
/* 0x000fc800078e4807 */
/*07e0*/ LOP3.LUT R0, R0, 0x7f800000, RZ, 0xfc, !PT ; /* 0x7f80000000007812 */ // 设置指数为全 1
/* 0x000fe200078efcff */
/*07f0*/ BRA 0x850 ; /* 0x0000000000147947 */ // 跳往返回准备
/* 0x000fec0003800000 */
// 0/0 或 Inf/Inf 情况：组合符号位
/*0800*/ LOP3.LUT R0, R8, 0x80000000, R7, 0x48, !PT ; /* 0x8000000008007812 */
/* 0x000fe200078e4807 */
/*0810*/ BRA 0x850 ; /* 0x00000000000c7947 */ // 跳转到共享尾部
/* 0x000fec0003800000 */
/*0820*/ MUFU.RSQ R0, -QNAN ; /* 0xffc0000000007908 */ // 通过对 QNaN 取倒数平方根生成 NaN
/* 0x000e220000001400 */
// 跳回统一收尾
/*0830*/ BRA 0x850 ; /* 0x0000000000047947 */
/* 0x000fea0003800000 */
/*0840*/ FADD.FTZ R0, R0, R3 ; /* 0x0000000300007221 */ // 将特殊情况构造的值与保留尾数叠加（保持符号）
/* 0x000fce0000010000 *
/////////
// 程序结束
/////////
/*0850*/ BSYNC.RECONVERGENT B1 ; /* 0x0000000000017941 */ // 子程序尾部 reconvergence，与主流程同步
/* 0x000fea0003800200 */
/*0860*/ IMAD.MOV.U32 R5, RZ, RZ, 0x0 ; /* 0x00000000ff057424 */ // 将返回值寄存器清零
/* 0x000fc800078e00ff */
/*0870*/ RET.REL.NODEC R4 0x0 ; /* 0xfffffff404e07950 */ // 返回主程序
/* 0x001fea0003c3ffff */
/*0880*/ BRA 0x880; /* 0xfffffffc00fc7947 */
/* 0x000fc0000383ffff */
/*0890*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*08a0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*08b0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*08c0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*08d0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*08e0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*08f0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0900*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0910*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0920*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0930*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0940*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0950*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0960*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
/*0970*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
..........
|
Albresky's Blog
支付宝
微信