Advertisement
Guest User

Untitled

a guest
Aug 28th, 2017
104
0
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
LLVM 25.82 KB | None | 0 0
  1.  
  2. ; Function Attrs: nounwind
  3. define void @__OpenCL_k02g_kernel(i32 addrspace(1)* noalias %g_output, %2 addrspace(2)* %g_header, %3 addrspace(1)* noalias %g_dag, i64 %start_nonce, i64 %target, i32 %isolate, i32 %dagsize, i32 %dm, i32 %dv, i32 %dmask, i32 addrspace(2)* noalias %d_header, i32 %d_ExtraNonce, i32 %DcrPerItem) #0 {
  4. entry:
  5.   %g_output.addr = alloca i32 addrspace(1)*, align 8
  6.   %g_header.addr = alloca %2 addrspace(2)*, align 8
  7.   %g_dag.addr = alloca %3 addrspace(1)*, align 8
  8.   %start_nonce.addr = alloca i64, align 8
  9.   %target.addr = alloca i64, align 8
  10.   %isolate.addr = alloca i32, align 4
  11.   %dagsize.addr = alloca i32, align 4
  12.   %dm.addr = alloca i32, align 4
  13.   %dv.addr = alloca i32, align 4
  14.   %dmask.addr = alloca i32, align 4
  15.   %d_header.addr = alloca i32 addrspace(2)*, align 8
  16.   %d_ExtraNonce.addr = alloca i32, align 4
  17.   %DcrPerItem.addr = alloca i32, align 4
  18.   %loc_arr = alloca [576 x i64] addrspace(3)*, align 8
  19.   %state = alloca [25 x i64], align 8
  20.   %gid = alloca i32, align 4
  21.   %lar = alloca i64 addrspace(3)*, align 8
  22.   %n = alloca i32, align 4
  23.   %dpi = alloca i32, align 4
  24.   %curstep = alloca i32, align 4
  25.   %i = alloca i32, align 4
  26.   %share = alloca %0 addrspace(3)*, align 8
  27.   %i49 = alloca i32, align 4
  28.   %i74 = alloca i32, align 4
  29.   %morebit = alloca i32, align 4
  30.   %thread_id = alloca i32, align 4
  31.   %hash_id = alloca i32, align 4
  32.   %i105 = alloca i32, align 4
  33.   %i117 = alloca i32, align 4
  34.   %mix = alloca <4 x i32>, align 16
  35.   %share0 = alloca i32 addrspace(3)*, align 8
  36.   %init0 = alloca i32, align 4
  37.   %a = alloca i32, align 4
  38.   %update_share = alloca i32, align 4
  39.   %i181 = alloca i32, align 4
  40.   %v = alloca i32, align 4
  41.   %ml = alloca i32, align 4
  42.   %aml = alloca i32, align 4
  43.   %res = alloca i32, align 4
  44.   %i266 = alloca i32, align 4
  45.   %i294 = alloca i32, align 4
  46.   %tmp319 = alloca <8 x i8>, align 8
  47.   %slot = alloca i32, align 4
  48.   store i32 addrspace(1)* %g_output, i32 addrspace(1)** %g_output.addr, align 8
  49.   store %2 addrspace(2)* %g_header, %2 addrspace(2)** %g_header.addr, align 8
  50.   store %3 addrspace(1)* %g_dag, %3 addrspace(1)** %g_dag.addr, align 8
  51.   store i64 %start_nonce, i64* %start_nonce.addr, align 8
  52.   store i64 %target, i64* %target.addr, align 8
  53.   store i32 %isolate, i32* %isolate.addr, align 4
  54.   store i32 %dagsize, i32* %dagsize.addr, align 4
  55.   store i32 %dm, i32* %dm.addr, align 4
  56.   store i32 %dv, i32* %dv.addr, align 4
  57.   store i32 %dmask, i32* %dmask.addr, align 4
  58.   store i32 addrspace(2)* %d_header, i32 addrspace(2)** %d_header.addr, align 8
  59.   store i32 %d_ExtraNonce, i32* %d_ExtraNonce.addr, align 4
  60.   store i32 %DcrPerItem, i32* %DcrPerItem.addr, align 4
  61.   store [576 x i64] addrspace(3)* @k02g_cllocal_loc_arr, [576 x i64] addrspace(3)** %loc_arr, align 8
  62.   call void @barrier(i32 2) #0
  63.   %tmp = load i32 addrspace(1)** %g_output.addr, align 8
  64.   %arrayidx = getelementptr i32 addrspace(1)* %tmp, i64 0
  65.   %tmp1 = load volatile i32 addrspace(1)* %arrayidx, align 4
  66.   %tmp2 = and i32 %tmp1, 15
  67.   %tobool = icmp ne i32 %tmp2, 0
  68.   br i1 %tobool, label %if.then, label %if.end
  69.  
  70. return:                                           ; preds = %if.end341, %for.exit, %if.then
  71.   ret void
  72.  
  73. if.end:                                           ; preds = %entry
  74.   %call = call i64 @get_global_id(i32 0) #0
  75.   %conv = trunc i64 %call to i32
  76.   store i32 %conv, i32* %gid, align 4
  77.   %tmp3 = load [576 x i64] addrspace(3)** %loc_arr, align 8
  78.   %ptr = getelementptr inbounds [576 x i64] addrspace(3)* %tmp3, i32 0, i32 0
  79.   %call4 = call i64 @get_local_id(i32 0) #0
  80.   %tmp5 = mul i64 9, %call4
  81.   %tmp6 = getelementptr i64 addrspace(3)* %ptr, i64 %tmp5
  82.   store i64 addrspace(3)* %tmp6, i64 addrspace(3)** %lar, align 8
  83.   %tmp7 = load i32* %gid, align 4
  84.   %tmp8 = udiv i32 %tmp7, 64
  85.   store i32 %tmp8, i32* %n, align 4
  86.   %tmp9 = load i32* %n, align 4
  87.   %tmp10 = urem i32 %tmp9, 3
  88.   %cmp = icmp eq i32 %tmp10, 0
  89.   br i1 %cmp, label %if.then12, label %if.end11
  90.  
  91. if.then:                                          ; preds = %entry
  92.   br label %return
  93.  
  94. if.end11:                                         ; preds = %if.end
  95.   %tmp46 = load [576 x i64] addrspace(3)** %loc_arr, align 8
  96.   %ptr47 = getelementptr inbounds [576 x i64] addrspace(3)* %tmp46, i32 0, i32 0
  97.   %conv48 = bitcast i64 addrspace(3)* %ptr47 to %0 addrspace(3)*
  98.   store %0 addrspace(3)* %conv48, %0 addrspace(3)** %share, align 8
  99.   store i32 0, i32* %i49, align 4
  100.   br label %for.cond50
  101.  
  102. if.then12:                                        ; preds = %if.end
  103.   %tmp13 = load i32* %DcrPerItem.addr, align 4
  104.   %tmp14 = add i32 %tmp13, 16
  105.   %tmp15 = udiv i32 %tmp14, 16
  106.   store i32 %tmp15, i32* %dpi, align 4
  107.   %tmp16 = load i32* %gid, align 4
  108.   %tmp17 = mul i32 16, %tmp16
  109.   %conv18 = zext i32 %tmp17 to i64
  110.   %call19 = call i64 @get_global_size(i32 0) #0
  111.   %tmp20 = udiv i64 %conv18, %call19
  112.   %conv21 = trunc i64 %tmp20 to i32
  113.   store i32 %conv21, i32* %curstep, align 4
  114.   %tmp22 = load i32* %curstep, align 4
  115.   %tmp23 = load i32* %DcrPerItem.addr, align 4
  116.   %tmp24 = and i32 %tmp23, 15
  117.   %cmp25 = icmp ult i32 %tmp22, %tmp24
  118.   br i1 %cmp25, label %if.then27, label %if.end26
  119.  
  120. if.end26:                                         ; preds = %if.then27, %if.then12
  121.   store i32 0, i32* %i, align 4
  122.   br label %for.cond
  123.  
  124. if.then27:                                        ; preds = %if.then12
  125.   %tmp28 = load i32* %dpi, align 4
  126.   %tmp29 = add i32 %tmp28, 1
  127.   store i32 %tmp29, i32* %dpi, align 4
  128.   br label %if.end26
  129.  
  130. for.cond:                                         ; preds = %for.inc, %if.end26
  131.   %tmp30 = load i32* %i, align 4
  132.   %tmp31 = load i32* %dpi, align 4
  133.   %cmp32 = icmp ult i32 %tmp30, %tmp31
  134.   br i1 %cmp32, label %for.body, label %for.exit
  135.  
  136. for.exit:                                         ; preds = %for.cond
  137.   br label %return
  138.  
  139. for.body:                                         ; preds = %for.cond
  140.   %tmp33 = load i32 addrspace(1)** %g_output.addr, align 8
  141.   %tmp34 = load i32 addrspace(2)** %d_header.addr, align 8
  142.   %tmp35 = load i32* %gid, align 4
  143.   %conv36 = zext i32 %tmp35 to i64
  144.   %tmp37 = load i32* %i, align 4
  145.   %conv38 = zext i32 %tmp37 to i64
  146.   %call39 = call i64 @get_global_size(i32 0) #0
  147.   %tmp40 = mul i64 %conv38, %call39
  148.   %tmp41 = add i64 %conv36, %tmp40
  149.   %conv42 = trunc i64 %tmp41 to i32
  150.   %tmp43 = load i64 addrspace(3)** %lar, align 8
  151.   call void @lbc_search(i32 addrspace(1)* noalias %tmp33, i32 addrspace(2)* %tmp34, i32 %conv42, i64 addrspace(3)* %tmp43) #0
  152.   br label %for.inc
  153.  
  154. for.inc:                                          ; preds = %for.body
  155.   %tmp44 = load i32* %i, align 4
  156.   %tmp45 = add i32 %tmp44, 1
  157.   store i32 %tmp45, i32* %i, align 4
  158.   br label %for.cond
  159.  
  160. for.cond50:                                       ; preds = %for.inc55, %if.end11
  161.   %tmp52 = load i32* %i49, align 4
  162.   %cmp53 = icmp ne i32 %tmp52, 4
  163.   br i1 %cmp53, label %for.body54, label %for.exit51
  164.  
  165. for.exit51:                                       ; preds = %for.cond50
  166.   %ptr68 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  167.   %arrayidx69 = getelementptr i64* %ptr68, i64 4
  168.   %tmp70 = load i64* %start_nonce.addr, align 8
  169.   %tmp71 = load i32* %gid, align 4
  170.   %conv72 = zext i32 %tmp71 to i64
  171.   %tmp73 = add i64 %tmp70, %conv72
  172.   store i64 %tmp73, i64* %arrayidx69, align 8
  173.   store i32 6, i32* %i74, align 4
  174.   br label %for.cond75
  175.  
  176. for.body54:                                       ; preds = %for.cond50
  177.   %ptr56 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  178.   %tmp57 = load i32* %i49, align 4
  179.   %conv58 = zext i32 %tmp57 to i64
  180.   %arrayidx59 = getelementptr i64* %ptr56, i64 %conv58
  181.   %tmp60 = load %2 addrspace(2)** %g_header.addr, align 8
  182.   %structele = getelementptr inbounds %2 addrspace(2)* %tmp60, i32 0, i32 0
  183.   %ptr61 = getelementptr inbounds [4 x i64] addrspace(2)* %structele, i32 0, i32 0
  184.   %tmp62 = load i32* %i49, align 4
  185.   %conv63 = zext i32 %tmp62 to i64
  186.   %arrayidx64 = getelementptr i64 addrspace(2)* %ptr61, i64 %conv63
  187.   %tmp65 = load i64 addrspace(2)* %arrayidx64, align 8
  188.   store i64 %tmp65, i64* %arrayidx59, align 8
  189.   br label %for.inc55
  190.  
  191. for.inc55:                                        ; preds = %for.body54
  192.   %tmp66 = load i32* %i49, align 4
  193.   %tmp67 = add i32 %tmp66, 1
  194.   store i32 %tmp67, i32* %i49, align 4
  195.   br label %for.cond50
  196.  
  197. for.cond75:                                       ; preds = %for.inc80, %for.exit51
  198.   %tmp77 = load i32* %i74, align 4
  199.   %cmp78 = icmp ne i32 %tmp77, 25
  200.   br i1 %cmp78, label %for.body79, label %for.exit76
  201.  
  202. for.exit76:                                       ; preds = %for.cond75
  203.   %ptr87 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  204.   %arrayidx88 = getelementptr i64* %ptr87, i64 5
  205.   store i64 1, i64* %arrayidx88, align 8
  206.   %ptr89 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  207.   %arrayidx90 = getelementptr i64* %ptr89, i64 8
  208.   store i64 -9223372036854775808, i64* %arrayidx90, align 8
  209.   %ptr91 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  210.   %conv92 = bitcast i64* %ptr91 to <2 x i32>*
  211.   %tmp93 = load i32* %isolate.addr, align 4
  212.   call void @keccak_f1600_no_absorb(<2 x i32>* %conv92, i32 8, i32 %tmp93) #0
  213.   %tmp94 = load i32* %dm.addr, align 4
  214.   %tmp95 = sub i32 32, %tmp94
  215.   %tmp96 = and i32 %tmp95, 63
  216.   %tmp97 = zext i32 %tmp96 to i64
  217.   %tmp98 = shl i64 1, %tmp97
  218.   %conv99 = trunc i64 %tmp98 to i32
  219.   store i32 %conv99, i32* %morebit, align 4
  220.   %tmp100 = load i32* %gid, align 4
  221.   %tmp101 = and i32 %tmp100, 7
  222.   store i32 %tmp101, i32* %thread_id, align 4
  223.   %tmp102 = load i32* %gid, align 4
  224.   %tmp103 = urem i32 %tmp102, 64
  225.   %tmp104 = lshr i32 %tmp103, 3
  226.   store i32 %tmp104, i32* %hash_id, align 4
  227.   store i32 0, i32* %i105, align 4
  228.   br label %for.cond106
  229.  
  230. for.body79:                                       ; preds = %for.cond75
  231.   %ptr81 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  232.   %tmp82 = load i32* %i74, align 4
  233.   %conv83 = zext i32 %tmp82 to i64
  234.   %arrayidx84 = getelementptr i64* %ptr81, i64 %conv83
  235.   store i64 0, i64* %arrayidx84, align 8
  236.   br label %for.inc80
  237.  
  238. for.inc80:                                        ; preds = %for.body79
  239.   %tmp85 = load i32* %i74, align 4
  240.   %tmp86 = add i32 %tmp85, 1
  241.   store i32 %tmp86, i32* %i74, align 4
  242.   br label %for.cond75
  243.  
  244. for.cond106:                                      ; preds = %for.inc111, %for.exit76
  245.   %tmp108 = load i32* %i105, align 4
  246.   %cmp109 = icmp slt i32 %tmp108, 8
  247.   br i1 %cmp109, label %for.body110, label %for.exit107
  248.  
  249. for.exit107:                                      ; preds = %for.cond106
  250.   store i32 13, i32* %i294, align 4
  251.   br label %for.cond295
  252.  
  253. for.body110:                                      ; preds = %for.cond106
  254.   %tmp112 = load i32* %i105, align 4
  255.   %tmp113 = load i32* %thread_id, align 4
  256.   %cmp114 = icmp eq i32 %tmp112, %tmp113
  257.   br i1 %cmp114, label %if.then116, label %if.end115
  258.  
  259. for.inc111:                                       ; preds = %if.end264
  260.   %tmp292 = load i32* %i105, align 4
  261.   %tmp293 = add nsw i32 %tmp292, 1
  262.   store i32 %tmp293, i32* %i105, align 4
  263.   br label %for.cond106
  264.  
  265. if.end115:                                        ; preds = %for.exit119, %for.body110
  266.   call void @barrier(i32 1) #0
  267.   %tmp141 = load %0 addrspace(3)** %share, align 8
  268.   %tmp142 = load i32* %hash_id, align 4
  269.   %conv143 = zext i32 %tmp142 to i64
  270.   %arrayidx144 = getelementptr %0 addrspace(3)* %tmp141, i64 %conv143
  271.   %structele145 = getelementptr inbounds %0 addrspace(3)* %arrayidx144, i32 0, i32 0
  272.   %ptr146 = getelementptr inbounds [4 x <4 x i32>] addrspace(3)* %structele145, i32 0, i32 0
  273.   %tmp147 = load i32* %thread_id, align 4
  274.   %tmp148 = and i32 %tmp147, 3
  275.   %conv149 = zext i32 %tmp148 to i64
  276.   %arrayidx150 = getelementptr <4 x i32> addrspace(3)* %ptr146, i64 %conv149
  277.   %tmp151 = load <4 x i32> addrspace(3)* %arrayidx150, align 16
  278.   store <4 x i32> %tmp151, <4 x i32>* %mix, align 16
  279.   call void @barrier(i32 1) #0
  280.   %tmp152 = load %0 addrspace(3)** %share, align 8
  281.   %tmp153 = load i32* %hash_id, align 4
  282.   %conv154 = zext i32 %tmp153 to i64
  283.   %arrayidx155 = getelementptr %0 addrspace(3)* %tmp152, i64 %conv154
  284.   %structele156 = getelementptr inbounds %0 addrspace(3)* %arrayidx155, i32 0, i32 0
  285.   %tmp157 = bitcast [4 x <4 x i32>] addrspace(3)* %structele156 to [16 x i32] addrspace(3)*
  286.   %ptr158 = getelementptr inbounds [16 x i32] addrspace(3)* %tmp157, i32 0, i32 0
  287.   store i32 addrspace(3)* %ptr158, i32 addrspace(3)** %share0, align 8
  288.   %tmp159 = load i32* %thread_id, align 4
  289.   %cmp160 = icmp eq i32 %tmp159, 0
  290.   br i1 %cmp160, label %if.then162, label %if.end161
  291.  
  292. if.then116:                                       ; preds = %for.body110
  293.   store i32 0, i32* %i117, align 4
  294.   br label %for.cond118
  295.  
  296. for.cond118:                                      ; preds = %for.inc123, %if.then116
  297.   %tmp120 = load i32* %i117, align 4
  298.   %cmp121 = icmp ne i32 %tmp120, 8
  299.   br i1 %cmp121, label %for.body122, label %for.exit119
  300.  
  301. for.exit119:                                      ; preds = %for.cond118
  302.   br label %if.end115
  303.  
  304. for.body122:                                      ; preds = %for.cond118
  305.   %tmp124 = load %0 addrspace(3)** %share, align 8
  306.   %tmp125 = load i32* %hash_id, align 4
  307.   %conv126 = zext i32 %tmp125 to i64
  308.   %arrayidx127 = getelementptr %0 addrspace(3)* %tmp124, i64 %conv126
  309.   %structele128 = getelementptr inbounds %0 addrspace(3)* %arrayidx127, i32 0, i32 0
  310.   %tmp129 = bitcast [4 x <4 x i32>] addrspace(3)* %structele128 to [8 x i64] addrspace(3)*
  311.   %ptr130 = getelementptr inbounds [8 x i64] addrspace(3)* %tmp129, i32 0, i32 0
  312.   %tmp131 = load i32* %i117, align 4
  313.   %conv132 = zext i32 %tmp131 to i64
  314.   %arrayidx133 = getelementptr i64 addrspace(3)* %ptr130, i64 %conv132
  315.   %ptr134 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  316.   %tmp135 = load i32* %i117, align 4
  317.   %conv136 = zext i32 %tmp135 to i64
  318.   %arrayidx137 = getelementptr i64* %ptr134, i64 %conv136
  319.   %tmp138 = load i64* %arrayidx137, align 8
  320.   store i64 %tmp138, i64 addrspace(3)* %arrayidx133, align 8
  321.   br label %for.inc123
  322.  
  323. for.inc123:                                       ; preds = %for.body122
  324.   %tmp139 = load i32* %i117, align 4
  325.   %tmp140 = add i32 %tmp139, 1
  326.   store i32 %tmp140, i32* %i117, align 4
  327.   br label %for.cond118
  328.  
  329. if.end161:                                        ; preds = %if.then162, %if.end115
  330.   call void @barrier(i32 1) #0
  331.   %tmp166 = load i32 addrspace(3)** %share0, align 8
  332.   %tmp167 = load i32 addrspace(3)* %tmp166, align 4
  333.   store i32 %tmp167, i32* %init0, align 4
  334.   store i32 0, i32* %a, align 4
  335.   br label %for.cond168
  336.  
  337. if.then162:                                       ; preds = %if.end115
  338.   %tmp163 = load i32 addrspace(3)** %share0, align 8
  339.   %tmp164 = load <4 x i32>* %mix, align 16
  340.   %tmp165 = extractelement <4 x i32> %tmp164, i32 0
  341.   store i32 %tmp165, i32 addrspace(3)* %tmp163, align 4
  342.   br label %if.end161
  343.  
  344. for.cond168:                                      ; preds = %for.inc173, %if.end161
  345.   %tmp170 = load i32* %a, align 4
  346.   %cmp171 = icmp ult i32 %tmp170, 64
  347.   br i1 %cmp171, label %for.body172, label %for.exit169
  348.  
  349. for.exit169:                                      ; preds = %for.cond168
  350.   %tmp248 = load %0 addrspace(3)** %share, align 8
  351.   %tmp249 = load i32* %hash_id, align 4
  352.   %conv250 = zext i32 %tmp249 to i64
  353.   %arrayidx251 = getelementptr %0 addrspace(3)* %tmp248, i64 %conv250
  354.   %structele252 = getelementptr inbounds %0 addrspace(3)* %arrayidx251, i32 0, i32 0
  355.   %tmp253 = bitcast [4 x <4 x i32>] addrspace(3)* %structele252 to [16 x i32] addrspace(3)*
  356.   %ptr254 = getelementptr inbounds [16 x i32] addrspace(3)* %tmp253, i32 0, i32 0
  357.   %tmp255 = load i32* %thread_id, align 4
  358.   %conv256 = zext i32 %tmp255 to i64
  359.   %arrayidx257 = getelementptr i32 addrspace(3)* %ptr254, i64 %conv256
  360.   %tmp258 = load <4 x i32>* %mix, align 16
  361.   %call259 = call i32 @fnv_reduce(<4 x i32> %tmp258) #0
  362.   store i32 %call259, i32 addrspace(3)* %arrayidx257, align 4
  363.   call void @barrier(i32 1) #0
  364.   %tmp260 = load i32* %i105, align 4
  365.   %tmp261 = load i32* %thread_id, align 4
  366.   %cmp262 = icmp eq i32 %tmp260, %tmp261
  367.   br i1 %cmp262, label %if.then265, label %if.end264
  368.  
  369. for.body172:                                      ; preds = %for.cond168
  370.   %tmp174 = load i32* %thread_id, align 4
  371.   %tmp175 = load i32* %a, align 4
  372.   %tmp176 = lshr i32 %tmp175, 2
  373.   %tmp177 = and i32 %tmp176, 7
  374.   %cmp178 = icmp eq i32 %tmp174, %tmp177
  375.   %cmp.ext = zext i1 %cmp178 to i32
  376.   %cmp179 = icmp ne i32 %cmp.ext, 0
  377.   %cmp.ext180 = zext i1 %cmp179 to i32
  378.   store i32 %cmp.ext180, i32* %update_share, align 4
  379.   store i32 0, i32* %i181, align 4
  380.   br label %for.cond182
  381.  
  382. for.inc173:                                       ; preds = %for.exit183
  383.   %tmp246 = load i32* %a, align 4
  384.   %tmp247 = add i32 %tmp246, 4
  385.   store i32 %tmp247, i32* %a, align 4
  386.   br label %for.cond168
  387.  
  388. for.cond182:                                      ; preds = %for.inc188, %for.body172
  389.   %tmp184 = load i32* %i181, align 4
  390.   %cmp185 = icmp ne i32 %tmp184, 4
  391.   br i1 %cmp185, label %for.body187, label %for.exit183
  392.  
  393. for.exit183:                                      ; preds = %for.cond182
  394.   br label %for.inc173
  395.  
  396. for.body187:                                      ; preds = %for.cond182
  397.   %tmp189 = load i32* %update_share, align 4
  398.   %tobool190 = icmp ne i32 %tmp189, 0
  399.   br i1 %tobool190, label %if.then192, label %if.end191
  400.  
  401. for.inc188:                                       ; preds = %if.end191
  402.   %tmp244 = load i32* %i181, align 4
  403.   %tmp245 = add i32 %tmp244, 1
  404.   store i32 %tmp245, i32* %i181, align 4
  405.   br label %for.cond182
  406.  
  407. if.end191:                                        ; preds = %if.end220, %for.body187
  408.   call void @barrier(i32 1) #0
  409.   %tmp231 = load <4 x i32>* %mix, align 16
  410.   %tmp232 = load %3 addrspace(1)** %g_dag.addr, align 8
  411.   %tmp233 = load i32 addrspace(3)** %share0, align 8
  412.   %tmp234 = load i32 addrspace(3)* %tmp233, align 4
  413.   %conv235 = zext i32 %tmp234 to i64
  414.   %arrayidx236 = getelementptr %3 addrspace(1)* %tmp232, i64 %conv235
  415.   %structele237 = getelementptr inbounds %3 addrspace(1)* %arrayidx236, i32 0, i32 0
  416.   %ptr238 = getelementptr inbounds [8 x <4 x i32>] addrspace(1)* %structele237, i32 0, i32 0
  417.   %tmp239 = load i32* %thread_id, align 4
  418.   %conv240 = zext i32 %tmp239 to i64
  419.   %arrayidx241 = getelementptr <4 x i32> addrspace(1)* %ptr238, i64 %conv240
  420.   %tmp242 = load <4 x i32> addrspace(1)* %arrayidx241, align 16
  421.   %call243 = call <4 x i32> @fnv4(<4 x i32> %tmp231, <4 x i32> %tmp242) #0
  422.   store <4 x i32> %call243, <4 x i32>* %mix, align 16
  423.   br label %for.inc188
  424.  
  425. if.then192:                                       ; preds = %for.body187
  426.   %tmp193 = load i32* %init0, align 4
  427.   %tmp194 = load i32* %a, align 4
  428.   %tmp195 = load i32* %i181, align 4
  429.   %tmp196 = add i32 %tmp194, %tmp195
  430.   %tmp197 = xor i32 %tmp193, %tmp196
  431.   %conv198 = bitcast <4 x i32>* %mix to i32*
  432.   %tmp199 = load i32* %i181, align 4
  433.   %conv200 = zext i32 %tmp199 to i64
  434.   %arrayidx201 = getelementptr i32* %conv198, i64 %conv200
  435.   %tmp202 = load i32* %arrayidx201, align 4
  436.   %call203 = call i32 @fnv(i32 %tmp197, i32 %tmp202) #0
  437.   store i32 %call203, i32* %v, align 4
  438.   %tmp204 = load i32* %v, align 4
  439.   %tmp205 = load i32* %dv.addr, align 4
  440.   %call206 = call i32 @__mul_hi_u32(i32 %tmp204, i32 %tmp205) #0
  441.   store i32 %call206, i32* %ml, align 4
  442.   %tmp207 = load i32* %ml, align 4
  443.   %tmp208 = load i32* %v, align 4
  444.   %tmp209 = load i32* %dmask.addr, align 4
  445.   %tmp210 = and i32 %tmp208, %tmp209
  446.   %tmp211 = add i32 %tmp207, %tmp210
  447.   store i32 %tmp211, i32* %aml, align 4
  448.   %tmp212 = load i32* %aml, align 4
  449.   %tmp213 = load i32* %dm.addr, align 4
  450.   %tmp214 = and i32 %tmp213, 31
  451.   %tmp215 = lshr i32 %tmp212, %tmp214
  452.   store i32 %tmp215, i32* %res, align 4
  453.   %tmp216 = load i32* %aml, align 4
  454.   %tmp217 = load i32* %ml, align 4
  455.   %cmp218 = icmp ult i32 %tmp216, %tmp217
  456.   br i1 %cmp218, label %if.then221, label %if.end220
  457.  
  458. if.end220:                                        ; preds = %if.then221, %if.then192
  459.   %tmp225 = load i32 addrspace(3)** %share0, align 8
  460.   %tmp226 = load i32* %v, align 4
  461.   %tmp227 = load i32* %dagsize.addr, align 4
  462.   %tmp228 = load i32* %res, align 4
  463.   %tmp229 = mul i32 %tmp227, %tmp228
  464.   %tmp230 = sub i32 %tmp226, %tmp229
  465.   store i32 %tmp230, i32 addrspace(3)* %tmp225, align 4
  466.   br label %if.end191
  467.  
  468. if.then221:                                       ; preds = %if.then192
  469.   %tmp222 = load i32* %res, align 4
  470.   %tmp223 = load i32* %morebit, align 4
  471.   %tmp224 = or i32 %tmp222, %tmp223
  472.   store i32 %tmp224, i32* %res, align 4
  473.   br label %if.end220
  474.  
  475. if.end264:                                        ; preds = %for.exit268, %for.exit169
  476.   call void @barrier(i32 1) #0
  477.   br label %for.inc111
  478.  
  479. if.then265:                                       ; preds = %for.exit169
  480.   store i32 0, i32* %i266, align 4
  481.   br label %for.cond267
  482.  
  483. for.cond267:                                      ; preds = %for.inc273, %if.then265
  484.   %tmp269 = load i32* %i266, align 4
  485.   %cmp270 = icmp ne i32 %tmp269, 4
  486.   br i1 %cmp270, label %for.body272, label %for.exit268
  487.  
  488. for.exit268:                                      ; preds = %for.cond267
  489.   br label %if.end264
  490.  
  491. for.body272:                                      ; preds = %for.cond267
  492.   %ptr274 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  493.   %tmp275 = getelementptr i64* %ptr274, i64 8
  494.   %tmp276 = load i32* %i266, align 4
  495.   %conv277 = zext i32 %tmp276 to i64
  496.   %arrayidx278 = getelementptr i64* %tmp275, i64 %conv277
  497.   %tmp279 = load %0 addrspace(3)** %share, align 8
  498.   %tmp280 = load i32* %hash_id, align 4
  499.   %conv281 = zext i32 %tmp280 to i64
  500.   %arrayidx282 = getelementptr %0 addrspace(3)* %tmp279, i64 %conv281
  501.   %structele283 = getelementptr inbounds %0 addrspace(3)* %arrayidx282, i32 0, i32 0
  502.   %tmp284 = bitcast [4 x <4 x i32>] addrspace(3)* %structele283 to [8 x i64] addrspace(3)*
  503.   %ptr285 = getelementptr inbounds [8 x i64] addrspace(3)* %tmp284, i32 0, i32 0
  504.   %tmp286 = load i32* %i266, align 4
  505.   %conv287 = zext i32 %tmp286 to i64
  506.   %arrayidx288 = getelementptr i64 addrspace(3)* %ptr285, i64 %conv287
  507.   %tmp289 = load i64 addrspace(3)* %arrayidx288, align 8
  508.   store i64 %tmp289, i64* %arrayidx278, align 8
  509.   br label %for.inc273
  510.  
  511. for.inc273:                                       ; preds = %for.body272
  512.   %tmp290 = load i32* %i266, align 4
  513.   %tmp291 = add i32 %tmp290, 1
  514.   store i32 %tmp291, i32* %i266, align 4
  515.   br label %for.cond267
  516.  
  517. for.cond295:                                      ; preds = %for.inc301, %for.exit107
  518.   %tmp297 = load i32* %i294, align 4
  519.   %cmp298 = icmp ne i32 %tmp297, 25
  520.   br i1 %cmp298, label %for.body300, label %for.exit296
  521.  
  522. for.exit296:                                      ; preds = %for.cond295
  523.   %ptr308 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  524.   %arrayidx309 = getelementptr i64* %ptr308, i64 12
  525.   store i64 1, i64* %arrayidx309, align 8
  526.   %ptr310 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  527.   %arrayidx311 = getelementptr i64* %ptr310, i64 16
  528.   store i64 -9223372036854775808, i64* %arrayidx311, align 8
  529.   %ptr312 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  530.   %conv313 = bitcast i64* %ptr312 to <2 x i32>*
  531.   %tmp314 = load i32* %isolate.addr, align 4
  532.   call void @keccak_f1600_no_absorb(<2 x i32>* %conv313, i32 1, i32 %tmp314) #0
  533.   %ptr315 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  534.   %arrayidx316 = getelementptr i64* %ptr315, i64 0
  535.   %tmp317 = load i64* %arrayidx316, align 8
  536.   %conv318 = bitcast i64 %tmp317 to <8 x i8>
  537.   store <8 x i8> %conv318, <8 x i8>* %tmp319, align 8
  538.   %tmp320 = load <8 x i8>* %tmp319, align 8
  539.   %tmp321 = extractelement <8 x i8> %tmp320, i32 7
  540.   %tmp322 = insertelement <8 x i8> undef, i8 %tmp321, i32 0
  541.   %tmp323 = extractelement <8 x i8> %tmp320, i32 6
  542.   %tmp324 = insertelement <8 x i8> %tmp322, i8 %tmp323, i32 1
  543.   %tmp325 = extractelement <8 x i8> %tmp320, i32 5
  544.   %tmp326 = insertelement <8 x i8> %tmp324, i8 %tmp325, i32 2
  545.   %tmp327 = extractelement <8 x i8> %tmp320, i32 4
  546.   %tmp328 = insertelement <8 x i8> %tmp326, i8 %tmp327, i32 3
  547.   %tmp329 = extractelement <8 x i8> %tmp320, i32 3
  548.   %tmp330 = insertelement <8 x i8> %tmp328, i8 %tmp329, i32 4
  549.   %tmp331 = extractelement <8 x i8> %tmp320, i32 2
  550.   %tmp332 = insertelement <8 x i8> %tmp330, i8 %tmp331, i32 5
  551.   %tmp333 = extractelement <8 x i8> %tmp320, i32 1
  552.   %tmp334 = insertelement <8 x i8> %tmp332, i8 %tmp333, i32 6
  553.   %tmp335 = extractelement <8 x i8> %tmp320, i32 0
  554.   %tmp336 = insertelement <8 x i8> %tmp334, i8 %tmp335, i32 7
  555.   %conv337 = bitcast <8 x i8> %tmp336 to i64
  556.   %tmp338 = load i64* %target.addr, align 8
  557.   %cmp339 = icmp ult i64 %conv337, %tmp338
  558.   br i1 %cmp339, label %if.then342, label %if.end341
  559.  
  560. for.body300:                                      ; preds = %for.cond295
  561.   %ptr302 = getelementptr inbounds [25 x i64]* %state, i32 0, i32 0
  562.   %tmp303 = load i32* %i294, align 4
  563.   %conv304 = zext i32 %tmp303 to i64
  564.   %arrayidx305 = getelementptr i64* %ptr302, i64 %conv304
  565.   store i64 0, i64* %arrayidx305, align 8
  566.   br label %for.inc301
  567.  
  568. for.inc301:                                       ; preds = %for.body300
  569.   %tmp306 = load i32* %i294, align 4
  570.   %tmp307 = add i32 %tmp306, 1
  571.   store i32 %tmp307, i32* %i294, align 4
  572.   br label %for.cond295
  573.  
  574. if.end341:                                        ; preds = %if.then342, %for.exit296
  575.   br label %return
  576.  
  577. if.then342:                                       ; preds = %for.exit296
  578.   %tmp343 = load i32 addrspace(1)** %g_output.addr, align 8
  579.   %tmp344 = getelementptr i32 addrspace(1)* %tmp343, i64 0
  580.   %call345 = call i32 @__atomic_inc_gu32(i32 addrspace(1)* %tmp344) #0
  581.   %call346 = call i32 @__min_u32(i32 3, i32 %call345) #0
  582.   store i32 %call346, i32* %slot, align 4
  583.   %tmp347 = load i32 addrspace(1)** %g_output.addr, align 8
  584.   %tmp348 = load i32* %slot, align 4
  585.   %tmp349 = mul i32 9, %tmp348
  586.   %tmp350 = add i32 1, %tmp349
  587.   %conv351 = zext i32 %tmp350 to i64
  588.   %arrayidx352 = getelementptr i32 addrspace(1)* %tmp347, i64 %conv351
  589.   %tmp353 = load i32* %gid, align 4
  590.   store volatile i32 %tmp353, i32 addrspace(1)* %arrayidx352, align 4
  591.   br label %if.end341
  592. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement