API tools faq
paste
Advertisement
Tristan9999

Untitled

Tristan9999
Jan 13th, 2024
443
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
text 54.14 KB | None | 0 0
raw download clone embed print report
  1. ---------------------------------------------------------------------------
  2. RuntimeError Traceback (most recent call last)
  3. Cell In[18], line 2
  4. 1 x = torch.ones(1, device=mps_device)
  5. ----> 2 print(x)
  6.  
  7. File ~/miniconda3/envs/pytorch/lib/python3.11/site-packages/torch/_tensor.py:461, in Tensor.__repr__(self, tensor_contents)
  8. 457 return handle_torch_function(
  9. 458 Tensor.__repr__, (self,), self, tensor_contents=tensor_contents
  10. 459 )
  11. 460 # All strings are unicode in Python 3.
  12. --> 461 return torch._tensor_str._str(self, tensor_contents=tensor_contents)
  13.  
  14. File ~/miniconda3/envs/pytorch/lib/python3.11/site-packages/torch/_tensor_str.py:677, in _str(self, tensor_contents)
  15. 675 with torch.no_grad(), torch.utils._python_dispatch._disable_current_modes():
  16. 676 guard = torch._C._DisableFuncTorch()
  17. --> 677 return _str_intern(self, tensor_contents=tensor_contents)
  18.  
  19. File ~/miniconda3/envs/pytorch/lib/python3.11/site-packages/torch/_tensor_str.py:597, in _str_intern(inp, tensor_contents)
  20. 595 tensor_str = _tensor_str(self.to_dense(), indent)
  21. 596 else:
  22. --> 597 tensor_str = _tensor_str(self, indent)
  23. 599 if self.layout != torch.strided:
  24. 600 suffixes.append("layout=" + str(self.layout))
  25.  
  26. File ~/miniconda3/envs/pytorch/lib/python3.11/site-packages/torch/_tensor_str.py:349, in _tensor_str(self, indent)
  27. 345 return _tensor_str_with_formatter(
  28. 346 self, indent, summarize, real_formatter, imag_formatter
  29. 347 )
  30. 348 else:
  31. --> 349 formatter = _Formatter(get_summarized_data(self) if summarize else self)
  32. 350 return _tensor_str_with_formatter(self, indent, summarize, formatter)
  33.  
  34. File ~/miniconda3/envs/pytorch/lib/python3.11/site-packages/torch/_tensor_str.py:137, in _Formatter.__init__(self, tensor)
  35. 134 self.max_width = max(self.max_width, len(value_str))
  36. 136 else:
  37. --> 137 nonzero_finite_vals = torch.masked_select(
  38. 138 tensor_view, torch.isfinite(tensor_view) & tensor_view.ne(0)
  39. 139 )
  40. 141 if nonzero_finite_vals.numel() == 0:
  41. 142 # no valid number, do nothing
  42. 143 return
  43.  
  44. RuntimeError: Failed to create indexing library, error: Error Domain=MTLLibraryErrorDomain Code=3 "program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  45. REGISTER_INDEX_OP_ALL_DTYPES(select);
  46. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  47. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  48. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  49. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  50. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  51. constant IDX_DTYPE * offsets [[buffer(3)]], \
  52. ^ ~~~~~~~~~
  53. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  54. program_source:160:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  55. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  56. ^
  57. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  58. REGISTER_INDEX_OP_ALL_DTYPES(select);
  59. ^
  60. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  61. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  62. ^
  63. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  64. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  65. ^
  66. <scratch space>:9:1: note: expanded from here
  67. index_select
  68. ^
  69. program_source:20:13: note: candidate template ignored: substitution failure [with T = char, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  70. kernel void index_select(
  71. ^
  72. program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  73. REGISTER_INDEX_OP_ALL_DTYPES(select);
  74. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  75. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  76. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  77. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  78. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  79. constant IDX_DTYPE * offsets [[buffer(3)]], \
  80. ^ ~~~~~~~~~
  81. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  82. program_source:162:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  83. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  84. ^
  85. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  86. REGISTER_INDEX_OP_ALL_DTYPES(select);
  87. ^
  88. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  89. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  90. ^
  91. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  92. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  93. ^
  94. <scratch space>:17:1: note: expanded from here
  95. index_select
  96. ^
  97. program_source:20:13: note: candidate template ignored: substitution failure [with T = short, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  98. kernel void index_select(
  99. ^
  100. program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  101. REGISTER_INDEX_OP_ALL_DTYPES(select);
  102. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  103. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  104. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  105. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  106. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  107. constant IDX_DTYPE * offsets [[buffer(3)]], \
  108. ^ ~~~~~~~~~
  109. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  110. program_source:164:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  111. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  112. ^
  113. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  114. REGISTER_INDEX_OP_ALL_DTYPES(select);
  115. ^
  116. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  117. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  118. ^
  119. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  120. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  121. ^
  122. <scratch space>:25:1: note: expanded from here
  123. index_select
  124. ^
  125. program_source:20:13: note: candidate template ignored: substitution failure [with T = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  126. kernel void index_select(
  127. ^
  128. program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  129. REGISTER_INDEX_OP_ALL_DTYPES(select);
  130. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  131. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  132. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  133. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  134. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  135. constant IDX_DTYPE * offsets [[buffer(3)]], \
  136. ^ ~~~~~~~~~
  137. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  138. program_source:166:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  139. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  140. ^
  141. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  142. REGISTER_INDEX_OP_ALL_DTYPES(select);
  143. ^
  144. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  145. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  146. ^
  147. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  148. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  149. ^
  150. <scratch space>:33:1: note: expanded from here
  151. index_select
  152. ^
  153. program_source:20:13: note: candidate template ignored: substitution failure [with T = long, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  154. kernel void index_select(
  155. ^
  156. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  157. REGISTER_INDEX_OP_ALL_DTYPES(put);
  158. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  159. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  160. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  161. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  162. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  163. constant IDX_DTYPE * offsets [[buffer(3)]], \
  164. ^ ~~~~~~~~~
  165. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  166. program_source:160:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  167. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  168. ^
  169. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  170. REGISTER_INDEX_OP_ALL_DTYPES(put);
  171. ^
  172. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  173. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  174. ^
  175. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  176. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  177. ^
  178. <scratch space>:41:1: note: expanded from here
  179. index_put
  180. ^
  181. program_source:111:13: note: candidate template ignored: substitution failure [with T = char, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  182. kernel void index_put(
  183. ^
  184. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  185. REGISTER_INDEX_OP_ALL_DTYPES(put);
  186. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  187. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  188. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  189. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  190. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  191. constant IDX_DTYPE * offsets [[buffer(3)]], \
  192. ^ ~~~~~~~~~
  193. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  194. program_source:162:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  195. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  196. ^
  197. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  198. REGISTER_INDEX_OP_ALL_DTYPES(put);
  199. ^
  200. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  201. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  202. ^
  203. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  204. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  205. ^
  206. <scratch space>:49:1: note: expanded from here
  207. index_put
  208. ^
  209. program_source:111:13: note: candidate template ignored: substitution failure [with T = short, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  210. kernel void index_put(
  211. ^
  212. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  213. REGISTER_INDEX_OP_ALL_DTYPES(put);
  214. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  215. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  216. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  217. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  218. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  219. constant IDX_DTYPE * offsets [[buffer(3)]], \
  220. ^ ~~~~~~~~~
  221. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  222. program_source:164:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  223. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  224. ^
  225. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  226. REGISTER_INDEX_OP_ALL_DTYPES(put);
  227. ^
  228. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  229. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  230. ^
  231. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  232. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  233. ^
  234. <scratch space>:57:1: note: expanded from here
  235. index_put
  236. ^
  237. program_source:111:13: note: candidate template ignored: substitution failure [with T = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  238. kernel void index_put(
  239. ^
  240. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  241. REGISTER_INDEX_OP_ALL_DTYPES(put);
  242. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  243. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  244. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  245. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  246. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  247. constant IDX_DTYPE * offsets [[buffer(3)]], \
  248. ^ ~~~~~~~~~
  249. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  250. program_source:166:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  251. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  252. ^
  253. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  254. REGISTER_INDEX_OP_ALL_DTYPES(put);
  255. ^
  256. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  257. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  258. ^
  259. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  260. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  261. ^
  262. <scratch space>:65:1: note: expanded from here
  263. index_put
  264. ^
  265. program_source:111:13: note: candidate template ignored: substitution failure [with T = long, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  266. kernel void index_put(
  267. ^
  268. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  269. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  270. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  271. program_source:203:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  272. REGISTER_SINGLE_THREADED_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  273. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  274. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  275. constant IDX_DTYPE * offsets [[buffer(3)]], \
  276. ^ ~~~~~~~~~
  277. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  278. program_source:203:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  279. REGISTER_SINGLE_THREADED_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  280. ^
  281. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  282. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  283. ^
  284. program_source:203:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  285. REGISTER_SINGLE_THREADED_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  286. ^
  287. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  288. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  289. ^
  290. <scratch space>:73:1: note: expanded from here
  291. index_put_serial
  292. ^
  293. program_source:87:13: note: candidate template ignored: substitution failure [with T = char, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  294. kernel void index_put_serial(
  295. ^
  296. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  297. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  298. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  299. program_source:205:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  300. REGISTER_SINGLE_THREADED_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  301. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  302. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  303. constant IDX_DTYPE * offsets [[buffer(3)]], \
  304. ^ ~~~~~~~~~
  305. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  306. program_source:205:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  307. REGISTER_SINGLE_THREADED_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  308. ^
  309. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  310. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  311. ^
  312. program_source:205:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  313. REGISTER_SINGLE_THREADED_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  314. ^
  315. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  316. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  317. ^
  318. <scratch space>:81:1: note: expanded from here
  319. index_put_serial
  320. ^
  321. program_source:87:13: note: candidate template ignored: substitution failure [with T = short, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  322. kernel void index_put_serial(
  323. ^
  324. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  325. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  326. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  327. program_source:207:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  328. REGISTER_SINGLE_THREADED_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  329. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  330. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  331. constant IDX_DTYPE * offsets [[buffer(3)]], \
  332. ^ ~~~~~~~~~
  333. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  334. program_source:207:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  335. REGISTER_SINGLE_THREADED_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  336. ^
  337. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  338. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  339. ^
  340. program_source:207:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  341. REGISTER_SINGLE_THREADED_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  342. ^
  343. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  344. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  345. ^
  346. <scratch space>:89:1: note: expanded from here
  347. index_put_serial
  348. ^
  349. program_source:87:13: note: candidate template ignored: substitution failure [with T = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  350. kernel void index_put_serial(
  351. ^
  352. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  353. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  354. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  355. program_source:209:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  356. REGISTER_SINGLE_THREADED_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  357. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  358. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  359. constant IDX_DTYPE * offsets [[buffer(3)]], \
  360. ^ ~~~~~~~~~
  361. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  362. program_source:209:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  363. REGISTER_SINGLE_THREADED_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  364. ^
  365. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  366. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  367. ^
  368. program_source:209:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  369. REGISTER_SINGLE_THREADED_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  370. ^
  371. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  372. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  373. ^
  374. <scratch space>:97:1: note: expanded from here
  375. index_put_serial
  376. ^
  377. program_source:87:13: note: candidate template ignored: substitution failure [with T = long, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  378. kernel void index_put_serial(
  379. ^
  380. program_source:242:17: error: type 'device ulong3 *' is not valid for attribute 'buffer'
  381. device ulong3 * data_offsets [[buffer(1)]],
  382. ^ ~~~~~~~~~
  383. program_source:242:24: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  384. device ulong3 * data_offsets [[buffer(1)]],
  385. ^
  386. program_source:240:13: error: explicit instantiation of 'kernel_index_offsets' does not refer to a function template, variable template, member function, member class, or static data member
  387. kernel void kernel_index_offsets<packed_uint3, ulong3>(
  388. ^
  389. program_source:214:13: note: candidate template ignored: substitution failure [with StridesT = unsigned int __attribute__((packed_vector_type(3))), DataT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long device * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  390. kernel void kernel_index_offsets(constant StridesT * strides [[buffer(0)]],
  391. ^
  392. program_source:351:5: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  393. constant ulong3 * offsets [[buffer(3)]],
  394. ^ ~~~~~~~~~
  395. program_source:351:14: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  396. constant ulong3 * offsets [[buffer(3)]],
  397. ^
  398. program_source:343:13: error: explicit instantiation of 'atomic_index_put_accumulate' does not refer to a function template, variable template, member function, member class, or static data member
  399. kernel void atomic_index_put_accumulate<float, ulong3>(
  400. ^
  401. program_source:292:13: note: candidate template ignored: substitution failure [with T = float, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  402. kernel void atomic_index_put_accumulate(
  403. ^
  404. program_source:383:5: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  405. constant ulong3 * offsets [[buffer(3)]],
  406. ^ ~~~~~~~~~
  407. program_source:383:14: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  408. constant ulong3 * offsets [[buffer(3)]],
  409. ^
  410. program_source:375:13: error: explicit instantiation of 'index_put_accumulate_native_dtypes' does not refer to a function template, variable template, member function, member class, or static data member
  411. kernel void index_put_accumulate_native_dtypes<atomic_int, int, ulong3>(
  412. ^
  413. program_source:248:13: note: candidate template ignored: substitution failure [with T = metal::_atomic<int>, E = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  414. kernel void index_put_accumulate_native_dtypes(
  415. ^
  416. " UserInfo={NSLocalizedDescription=program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  417. REGISTER_INDEX_OP_ALL_DTYPES(select);
  418. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  419. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  420. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  421. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  422. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  423. constant IDX_DTYPE * offsets [[buffer(3)]], \
  424. ^ ~~~~~~~~~
  425. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  426. program_source:160:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  427. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  428. ^
  429. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  430. REGISTER_INDEX_OP_ALL_DTYPES(select);
  431. ^
  432. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  433. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  434. ^
  435. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  436. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  437. ^
  438. <scratch space>:9:1: note: expanded from here
  439. index_select
  440. ^
  441. program_source:20:13: note: candidate template ignored: substitution failure [with T = char, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  442. kernel void index_select(
  443. ^
  444. program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  445. REGISTER_INDEX_OP_ALL_DTYPES(select);
  446. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  447. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  448. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  449. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  450. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  451. constant IDX_DTYPE * offsets [[buffer(3)]], \
  452. ^ ~~~~~~~~~
  453. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  454. program_source:162:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  455. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  456. ^
  457. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  458. REGISTER_INDEX_OP_ALL_DTYPES(select);
  459. ^
  460. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  461. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  462. ^
  463. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  464. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  465. ^
  466. <scratch space>:17:1: note: expanded from here
  467. index_select
  468. ^
  469. program_source:20:13: note: candidate template ignored: substitution failure [with T = short, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  470. kernel void index_select(
  471. ^
  472. program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  473. REGISTER_INDEX_OP_ALL_DTYPES(select);
  474. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  475. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  476. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  477. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  478. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  479. constant IDX_DTYPE * offsets [[buffer(3)]], \
  480. ^ ~~~~~~~~~
  481. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  482. program_source:164:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  483. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  484. ^
  485. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  486. REGISTER_INDEX_OP_ALL_DTYPES(select);
  487. ^
  488. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  489. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  490. ^
  491. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  492. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  493. ^
  494. <scratch space>:25:1: note: expanded from here
  495. index_select
  496. ^
  497. program_source:20:13: note: candidate template ignored: substitution failure [with T = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  498. kernel void index_select(
  499. ^
  500. program_source:168:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  501. REGISTER_INDEX_OP_ALL_DTYPES(select);
  502. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  503. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  504. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  505. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  506. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  507. constant IDX_DTYPE * offsets [[buffer(3)]], \
  508. ^ ~~~~~~~~~
  509. program_source:168:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  510. program_source:166:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  511. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  512. ^
  513. program_source:168:1: error: explicit instantiation of 'index_select' does not refer to a function template, variable template, member function, member class, or static data member
  514. REGISTER_INDEX_OP_ALL_DTYPES(select);
  515. ^
  516. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  517. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  518. ^
  519. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  520. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  521. ^
  522. <scratch space>:33:1: note: expanded from here
  523. index_select
  524. ^
  525. program_source:20:13: note: candidate template ignored: substitution failure [with T = long, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  526. kernel void index_select(
  527. ^
  528. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  529. REGISTER_INDEX_OP_ALL_DTYPES(put);
  530. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  531. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  532. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  533. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  534. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  535. constant IDX_DTYPE * offsets [[buffer(3)]], \
  536. ^ ~~~~~~~~~
  537. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  538. program_source:160:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  539. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  540. ^
  541. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  542. REGISTER_INDEX_OP_ALL_DTYPES(put);
  543. ^
  544. program_source:160:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  545. REGISTER_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  546. ^
  547. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  548. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  549. ^
  550. <scratch space>:41:1: note: expanded from here
  551. index_put
  552. ^
  553. program_source:111:13: note: candidate template ignored: substitution failure [with T = char, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  554. kernel void index_put(
  555. ^
  556. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  557. REGISTER_INDEX_OP_ALL_DTYPES(put);
  558. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  559. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  560. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  561. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  562. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  563. constant IDX_DTYPE * offsets [[buffer(3)]], \
  564. ^ ~~~~~~~~~
  565. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  566. program_source:162:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  567. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  568. ^
  569. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  570. REGISTER_INDEX_OP_ALL_DTYPES(put);
  571. ^
  572. program_source:162:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  573. REGISTER_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  574. ^
  575. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  576. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  577. ^
  578. <scratch space>:49:1: note: expanded from here
  579. index_put
  580. ^
  581. program_source:111:13: note: candidate template ignored: substitution failure [with T = short, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  582. kernel void index_put(
  583. ^
  584. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  585. REGISTER_INDEX_OP_ALL_DTYPES(put);
  586. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  587. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  588. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  589. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  590. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  591. constant IDX_DTYPE * offsets [[buffer(3)]], \
  592. ^ ~~~~~~~~~
  593. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  594. program_source:164:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  595. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  596. ^
  597. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  598. REGISTER_INDEX_OP_ALL_DTYPES(put);
  599. ^
  600. program_source:164:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  601. REGISTER_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  602. ^
  603. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  604. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  605. ^
  606. <scratch space>:57:1: note: expanded from here
  607. index_put
  608. ^
  609. program_source:111:13: note: candidate template ignored: substitution failure [with T = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  610. kernel void index_put(
  611. ^
  612. program_source:169:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  613. REGISTER_INDEX_OP_ALL_DTYPES(put);
  614. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  615. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  616. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  617. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  618. program_source:138:5: note: expanded from macro 'REGISTER_INDEX_OP'
  619. constant IDX_DTYPE * offsets [[buffer(3)]], \
  620. ^ ~~~~~~~~~
  621. program_source:169:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  622. program_source:166:59: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  623. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  624. ^
  625. program_source:169:1: error: explicit instantiation of 'index_put' does not refer to a function template, variable template, member function, member class, or static data member
  626. REGISTER_INDEX_OP_ALL_DTYPES(put);
  627. ^
  628. program_source:166:5: note: expanded from macro 'REGISTER_INDEX_OP_ALL_DTYPES'
  629. REGISTER_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  630. ^
  631. program_source:134:13: note: expanded from macro 'REGISTER_INDEX_OP'
  632. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  633. ^
  634. <scratch space>:65:1: note: expanded from here
  635. index_put
  636. ^
  637. program_source:111:13: note: candidate template ignored: substitution failure [with T = long, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  638. kernel void index_put(
  639. ^
  640. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  641. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  642. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  643. program_source:203:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  644. REGISTER_SINGLE_THREADED_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  645. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  646. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  647. constant IDX_DTYPE * offsets [[buffer(3)]], \
  648. ^ ~~~~~~~~~
  649. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  650. program_source:203:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  651. REGISTER_SINGLE_THREADED_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  652. ^
  653. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  654. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  655. ^
  656. program_source:203:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  657. REGISTER_SINGLE_THREADED_INDEX_OP(8bit, idx64, char, INDEX_OP_TYPE, ulong3); \
  658. ^
  659. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  660. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  661. ^
  662. <scratch space>:73:1: note: expanded from here
  663. index_put_serial
  664. ^
  665. program_source:87:13: note: candidate template ignored: substitution failure [with T = char, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  666. kernel void index_put_serial(
  667. ^
  668. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  669. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  670. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  671. program_source:205:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  672. REGISTER_SINGLE_THREADED_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  673. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  674. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  675. constant IDX_DTYPE * offsets [[buffer(3)]], \
  676. ^ ~~~~~~~~~
  677. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  678. program_source:205:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  679. REGISTER_SINGLE_THREADED_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  680. ^
  681. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  682. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  683. ^
  684. program_source:205:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  685. REGISTER_SINGLE_THREADED_INDEX_OP(16bit, idx64, short, INDEX_OP_TYPE, ulong3); \
  686. ^
  687. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  688. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  689. ^
  690. <scratch space>:81:1: note: expanded from here
  691. index_put_serial
  692. ^
  693. program_source:87:13: note: candidate template ignored: substitution failure [with T = short, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  694. kernel void index_put_serial(
  695. ^
  696. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  697. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  698. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  699. program_source:207:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  700. REGISTER_SINGLE_THREADED_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  701. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  702. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  703. constant IDX_DTYPE * offsets [[buffer(3)]], \
  704. ^ ~~~~~~~~~
  705. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  706. program_source:207:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  707. REGISTER_SINGLE_THREADED_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  708. ^
  709. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  710. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  711. ^
  712. program_source:207:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  713. REGISTER_SINGLE_THREADED_INDEX_OP(32bit, idx64, int, INDEX_OP_TYPE, ulong3); \
  714. ^
  715. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  716. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  717. ^
  718. <scratch space>:89:1: note: expanded from here
  719. index_put_serial
  720. ^
  721. program_source:87:13: note: candidate template ignored: substitution failure [with T = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  722. kernel void index_put_serial(
  723. ^
  724. program_source:211:1: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  725. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  726. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  727. program_source:209:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  728. REGISTER_SINGLE_THREADED_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  729. ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
  730. program_source:179:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  731. constant IDX_DTYPE * offsets [[buffer(3)]], \
  732. ^ ~~~~~~~~~
  733. program_source:211:1: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  734. program_source:209:75: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  735. REGISTER_SINGLE_THREADED_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  736. ^
  737. program_source:211:1: error: explicit instantiation of 'index_put_serial' does not refer to a function template, variable template, member function, member class, or static data member
  738. REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES(put_serial);
  739. ^
  740. program_source:209:5: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP_ALL_DTYPES'
  741. REGISTER_SINGLE_THREADED_INDEX_OP(64bit, idx64, long, INDEX_OP_TYPE, ulong3);
  742. ^
  743. program_source:175:13: note: expanded from macro 'REGISTER_SINGLE_THREADED_INDEX_OP'
  744. kernel void index_ ## INDEX_OP_TYPE<DTYPE, IDX_DTYPE>( \
  745. ^
  746. <scratch space>:97:1: note: expanded from here
  747. index_put_serial
  748. ^
  749. program_source:87:13: note: candidate template ignored: substitution failure [with T = long, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  750. kernel void index_put_serial(
  751. ^
  752. program_source:242:17: error: type 'device ulong3 *' is not valid for attribute 'buffer'
  753. device ulong3 * data_offsets [[buffer(1)]],
  754. ^ ~~~~~~~~~
  755. program_source:242:24: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  756. device ulong3 * data_offsets [[buffer(1)]],
  757. ^
  758. program_source:240:13: error: explicit instantiation of 'kernel_index_offsets' does not refer to a function template, variable template, member function, member class, or static data member
  759. kernel void kernel_index_offsets<packed_uint3, ulong3>(
  760. ^
  761. program_source:214:13: note: candidate template ignored: substitution failure [with StridesT = unsigned int __attribute__((packed_vector_type(3))), DataT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long device * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  762. kernel void kernel_index_offsets(constant StridesT * strides [[buffer(0)]],
  763. ^
  764. program_source:351:5: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  765. constant ulong3 * offsets [[buffer(3)]],
  766. ^ ~~~~~~~~~
  767. program_source:351:14: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  768. constant ulong3 * offsets [[buffer(3)]],
  769. ^
  770. program_source:343:13: error: explicit instantiation of 'atomic_index_put_accumulate' does not refer to a function template, variable template, member function, member class, or static data member
  771. kernel void atomic_index_put_accumulate<float, ulong3>(
  772. ^
  773. program_source:292:13: note: candidate template ignored: substitution failure [with T = float, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  774. kernel void atomic_index_put_accumulate(
  775. ^
  776. program_source:383:5: error: type 'const constant ulong3 *' is not valid for attribute 'buffer'
  777. constant ulong3 * offsets [[buffer(3)]],
  778. ^ ~~~~~~~~~
  779. program_source:383:14: note: type 'ulong3' (vector of 3 'unsigned long' values) cannot be used in buffer pointee type
  780. constant ulong3 * offsets [[buffer(3)]],
  781. ^
  782. program_source:375:13: error: explicit instantiation of 'index_put_accumulate_native_dtypes' does not refer to a function template, variable template, member function, member class, or static data member
  783. kernel void index_put_accumulate_native_dtypes<atomic_int, int, ulong3>(
  784. ^
  785. program_source:248:13: note: candidate template ignored: substitution failure [with T = metal::_atomic<int>, E = int, OffsetsT = unsigned long __attribute__((ext_vector_type(3)))]: type 'unsigned long const constant * __attribute__((ext_vector_type(3)))' is not valid for attribute 'buffer'
  786. kernel void index_put_accumulate_native_dtypes(
  787. ^
  788. }
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!