Guest User

Untitled

a guest
Dec 10th, 2017
80
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. AMD Kernel Code for &__OpenCL_search_kernel:
  2. amd_kernel_code_version_major = 1
  3. amd_kernel_code_version_minor = 1
  4. amd_machine_kind = AMDGPU
  5. amd_machine_version_major = 0
  6. amd_machine_version_minor = 0
  7. amd_machine_version_stepping = 0
  8. kernel_code_entry_byte_offset = 256
  9. max_scratch_backing_memory_byte_size = 0
  10. COMPUTE_PGM_RSRC1 (0xac02d800):
  11. granulated_workitem_vgpr_count = 24
  12. granulated_wavefront_sgpr_count = 11
  13. priority = 0
  14. float_round_mode_32 = NEAREST_EVEN
  15. float_round_mode_16_64 = NEAREST_EVEN
  16. float_denorm_mode_32 = FLUSH_SOURCE_OUTPUT
  17. float_denorm_mode_16_64 = FLUSH_NONE
  18. enable_dx10_clamp = TRUE
  19. enable_ieee_mode = TRUE
  20. COMPUTE_PGM_RSRC2 (0x8c000000):
  21. user_sgpr_count = 6
  22. enable_sgpr_workgroup_id_x = TRUE
  23. enable_vgpr_workitem_id = X
  24. granulated_lds_size = 0
  25. KERNEL_CODE_PROPERTIES (0xa0009000):
  26. enable_sgpr_private_segment_buffer = TRUE
  27. enable_sgpr_kernarg_segment_ptr = TRUE
  28. private_element_size = DWORD (4 bytes)
  29. is_ptr64 = TRUE
  30. kernarg_segment_byte_size = 112
  31. wavefront_sgpr_count = 95
  32. workitem_vgpr_count = 99
  33. kernarg_segment_alignment = 16 (4)
  34. group_segment_alignment = 16 (4)
  35. private_segment_alignment = 16 (4)
  36. wavefront_size = 64 (6)
  37. CONTROL_DIRECTIVES:
  38. required_workgroup_size = (64, 1, 1)
  39. require_no_partial_workgroups = TRUE
  40.  
  41. Disassembly for &__OpenCL_search_kernel:
  42. asic(CI)
  43. type(CS)
  44.  
  45. //
  46. // &__OpenCL_search_kernel:
  47. //
  48. s_lshl_b32 s6, s6, 6 // 000000000100: 8F068606
  49. v_add_i32 v0, vcc, s6, v0 // 000000000104: 4A000006
  50. s_load_dwordx2 s[8:9], s[4:5], 0x00 // 000000000108: C0440500
  51. s_mov_b32 s11, 0x55555555 // 00000000010C: BE8B03FF 55555555
  52. s_mov_b32 s10, 0x55555555 // 000000000114: BE8A03FF 55555555
  53. s_load_dwordx8 s[12:19], s[10:11], 0x00 // 00000000011C: C0C60B00
  54. s_waitcnt lgkmcnt(0) // 000000000120: BF8C007F
  55. v_add_i32 v0, vcc, s8, v0 // 000000000124: 4A000008
  56. s_load_dwordx2 s[8:9], s[4:5], 0x16 // 000000000128: C0440516
  57. s_load_dword s6, s[4:5], 0x18 // 00000000012C: C0030518
  58. v_lshrrev_b32 v1, 8, v0 // 000000000130: 2C020088
  59. v_lshlrev_b32 v2, 8, v0 // 000000000134: 34040088
  60. s_movk_i32 s20, 0x00ff // 000000000138: B01400FF
  61. v_bfi_b32 v1, s20, v1, v2 // 00000000013C: D2940001 040A0214
  62. v_lshlrev_b32 v1, 8, v1 // 000000000144: 34020288
  63. v_lshrrev_b32 v2, 16, v0 // 000000000148: 2C040090
  64. v_bfi_b32 v1, s20, v2, v1 // 00000000014C: D2940001 04060414
  65. v_lshlrev_b32 v1, 8, v1 // 000000000154: 34020288
  66. v_lshrrev_b32 v2, 24, v0 // 000000000158: 2C040098
  67. s_lshl_b32 s21, s18, 2 // 00000000015C: 8F158212
  68. s_lshl_b32 s22, s14, 2 // 000000000160: 8F16820E
  69. s_lshl_b32 s23, s12, 2 // 000000000164: 8F17820C
  70. s_lshl_b32 s24, s16, 2 // 000000000168: 8F188210
  71. v_bfi_b32 v1, s20, v2, v1 // 00000000016C: D2940001 04060414
  72. s_waitcnt lgkmcnt(0) // 000000000174: BF8C007F
  73. v_mov_b32 v83, s8 // 000000000178: 7EA60208
  74. v_mov_b32 v84, s9 // 00000000017C: 7EA80209
  75. v_mov_b32 v85, s6 // 000000000180: 7EAA0206
  76. v_mov_b32 v86, v1 // 000000000184: 7EAC0301
  77. v_mov_b32 v87, 0x80000000 // 000000000188: 7EAE02FF 80000000
  78. v_mov_b32 v88, 0 // 000000000190: 7EB00280
  79. v_mov_b32 v89, 0 // 000000000194: 7EB20280
  80. v_mov_b32 v90, 0 // 000000000198: 7EB40280
  81. v_mov_b32 v91, 0 // 00000000019C: 7EB60280
  82. v_mov_b32 v92, 0 // 0000000001A0: 7EB80280
  83. v_mov_b32 v93, 0 // 0000000001A4: 7EBA0280
  84. v_mov_b32 v94, 0 // 0000000001A8: 7EBC0280
  85. v_mov_b32 v95, 0 // 0000000001AC: 7EBE0280
  86. v_mov_b32 v96, 1 // 0000000001B0: 7EC00281
  87. v_mov_b32 v97, 0 // 0000000001B4: 7EC20280
  88. v_mov_b32 v98, 0x00000280 // 0000000001B8: 7EC402FF 00000280
  89. s_lshr_b32 s92, s22, 2 // 0000000001C0: 905C8216
  90. s_mov_b32 m0, s92 // 0000000001C4: BEFC035C
  91. v_movrels_b32 v1, v83 // 0000000001C8: 7E028753
  92. s_lshr_b32 s92, s21, 2 // 0000000001CC: 905C8215
  93. s_mov_b32 m0, s92 // 0000000001D0: BEFC035C
  94. v_movrels_b32 v2, v83 // 0000000001D4: 7E048753
  95. s_lshr_b32 s92, s24, 2 // 0000000001D8: 905C8218
  96. s_mov_b32 m0, s92 // 0000000001DC: BEFC035C
  97. v_movrels_b32 v3, v83 // 0000000001E0: 7E068753
  98. s_lshr_b32 s92, s23, 2 // 0000000001E4: 905C8217
  99. s_mov_b32 m0, s92 // 0000000001E8: BEFC035C
  100. v_movrels_b32 v4, v83 // 0000000001EC: 7E088753
  101. s_lshl_b32 s6, s19, 2 // 0000000001F0: 8F068213
  102. s_lshl_b32 s8, s17, 2 // 0000000001F4: 8F088211
  103. s_lshl_b32 s9, s15, 2 // 0000000001F8: 8F09820F
  104. s_lshl_b32 s20, s13, 2 // 0000000001FC: 8F14820D
  105. s_load_dwordx8 s[24:31], s[10:11], 0x08 // 000000000200: C0CC0B08
  106. s_lshr_b32 s92, s6, 2 // 000000000204: 905C8206
  107. s_mov_b32 m0, s92 // 000000000208: BEFC035C
  108. v_movrels_b32 v5, v83 // 00000000020C: 7E0A8753
  109. s_lshr_b32 s92, s8, 2 // 000000000210: 905C8208
  110. s_mov_b32 m0, s92 // 000000000214: BEFC035C
  111. v_movrels_b32 v6, v83 // 000000000218: 7E0C8753
  112. s_lshr_b32 s92, s9, 2 // 00000000021C: 905C8209
  113. s_mov_b32 m0, s92 // 000000000220: BEFC035C
  114. v_movrels_b32 v7, v83 // 000000000224: 7E0E8753
  115. s_lshr_b32 s92, s20, 2 // 000000000228: 905C8214
  116. s_mov_b32 m0, s92 // 00000000022C: BEFC035C
  117. v_movrels_b32 v8, v83 // 000000000230: 7E108753
  118. s_waitcnt lgkmcnt(0) // 000000000234: BF8C007F
  119. s_lshl_b32 s6, s28, 2 // 000000000238: 8F06821C
  120. s_lshl_b32 s8, s30, 2 // 00000000023C: 8F08821E
  121. s_lshl_b32 s9, s24, 2 // 000000000240: 8F098218
  122. s_lshl_b32 s20, s26, 2 // 000000000244: 8F14821A
  123. s_lshr_b32 s92, s6, 2 // 000000000248: 905C8206
  124. s_mov_b32 m0, s92 // 00000000024C: BEFC035C
  125. v_movrels_b32 v9, v83 // 000000000250: 7E128753
  126. s_lshr_b32 s92, s8, 2 // 000000000254: 905C8208
  127. s_mov_b32 m0, s92 // 000000000258: BEFC035C
  128. v_movrels_b32 v10, v83 // 00000000025C: 7E148753
  129. s_lshr_b32 s92, s20, 2 // 000000000260: 905C8214
  130. s_mov_b32 m0, s92 // 000000000264: BEFC035C
  131. v_movrels_b32 v11, v83 // 000000000268: 7E168753
  132. s_lshr_b32 s92, s9, 2 // 00000000026C: 905C8209
  133. s_mov_b32 m0, s92 // 000000000270: BEFC035C
  134. v_movrels_b32 v12, v83 // 000000000274: 7E188753
  135. s_lshl_b32 s6, s25, 2 // 000000000278: 8F068219
  136. s_lshl_b32 s8, s27, 2 // 00000000027C: 8F08821B
  137. s_lshl_b32 s9, s31, 2 // 000000000280: 8F09821F
  138. s_lshl_b32 s20, s29, 2 // 000000000284: 8F14821D
  139. s_load_dwordx8 s[32:39], s[10:11], 0x10 // 000000000288: C0D00B10
  140. s_lshr_b32 s92, s9, 2 // 00000000028C: 905C8209
  141. s_mov_b32 m0, s92 // 000000000290: BEFC035C
  142. v_movrels_b32 v13, v83 // 000000000294: 7E1A8753
  143. s_lshr_b32 s92, s8, 2 // 000000000298: 905C8208
  144. s_mov_b32 m0, s92 // 00000000029C: BEFC035C
  145. v_movrels_b32 v14, v83 // 0000000002A0: 7E1C8753
  146. s_lshr_b32 s92, s6, 2 // 0000000002A4: 905C8206
  147. s_mov_b32 m0, s92 // 0000000002A8: BEFC035C
  148. v_movrels_b32 v15, v83 // 0000000002AC: 7E1E8753
  149. s_lshr_b32 s92, s20, 2 // 0000000002B0: 905C8214
  150. s_mov_b32 m0, s92 // 0000000002B4: BEFC035C
  151. v_movrels_b32 v16, v83 // 0000000002B8: 7E208753
  152. s_waitcnt lgkmcnt(0) // 0000000002BC: BF8C007F
  153. s_lshl_b32 s6, s38, 2 // 0000000002C0: 8F068226
  154. s_lshl_b32 s8, s34, 2 // 0000000002C4: 8F088222
  155. s_lshl_b32 s9, s32, 2 // 0000000002C8: 8F098220
  156. s_lshl_b32 s20, s36, 2 // 0000000002CC: 8F148224
  157. s_lshr_b32 s92, s8, 2 // 0000000002D0: 905C8208
  158. s_mov_b32 m0, s92 // 0000000002D4: BEFC035C
  159. v_movrels_b32 v17, v83 // 0000000002D8: 7E228753
  160. s_lshr_b32 s92, s6, 2 // 0000000002DC: 905C8206
  161. s_mov_b32 m0, s92 // 0000000002E0: BEFC035C
  162. v_movrels_b32 v18, v83 // 0000000002E4: 7E248753
  163. s_lshr_b32 s92, s20, 2 // 0000000002E8: 905C8214
  164. s_mov_b32 m0, s92 // 0000000002EC: BEFC035C
  165. v_movrels_b32 v19, v83 // 0000000002F0: 7E268753
  166. s_lshr_b32 s92, s9, 2 // 0000000002F4: 905C8209
  167. s_mov_b32 m0, s92 // 0000000002F8: BEFC035C
  168. v_movrels_b32 v20, v83 // 0000000002FC: 7E288753
  169. s_lshl_b32 s6, s39, 2 // 000000000300: 8F068227
  170. s_lshl_b32 s8, s37, 2 // 000000000304: 8F088225
  171. s_lshl_b32 s9, s35, 2 // 000000000308: 8F098223
  172. s_lshl_b32 s20, s33, 2 // 00000000030C: 8F148221
  173. s_load_dwordx8 s[40:47], s[10:11], 0x18 // 000000000310: C0D40B18
  174. s_lshr_b32 s92, s6, 2 // 000000000314: 905C8206
  175. s_mov_b32 m0, s92 // 000000000318: BEFC035C
  176. v_movrels_b32 v21, v83 // 00000000031C: 7E2A8753
  177. s_lshr_b32 s92, s8, 2 // 000000000320: 905C8208
  178. s_mov_b32 m0, s92 // 000000000324: BEFC035C
  179. v_movrels_b32 v22, v83 // 000000000328: 7E2C8753
  180. s_lshr_b32 s92, s9, 2 // 00000000032C: 905C8209
  181. s_mov_b32 m0, s92 // 000000000330: BEFC035C
  182. v_movrels_b32 v23, v83 // 000000000334: 7E2E8753
  183. s_lshr_b32 s92, s20, 2 // 000000000338: 905C8214
  184. s_mov_b32 m0, s92 // 00000000033C: BEFC035C
  185. v_movrels_b32 v24, v83 // 000000000340: 7E308753
  186. s_waitcnt lgkmcnt(0) // 000000000344: BF8C007F
  187. s_lshl_b32 s6, s44, 2 // 000000000348: 8F06822C
  188. s_lshl_b32 s8, s46, 2 // 00000000034C: 8F08822E
  189. s_lshl_b32 s9, s40, 2 // 000000000350: 8F098228
  190. s_lshl_b32 s20, s42, 2 // 000000000354: 8F14822A
  191. s_lshr_b32 s92, s6, 2 // 000000000358: 905C8206
  192. s_mov_b32 m0, s92 // 00000000035C: BEFC035C
  193. v_movrels_b32 v25, v83 // 000000000360: 7E328753
  194. s_lshr_b32 s92, s8, 2 // 000000000364: 905C8208
  195. s_mov_b32 m0, s92 // 000000000368: BEFC035C
  196. v_movrels_b32 v26, v83 // 00000000036C: 7E348753
  197. s_lshr_b32 s92, s20, 2 // 000000000370: 905C8214
  198. s_mov_b32 m0, s92 // 000000000374: BEFC035C
  199. v_movrels_b32 v27, v83 // 000000000378: 7E368753
  200. s_lshr_b32 s92, s9, 2 // 00000000037C: 905C8209
  201. s_mov_b32 m0, s92 // 000000000380: BEFC035C
  202. v_movrels_b32 v28, v83 // 000000000384: 7E388753
  203. s_lshl_b32 s6, s41, 2 // 000000000388: 8F068229
  204. s_lshl_b32 s8, s43, 2 // 00000000038C: 8F08822B
  205. s_lshl_b32 s9, s47, 2 // 000000000390: 8F09822F
  206. s_lshl_b32 s20, s45, 2 // 000000000394: 8F14822D
  207. s_load_dwordx8 s[48:55], s[10:11], 0x20 // 000000000398: C0D80B20
  208. s_lshr_b32 s92, s9, 2 // 00000000039C: 905C8209
  209. s_mov_b32 m0, s92 // 0000000003A0: BEFC035C
  210. v_movrels_b32 v29, v83 // 0000000003A4: 7E3A8753
  211. s_lshr_b32 s92, s8, 2 // 0000000003A8: 905C8208
  212. s_mov_b32 m0, s92 // 0000000003AC: BEFC035C
  213. v_movrels_b32 v30, v83 // 0000000003B0: 7E3C8753
  214. s_lshr_b32 s92, s6, 2 // 0000000003B4: 905C8206
  215. s_mov_b32 m0, s92 // 0000000003B8: BEFC035C
  216. v_movrels_b32 v31, v83 // 0000000003BC: 7E3E8753
  217. s_lshr_b32 s92, s20, 2 // 0000000003C0: 905C8214
  218. s_mov_b32 m0, s92 // 0000000003C4: BEFC035C
  219. v_movrels_b32 v32, v83 // 0000000003C8: 7E408753
  220. s_waitcnt lgkmcnt(0) // 0000000003CC: BF8C007F
  221. s_lshl_b32 s6, s54, 2 // 0000000003D0: 8F068236
  222. s_lshl_b32 s8, s50, 2 // 0000000003D4: 8F088232
  223. s_lshl_b32 s9, s48, 2 // 0000000003D8: 8F098230
  224. s_lshl_b32 s20, s52, 2 // 0000000003DC: 8F148234
  225. s_lshr_b32 s92, s8, 2 // 0000000003E0: 905C8208
  226. s_mov_b32 m0, s92 // 0000000003E4: BEFC035C
  227. v_movrels_b32 v33, v83 // 0000000003E8: 7E428753
  228. s_lshr_b32 s92, s6, 2 // 0000000003EC: 905C8206
  229. s_mov_b32 m0, s92 // 0000000003F0: BEFC035C
  230. v_movrels_b32 v34, v83 // 0000000003F4: 7E448753
  231. s_lshr_b32 s92, s20, 2 // 0000000003F8: 905C8214
  232. s_mov_b32 m0, s92 // 0000000003FC: BEFC035C
  233. v_movrels_b32 v35, v83 // 000000000400: 7E468753
  234. s_lshr_b32 s92, s9, 2 // 000000000404: 905C8209
  235. s_mov_b32 m0, s92 // 000000000408: BEFC035C
  236. v_movrels_b32 v36, v83 // 00000000040C: 7E488753
  237. s_lshl_b32 s6, s55, 2 // 000000000410: 8F068237
  238. s_lshl_b32 s8, s53, 2 // 000000000414: 8F088235
  239. s_lshl_b32 s9, s51, 2 // 000000000418: 8F098233
  240. s_lshl_b32 s20, s49, 2 // 00000000041C: 8F148231
  241. s_load_dwordx8 s[56:63], s[10:11], 0x28 // 000000000420: C0DC0B28
  242. s_lshr_b32 s92, s6, 2 // 000000000424: 905C8206
  243. s_mov_b32 m0, s92 // 000000000428: BEFC035C
  244. v_movrels_b32 v37, v83 // 00000000042C: 7E4A8753
  245. s_lshr_b32 s92, s8, 2 // 000000000430: 905C8208
  246. s_mov_b32 m0, s92 // 000000000434: BEFC035C
  247. v_movrels_b32 v38, v83 // 000000000438: 7E4C8753
  248. s_lshr_b32 s92, s9, 2 // 00000000043C: 905C8209
  249. s_mov_b32 m0, s92 // 000000000440: BEFC035C
  250. v_movrels_b32 v39, v83 // 000000000444: 7E4E8753
  251. s_lshr_b32 s92, s20, 2 // 000000000448: 905C8214
  252. s_mov_b32 m0, s92 // 00000000044C: BEFC035C
  253. v_movrels_b32 v40, v83 // 000000000450: 7E508753
  254. s_waitcnt lgkmcnt(0) // 000000000454: BF8C007F
  255. s_lshl_b32 s6, s60, 2 // 000000000458: 8F06823C
  256. s_lshl_b32 s8, s62, 2 // 00000000045C: 8F08823E
  257. s_lshl_b32 s9, s56, 2 // 000000000460: 8F098238
  258. s_lshl_b32 s20, s58, 2 // 000000000464: 8F14823A
  259. s_lshr_b32 s92, s6, 2 // 000000000468: 905C8206
  260. s_mov_b32 m0, s92 // 00000000046C: BEFC035C
  261. v_movrels_b32 v41, v83 // 000000000470: 7E528753
  262. s_lshr_b32 s92, s8, 2 // 000000000474: 905C8208
  263. s_mov_b32 m0, s92 // 000000000478: BEFC035C
  264. v_movrels_b32 v42, v83 // 00000000047C: 7E548753
  265. s_lshr_b32 s92, s20, 2 // 000000000480: 905C8214
  266. s_mov_b32 m0, s92 // 000000000484: BEFC035C
  267. v_movrels_b32 v43, v83 // 000000000488: 7E568753
  268. s_lshr_b32 s92, s9, 2 // 00000000048C: 905C8209
  269. s_mov_b32 m0, s92 // 000000000490: BEFC035C
  270. v_movrels_b32 v44, v83 // 000000000494: 7E588753
  271. s_mov_b32 s8, s19 // 000000000498: BE880313
  272. s_movk_i32 s9, 0x0000 // 00000000049C: B0090000
  273. s_lshl_b64 s[8:9], s[8:9], 2 // 0000000004A0: 8F888208
  274. s_mov_b32 s20, s13 // 0000000004A4: BE94030D
  275. s_movk_i32 s21, 0x0000 // 0000000004A8: B0150000
  276. s_lshl_b64 s[20:21], s[20:21], 2 // 0000000004AC: 8F948214
  277. s_mov_b32 s22, s15 // 0000000004B0: BE96030F
  278. s_movk_i32 s23, 0x0000 // 0000000004B4: B0170000
  279. s_lshl_b64 s[22:23], s[22:23], 2 // 0000000004B8: 8F968216
  280. s_mov_b32 s64, s17 // 0000000004BC: BEC00311
  281. s_movk_i32 s65, 0x0000 // 0000000004C0: B0410000
  282. s_lshl_b64 s[64:65], s[64:65], 2 // 0000000004C4: 8FC08240
  283. s_mov_b32 s67, 0x55555555 // 0000000004C8: BEC303FF 55555555
  284. s_mov_b32 s66, 0x55555555 // 0000000004D0: BEC203FF 55555555
  285. s_lshl_b32 s6, s57, 2 // 0000000004D8: 8F068239
  286. s_lshl_b32 s13, s59, 2 // 0000000004DC: 8F0D823B
  287. s_lshl_b32 s15, s63, 2 // 0000000004E0: 8F0F823F
  288. s_lshl_b32 s17, s61, 2 // 0000000004E4: 8F11823D
  289. s_load_dwordx8 s[68:75], s[10:11], 0x30 // 0000000004E8: C0E20B30
  290. s_add_u32 s20, s66, s20 // 0000000004EC: 80141442
  291. s_addc_u32 s21, s67, s21 // 0000000004F0: 82151543
  292. s_add_u32 s8, s66, s8 // 0000000004F4: 80080842
  293. s_addc_u32 s9, s67, s9 // 0000000004F8: 82090943
  294. s_add_u32 s64, s66, s64 // 0000000004FC: 80404042
  295. s_addc_u32 s65, s67, s65 // 000000000500: 82414143
  296. s_add_u32 s22, s66, s22 // 000000000504: 80161642
  297. s_addc_u32 s23, s67, s23 // 000000000508: 82171743
  298. v_mov_b32 v46, s13 // 00000000050C: 7E5C020D
  299. v_mov_b32 v47, s6 // 000000000510: 7E5E0206
  300. v_mov_b32 v48, s17 // 000000000514: 7E600211
  301. s_load_dwordx8 s[76:83], s[4:5], 0x0c // 000000000518: C0E6050C
  302. s_load_dwordx2 s[4:5], s[4:5], 0x14 // 00000000051C: C0420514
  303. s_load_dword s6, s[20:21], 0x00 // 000000000520: C0031500
  304. s_load_dword s8, s[8:9], 0x00 // 000000000524: C0040900
  305. s_load_dword s9, s[64:65], 0x00 // 000000000528: C004C100
  306. s_load_dword s13, s[22:23], 0x00 // 00000000052C: C0069700
  307. s_lshr_b32 s92, s15, 2 // 000000000530: 905C820F
  308. s_mov_b32 m0, s92 // 000000000534: BEFC035C
  309. v_movrels_b32 v45, v83 // 000000000538: 7E5A8753
  310. v_readfirstlane_b32 s92, v46 // 00000000053C: 7EB8052E
  311. s_lshr_b32 s92, s92, 2 // 000000000540: 905C825C
  312. s_mov_b32 m0, s92 // 000000000544: BEFC035C
  313. v_movrels_b32 v46, v83 // 000000000548: 7E5C8753
  314. v_readfirstlane_b32 s92, v47 // 00000000054C: 7EB8052F
  315. s_lshr_b32 s92, s92, 2 // 000000000550: 905C825C
  316. s_mov_b32 m0, s92 // 000000000554: BEFC035C
  317. v_movrels_b32 v47, v83 // 000000000558: 7E5E8753
  318. v_readfirstlane_b32 s92, v48 // 00000000055C: 7EB80530
  319. s_lshr_b32 s92, s92, 2 // 000000000560: 905C825C
  320. s_mov_b32 m0, s92 // 000000000564: BEFC035C
  321. v_movrels_b32 v48, v83 // 000000000568: 7E608753
  322. s_mov_b32 s20, s12 // 00000000056C: BE94030C
  323. s_movk_i32 s21, 0x0000 // 000000000570: B0150000
  324. s_lshl_b64 s[20:21], s[20:21], 2 // 000000000574: 8F948214
  325. s_movk_i32 s19, 0x0000 // 000000000578: B0130000
  326. s_lshl_b64 s[18:19], s[18:19], 2 // 00000000057C: 8F928212
  327. s_movk_i32 s15, 0x0000 // 000000000580: B00F0000
  328. s_lshl_b64 s[14:15], s[14:15], 2 // 000000000584: 8F8E820E
  329. s_movk_i32 s17, 0x0000 // 000000000588: B0110000
  330. s_lshl_b64 s[16:17], s[16:17], 2 // 00000000058C: 8F908210
  331. s_add_u32 s16, s66, s16 // 000000000590: 80101042
  332. s_addc_u32 s17, s67, s17 // 000000000594: 82111143
  333. s_add_u32 s14, s66, s14 // 000000000598: 800E0E42
  334. s_addc_u32 s15, s67, s15 // 00000000059C: 820F0F43
  335. s_add_u32 s18, s66, s18 // 0000000005A0: 80121242
  336. s_addc_u32 s19, s67, s19 // 0000000005A4: 82131343
  337. s_add_u32 s20, s66, s20 // 0000000005A8: 80141442
  338. s_addc_u32 s21, s67, s21 // 0000000005AC: 82151543
  339. s_waitcnt lgkmcnt(0) // 0000000005B0: BF8C007F
  340. s_lshl_b32 s12, s74, 2 // 0000000005B4: 8F0C824A
  341. s_lshl_b32 s22, s70, 2 // 0000000005B8: 8F168246
  342. s_lshl_b32 s23, s68, 2 // 0000000005BC: 8F178244
  343. s_lshl_b32 s64, s72, 2 // 0000000005C0: 8F408248
  344. v_xor_b32 v3, s9, v3 // 0000000005C4: 3A060609
  345. v_xor_b32 v2, s8, v2 // 0000000005C8: 3A040408
  346. v_xor_b32 v1, s13, v1 // 0000000005CC: 3A02020D
  347. v_xor_b32 v4, s6, v4 // 0000000005D0: 3A080806
  348. s_add_u32 s6, s4, s80 // 0000000005D4: 80065004
  349. s_add_u32 s8, s5, s81 // 0000000005D8: 80085105
  350. s_add_u32 s9, s83, s79 // 0000000005DC: 80094F53
  351. s_add_u32 s13, s82, s78 // 0000000005E0: 800D4E52
  352. v_mov_b32 v50, s12 // 0000000005E4: 7E64020C
  353. v_mov_b32 v51, s64 // 0000000005E8: 7E660240
  354. v_mov_b32 v52, s23 // 0000000005EC: 7E680217
  355. s_load_dword s12, s[16:17], 0x00 // 0000000005F0: C0061100
  356. s_load_dword s14, s[14:15], 0x00 // 0000000005F4: C0070F00
  357. s_load_dword s15, s[18:19], 0x00 // 0000000005F8: C0079300
  358. s_load_dword s16, s[20:21], 0x00 // 0000000005FC: C0081500
  359. v_add_i32 v3, vcc, s6, v3 // 000000000600: 4A060606
  360. v_add_i32 v2, vcc, s8, v2 // 000000000604: 4A040408
  361. v_add_i32 v1, vcc, s9, v1 // 000000000608: 4A020209
  362. v_add_i32 v4, vcc, s13, v4 // 00000000060C: 4A08080D
  363. v_xor_b32 v53, 0x082efa98, v3 // 000000000610: 3A6A06FF 082EFA98
  364. v_xor_b32 v54, 0xec4e6c89, v2 // 000000000618: 3A6C04FF EC4E6C89
  365. v_xor_b32 v55, 0x299f3350, v1 // 000000000620: 3A6E02FF 299F3350
  366. v_xor_b32 v56, 0xa4093aa2, v4 // 000000000628: 3A7008FF A4093AA2
  367. s_lshr_b32 s92, s22, 2 // 000000000630: 905C8216
  368. s_mov_b32 m0, s92 // 000000000634: BEFC035C
  369. v_movrels_b32 v49, v83 // 000000000638: 7E628753
  370. v_readfirstlane_b32 s92, v50 // 00000000063C: 7EB80532
  371. s_lshr_b32 s92, s92, 2 // 000000000640: 905C825C
  372. s_mov_b32 m0, s92 // 000000000644: BEFC035C
  373. v_movrels_b32 v50, v83 // 000000000648: 7E648753
  374. v_readfirstlane_b32 s92, v51 // 00000000064C: 7EB80533
  375. s_lshr_b32 s92, s92, 2 // 000000000650: 905C825C
  376. s_mov_b32 m0, s92 // 000000000654: BEFC035C
  377. v_movrels_b32 v51, v83 // 000000000658: 7E668753
  378. v_readfirstlane_b32 s92, v52 // 00000000065C: 7EB80534
  379. s_lshr_b32 s92, s92, 2 // 000000000660: 905C825C
  380. s_mov_b32 m0, s92 // 000000000664: BEFC035C
  381. v_movrels_b32 v52, v83 // 000000000668: 7E688753
  382. v_alignbyte_b32 v53, v53, v53, 2 // 00000000066C: D29E0035 020A6B35
  383. v_alignbyte_b32 v54, v54, v54, 2 // 000000000674: D29E0036 020A6D36
  384. v_alignbyte_b32 v55, v55, v55, 2 // 00000000067C: D29E0037 020A6F37
  385. v_alignbyte_b32 v56, v56, v56, 2 // 000000000684: D29E0038 020A7138
  386. s_mov_b32 s8, s29 // 00000000068C: BE88031D
  387. s_movk_i32 s9, 0x0000 // 000000000690: B0090000
  388. s_lshl_b64 s[8:9], s[8:9], 2 // 000000000694: 8F888208
  389. s_mov_b32 s18, s31 // 000000000698: BE92031F
  390. s_movk_i32 s19, 0x0000 // 00000000069C: B0130000
  391. s_lshl_b64 s[18:19], s[18:19], 2 // 0000000006A0: 8F928212
  392. s_mov_b32 s20, s27 // 0000000006A4: BE94031B
  393. s_movk_i32 s21, 0x0000 // 0000000006A8: B0150000
  394. s_lshl_b64 s[20:21], s[20:21], 2 // 0000000006AC: 8F948214
  395. s_mov_b32 s22, s25 // 0000000006B0: BE960319
  396. s_movk_i32 s23, 0x0000 // 0000000006B4: B0170000
  397. s_lshl_b64 s[22:23], s[22:23], 2 // 0000000006B8: 8F968216
  398. v_add_i32 v57, vcc, 0x13198a2e, v53 // 0000000006BC: 4A726AFF 13198A2E
  399. v_add_i32 v58, vcc, 0x03707344, v54 // 0000000006C4: 4A746CFF 03707344
  400. v_add_i32 v59, vcc, 0x85a308d3, v55 // 0000000006CC: 4A766EFF 85A308D3
  401. v_add_i32 v60, vcc, 0x243f6a88, v56 // 0000000006D4: 4A7870FF 243F6A88
  402. v_xor_b32 v61, s4, v57 // 0000000006DC: 3A7A7204
  403. v_xor_b32 v62, s5, v58 // 0000000006E0: 3A7C7405
  404. v_xor_b32 v63, s83, v59 // 0000000006E4: 3A7E7653
  405. v_xor_b32 v64, s82, v60 // 0000000006E8: 3A807852
  406. s_waitcnt lgkmcnt(0) // 0000000006EC: BF8C007F
  407. v_xor_b32 v6, s12, v6
RAW Paste Data