SHARE
TWEET

Untitled

a guest Dec 10th, 2017 68 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
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy. OK, I Understand
 
Top