API tools faq
paste
Advertisement
TheGhastModding

Error log

TheGhastModding
Aug 29th, 2019
267
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 73.84 KB | None | 0 0
raw download clone embed print report
  1. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  2. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  3. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  4. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  5. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  6. MIOpenConvD3x3 -DMLO_DIR_FORWARD=1 -DMLO_GRP_SZ=256 -DMLO_GRP_SZ0=256 -DMLO_GRP_SZ1=1 -DMLO_GRP_SZ2=1 -DMLO_FILTER_SIZE0=3 -DMLO_FILTER_SIZE1=3 -DMLO_FILTER_PAD0=1 -DMLO_FILTER_PAD1=1 -DMLO_N_OUTPUTS=128 -DMLO_N_INPUTS=64 -DMLO_BATCH_SZ=32 -DMLO_OUT_BATCH_STRIDE=524288 -DMLO_OUT_CHANNEL_STRIDE=4096 -DMLO_OUT_STRIDE=64 -DMLO_IN_BATCH_STRIDE=262144 -DMLO_IN_CHANNEL_STRIDE=4096 -DMLO_IN_STRIDE=64 -DMLO_WEI_BATCH_STRIDE=576 -DMLO_WEI_CHANNEL_STRIDE=9 -DMLO_IN_WIDTH=64 -DMLO_IN_HEIGHT=64 -DMLO_N_LCL_BATCHS=1 -DMLO_N_LCL_OUT_MAPS=4 -DMLO_N_LCL_IN_MAPS=1 -DMLO_OUT_TILE0=4 -DMLO_OUT_TILE1=2 -DMLO_ALU_EXTENT_X=16 -DMLO_LG2ALU_EXTENT_X=4 -DMLO_ALU_EXTENT_Y=16 -DMLO_LG2ALU_EXTENT_Y=4 -DMLO_OUT_EXTENT1=8 -DMLO_OUT_EXTENT0=64 -DMLO_N_WAVES=4 -DMLO_N_WAVES_MASK=3 -DMLO_LG2_WAVE_SZ=6 -DMLO_LG2_WAVE_SZ0=4 -DMLO_READ_TYPE=_FLOAT4 -DMLO_READ_UNIT=4 -DMLO_CONV_BIAS=0 -DMIOPEN_USE_FP16=0 -DMIOPEN_USE_FP32=1 -DMIOPEN_USE_INT8=0 -DMIOPEN_USE_INT8x4=0 -DMIOPEN_USE_BFP16=0 -DMIOPEN_USE_INT32=0 -DMIOPEN_USE_RNE_BFLOAT16=1
  7. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  8. MIOpenConvD3x3 -DMLO_DIR_FORWARD=1 -DMLO_GRP_SZ=256 -DMLO_GRP_SZ0=256 -DMLO_GRP_SZ1=1 -DMLO_GRP_SZ2=1 -DMLO_FILTER_SIZE0=3 -DMLO_FILTER_SIZE1=3 -DMLO_FILTER_PAD0=1 -DMLO_FILTER_PAD1=1 -DMLO_N_OUTPUTS=1 -DMLO_N_INPUTS=128 -DMLO_BATCH_SZ=32 -DMLO_OUT_BATCH_STRIDE=16384 -DMLO_OUT_CHANNEL_STRIDE=16384 -DMLO_OUT_STRIDE=128 -DMLO_IN_BATCH_STRIDE=2097152 -DMLO_IN_CHANNEL_STRIDE=16384 -DMLO_IN_STRIDE=128 -DMLO_WEI_BATCH_STRIDE=1152 -DMLO_WEI_CHANNEL_STRIDE=9 -DMLO_IN_WIDTH=128 -DMLO_IN_HEIGHT=128 -DMLO_N_LCL_BATCHS=1 -DMLO_N_LCL_OUT_MAPS=4 -DMLO_N_LCL_IN_MAPS=1 -DMLO_OUT_TILE0=4 -DMLO_OUT_TILE1=2 -DMLO_ALU_EXTENT_X=32 -DMLO_LG2ALU_EXTENT_X=5 -DMLO_ALU_EXTENT_Y=8 -DMLO_LG2ALU_EXTENT_Y=3 -DMLO_OUT_EXTENT1=4 -DMLO_OUT_EXTENT0=128 -DMLO_N_WAVES=4 -DMLO_N_WAVES_MASK=3 -DMLO_LG2_WAVE_SZ=6 -DMLO_LG2_WAVE_SZ0=5 -DMLO_READ_TYPE=_FLOAT4 -DMLO_READ_UNIT=4 -DMLO_CONV_BIAS=0 -DMIOPEN_USE_FP16=0 -DMIOPEN_USE_FP32=1 -DMIOPEN_USE_INT8=0 -DMIOPEN_USE_INT8x4=0 -DMIOPEN_USE_BFP16=0 -DMIOPEN_USE_INT32=0 -DMIOPEN_USE_RNE_BFLOAT16=1
  9. MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvAsm3x3U, enforce: SEARCH_DB_UPDATE(4), ALL(1)
  10. MIOpen(HIP): Warning [GenericSearch] ConvAsm3x3U: Searching the best solution among 640...
  11. MIOpen(HIP): Warning [GenericSearch] Done: 640/0/640, best #111 1.33233 1,4,2
  12. MIOpen(HIP): Warning [GenericSearch] ...Score: 1.88003 (default time 2.50481)
  13. MIOpenConvD3x3 -DMLO_DIR_FORWARD=1 -DMLO_GRP_SZ=256 -DMLO_GRP_SZ0=256 -DMLO_GRP_SZ1=1 -DMLO_GRP_SZ2=1 -DMLO_FILTER_SIZE0=3 -DMLO_FILTER_SIZE1=3 -DMLO_FILTER_PAD0=1 -DMLO_FILTER_PAD1=1 -DMLO_N_OUTPUTS=1 -DMLO_N_INPUTS=128 -DMLO_BATCH_SZ=32 -DMLO_OUT_BATCH_STRIDE=16384 -DMLO_OUT_CHANNEL_STRIDE=16384 -DMLO_OUT_STRIDE=128 -DMLO_IN_BATCH_STRIDE=2097152 -DMLO_IN_CHANNEL_STRIDE=16384 -DMLO_IN_STRIDE=128 -DMLO_WEI_BATCH_STRIDE=1152 -DMLO_WEI_CHANNEL_STRIDE=9 -DMLO_IN_WIDTH=128 -DMLO_IN_HEIGHT=128 -DMLO_N_LCL_BATCHS=1 -DMLO_N_LCL_OUT_MAPS=4 -DMLO_N_LCL_IN_MAPS=1 -DMLO_OUT_TILE0=4 -DMLO_OUT_TILE1=2 -DMLO_ALU_EXTENT_X=32 -DMLO_LG2ALU_EXTENT_X=5 -DMLO_ALU_EXTENT_Y=8 -DMLO_LG2ALU_EXTENT_Y=3 -DMLO_OUT_EXTENT1=4 -DMLO_OUT_EXTENT0=128 -DMLO_N_WAVES=4 -DMLO_N_WAVES_MASK=3 -DMLO_LG2_WAVE_SZ=6 -DMLO_LG2_WAVE_SZ0=5 -DMLO_READ_TYPE=_FLOAT4 -DMLO_READ_UNIT=4 -DMLO_CONV_BIAS=0 -DMIOPEN_USE_FP16=0 -DMIOPEN_USE_FP32=1 -DMIOPEN_USE_INT8=0 -DMIOPEN_USE_INT8x4=0 -DMIOPEN_USE_BFP16=0 -DMIOPEN_USE_INT32=0 -DMIOPEN_USE_RNE_BFLOAT16=1
  14. MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvOclDirectFwd, enforce: SEARCH_DB_UPDATE(4), ALL(1)
  15. MIOpen(HIP): Warning [SearchImpl] Searching the best solution in the 9 dim space. Please, be patient...
  16. MIOpen(HIP): Warning [SearchImpl] Runs left: 863, min time so far: 2.7725, curr time: 2.7725 16,16,16,16,1,1,1,1,1
  17. MIOpen(HIP): Warning [SearchImpl] Runs left: 814, min time so far: 2.11793, curr time: 2.11793 16,16,16,32,1,2,1,1,2
  18. MIOpen(HIP): Warning [SearchImpl] Runs left: 777, min time so far: 2.04033, curr time: 3.12978 8,8,16,32,2,4,1,2,1
  19. MIOpen(HIP): Warning [SearchImpl] Runs left: 711, min time so far: 1.96545, curr time: 2.02977 8,32,32,32,4,1,1,2,2
  20. MIOpen(HIP): Warning [SearchImpl] Default run, min time so far: 1.96545, default time: 2.00721 16,16,32,32,2,2,8,2,1
  21. MIOpen(HIP): Warning [SearchImpl] ...Score: 1.02125
  22. MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvOclDirectFwd, enforce: SEARCH_DB_UPDATE(4), ALL(1)
  23. MIOpen(HIP): Warning [SearchImpl] Searching the best solution in the 9 dim space. Please, be patient...
  24. MIOpen(HIP): Warning [SearchImpl] Runs left: 863, min time so far: 8.72165, curr time: 8.72165 16,16,16,16,1,1,1,1,1
  25. MIOpen(HIP): Warning [SearchImpl] Runs left: 813, min time so far: 2.46433, curr time: 2.46433 8,8,16,16,2,2,1,1,2
  26. MIOpen(HIP): Warning [SearchImpl] Runs left: 744, min time so far: 1.50913, curr time: 3.45074 8,32,16,32,2,1,2,1,1
  27. MIOpen(HIP): Warning [SearchImpl] Runs left: 658, min time so far: 1.50913, curr time: 2.26433 32,8,32,32,1,4,4,1,1
  28. MIOpen(HIP): Warning [SearchImpl] Runs left: 590, min time so far: 1.50913, curr time: 9.71878 8,32,32,32,4,1,8,1,1
  29. MIOpen(HIP): Warning [SearchImpl] Default run, min time so far: 1.50913, default time: 1.83985 16,16,32,32,2,2,8,2,1
  30. MIOpen(HIP): Warning [SearchImpl] ...Score: 1.21915
  31. MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvAsmBwdWrW3x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
  32. MIOpen(HIP): Warning [GenericSearch] ConvAsmBwdWrW3x3: Searching the best solution among 96...
  33. MIOpen(HIP): Warning [GenericSearch] Done: 96/0/96, best #9 1.33668 0,0,16,1,5,1
  34. MIOpen(HIP): Warning [GenericSearch] ...Score: 1.24345 (default time 1.66209)
  35. <instantiation>:4:3: error: instruction not supported on this GPU
  36. s_mul_hi_u32 s[tiles_w], div_const_1_4, s[tiles_w]
  37. ^
  38. <instantiation>:2:2: note: while in macro instantiation
  39. _s_div_const_u32_u16 s[tiles_w], s[tiles_w], 4
  40. ^
  41. <stdin>:796:2: note: while in macro instantiation
  42. _s_ceil_u32 s[tiles_w], s[S], %xformx_f_size
  43. ^
  44. <instantiation>:4:3: error: instruction not supported on this GPU
  45. s_mul_hi_u32 s[tiles_h], div_const_1_4, s[tiles_h]
  46. ^
  47. <instantiation>:2:2: note: while in macro instantiation
  48. _s_div_const_u32_u16 s[tiles_h], s[tiles_h], 4
  49. ^
  50. <stdin>:797:2: note: while in macro instantiation
  51. _s_ceil_u32 s[tiles_h], s[R], %xformy_f_size
  52. ^
  53. <instantiation>:8:2: error: instruction not supported on this GPU
  54. v_sub_u32 v[vtmp+1+3], 0, v[vtmp+1+1]
  55. ^
  56. <stdin>:857:2: note: while in macro instantiation
  57. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  58. ^
  59. <instantiation>:12:2: error: instruction not supported on this GPU
  60. v_sub_u32 v[vtmp+1+2], v[vtmp+1], v[vtmp+1+1]
  61. ^
  62. <stdin>:857:2: note: while in macro instantiation
  63. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  64. ^
  65. <instantiation>:13:2: error: instruction not supported on this GPU
  66. v_add_co_u32 v[vtmp+1], vcc, v[vtmp+1], v[vtmp+1+1]
  67. ^
  68. <stdin>:857:2: note: while in macro instantiation
  69. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  70. ^
  71. <instantiation>:17:2: error: instruction not supported on this GPU
  72. v_sub_u32 v[vtmp+1+2], -1, v[vtmp+1+1]
  73. ^
  74. <stdin>:857:2: note: while in macro instantiation
  75. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  76. ^
  77. <instantiation>:19:2: error: instruction not supported on this GPU
  78. v_add_u32 v[vtmp+1+2], 1, v[vtmp+1]
  79. ^
  80. <stdin>:857:2: note: while in macro instantiation
  81. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  82. ^
  83. <instantiation>:20:2: error: instruction not supported on this GPU
  84. v_add_co_u32 v[vtmp+1+1], vcc, -1, v[vtmp+1]
  85. ^
  86. <stdin>:857:2: note: while in macro instantiation
  87. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  88. ^
  89. <instantiation>:22:2: error: instruction not supported on this GPU
  90. v_add_u32 v[vtmp+1+2], 1, v[vtmp+1+2]
  91. ^
  92. <stdin>:857:2: note: while in macro instantiation
  93. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  94. ^
  95. <stdin>:866:2: error: instruction not supported on this GPU
  96. v_sub_u32 v[vlds_waddr], v[tid], v[vtmp]
  97. ^
  98. <stdin>:867:2: error: instruction not supported on this GPU
  99. v_add_u32 v[vcur_tw], s[base_tile], v[vlds_waddr]
  100. ^
  101. <stdin>:877:2: error: instruction not supported on this GPU
  102. v_add_u32 v[vlds_waddr], v[vlds_waddr], v[vtmp]
  103. ^
  104. <instantiation>:20:3: error: instruction not supported on this GPU
  105. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  106. ^
  107. <stdin>:1251:2: note: while in macro instantiation
  108. .rept pipe_depth
  109. ^
  110. <instantiation>:6:3: error: instruction not supported on this GPU
  111. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  112. ^
  113. <instantiation>:24:3: note: while in macro instantiation
  114. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  115. ^
  116. <stdin>:1251:2: note: while in macro instantiation
  117. .rept pipe_depth
  118. ^
  119. <instantiation>:12:3: error: instruction not supported on this GPU
  120. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  121. ^
  122. <instantiation>:24:3: note: while in macro instantiation
  123. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  124. ^
  125. <stdin>:1251:2: note: while in macro instantiation
  126. .rept pipe_depth
  127. ^
  128. <instantiation>:18:3: error: instruction not supported on this GPU
  129. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  130. ^
  131. <instantiation>:24:3: note: while in macro instantiation
  132. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  133. ^
  134. <stdin>:1251:2: note: while in macro instantiation
  135. .rept pipe_depth
  136. ^
  137. <instantiation>:5:3: error: instruction not supported on this GPU
  138. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  139. ^
  140. <instantiation>:27:3: note: while in macro instantiation
  141. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  142. ^
  143. <stdin>:1251:2: note: while in macro instantiation
  144. .rept pipe_depth
  145. ^
  146. <instantiation>:7:3: error: instruction not supported on this GPU
  147. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  148. ^
  149. <instantiation>:27:3: note: while in macro instantiation
  150. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  151. ^
  152. <stdin>:1251:2: note: while in macro instantiation
  153. .rept pipe_depth
  154. ^
  155. <instantiation>:20:4: error: instruction not supported on this GPU
  156. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  157. ^
  158. <instantiation>:27:3: note: while in macro instantiation
  159. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  160. ^
  161. <stdin>:1251:2: note: while in macro instantiation
  162. .rept pipe_depth
  163. ^
  164. <instantiation>:21:4: error: instruction not supported on this GPU
  165. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  166. ^
  167. <instantiation>:27:3: note: while in macro instantiation
  168. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  169. ^
  170. <stdin>:1251:2: note: while in macro instantiation
  171. .rept pipe_depth
  172. ^
  173. <instantiation>:30:3: error: instruction not supported on this GPU
  174. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  175. ^
  176. <instantiation>:27:3: note: while in macro instantiation
  177. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  178. ^
  179. <stdin>:1251:2: note: while in macro instantiation
  180. .rept pipe_depth
  181. ^
  182. <instantiation>:32:3: error: instruction not supported on this GPU
  183. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  184. ^
  185. <instantiation>:27:3: note: while in macro instantiation
  186. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  187. ^
  188. <stdin>:1251:2: note: while in macro instantiation
  189. .rept pipe_depth
  190. ^
  191. <instantiation>:34:3: error: instruction not supported on this GPU
  192. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  193. ^
  194. <instantiation>:27:3: note: while in macro instantiation
  195. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  196. ^
  197. <stdin>:1251:2: note: while in macro instantiation
  198. .rept pipe_depth
  199. ^
  200. <instantiation>:51:19: error: too large value for vmcnt
  201. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  202. ^
  203. <stdin>:1251:2: note: while in macro instantiation
  204. .rept pipe_depth
  205. ^
  206. <instantiation>:51:70: error: unknown token in expression
  207. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  208. ^
  209. <stdin>:1251:2: note: while in macro instantiation
  210. .rept pipe_depth
  211. ^
  212. <instantiation>:51:70: error: failed parsing operand.
  213. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  214. ^
  215. <stdin>:1251:2: note: while in macro instantiation
  216. .rept pipe_depth
  217. ^
  218. <instantiation>:5:5: error: instruction not supported on this GPU
  219. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  220. ^
  221. <instantiation>:56:3: note: while in macro instantiation
  222. .rept in_tile_width
  223. ^
  224. <stdin>:1251:2: note: while in macro instantiation
  225. .rept pipe_depth
  226. ^
  227. <instantiation>:11:5: error: instruction not supported on this GPU
  228. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  229. ^
  230. <instantiation>:56:3: note: while in macro instantiation
  231. .rept in_tile_width
  232. ^
  233. <stdin>:1251:2: note: while in macro instantiation
  234. .rept pipe_depth
  235. ^
  236. <instantiation>:17:5: error: instruction not supported on this GPU
  237. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  238. ^
  239. <instantiation>:56:3: note: while in macro instantiation
  240. .rept in_tile_width
  241. ^
  242. <stdin>:1251:2: note: while in macro instantiation
  243. .rept pipe_depth
  244. ^
  245. <instantiation>:23:5: error: instruction not supported on this GPU
  246. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  247. ^
  248. <instantiation>:56:3: note: while in macro instantiation
  249. .rept in_tile_width
  250. ^
  251. <stdin>:1251:2: note: while in macro instantiation
  252. .rept pipe_depth
  253. ^
  254. <instantiation>:29:5: error: instruction not supported on this GPU
  255. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  256. ^
  257. <instantiation>:56:3: note: while in macro instantiation
  258. .rept in_tile_width
  259. ^
  260. <stdin>:1251:2: note: while in macro instantiation
  261. .rept pipe_depth
  262. ^
  263. <instantiation>:104:3: error: instruction not supported on this GPU
  264. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  265. ^
  266. <stdin>:1251:2: note: while in macro instantiation
  267. .rept pipe_depth
  268. ^
  269. <instantiation>:6:3: error: instruction not supported on this GPU
  270. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  271. ^
  272. <instantiation>:108:3: note: while in macro instantiation
  273. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  274. ^
  275. <stdin>:1251:2: note: while in macro instantiation
  276. .rept pipe_depth
  277. ^
  278. <instantiation>:12:3: error: instruction not supported on this GPU
  279. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  280. ^
  281. <instantiation>:108:3: note: while in macro instantiation
  282. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  283. ^
  284. <stdin>:1251:2: note: while in macro instantiation
  285. .rept pipe_depth
  286. ^
  287. <instantiation>:18:3: error: instruction not supported on this GPU
  288. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  289. ^
  290. <instantiation>:108:3: note: while in macro instantiation
  291. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  292. ^
  293. <stdin>:1251:2: note: while in macro instantiation
  294. .rept pipe_depth
  295. ^
  296. <instantiation>:5:3: error: instruction not supported on this GPU
  297. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  298. ^
  299. <instantiation>:111:3: note: while in macro instantiation
  300. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  301. ^
  302. <stdin>:1251:2: note: while in macro instantiation
  303. .rept pipe_depth
  304. ^
  305. <instantiation>:7:3: error: instruction not supported on this GPU
  306. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  307. ^
  308. <instantiation>:111:3: note: while in macro instantiation
  309. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  310. ^
  311. <stdin>:1251:2: note: while in macro instantiation
  312. .rept pipe_depth
  313. ^
  314. <instantiation>:20:4: error: instruction not supported on this GPU
  315. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  316. ^
  317. <instantiation>:111:3: note: while in macro instantiation
  318. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  319. ^
  320. <stdin>:1251:2: note: while in macro instantiation
  321. .rept pipe_depth
  322. ^
  323. <instantiation>:21:4: error: instruction not supported on this GPU
  324. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  325. ^
  326. <instantiation>:111:3: note: while in macro instantiation
  327. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  328. ^
  329. <stdin>:1251:2: note: while in macro instantiation
  330. .rept pipe_depth
  331. ^
  332. <instantiation>:30:3: error: instruction not supported on this GPU
  333. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  334. ^
  335. <instantiation>:111:3: note: while in macro instantiation
  336. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  337. ^
  338. <stdin>:1251:2: note: while in macro instantiation
  339. .rept pipe_depth
  340. ^
  341. <instantiation>:32:3: error: instruction not supported on this GPU
  342. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  343. ^
  344. <instantiation>:111:3: note: while in macro instantiation
  345. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  346. ^
  347. <stdin>:1251:2: note: while in macro instantiation
  348. .rept pipe_depth
  349. ^
  350. <instantiation>:34:3: error: instruction not supported on this GPU
  351. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  352. ^
  353. <instantiation>:111:3: note: while in macro instantiation
  354. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  355. ^
  356. <stdin>:1251:2: note: while in macro instantiation
  357. .rept pipe_depth
  358. ^
  359. <instantiation>:135:19: error: too large value for vmcnt
  360. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  361. ^
  362. <stdin>:1251:2: note: while in macro instantiation
  363. .rept pipe_depth
  364. ^
  365. <instantiation>:135:70: error: unknown token in expression
  366. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  367. ^
  368. <stdin>:1251:2: note: while in macro instantiation
  369. .rept pipe_depth
  370. ^
  371. <instantiation>:135:70: error: failed parsing operand.
  372. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  373. ^
  374. <stdin>:1251:2: note: while in macro instantiation
  375. .rept pipe_depth
  376. ^
  377. <instantiation>:5:5: error: instruction not supported on this GPU
  378. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  379. ^
  380. <instantiation>:140:3: note: while in macro instantiation
  381. .rept in_tile_width
  382. ^
  383. <stdin>:1251:2: note: while in macro instantiation
  384. .rept pipe_depth
  385. ^
  386. <instantiation>:11:5: error: instruction not supported on this GPU
  387. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  388. ^
  389. <instantiation>:140:3: note: while in macro instantiation
  390. .rept in_tile_width
  391. ^
  392. <stdin>:1251:2: note: while in macro instantiation
  393. .rept pipe_depth
  394. ^
  395. <instantiation>:17:5: error: instruction not supported on this GPU
  396. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  397. ^
  398. <instantiation>:140:3: note: while in macro instantiation
  399. .rept in_tile_width
  400. ^
  401. <stdin>:1251:2: note: while in macro instantiation
  402. .rept pipe_depth
  403. ^
  404. <instantiation>:23:5: error: instruction not supported on this GPU
  405. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  406. ^
  407. <instantiation>:140:3: note: while in macro instantiation
  408. .rept in_tile_width
  409. ^
  410. <stdin>:1251:2: note: while in macro instantiation
  411. .rept pipe_depth
  412. ^
  413. <instantiation>:29:5: error: instruction not supported on this GPU
  414. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  415. ^
  416. <instantiation>:140:3: note: while in macro instantiation
  417. .rept in_tile_width
  418. ^
  419. <stdin>:1251:2: note: while in macro instantiation
  420. .rept pipe_depth
  421. ^
  422. <instantiation>:188:3: error: instruction not supported on this GPU
  423. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  424. ^
  425. <stdin>:1251:2: note: while in macro instantiation
  426. .rept pipe_depth
  427. ^
  428. <instantiation>:6:3: error: instruction not supported on this GPU
  429. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  430. ^
  431. <instantiation>:192:3: note: while in macro instantiation
  432. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  433. ^
  434. <stdin>:1251:2: note: while in macro instantiation
  435. .rept pipe_depth
  436. ^
  437. <instantiation>:12:3: error: instruction not supported on this GPU
  438. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  439. ^
  440. <instantiation>:192:3: note: while in macro instantiation
  441. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  442. ^
  443. <stdin>:1251:2: note: while in macro instantiation
  444. .rept pipe_depth
  445. ^
  446. <instantiation>:18:3: error: instruction not supported on this GPU
  447. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  448. ^
  449. <instantiation>:192:3: note: while in macro instantiation
  450. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  451. ^
  452. <stdin>:1251:2: note: while in macro instantiation
  453. .rept pipe_depth
  454. ^
  455. <instantiation>:5:3: error: instruction not supported on this GPU
  456. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  457. ^
  458. <instantiation>:195:3: note: while in macro instantiation
  459. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  460. ^
  461. <stdin>:1251:2: note: while in macro instantiation
  462. .rept pipe_depth
  463. ^
  464. <instantiation>:7:3: error: instruction not supported on this GPU
  465. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  466. ^
  467. <instantiation>:195:3: note: while in macro instantiation
  468. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  469. ^
  470. <stdin>:1251:2: note: while in macro instantiation
  471. .rept pipe_depth
  472. ^
  473. <instantiation>:20:4: error: instruction not supported on this GPU
  474. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  475. ^
  476. <instantiation>:195:3: note: while in macro instantiation
  477. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  478. ^
  479. <stdin>:1251:2: note: while in macro instantiation
  480. .rept pipe_depth
  481. ^
  482. <instantiation>:21:4: error: instruction not supported on this GPU
  483. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  484. ^
  485. <instantiation>:195:3: note: while in macro instantiation
  486. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  487. ^
  488. <stdin>:1251:2: note: while in macro instantiation
  489. .rept pipe_depth
  490. ^
  491. <instantiation>:30:3: error: instruction not supported on this GPU
  492. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  493. ^
  494. <instantiation>:195:3: note: while in macro instantiation
  495. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  496. ^
  497. <stdin>:1251:2: note: while in macro instantiation
  498. .rept pipe_depth
  499. ^
  500. <instantiation>:32:3: error: instruction not supported on this GPU
  501. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  502. ^
  503. <instantiation>:195:3: note: while in macro instantiation
  504. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  505. ^
  506. <stdin>:1251:2: note: while in macro instantiation
  507. .rept pipe_depth
  508. ^
  509. <instantiation>:34:3: error: instruction not supported on this GPU
  510. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  511. ^
  512. <instantiation>:195:3: note: while in macro instantiation
  513. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  514. ^
  515. <stdin>:1251:2: note: while in macro instantiation
  516. .rept pipe_depth
  517. ^
  518. <instantiation>:219:19: error: too large value for vmcnt
  519. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  520. ^
  521. <stdin>:1251:2: note: while in macro instantiation
  522. .rept pipe_depth
  523. ^
  524. <instantiation>:219:70: error: unknown token in expression
  525. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  526. ^
  527. <stdin>:1251:2: note: while in macro instantiation
  528. .rept pipe_depth
  529. ^
  530. <instantiation>:219:70: error: failed parsing operand.
  531. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  532. ^
  533. <stdin>:1251:2: note: while in macro instantiation
  534. .rept pipe_depth
  535. ^
  536. <instantiation>:5:5: error: instruction not supported on this GPU
  537. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  538. ^
  539. <instantiation>:224:3: note: while in macro instantiation
  540. .rept in_tile_width
  541. ^
  542. <stdin>:1251:2: note: while in macro instantiation
  543. .rept pipe_depth
  544. ^
  545. <instantiation>:11:5: error: instruction not supported on this GPU
  546. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  547. ^
  548. <instantiation>:224:3: note: while in macro instantiation
  549. .rept in_tile_width
  550. ^
  551. <stdin>:1251:2: note: while in macro instantiation
  552. .rept pipe_depth
  553. ^
  554. <instantiation>:17:5: error: instruction not supported on this GPU
  555. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  556. ^
  557. <instantiation>:224:3: note: while in macro instantiation
  558. .rept in_tile_width
  559. ^
  560. <stdin>:1251:2: note: while in macro instantiation
  561. .rept pipe_depth
  562. ^
  563. <instantiation>:23:5: error: instruction not supported on this GPU
  564. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  565. ^
  566. <instantiation>:224:3: note: while in macro instantiation
  567. .rept in_tile_width
  568. ^
  569. <stdin>:1251:2: note: while in macro instantiation
  570. .rept pipe_depth
  571. ^
  572. <instantiation>:29:5: error: instruction not supported on this GPU
  573. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  574. ^
  575. <instantiation>:224:3: note: while in macro instantiation
  576. .rept in_tile_width
  577. ^
  578. <stdin>:1251:2: note: while in macro instantiation
  579. .rept pipe_depth
  580. ^
  581. <instantiation>:272:3: error: instruction not supported on this GPU
  582. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  583. ^
  584. <stdin>:1251:2: note: while in macro instantiation
  585. .rept pipe_depth
  586. ^
  587. <instantiation>:6:3: error: instruction not supported on this GPU
  588. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  589. ^
  590. <instantiation>:276:3: note: while in macro instantiation
  591. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  592. ^
  593. <stdin>:1251:2: note: while in macro instantiation
  594. .rept pipe_depth
  595. ^
  596. <instantiation>:12:3: error: instruction not supported on this GPU
  597. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  598. ^
  599. <instantiation>:276:3: note: while in macro instantiation
  600. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  601. ^
  602. <stdin>:1251:2: note: while in macro instantiation
  603. .rept pipe_depth
  604. ^
  605. <instantiation>:18:3: error: instruction not supported on this GPU
  606. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  607. ^
  608. <instantiation>:276:3: note: while in macro instantiation
  609. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  610. ^
  611. <stdin>:1251:2: note: while in macro instantiation
  612. .rept pipe_depth
  613. ^
  614. <instantiation>:5:3: error: instruction not supported on this GPU
  615. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  616. ^
  617. <instantiation>:279:3: note: while in macro instantiation
  618. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  619. ^
  620. <stdin>:1251:2: note: while in macro instantiation
  621. .rept pipe_depth
  622. ^
  623. <instantiation>:7:3: error: instruction not supported on this GPU
  624. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  625. ^
  626. <instantiation>:279:3: note: while in macro instantiation
  627. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  628. ^
  629. <stdin>:1251:2: note: while in macro instantiation
  630. .rept pipe_depth
  631. ^
  632. <instantiation>:20:4: error: instruction not supported on this GPU
  633. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  634. ^
  635. <instantiation>:279:3: note: while in macro instantiation
  636. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  637. ^
  638. <stdin>:1251:2: note: while in macro instantiation
  639. .rept pipe_depth
  640. ^
  641. <instantiation>:21:4: error: instruction not supported on this GPU
  642. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  643. ^
  644. <instantiation>:279:3: note: while in macro instantiation
  645. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  646. ^
  647. <stdin>:1251:2: note: while in macro instantiation
  648. .rept pipe_depth
  649. ^
  650. <instantiation>:30:3: error: instruction not supported on this GPU
  651. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  652. ^
  653. <instantiation>:279:3: note: while in macro instantiation
  654. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  655. ^
  656. <stdin>:1251:2: note: while in macro instantiation
  657. .rept pipe_depth
  658. ^
  659. <instantiation>:32:3: error: instruction not supported on this GPU
  660. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  661. ^
  662. <instantiation>:279:3: note: while in macro instantiation
  663. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  664. ^
  665. <stdin>:1251:2: note: while in macro instantiation
  666. .rept pipe_depth
  667. ^
  668. <instantiation>:34:3: error: instruction not supported on this GPU
  669. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  670. ^
  671. <instantiation>:279:3: note: while in macro instantiation
  672. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  673. ^
  674. <stdin>:1251:2: note: while in macro instantiation
  675. .rept pipe_depth
  676. ^
  677. <instantiation>:303:19: error: too large value for vmcnt
  678. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  679. ^
  680. <stdin>:1251:2: note: while in macro instantiation
  681. .rept pipe_depth
  682. ^
  683. <instantiation>:303:70: error: unknown token in expression
  684. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  685. ^
  686. <stdin>:1251:2: note: while in macro instantiation
  687. .rept pipe_depth
  688. ^
  689. <instantiation>:303:70: error: failed parsing operand.
  690. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  691. ^
  692. <stdin>:1251:2: note: while in macro instantiation
  693. .rept pipe_depth
  694. ^
  695. <instantiation>:5:5: error: instruction not supported on this GPU
  696. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  697. ^
  698. <instantiation>:308:3: note: while in macro instantiation
  699. .rept in_tile_width
  700. ^
  701. <stdin>:1251:2: note: while in macro instantiation
  702. .rept pipe_depth
  703. ^
  704. <instantiation>:11:5: error: instruction not supported on this GPU
  705. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  706. ^
  707. <instantiation>:308:3: note: while in macro instantiation
  708. .rept in_tile_width
  709. ^
  710. <stdin>:1251:2: note: while in macro instantiation
  711. .rept pipe_depth
  712. ^
  713. <instantiation>:17:5: error: instruction not supported on this GPU
  714. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  715. ^
  716. <instantiation>:308:3: note: while in macro instantiation
  717. .rept in_tile_width
  718. ^
  719. <stdin>:1251:2: note: while in macro instantiation
  720. .rept pipe_depth
  721. ^
  722. <instantiation>:23:5: error: instruction not supported on this GPU
  723. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  724. ^
  725. <instantiation>:308:3: note: while in macro instantiation
  726. .rept in_tile_width
  727. ^
  728. <stdin>:1251:2: note: while in macro instantiation
  729. .rept pipe_depth
  730. ^
  731. <instantiation>:29:5: error: instruction not supported on this GPU
  732. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  733. ^
  734. <instantiation>:308:3: note: while in macro instantiation
  735. .rept in_tile_width
  736. ^
  737. <stdin>:1251:2: note: while in macro instantiation
  738. .rept pipe_depth
  739. ^
  740. MIOpen(HIP): Warning [AmdgcnAssemble] -x assembler -target amdgcn--amdhsa -mno-code-object-v3 -Wa,-defsym,ROCM_METADATA_VERSION=4 -mcpu=gfx803 - -o /tmp/miopen-tmp-d0bc-622f-61e6-923d/amdgcn-asm-out-XXXXXX
  741. MIOpen(HIP): Warning [FindConvBwdWeightsAlgorithm] Find Winograd WrW failed:/root/driver/MLOpen/src/ocl/gcn_asm_utils.cpp:228: Assembly error(1)
  742. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  743. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  744. MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvAsmBwdWrW3x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
  745. MIOpen(HIP): Warning [GenericSearch] ConvAsmBwdWrW3x3: Searching the best solution among 896...
  746. MIOpen(HIP): Warning [Monitor] 176/0/896 4.55788, best within recent 177: 4.55788 #38 0,0,16,4,3,1, ETA:12.3311 sec.
  747. MIOpen(HIP): Warning [Monitor] 326/0/896 4.55788, best within recent 150: 6.08644 #261 0,1,16,2,14,2, ETA:10.5251 sec.
  748. MIOpen(HIP): Warning [Monitor] 471/0/896 4.55788, best within recent 145: 5.33091 #425 0,1,8,4,2,4, ETA:8.14203 sec.
  749. MIOpen(HIP): Warning [Monitor] 614/0/896 4.55788, best within recent 143: 6.54164 #525 0,1,16,2,13,4, ETA:5.52234 sec.
  750. MIOpen(HIP): Warning [Monitor] 772/0/896 4.55788, best within recent 158: 7.1362 #733 0,1,8,2,2,7, ETA:2.4138 sec.
  751. ^CMIOpen(HIP): Warning [GenericSearch] Done: 896/0/896, best #38 4.55788 0,0,16,4,3,1
  752. MIOpen(HIP): Warning [GenericSearch] ...Score: 4.76417 (default time 21.7145)
  753. <instantiation>:4:3: error: instruction not supported on this GPU
  754. s_mul_hi_u32 s[tiles_w], div_const_1_4, s[tiles_w]
  755. ^
  756. <instantiation>:2:2: note: while in macro instantiation
  757. _s_div_const_u32_u16 s[tiles_w], s[tiles_w], 4
  758. ^
  759. <stdin>:796:2: note: while in macro instantiation
  760. _s_ceil_u32 s[tiles_w], s[S], %xformx_f_size
  761. ^
  762. <instantiation>:4:3: error: instruction not supported on this GPU
  763. s_mul_hi_u32 s[tiles_h], div_const_1_4, s[tiles_h]
  764. ^
  765. <instantiation>:2:2: note: while in macro instantiation
  766. _s_div_const_u32_u16 s[tiles_h], s[tiles_h], 4
  767. ^
  768. <stdin>:797:2: note: while in macro instantiation
  769. _s_ceil_u32 s[tiles_h], s[R], %xformy_f_size
  770. ^
  771. <instantiation>:8:2: error: instruction not supported on this GPU
  772. v_sub_u32 v[vtmp+1+3], 0, v[vtmp+1+1]
  773. ^
  774. <stdin>:857:2: note: while in macro instantiation
  775. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  776. ^
  777. <instantiation>:12:2: error: instruction not supported on this GPU
  778. v_sub_u32 v[vtmp+1+2], v[vtmp+1], v[vtmp+1+1]
  779. ^
  780. <stdin>:857:2: note: while in macro instantiation
  781. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  782. ^
  783. <instantiation>:13:2: error: instruction not supported on this GPU
  784. v_add_co_u32 v[vtmp+1], vcc, v[vtmp+1], v[vtmp+1+1]
  785. ^
  786. <stdin>:857:2: note: while in macro instantiation
  787. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  788. ^
  789. <instantiation>:17:2: error: instruction not supported on this GPU
  790. v_sub_u32 v[vtmp+1+2], -1, v[vtmp+1+1]
  791. ^
  792. <stdin>:857:2: note: while in macro instantiation
  793. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  794. ^
  795. <instantiation>:19:2: error: instruction not supported on this GPU
  796. v_add_u32 v[vtmp+1+2], 1, v[vtmp+1]
  797. ^
  798. <stdin>:857:2: note: while in macro instantiation
  799. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  800. ^
  801. <instantiation>:20:2: error: instruction not supported on this GPU
  802. v_add_co_u32 v[vtmp+1+1], vcc, -1, v[vtmp+1]
  803. ^
  804. <stdin>:857:2: note: while in macro instantiation
  805. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  806. ^
  807. <instantiation>:22:2: error: instruction not supported on this GPU
  808. v_add_u32 v[vtmp+1+2], 1, v[vtmp+1+2]
  809. ^
  810. <stdin>:857:2: note: while in macro instantiation
  811. ceil_2_32_div_u16 v[vtmp], v[vtmp], vtmp+1, stmp
  812. ^
  813. <stdin>:866:2: error: instruction not supported on this GPU
  814. v_sub_u32 v[vlds_waddr], v[tid], v[vtmp]
  815. ^
  816. <stdin>:867:2: error: instruction not supported on this GPU
  817. v_add_u32 v[vcur_tw], s[base_tile], v[vlds_waddr]
  818. ^
  819. <stdin>:877:2: error: instruction not supported on this GPU
  820. v_add_u32 v[vlds_waddr], v[vlds_waddr], v[vtmp]
  821. ^
  822. <instantiation>:20:3: error: instruction not supported on this GPU
  823. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  824. ^
  825. <stdin>:1251:2: note: while in macro instantiation
  826. .rept pipe_depth
  827. ^
  828. <instantiation>:6:3: error: instruction not supported on this GPU
  829. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  830. ^
  831. <instantiation>:24:3: note: while in macro instantiation
  832. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  833. ^
  834. <stdin>:1251:2: note: while in macro instantiation
  835. .rept pipe_depth
  836. ^
  837. <instantiation>:12:3: error: instruction not supported on this GPU
  838. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  839. ^
  840. <instantiation>:24:3: note: while in macro instantiation
  841. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  842. ^
  843. <stdin>:1251:2: note: while in macro instantiation
  844. .rept pipe_depth
  845. ^
  846. <instantiation>:18:3: error: instruction not supported on this GPU
  847. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  848. ^
  849. <instantiation>:24:3: note: while in macro instantiation
  850. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  851. ^
  852. <stdin>:1251:2: note: while in macro instantiation
  853. .rept pipe_depth
  854. ^
  855. <instantiation>:5:3: error: instruction not supported on this GPU
  856. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  857. ^
  858. <instantiation>:27:3: note: while in macro instantiation
  859. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  860. ^
  861. <stdin>:1251:2: note: while in macro instantiation
  862. .rept pipe_depth
  863. ^
  864. <instantiation>:7:3: error: instruction not supported on this GPU
  865. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  866. ^
  867. <instantiation>:27:3: note: while in macro instantiation
  868. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  869. ^
  870. <stdin>:1251:2: note: while in macro instantiation
  871. .rept pipe_depth
  872. ^
  873. <instantiation>:20:4: error: instruction not supported on this GPU
  874. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  875. ^
  876. <instantiation>:27:3: note: while in macro instantiation
  877. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  878. ^
  879. <stdin>:1251:2: note: while in macro instantiation
  880. .rept pipe_depth
  881. ^
  882. <instantiation>:21:4: error: instruction not supported on this GPU
  883. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  884. ^
  885. <instantiation>:27:3: note: while in macro instantiation
  886. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  887. ^
  888. <stdin>:1251:2: note: while in macro instantiation
  889. .rept pipe_depth
  890. ^
  891. <instantiation>:30:3: error: instruction not supported on this GPU
  892. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  893. ^
  894. <instantiation>:27:3: note: while in macro instantiation
  895. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  896. ^
  897. <stdin>:1251:2: note: while in macro instantiation
  898. .rept pipe_depth
  899. ^
  900. <instantiation>:32:3: error: instruction not supported on this GPU
  901. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  902. ^
  903. <instantiation>:27:3: note: while in macro instantiation
  904. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  905. ^
  906. <stdin>:1251:2: note: while in macro instantiation
  907. .rept pipe_depth
  908. ^
  909. <instantiation>:34:3: error: instruction not supported on this GPU
  910. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  911. ^
  912. <instantiation>:27:3: note: while in macro instantiation
  913. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  914. ^
  915. <stdin>:1251:2: note: while in macro instantiation
  916. .rept pipe_depth
  917. ^
  918. <instantiation>:51:19: error: too large value for vmcnt
  919. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  920. ^
  921. <stdin>:1251:2: note: while in macro instantiation
  922. .rept pipe_depth
  923. ^
  924. <instantiation>:51:70: error: unknown token in expression
  925. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  926. ^
  927. <stdin>:1251:2: note: while in macro instantiation
  928. .rept pipe_depth
  929. ^
  930. <instantiation>:51:70: error: failed parsing operand.
  931. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  932. ^
  933. <stdin>:1251:2: note: while in macro instantiation
  934. .rept pipe_depth
  935. ^
  936. <instantiation>:5:5: error: instruction not supported on this GPU
  937. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  938. ^
  939. <instantiation>:56:3: note: while in macro instantiation
  940. .rept in_tile_width
  941. ^
  942. <stdin>:1251:2: note: while in macro instantiation
  943. .rept pipe_depth
  944. ^
  945. <instantiation>:11:5: error: instruction not supported on this GPU
  946. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  947. ^
  948. <instantiation>:56:3: note: while in macro instantiation
  949. .rept in_tile_width
  950. ^
  951. <stdin>:1251:2: note: while in macro instantiation
  952. .rept pipe_depth
  953. ^
  954. <instantiation>:17:5: error: instruction not supported on this GPU
  955. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  956. ^
  957. <instantiation>:56:3: note: while in macro instantiation
  958. .rept in_tile_width
  959. ^
  960. <stdin>:1251:2: note: while in macro instantiation
  961. .rept pipe_depth
  962. ^
  963. <instantiation>:23:5: error: instruction not supported on this GPU
  964. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  965. ^
  966. <instantiation>:56:3: note: while in macro instantiation
  967. .rept in_tile_width
  968. ^
  969. <stdin>:1251:2: note: while in macro instantiation
  970. .rept pipe_depth
  971. ^
  972. <instantiation>:29:5: error: instruction not supported on this GPU
  973. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  974. ^
  975. <instantiation>:56:3: note: while in macro instantiation
  976. .rept in_tile_width
  977. ^
  978. <stdin>:1251:2: note: while in macro instantiation
  979. .rept pipe_depth
  980. ^
  981. <instantiation>:104:3: error: instruction not supported on this GPU
  982. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  983. ^
  984. <stdin>:1251:2: note: while in macro instantiation
  985. .rept pipe_depth
  986. ^
  987. <instantiation>:6:3: error: instruction not supported on this GPU
  988. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  989. ^
  990. <instantiation>:108:3: note: while in macro instantiation
  991. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  992. ^
  993. <stdin>:1251:2: note: while in macro instantiation
  994. .rept pipe_depth
  995. ^
  996. <instantiation>:12:3: error: instruction not supported on this GPU
  997. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  998. ^
  999. <instantiation>:108:3: note: while in macro instantiation
  1000. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1001. ^
  1002. <stdin>:1251:2: note: while in macro instantiation
  1003. .rept pipe_depth
  1004. ^
  1005. <instantiation>:18:3: error: instruction not supported on this GPU
  1006. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  1007. ^
  1008. <instantiation>:108:3: note: while in macro instantiation
  1009. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1010. ^
  1011. <stdin>:1251:2: note: while in macro instantiation
  1012. .rept pipe_depth
  1013. ^
  1014. <instantiation>:5:3: error: instruction not supported on this GPU
  1015. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  1016. ^
  1017. <instantiation>:111:3: note: while in macro instantiation
  1018. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1019. ^
  1020. <stdin>:1251:2: note: while in macro instantiation
  1021. .rept pipe_depth
  1022. ^
  1023. <instantiation>:7:3: error: instruction not supported on this GPU
  1024. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  1025. ^
  1026. <instantiation>:111:3: note: while in macro instantiation
  1027. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1028. ^
  1029. <stdin>:1251:2: note: while in macro instantiation
  1030. .rept pipe_depth
  1031. ^
  1032. <instantiation>:20:4: error: instruction not supported on this GPU
  1033. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  1034. ^
  1035. <instantiation>:111:3: note: while in macro instantiation
  1036. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1037. ^
  1038. <stdin>:1251:2: note: while in macro instantiation
  1039. .rept pipe_depth
  1040. ^
  1041. <instantiation>:21:4: error: instruction not supported on this GPU
  1042. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  1043. ^
  1044. <instantiation>:111:3: note: while in macro instantiation
  1045. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1046. ^
  1047. <stdin>:1251:2: note: while in macro instantiation
  1048. .rept pipe_depth
  1049. ^
  1050. <instantiation>:30:3: error: instruction not supported on this GPU
  1051. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1052. ^
  1053. <instantiation>:111:3: note: while in macro instantiation
  1054. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1055. ^
  1056. <stdin>:1251:2: note: while in macro instantiation
  1057. .rept pipe_depth
  1058. ^
  1059. <instantiation>:32:3: error: instruction not supported on this GPU
  1060. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1061. ^
  1062. <instantiation>:111:3: note: while in macro instantiation
  1063. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1064. ^
  1065. <stdin>:1251:2: note: while in macro instantiation
  1066. .rept pipe_depth
  1067. ^
  1068. <instantiation>:34:3: error: instruction not supported on this GPU
  1069. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1070. ^
  1071. <instantiation>:111:3: note: while in macro instantiation
  1072. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1073. ^
  1074. <stdin>:1251:2: note: while in macro instantiation
  1075. .rept pipe_depth
  1076. ^
  1077. <instantiation>:135:19: error: too large value for vmcnt
  1078. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1079. ^
  1080. <stdin>:1251:2: note: while in macro instantiation
  1081. .rept pipe_depth
  1082. ^
  1083. <instantiation>:135:70: error: unknown token in expression
  1084. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1085. ^
  1086. <stdin>:1251:2: note: while in macro instantiation
  1087. .rept pipe_depth
  1088. ^
  1089. <instantiation>:135:70: error: failed parsing operand.
  1090. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1091. ^
  1092. <stdin>:1251:2: note: while in macro instantiation
  1093. .rept pipe_depth
  1094. ^
  1095. <instantiation>:5:5: error: instruction not supported on this GPU
  1096. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1097. ^
  1098. <instantiation>:140:3: note: while in macro instantiation
  1099. .rept in_tile_width
  1100. ^
  1101. <stdin>:1251:2: note: while in macro instantiation
  1102. .rept pipe_depth
  1103. ^
  1104. <instantiation>:11:5: error: instruction not supported on this GPU
  1105. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1106. ^
  1107. <instantiation>:140:3: note: while in macro instantiation
  1108. .rept in_tile_width
  1109. ^
  1110. <stdin>:1251:2: note: while in macro instantiation
  1111. .rept pipe_depth
  1112. ^
  1113. <instantiation>:17:5: error: instruction not supported on this GPU
  1114. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1115. ^
  1116. <instantiation>:140:3: note: while in macro instantiation
  1117. .rept in_tile_width
  1118. ^
  1119. <stdin>:1251:2: note: while in macro instantiation
  1120. .rept pipe_depth
  1121. ^
  1122. <instantiation>:23:5: error: instruction not supported on this GPU
  1123. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1124. ^
  1125. <instantiation>:140:3: note: while in macro instantiation
  1126. .rept in_tile_width
  1127. ^
  1128. <stdin>:1251:2: note: while in macro instantiation
  1129. .rept pipe_depth
  1130. ^
  1131. <instantiation>:29:5: error: instruction not supported on this GPU
  1132. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1133. ^
  1134. <instantiation>:140:3: note: while in macro instantiation
  1135. .rept in_tile_width
  1136. ^
  1137. <stdin>:1251:2: note: while in macro instantiation
  1138. .rept pipe_depth
  1139. ^
  1140. <instantiation>:188:3: error: instruction not supported on this GPU
  1141. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  1142. ^
  1143. <stdin>:1251:2: note: while in macro instantiation
  1144. .rept pipe_depth
  1145. ^
  1146. <instantiation>:6:3: error: instruction not supported on this GPU
  1147. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  1148. ^
  1149. <instantiation>:192:3: note: while in macro instantiation
  1150. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1151. ^
  1152. <stdin>:1251:2: note: while in macro instantiation
  1153. .rept pipe_depth
  1154. ^
  1155. <instantiation>:12:3: error: instruction not supported on this GPU
  1156. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  1157. ^
  1158. <instantiation>:192:3: note: while in macro instantiation
  1159. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1160. ^
  1161. <stdin>:1251:2: note: while in macro instantiation
  1162. .rept pipe_depth
  1163. ^
  1164. <instantiation>:18:3: error: instruction not supported on this GPU
  1165. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  1166. ^
  1167. <instantiation>:192:3: note: while in macro instantiation
  1168. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1169. ^
  1170. <stdin>:1251:2: note: while in macro instantiation
  1171. .rept pipe_depth
  1172. ^
  1173. <instantiation>:5:3: error: instruction not supported on this GPU
  1174. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  1175. ^
  1176. <instantiation>:195:3: note: while in macro instantiation
  1177. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1178. ^
  1179. <stdin>:1251:2: note: while in macro instantiation
  1180. .rept pipe_depth
  1181. ^
  1182. <instantiation>:7:3: error: instruction not supported on this GPU
  1183. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  1184. ^
  1185. <instantiation>:195:3: note: while in macro instantiation
  1186. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1187. ^
  1188. <stdin>:1251:2: note: while in macro instantiation
  1189. .rept pipe_depth
  1190. ^
  1191. <instantiation>:20:4: error: instruction not supported on this GPU
  1192. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  1193. ^
  1194. <instantiation>:195:3: note: while in macro instantiation
  1195. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1196. ^
  1197. <stdin>:1251:2: note: while in macro instantiation
  1198. .rept pipe_depth
  1199. ^
  1200. <instantiation>:21:4: error: instruction not supported on this GPU
  1201. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  1202. ^
  1203. <instantiation>:195:3: note: while in macro instantiation
  1204. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1205. ^
  1206. <stdin>:1251:2: note: while in macro instantiation
  1207. .rept pipe_depth
  1208. ^
  1209. <instantiation>:30:3: error: instruction not supported on this GPU
  1210. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1211. ^
  1212. <instantiation>:195:3: note: while in macro instantiation
  1213. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1214. ^
  1215. <stdin>:1251:2: note: while in macro instantiation
  1216. .rept pipe_depth
  1217. ^
  1218. <instantiation>:32:3: error: instruction not supported on this GPU
  1219. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1220. ^
  1221. <instantiation>:195:3: note: while in macro instantiation
  1222. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1223. ^
  1224. <stdin>:1251:2: note: while in macro instantiation
  1225. .rept pipe_depth
  1226. ^
  1227. <instantiation>:34:3: error: instruction not supported on this GPU
  1228. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1229. ^
  1230. <instantiation>:195:3: note: while in macro instantiation
  1231. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1232. ^
  1233. <stdin>:1251:2: note: while in macro instantiation
  1234. .rept pipe_depth
  1235. ^
  1236. <instantiation>:219:19: error: too large value for vmcnt
  1237. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1238. ^
  1239. <stdin>:1251:2: note: while in macro instantiation
  1240. .rept pipe_depth
  1241. ^
  1242. <instantiation>:219:70: error: unknown token in expression
  1243. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1244. ^
  1245. <stdin>:1251:2: note: while in macro instantiation
  1246. .rept pipe_depth
  1247. ^
  1248. <instantiation>:219:70: error: failed parsing operand.
  1249. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1250. ^
  1251. <stdin>:1251:2: note: while in macro instantiation
  1252. .rept pipe_depth
  1253. ^
  1254. <instantiation>:5:5: error: instruction not supported on this GPU
  1255. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1256. ^
  1257. <instantiation>:224:3: note: while in macro instantiation
  1258. .rept in_tile_width
  1259. ^
  1260. <stdin>:1251:2: note: while in macro instantiation
  1261. .rept pipe_depth
  1262. ^
  1263. <instantiation>:11:5: error: instruction not supported on this GPU
  1264. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1265. ^
  1266. <instantiation>:224:3: note: while in macro instantiation
  1267. .rept in_tile_width
  1268. ^
  1269. <stdin>:1251:2: note: while in macro instantiation
  1270. .rept pipe_depth
  1271. ^
  1272. <instantiation>:17:5: error: instruction not supported on this GPU
  1273. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1274. ^
  1275. <instantiation>:224:3: note: while in macro instantiation
  1276. .rept in_tile_width
  1277. ^
  1278. <stdin>:1251:2: note: while in macro instantiation
  1279. .rept pipe_depth
  1280. ^
  1281. <instantiation>:23:5: error: instruction not supported on this GPU
  1282. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1283. ^
  1284. <instantiation>:224:3: note: while in macro instantiation
  1285. .rept in_tile_width
  1286. ^
  1287. <stdin>:1251:2: note: while in macro instantiation
  1288. .rept pipe_depth
  1289. ^
  1290. <instantiation>:29:5: error: instruction not supported on this GPU
  1291. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1292. ^
  1293. <instantiation>:224:3: note: while in macro instantiation
  1294. .rept in_tile_width
  1295. ^
  1296. <stdin>:1251:2: note: while in macro instantiation
  1297. .rept pipe_depth
  1298. ^
  1299. <instantiation>:272:3: error: instruction not supported on this GPU
  1300. v_add_u32 v[vcur_tw], s[chw_step], v[vcur_tw]
  1301. ^
  1302. <stdin>:1251:2: note: while in macro instantiation
  1303. .rept pipe_depth
  1304. ^
  1305. <instantiation>:6:3: error: instruction not supported on this GPU
  1306. v_add_u32 v[vcur_th], v[vtmp], v[vcur_th]
  1307. ^
  1308. <instantiation>:276:3: note: while in macro instantiation
  1309. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1310. ^
  1311. <stdin>:1251:2: note: while in macro instantiation
  1312. .rept pipe_depth
  1313. ^
  1314. <instantiation>:12:3: error: instruction not supported on this GPU
  1315. v_add_u32 v[vcur_c], v[vtmp], v[vcur_c]
  1316. ^
  1317. <instantiation>:276:3: note: while in macro instantiation
  1318. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1319. ^
  1320. <stdin>:1251:2: note: while in macro instantiation
  1321. .rept pipe_depth
  1322. ^
  1323. <instantiation>:18:3: error: instruction not supported on this GPU
  1324. v_add_u32 v[vcur_n], v[vtmp], v[vcur_n]
  1325. ^
  1326. <instantiation>:276:3: note: while in macro instantiation
  1327. normalize_nchw_idx_u16 v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1328. ^
  1329. <stdin>:1251:2: note: while in macro instantiation
  1330. .rept pipe_depth
  1331. ^
  1332. <instantiation>:5:3: error: instruction not supported on this GPU
  1333. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  1334. ^
  1335. <instantiation>:279:3: note: while in macro instantiation
  1336. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1337. ^
  1338. <stdin>:1251:2: note: while in macro instantiation
  1339. .rept pipe_depth
  1340. ^
  1341. <instantiation>:7:3: error: instruction not supported on this GPU
  1342. v_add_u32 v[voff_o], v[voff_o], v[vtmp]
  1343. ^
  1344. <instantiation>:279:3: note: while in macro instantiation
  1345. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1346. ^
  1347. <stdin>:1251:2: note: while in macro instantiation
  1348. .rept pipe_depth
  1349. ^
  1350. <instantiation>:20:4: error: instruction not supported on this GPU
  1351. v_subrev_u32 v[vcur_w], s[pad_w], v[vcur_w]
  1352. ^
  1353. <instantiation>:279:3: note: while in macro instantiation
  1354. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1355. ^
  1356. <stdin>:1251:2: note: while in macro instantiation
  1357. .rept pipe_depth
  1358. ^
  1359. <instantiation>:21:4: error: instruction not supported on this GPU
  1360. v_subrev_u32 v[vtmp], s[pad_h], v[vtmp]
  1361. ^
  1362. <instantiation>:279:3: note: while in macro instantiation
  1363. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1364. ^
  1365. <stdin>:1251:2: note: while in macro instantiation
  1366. .rept pipe_depth
  1367. ^
  1368. <instantiation>:30:3: error: instruction not supported on this GPU
  1369. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1370. ^
  1371. <instantiation>:279:3: note: while in macro instantiation
  1372. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1373. ^
  1374. <stdin>:1251:2: note: while in macro instantiation
  1375. .rept pipe_depth
  1376. ^
  1377. <instantiation>:32:3: error: instruction not supported on this GPU
  1378. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1379. ^
  1380. <instantiation>:279:3: note: while in macro instantiation
  1381. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1382. ^
  1383. <stdin>:1251:2: note: while in macro instantiation
  1384. .rept pipe_depth
  1385. ^
  1386. <instantiation>:34:3: error: instruction not supported on this GPU
  1387. v_add_u32 v[voff_d], v[voff_d], v[vtmp]
  1388. ^
  1389. <instantiation>:279:3: note: while in macro instantiation
  1390. compute_voff v[voff_d], v[voff_o], v[vcur_w], v[vcur_n], v[vcur_c], v[vcur_th], v[vcur_tw], v[vtmp]
  1391. ^
  1392. <stdin>:1251:2: note: while in macro instantiation
  1393. .rept pipe_depth
  1394. ^
  1395. <instantiation>:303:19: error: too large value for vmcnt
  1396. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1397. ^
  1398. <stdin>:1251:2: note: while in macro instantiation
  1399. .rept pipe_depth
  1400. ^
  1401. <instantiation>:303:70: error: unknown token in expression
  1402. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1403. ^
  1404. <stdin>:1251:2: note: while in macro instantiation
  1405. .rept pipe_depth
  1406. ^
  1407. <instantiation>:303:70: error: failed parsing operand.
  1408. s_waitcnt vmcnt(0+(pipe_depth - 1) * (reads_per_slot+xformy_d_size))
  1409. ^
  1410. <stdin>:1251:2: note: while in macro instantiation
  1411. .rept pipe_depth
  1412. ^
  1413. <instantiation>:5:5: error: instruction not supported on this GPU
  1414. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1415. ^
  1416. <instantiation>:308:3: note: while in macro instantiation
  1417. .rept in_tile_width
  1418. ^
  1419. <stdin>:1251:2: note: while in macro instantiation
  1420. .rept pipe_depth
  1421. ^
  1422. <instantiation>:11:5: error: instruction not supported on this GPU
  1423. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1424. ^
  1425. <instantiation>:308:3: note: while in macro instantiation
  1426. .rept in_tile_width
  1427. ^
  1428. <stdin>:1251:2: note: while in macro instantiation
  1429. .rept pipe_depth
  1430. ^
  1431. <instantiation>:17:5: error: instruction not supported on this GPU
  1432. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1433. ^
  1434. <instantiation>:308:3: note: while in macro instantiation
  1435. .rept in_tile_width
  1436. ^
  1437. <stdin>:1251:2: note: while in macro instantiation
  1438. .rept pipe_depth
  1439. ^
  1440. <instantiation>:23:5: error: instruction not supported on this GPU
  1441. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1442. ^
  1443. <instantiation>:308:3: note: while in macro instantiation
  1444. .rept in_tile_width
  1445. ^
  1446. <stdin>:1251:2: note: while in macro instantiation
  1447. .rept pipe_depth
  1448. ^
  1449. <instantiation>:29:5: error: instruction not supported on this GPU
  1450. v_add_u32 v[waddrbuf+xf_slot], 1, v[waddrbuf+xf_slot]
  1451. ^
  1452. <instantiation>:308:3: note: while in macro instantiation
  1453. .rept in_tile_width
  1454. ^
  1455. <stdin>:1251:2: note: while in macro instantiation
  1456. .rept pipe_depth
  1457. ^
  1458. MIOpen(HIP): Warning [AmdgcnAssemble] -x assembler -target amdgcn--amdhsa -mno-code-object-v3 -Wa,-defsym,ROCM_METADATA_VERSION=4 -mcpu=gfx803 - -o /tmp/miopen-tmp-6860-6767-379b-2573/amdgcn-asm-out-XXXXXX
  1459. MIOpen(HIP): Warning [FindConvBwdWeightsAlgorithm] Find Winograd WrW failed:/root/driver/MLOpen/src/ocl/gcn_asm_utils.cpp:228: Assembly error(1)
  1460. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  1461. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  1462. MIOpen(HIP): Warning [FindRecordUnsafe] File is unreadable: /opt/rocm/miopen/share/miopen/db/gfx803_32.cd.pdb.txt
  1463. MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvAsmBwdWrW3x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
  1464. MIOpen(HIP): Warning [GenericSearch] ConvAsmBwdWrW3x3: Searching the best solution among 1288...
  1465. MIOpen(HIP): Warning [Monitor] 517/0/1288 0.331394, best within recent 518: 0.331394 #228 0,0,16,4,4,2, ETA:4.5566 sec.
  1466. MIOpen(HIP): Warning [Monitor] 531/0/1288 0.331394, best within recent 14: 0.483203 #520 0,0,16,4,15,3, ETA:8.6853 sec.
  1467. MIOpen(HIP): Warning [Monitor] 546/0/1288 0.331394, best within recent 15: 0.461282 #533 0,1,8,2,1,4, ETA:12.5302 sec.
  1468. MIOpen(HIP): Warning [Monitor] 555/0/1288 0.331394, best within recent 9: 0.396962 #552 0,0,16,4,2,4, ETA:16.1869 sec.
  1469. MIOpen(HIP): Warning [Monitor] 567/0/1288 0.331394, best within recent 12: 0.387682 #567 0,1,16,4,3,4, ETA:19.9995 sec.
  1470. MIOpen(HIP): Warning [Monitor] 576/0/1288 0.331394, best within recent 9: 0.435522 #574 0,0,8,2,4,4, ETA:23.2246 sec.
  1471. MIOpen(HIP): Warning [Monitor] 583/0/1288 0.331394, best within recent 7: 0.382402 #581 0,1,16,4,4,4, ETA:26.3985 sec.
  1472. MIOpen(HIP): Warning [Monitor] 595/0/1288 0.331394, best within recent 12: 0.381442 #595 0,1,16,4,5,4, ETA:29.2122 sec.
  1473. ^C^CMIOpen Error: basic_string::_M_construct null not valid
  1474. Exception ignored in: <bound method _MultiProcessingDataLoaderIter.__del__ of <torch.utils.data.dataloader._MultiProcessingDataLoaderIter object at 0x7fe7fa9f2908>>
  1475. Traceback (most recent call last):
  1476. File "/home/lucah/.local/lib/python3.6/site-packages/torch/utils/data/dataloader.py", line 925, in __del__
  1477. def __del__(self):
  1478. KeyboardInterrupt
  1479. Traceback (most recent call last):
  1480. File "BEGAN.py", line 178, in <module>
  1481. loss.backward()
  1482. File "/home/lucah/.local/lib/python3.6/site-packages/torch/tensor.py", line 119, in backward
  1483. torch.autograd.backward(self, gradient, retain_graph, create_graph)
  1484. File "/home/lucah/.local/lib/python3.6/site-packages/torch/autograd/__init__.py", line 93, in backward
  1485. allow_unreachable=True) # allow_unreachable flag
  1486. RuntimeError: miopenStatusUnknownError
  1487. ^C^C^C
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement
Public Pastes
Advertisement
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
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!