libdevice.py 55 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735736737738739740741742743744745746747748749750751752753754755756757758759760761762763764765766767768769770771772773774775776777778779780781782783784785786787788789790791792793794795796797798799800801802803804805806807808809810811812813814815816817818819820821822823824825826827828829830831832833834835836837838839840841842843844845846847848849850851852853854855856857858859860861862863864865866867868869870871872873874875876877878879880881882883884885886887888889890891892893894895896897898899900901902903904905906907908909910911912913914915916917918919920921922923924925926927928929930931932933934935936937938939940941942943944945946947948949950951952953954955956957958959960961962963964965966967968969970971972973974975976977978979980981982983984985986987988989990991992993994995996997998999100010011002100310041005100610071008100910101011101210131014101510161017101810191020102110221023102410251026102710281029103010311032103310341035103610371038103910401041104210431044104510461047104810491050105110521053105410551056105710581059106010611062106310641065106610671068106910701071107210731074107510761077107810791080108110821083108410851086108710881089109010911092109310941095109610971098109911001101110211031104110511061107110811091110111111121113111411151116111711181119112011211122112311241125112611271128112911301131113211331134113511361137113811391140114111421143114411451146114711481149115011511152115311541155115611571158115911601161116211631164116511661167116811691170117111721173117411751176117711781179118011811182118311841185118611871188118911901191119211931194119511961197119811991200120112021203120412051206120712081209121012111212121312141215121612171218121912201221122212231224122512261227122812291230123112321233123412351236123712381239124012411242124312441245124612471248124912501251125212531254125512561257125812591260126112621263126412651266126712681269127012711272127312741275127612771278127912801281128212831284128512861287128812891290129112921293129412951296129712981299130013011302130313041305130613071308130913101311131213131314131513161317131813191320132113221323132413251326132713281329133013311332133313341335133613371338133913401341134213431344134513461347134813491350135113521353135413551356135713581359136013611362136313641365136613671368136913701371137213731374137513761377137813791380138113821383138413851386138713881389139013911392139313941395139613971398139914001401140214031404140514061407140814091410141114121413141414151416141714181419142014211422142314241425142614271428142914301431143214331434143514361437143814391440144114421443144414451446144714481449145014511452145314541455145614571458145914601461146214631464146514661467146814691470147114721473147414751476147714781479148014811482148314841485148614871488148914901491149214931494149514961497149814991500150115021503150415051506150715081509151015111512151315141515151615171518151915201521152215231524152515261527152815291530153115321533153415351536153715381539154015411542154315441545154615471548154915501551155215531554155515561557155815591560156115621563156415651566156715681569157015711572157315741575157615771578157915801581158215831584158515861587158815891590159115921593159415951596159715981599160016011602160316041605160616071608160916101611161216131614161516161617161816191620162116221623162416251626162716281629
  1. from triton.language import core
  2. @core.extern
  3. def clz(arg0, _semantic=None):
  4. return core.extern_elementwise(
  5. "", "", [arg0], {
  6. (core.dtype("int32"), ): ("__nv_clz", core.dtype("int32")),
  7. (core.dtype("int64"), ): ("__nv_clzll", core.dtype("int32")),
  8. }, is_pure=True, _semantic=_semantic)
  9. @core.extern
  10. def popc(arg0, _semantic=None):
  11. return core.extern_elementwise(
  12. "", "", [arg0], {
  13. (core.dtype("int32"), ): ("__nv_popc", core.dtype("int32")),
  14. (core.dtype("int64"), ): ("__nv_popcll", core.dtype("int32")),
  15. }, is_pure=True, _semantic=_semantic)
  16. @core.extern
  17. def byte_perm(arg0, arg1, arg2, _semantic=None):
  18. return core.extern_elementwise("", "", [arg0, arg1, arg2], {
  19. (core.dtype("int32"), core.dtype("int32"), core.dtype("int32")): ("__nv_byte_perm", core.dtype("int32")),
  20. }, is_pure=True, _semantic=_semantic)
  21. @core.extern
  22. def mulhi(arg0, arg1, _semantic=None):
  23. return core.extern_elementwise(
  24. "", "", [arg0, arg1], {
  25. (core.dtype("int32"), core.dtype("int32")): ("__nv_mulhi", core.dtype("int32")),
  26. (core.dtype("uint32"), core.dtype("uint32")): ("__nv_umulhi", core.dtype("uint32")),
  27. (core.dtype("int64"), core.dtype("int64")): ("__nv_mul64hi", core.dtype("int64")),
  28. (core.dtype("uint64"), core.dtype("uint64")): ("__nv_umul64hi", core.dtype("uint64")),
  29. }, is_pure=True, _semantic=_semantic)
  30. @core.extern
  31. def mul24(arg0, arg1, _semantic=None):
  32. return core.extern_elementwise(
  33. "", "", [arg0, arg1], {
  34. (core.dtype("int32"), core.dtype("int32")): ("__nv_mul24", core.dtype("int32")),
  35. (core.dtype("uint32"), core.dtype("uint32")): ("__nv_umul24", core.dtype("uint32")),
  36. }, is_pure=True, _semantic=_semantic)
  37. @core.extern
  38. def brev(arg0, _semantic=None):
  39. return core.extern_elementwise(
  40. "", "", [arg0], {
  41. (core.dtype("int32"), ): ("__nv_brev", core.dtype("int32")),
  42. (core.dtype("int64"), ): ("__nv_brevll", core.dtype("int64")),
  43. }, is_pure=True, _semantic=_semantic)
  44. @core.extern
  45. def sad(arg0, arg1, arg2, _semantic=None):
  46. return core.extern_elementwise(
  47. "", "", [arg0, arg1, arg2], {
  48. (core.dtype("int32"), core.dtype("int32"), core.dtype("uint32")): ("__nv_sad", core.dtype("int32")),
  49. (core.dtype("uint32"), core.dtype("uint32"), core.dtype("uint32")): ("__nv_usad", core.dtype("uint32")),
  50. }, is_pure=True, _semantic=_semantic)
  51. @core.extern
  52. def abs(arg0, _semantic=None):
  53. return core.extern_elementwise(
  54. "", "", [arg0], {
  55. (core.dtype("int32"), ): ("__nv_abs", core.dtype("int32")),
  56. (core.dtype("int64"), ): ("__nv_llabs", core.dtype("int64")),
  57. (core.dtype("fp32"), ): ("__nv_fabsf", core.dtype("fp32")),
  58. (core.dtype("fp64"), ): ("__nv_fabs", core.dtype("fp64")),
  59. }, is_pure=True, _semantic=_semantic)
  60. @core.extern
  61. def floor(arg0, _semantic=None):
  62. return core.extern_elementwise(
  63. "", "", [arg0], {
  64. (core.dtype("fp32"), ): ("__nv_floorf", core.dtype("fp32")),
  65. (core.dtype("fp64"), ): ("__nv_floor", core.dtype("fp64")),
  66. }, is_pure=True, _semantic=_semantic)
  67. @core.extern
  68. def rcp64h(arg0, _semantic=None):
  69. return core.extern_elementwise("", "", [arg0], {
  70. (core.dtype("fp64"), ): ("__nv_rcp64h", core.dtype("fp64")),
  71. }, is_pure=True, _semantic=_semantic)
  72. @core.extern
  73. def rsqrt(arg0, _semantic=None):
  74. return core.extern_elementwise(
  75. "", "", [arg0], {
  76. (core.dtype("fp32"), ): ("__nv_rsqrtf", core.dtype("fp32")),
  77. (core.dtype("fp64"), ): ("__nv_rsqrt", core.dtype("fp64")),
  78. }, is_pure=True, _semantic=_semantic)
  79. @core.extern
  80. def ceil(arg0, _semantic=None):
  81. return core.extern_elementwise(
  82. "", "", [arg0], {
  83. (core.dtype("fp64"), ): ("__nv_ceil", core.dtype("fp64")),
  84. (core.dtype("fp32"), ): ("__nv_ceilf", core.dtype("fp32")),
  85. }, is_pure=True, _semantic=_semantic)
  86. @core.extern
  87. def trunc(arg0, _semantic=None):
  88. return core.extern_elementwise(
  89. "", "", [arg0], {
  90. (core.dtype("fp64"), ): ("__nv_trunc", core.dtype("fp64")),
  91. (core.dtype("fp32"), ): ("__nv_truncf", core.dtype("fp32")),
  92. }, is_pure=True, _semantic=_semantic)
  93. @core.extern
  94. def exp2(arg0, _semantic=None):
  95. return core.extern_elementwise(
  96. "", "", [arg0], {
  97. (core.dtype("fp32"), ): ("__nv_exp2f", core.dtype("fp32")),
  98. (core.dtype("fp64"), ): ("__nv_exp2", core.dtype("fp64")),
  99. }, is_pure=True, _semantic=_semantic)
  100. @core.extern
  101. def saturatef(arg0, _semantic=None):
  102. return core.extern_elementwise("", "", [arg0], {
  103. (core.dtype("fp32"), ): ("__nv_saturatef", core.dtype("fp32")),
  104. }, is_pure=True, _semantic=_semantic)
  105. @core.extern
  106. def fma_rn(arg0, arg1, arg2, _semantic=None):
  107. return core.extern_elementwise(
  108. "", "", [arg0, arg1, arg2], {
  109. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_rn", core.dtype("fp32")),
  110. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_rn", core.dtype("fp64")),
  111. }, is_pure=True, _semantic=_semantic)
  112. @core.extern
  113. def fma_rz(arg0, arg1, arg2, _semantic=None):
  114. return core.extern_elementwise(
  115. "", "", [arg0, arg1, arg2], {
  116. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_rz", core.dtype("fp32")),
  117. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_rz", core.dtype("fp64")),
  118. }, is_pure=True, _semantic=_semantic)
  119. @core.extern
  120. def fma_rd(arg0, arg1, arg2, _semantic=None):
  121. return core.extern_elementwise(
  122. "", "", [arg0, arg1, arg2], {
  123. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_rd", core.dtype("fp32")),
  124. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_rd", core.dtype("fp64")),
  125. }, is_pure=True, _semantic=_semantic)
  126. @core.extern
  127. def fma_ru(arg0, arg1, arg2, _semantic=None):
  128. return core.extern_elementwise(
  129. "", "", [arg0, arg1, arg2], {
  130. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_ru", core.dtype("fp32")),
  131. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_ru", core.dtype("fp64")),
  132. }, is_pure=True, _semantic=_semantic)
  133. @core.extern
  134. def fast_dividef(arg0, arg1, _semantic=None):
  135. return core.extern_elementwise("", "", [arg0, arg1], {
  136. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fast_fdividef", core.dtype("fp32")),
  137. }, is_pure=True, _semantic=_semantic)
  138. @core.extern
  139. def div_rn(arg0, arg1, _semantic=None):
  140. return core.extern_elementwise(
  141. "", "", [arg0, arg1], {
  142. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_rn", core.dtype("fp32")),
  143. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_rn", core.dtype("fp64")),
  144. }, is_pure=True, _semantic=_semantic)
  145. @core.extern
  146. def div_rz(arg0, arg1, _semantic=None):
  147. return core.extern_elementwise(
  148. "", "", [arg0, arg1], {
  149. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_rz", core.dtype("fp32")),
  150. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_rz", core.dtype("fp64")),
  151. }, is_pure=True, _semantic=_semantic)
  152. @core.extern
  153. def div_rd(arg0, arg1, _semantic=None):
  154. return core.extern_elementwise(
  155. "", "", [arg0, arg1], {
  156. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_rd", core.dtype("fp32")),
  157. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_rd", core.dtype("fp64")),
  158. }, is_pure=True, _semantic=_semantic)
  159. @core.extern
  160. def div_ru(arg0, arg1, _semantic=None):
  161. return core.extern_elementwise(
  162. "", "", [arg0, arg1], {
  163. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_ru", core.dtype("fp32")),
  164. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_ru", core.dtype("fp64")),
  165. }, is_pure=True, _semantic=_semantic)
  166. @core.extern
  167. def rcp_rn(arg0, _semantic=None):
  168. return core.extern_elementwise(
  169. "", "", [arg0], {
  170. (core.dtype("fp32"), ): ("__nv_frcp_rn", core.dtype("fp32")),
  171. (core.dtype("fp64"), ): ("__nv_drcp_rn", core.dtype("fp64")),
  172. }, is_pure=True, _semantic=_semantic)
  173. @core.extern
  174. def rcp_rz(arg0, _semantic=None):
  175. return core.extern_elementwise(
  176. "", "", [arg0], {
  177. (core.dtype("fp32"), ): ("__nv_frcp_rz", core.dtype("fp32")),
  178. (core.dtype("fp64"), ): ("__nv_drcp_rz", core.dtype("fp64")),
  179. }, is_pure=True, _semantic=_semantic)
  180. @core.extern
  181. def rcp_rd(arg0, _semantic=None):
  182. return core.extern_elementwise(
  183. "", "", [arg0], {
  184. (core.dtype("fp32"), ): ("__nv_frcp_rd", core.dtype("fp32")),
  185. (core.dtype("fp64"), ): ("__nv_drcp_rd", core.dtype("fp64")),
  186. }, is_pure=True, _semantic=_semantic)
  187. @core.extern
  188. def rcp_ru(arg0, _semantic=None):
  189. return core.extern_elementwise(
  190. "", "", [arg0], {
  191. (core.dtype("fp32"), ): ("__nv_frcp_ru", core.dtype("fp32")),
  192. (core.dtype("fp64"), ): ("__nv_drcp_ru", core.dtype("fp64")),
  193. }, is_pure=True, _semantic=_semantic)
  194. @core.extern
  195. def sqrt_rn(arg0, _semantic=None):
  196. return core.extern_elementwise(
  197. "", "", [arg0], {
  198. (core.dtype("fp32"), ): ("__nv_fsqrt_rn", core.dtype("fp32")),
  199. (core.dtype("fp64"), ): ("__nv_dsqrt_rn", core.dtype("fp64")),
  200. }, is_pure=True, _semantic=_semantic)
  201. @core.extern
  202. def sqrt_rz(arg0, _semantic=None):
  203. return core.extern_elementwise(
  204. "", "", [arg0], {
  205. (core.dtype("fp32"), ): ("__nv_fsqrt_rz", core.dtype("fp32")),
  206. (core.dtype("fp64"), ): ("__nv_dsqrt_rz", core.dtype("fp64")),
  207. }, is_pure=True, _semantic=_semantic)
  208. @core.extern
  209. def sqrt_rd(arg0, _semantic=None):
  210. return core.extern_elementwise(
  211. "", "", [arg0], {
  212. (core.dtype("fp32"), ): ("__nv_fsqrt_rd", core.dtype("fp32")),
  213. (core.dtype("fp64"), ): ("__nv_dsqrt_rd", core.dtype("fp64")),
  214. }, is_pure=True, _semantic=_semantic)
  215. @core.extern
  216. def sqrt_ru(arg0, _semantic=None):
  217. return core.extern_elementwise(
  218. "", "", [arg0], {
  219. (core.dtype("fp32"), ): ("__nv_fsqrt_ru", core.dtype("fp32")),
  220. (core.dtype("fp64"), ): ("__nv_dsqrt_ru", core.dtype("fp64")),
  221. }, is_pure=True, _semantic=_semantic)
  222. @core.extern
  223. def sqrt(arg0, _semantic=None):
  224. return core.extern_elementwise(
  225. "", "", [arg0], {
  226. (core.dtype("fp32"), ): ("__nv_sqrtf", core.dtype("fp32")),
  227. (core.dtype("fp64"), ): ("__nv_sqrt", core.dtype("fp64")),
  228. }, is_pure=True, _semantic=_semantic)
  229. @core.extern
  230. def add_rn(arg0, arg1, _semantic=None):
  231. return core.extern_elementwise(
  232. "", "", [arg0, arg1], {
  233. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_rn", core.dtype("fp64")),
  234. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_rn", core.dtype("fp32")),
  235. }, is_pure=True, _semantic=_semantic)
  236. @core.extern
  237. def add_rz(arg0, arg1, _semantic=None):
  238. return core.extern_elementwise(
  239. "", "", [arg0, arg1], {
  240. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_rz", core.dtype("fp64")),
  241. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_rz", core.dtype("fp32")),
  242. }, is_pure=True, _semantic=_semantic)
  243. @core.extern
  244. def add_rd(arg0, arg1, _semantic=None):
  245. return core.extern_elementwise(
  246. "", "", [arg0, arg1], {
  247. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_rd", core.dtype("fp64")),
  248. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_rd", core.dtype("fp32")),
  249. }, is_pure=True, _semantic=_semantic)
  250. @core.extern
  251. def add_ru(arg0, arg1, _semantic=None):
  252. return core.extern_elementwise(
  253. "", "", [arg0, arg1], {
  254. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_ru", core.dtype("fp64")),
  255. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_ru", core.dtype("fp32")),
  256. }, is_pure=True, _semantic=_semantic)
  257. @core.extern
  258. def mul_rn(arg0, arg1, _semantic=None):
  259. return core.extern_elementwise(
  260. "", "", [arg0, arg1], {
  261. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dmul_rn", core.dtype("fp64")),
  262. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmul_rn", core.dtype("fp32")),
  263. }, is_pure=True, _semantic=_semantic)
  264. @core.extern
  265. def mul_rz(arg0, arg1, _semantic=None):
  266. return core.extern_elementwise(
  267. "", "", [arg0, arg1], {
  268. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dmul_rz", core.dtype("fp64")),
  269. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmul_rz", core.dtype("fp32")),
  270. }, is_pure=True, _semantic=_semantic)
  271. @core.extern
  272. def mul_rd(arg0, arg1, _semantic=None):
  273. return core.extern_elementwise(
  274. "", "", [arg0, arg1], {
  275. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dmul_rd", core.dtype("fp64")),
  276. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmul_rd", core.dtype("fp32")),
  277. }, is_pure=True, _semantic=_semantic)
  278. @core.extern
  279. def mul_ru(arg0, arg1, _semantic=None):
  280. return core.extern_elementwise(
  281. "", "", [
  282. arg0,
  283. arg1,
  284. ], {
  285. (
  286. core.dtype("fp64"),
  287. core.dtype("fp64"),
  288. ): ("__nv_dmul_ru", core.dtype("fp64")),
  289. (
  290. core.dtype("fp32"),
  291. core.dtype("fp32"),
  292. ): ("__nv_fmul_ru", core.dtype("fp32")),
  293. }, is_pure=True, _semantic=_semantic)
  294. @core.extern
  295. def double2float_rn(arg0, _semantic=None):
  296. return core.extern_elementwise("", "", [arg0], {
  297. (core.dtype("fp64"), ): ("__nv_double2float_rn", core.dtype("fp32")),
  298. }, is_pure=True, _semantic=_semantic)
  299. @core.extern
  300. def double2float_rz(arg0, _semantic=None):
  301. return core.extern_elementwise("", "", [arg0], {
  302. (core.dtype("fp64"), ): ("__nv_double2float_rz", core.dtype("fp32")),
  303. }, is_pure=True, _semantic=_semantic)
  304. @core.extern
  305. def double2float_rd(arg0, _semantic=None):
  306. return core.extern_elementwise("", "", [arg0], {
  307. (core.dtype("fp64"), ): ("__nv_double2float_rd", core.dtype("fp32")),
  308. }, is_pure=True, _semantic=_semantic)
  309. @core.extern
  310. def double2float_ru(arg0, _semantic=None):
  311. return core.extern_elementwise("", "", [arg0], {
  312. (core.dtype("fp64"), ): ("__nv_double2float_ru", core.dtype("fp32")),
  313. }, is_pure=True, _semantic=_semantic)
  314. @core.extern
  315. def double2int_rn(arg0, _semantic=None):
  316. return core.extern_elementwise("", "", [arg0], {
  317. (core.dtype("fp64"), ): ("__nv_double2int_rn", core.dtype("int32")),
  318. }, is_pure=True, _semantic=_semantic)
  319. @core.extern
  320. def double2int_rz(arg0, _semantic=None):
  321. return core.extern_elementwise("", "", [arg0], {
  322. (core.dtype("fp64"), ): ("__nv_double2int_rz", core.dtype("int32")),
  323. }, is_pure=True, _semantic=_semantic)
  324. @core.extern
  325. def double2int_rd(arg0, _semantic=None):
  326. return core.extern_elementwise("", "", [arg0], {
  327. (core.dtype("fp64"), ): ("__nv_double2int_rd", core.dtype("int32")),
  328. }, is_pure=True, _semantic=_semantic)
  329. @core.extern
  330. def double2int_ru(arg0, _semantic=None):
  331. return core.extern_elementwise("", "", [arg0], {
  332. (core.dtype("fp64"), ): ("__nv_double2int_ru", core.dtype("int32")),
  333. }, is_pure=True, _semantic=_semantic)
  334. @core.extern
  335. def double2uint_rn(arg0, _semantic=None):
  336. return core.extern_elementwise("", "", [arg0], {
  337. (core.dtype("fp64"), ): ("__nv_double2uint_rn", core.dtype("int32")),
  338. }, is_pure=True, _semantic=_semantic)
  339. @core.extern
  340. def double2uint_rz(arg0, _semantic=None):
  341. return core.extern_elementwise("", "", [arg0], {
  342. (core.dtype("fp64"), ): ("__nv_double2uint_rz", core.dtype("int32")),
  343. }, is_pure=True, _semantic=_semantic)
  344. @core.extern
  345. def double2uint_rd(arg0, _semantic=None):
  346. return core.extern_elementwise("", "", [arg0], {
  347. (core.dtype("fp64"), ): ("__nv_double2uint_rd", core.dtype("int32")),
  348. }, is_pure=True, _semantic=_semantic)
  349. @core.extern
  350. def double2uint_ru(arg0, _semantic=None):
  351. return core.extern_elementwise("", "", [arg0], {
  352. (core.dtype("fp64"), ): ("__nv_double2uint_ru", core.dtype("int32")),
  353. }, is_pure=True, _semantic=_semantic)
  354. @core.extern
  355. def int2double_rn(arg0, _semantic=None):
  356. return core.extern_elementwise("", "", [arg0], {
  357. (core.dtype("int32"), ): ("__nv_int2double_rn", core.dtype("fp64")),
  358. }, is_pure=True, _semantic=_semantic)
  359. @core.extern
  360. def uint2double_rn(arg0, _semantic=None):
  361. return core.extern_elementwise("", "", [arg0], {
  362. (core.dtype("uint32"), ): ("__nv_uint2double_rn", core.dtype("fp64")),
  363. }, is_pure=True, _semantic=_semantic)
  364. @core.extern
  365. def float2int_rn(arg0, _semantic=None):
  366. return core.extern_elementwise("", "", [arg0], {
  367. (core.dtype("fp32"), ): ("__nv_float2int_rn", core.dtype("int32")),
  368. }, is_pure=True, _semantic=_semantic)
  369. @core.extern
  370. def float2int_rz(arg0, _semantic=None):
  371. return core.extern_elementwise("", "", [arg0], {
  372. (core.dtype("fp32"), ): ("__nv_float2int_rz", core.dtype("int32")),
  373. }, is_pure=True, _semantic=_semantic)
  374. @core.extern
  375. def float2int_rd(arg0, _semantic=None):
  376. return core.extern_elementwise("", "", [arg0], {
  377. (core.dtype("fp32"), ): ("__nv_float2int_rd", core.dtype("int32")),
  378. }, is_pure=True, _semantic=_semantic)
  379. @core.extern
  380. def float2int_ru(arg0, _semantic=None):
  381. return core.extern_elementwise("", "", [arg0], {
  382. (core.dtype("fp32"), ): ("__nv_float2int_ru", core.dtype("int32")),
  383. }, is_pure=True, _semantic=_semantic)
  384. @core.extern
  385. def float2uint_rn(arg0, _semantic=None):
  386. return core.extern_elementwise("", "", [arg0], {
  387. (core.dtype("fp32"), ): ("__nv_float2uint_rn", core.dtype("int32")),
  388. }, is_pure=True, _semantic=_semantic)
  389. @core.extern
  390. def float2uint_rz(arg0, _semantic=None):
  391. return core.extern_elementwise("", "", [arg0], {
  392. (core.dtype("fp32"), ): ("__nv_float2uint_rz", core.dtype("int32")),
  393. }, is_pure=True, _semantic=_semantic)
  394. @core.extern
  395. def float2uint_rd(arg0, _semantic=None):
  396. return core.extern_elementwise("", "", [arg0], {
  397. (core.dtype("fp32"), ): ("__nv_float2uint_rd", core.dtype("int32")),
  398. }, is_pure=True, _semantic=_semantic)
  399. @core.extern
  400. def float2uint_ru(arg0, _semantic=None):
  401. return core.extern_elementwise("", "", [arg0], {
  402. (core.dtype("fp32"), ): ("__nv_float2uint_ru", core.dtype("int32")),
  403. }, is_pure=True, _semantic=_semantic)
  404. @core.extern
  405. def int2float_rn(arg0, _semantic=None):
  406. return core.extern_elementwise("", "", [arg0], {
  407. (core.dtype("int32"), ): ("__nv_int2float_rn", core.dtype("fp32")),
  408. }, is_pure=True, _semantic=_semantic)
  409. @core.extern
  410. def int2float_rz(arg0, _semantic=None):
  411. return core.extern_elementwise("", "", [arg0], {
  412. (core.dtype("int32"), ): ("__nv_int2float_rz", core.dtype("fp32")),
  413. }, is_pure=True, _semantic=_semantic)
  414. @core.extern
  415. def int2float_rd(arg0, _semantic=None):
  416. return core.extern_elementwise("", "", [arg0], {
  417. (core.dtype("int32"), ): ("__nv_int2float_rd", core.dtype("fp32")),
  418. }, is_pure=True, _semantic=_semantic)
  419. @core.extern
  420. def int2float_ru(arg0, _semantic=None):
  421. return core.extern_elementwise("", "", [arg0], {
  422. (core.dtype("int32"), ): ("__nv_int2float_ru", core.dtype("fp32")),
  423. }, is_pure=True, _semantic=_semantic)
  424. @core.extern
  425. def uint2float_rn(arg0, _semantic=None):
  426. return core.extern_elementwise("", "", [arg0], {
  427. (core.dtype("uint32"), ): ("__nv_uint2float_rn", core.dtype("fp32")),
  428. }, is_pure=True, _semantic=_semantic)
  429. @core.extern
  430. def uint2float_rz(arg0, _semantic=None):
  431. return core.extern_elementwise("", "", [arg0], {
  432. (core.dtype("uint32"), ): ("__nv_uint2float_rz", core.dtype("fp32")),
  433. }, is_pure=True, _semantic=_semantic)
  434. @core.extern
  435. def uint2float_rd(arg0, _semantic=None):
  436. return core.extern_elementwise("", "", [arg0], {
  437. (core.dtype("uint32"), ): ("__nv_uint2float_rd", core.dtype("fp32")),
  438. }, is_pure=True, _semantic=_semantic)
  439. @core.extern
  440. def uint2float_ru(arg0, _semantic=None):
  441. return core.extern_elementwise("", "", [arg0], {
  442. (core.dtype("uint32"), ): ("__nv_uint2float_ru", core.dtype("fp32")),
  443. }, is_pure=True, _semantic=_semantic)
  444. @core.extern
  445. def hiloint2double(arg0, arg1, _semantic=None):
  446. return core.extern_elementwise("", "", [arg0, arg1], {
  447. (core.dtype("int32"), core.dtype("int32")): ("__nv_hiloint2double", core.dtype("fp64")),
  448. }, is_pure=True, _semantic=_semantic)
  449. @core.extern
  450. def double2loint(arg0, _semantic=None):
  451. return core.extern_elementwise("", "", [arg0], {
  452. (core.dtype("fp64"), ): ("__nv_double2loint", core.dtype("int32")),
  453. }, is_pure=True, _semantic=_semantic)
  454. @core.extern
  455. def double2hiint(arg0, _semantic=None):
  456. return core.extern_elementwise("", "", [arg0], {
  457. (core.dtype("fp64"), ): ("__nv_double2hiint", core.dtype("int32")),
  458. }, is_pure=True, _semantic=_semantic)
  459. @core.extern
  460. def float2ll_rn(arg0, _semantic=None):
  461. return core.extern_elementwise("", "", [arg0], {
  462. (core.dtype("fp32"), ): ("__nv_float2ll_rn", core.dtype("int64")),
  463. }, is_pure=True, _semantic=_semantic)
  464. @core.extern
  465. def float2ll_rz(arg0, _semantic=None):
  466. return core.extern_elementwise("", "", [arg0], {
  467. (core.dtype("fp32"), ): ("__nv_float2ll_rz", core.dtype("int64")),
  468. }, is_pure=True, _semantic=_semantic)
  469. @core.extern
  470. def float2ll_rd(arg0, _semantic=None):
  471. return core.extern_elementwise("", "", [arg0], {
  472. (core.dtype("fp32"), ): ("__nv_float2ll_rd", core.dtype("int64")),
  473. }, is_pure=True, _semantic=_semantic)
  474. @core.extern
  475. def float2ll_ru(arg0, _semantic=None):
  476. return core.extern_elementwise("", "", [arg0], {
  477. (core.dtype("fp32"), ): ("__nv_float2ll_ru", core.dtype("int64")),
  478. }, is_pure=True, _semantic=_semantic)
  479. @core.extern
  480. def float2ull_rn(arg0, _semantic=None):
  481. return core.extern_elementwise("", "", [arg0], {
  482. (core.dtype("fp32"), ): ("__nv_float2ull_rn", core.dtype("int64")),
  483. }, is_pure=True, _semantic=_semantic)
  484. @core.extern
  485. def float2ull_rz(arg0, _semantic=None):
  486. return core.extern_elementwise("", "", [arg0], {
  487. (core.dtype("fp32"), ): ("__nv_float2ull_rz", core.dtype("int64")),
  488. }, is_pure=True, _semantic=_semantic)
  489. @core.extern
  490. def float2ull_rd(arg0, _semantic=None):
  491. return core.extern_elementwise("", "", [arg0], {
  492. (core.dtype("fp32"), ): ("__nv_float2ull_rd", core.dtype("int64")),
  493. }, is_pure=True, _semantic=_semantic)
  494. @core.extern
  495. def float2ull_ru(arg0, _semantic=None):
  496. return core.extern_elementwise("", "", [arg0], {
  497. (core.dtype("fp32"), ): ("__nv_float2ull_ru", core.dtype("int64")),
  498. }, is_pure=True, _semantic=_semantic)
  499. @core.extern
  500. def double2ll_rn(arg0, _semantic=None):
  501. return core.extern_elementwise("", "", [arg0], {
  502. (core.dtype("fp64"), ): ("__nv_double2ll_rn", core.dtype("int64")),
  503. }, is_pure=True, _semantic=_semantic)
  504. @core.extern
  505. def double2ll_rz(arg0, _semantic=None):
  506. return core.extern_elementwise("", "", [arg0], {
  507. (core.dtype("fp64"), ): ("__nv_double2ll_rz", core.dtype("int64")),
  508. }, is_pure=True, _semantic=_semantic)
  509. @core.extern
  510. def double2ll_rd(arg0, _semantic=None):
  511. return core.extern_elementwise("", "", [arg0], {
  512. (core.dtype("fp64"), ): ("__nv_double2ll_rd", core.dtype("int64")),
  513. }, is_pure=True, _semantic=_semantic)
  514. @core.extern
  515. def double2ll_ru(arg0, _semantic=None):
  516. return core.extern_elementwise("", "", [arg0], {
  517. (core.dtype("fp64"), ): ("__nv_double2ll_ru", core.dtype("int64")),
  518. }, is_pure=True, _semantic=_semantic)
  519. @core.extern
  520. def double2ull_rn(arg0, _semantic=None):
  521. return core.extern_elementwise("", "", [arg0], {
  522. (core.dtype("fp64"), ): ("__nv_double2ull_rn", core.dtype("int64")),
  523. }, is_pure=True, _semantic=_semantic)
  524. @core.extern
  525. def double2ull_rz(arg0, _semantic=None):
  526. return core.extern_elementwise("", "", [arg0], {
  527. (core.dtype("fp64"), ): ("__nv_double2ull_rz", core.dtype("int64")),
  528. }, is_pure=True, _semantic=_semantic)
  529. @core.extern
  530. def double2ull_rd(arg0, _semantic=None):
  531. return core.extern_elementwise("", "", [arg0], {
  532. (core.dtype("fp64"), ): ("__nv_double2ull_rd", core.dtype("int64")),
  533. }, is_pure=True, _semantic=_semantic)
  534. @core.extern
  535. def double2ull_ru(arg0, _semantic=None):
  536. return core.extern_elementwise("", "", [arg0], {
  537. (core.dtype("fp64"), ): ("__nv_double2ull_ru", core.dtype("int64")),
  538. }, is_pure=True, _semantic=_semantic)
  539. @core.extern
  540. def ll2float_rn(arg0, _semantic=None):
  541. return core.extern_elementwise("", "", [arg0], {
  542. (core.dtype("int64"), ): ("__nv_ll2float_rn", core.dtype("fp32")),
  543. }, is_pure=True, _semantic=_semantic)
  544. @core.extern
  545. def ll2float_rz(arg0, _semantic=None):
  546. return core.extern_elementwise("", "", [arg0], {
  547. (core.dtype("int64"), ): ("__nv_ll2float_rz", core.dtype("fp32")),
  548. }, is_pure=True, _semantic=_semantic)
  549. @core.extern
  550. def ll2float_rd(arg0, _semantic=None):
  551. return core.extern_elementwise("", "", [arg0], {
  552. (core.dtype("int64"), ): ("__nv_ll2float_rd", core.dtype("fp32")),
  553. }, is_pure=True, _semantic=_semantic)
  554. @core.extern
  555. def ll2float_ru(arg0, _semantic=None):
  556. return core.extern_elementwise("", "", [arg0], {
  557. (core.dtype("int64"), ): ("__nv_ll2float_ru", core.dtype("fp32")),
  558. }, is_pure=True, _semantic=_semantic)
  559. @core.extern
  560. def ull2float_rn(arg0, _semantic=None):
  561. return core.extern_elementwise("", "", [arg0], {
  562. (core.dtype("uint64"), ): ("__nv_ull2float_rn", core.dtype("fp32")),
  563. }, is_pure=True, _semantic=_semantic)
  564. @core.extern
  565. def ull2float_rz(arg0, _semantic=None):
  566. return core.extern_elementwise("", "", [arg0], {
  567. (core.dtype("uint64"), ): ("__nv_ull2float_rz", core.dtype("fp32")),
  568. }, is_pure=True, _semantic=_semantic)
  569. @core.extern
  570. def ull2float_rd(arg0, _semantic=None):
  571. return core.extern_elementwise("", "", [arg0], {
  572. (core.dtype("uint64"), ): ("__nv_ull2float_rd", core.dtype("fp32")),
  573. }, is_pure=True, _semantic=_semantic)
  574. @core.extern
  575. def ull2float_ru(arg0, _semantic=None):
  576. return core.extern_elementwise("", "", [arg0], {
  577. (core.dtype("uint64"), ): ("__nv_ull2float_ru", core.dtype("fp32")),
  578. }, is_pure=True, _semantic=_semantic)
  579. @core.extern
  580. def ll2double_rn(arg0, _semantic=None):
  581. return core.extern_elementwise("", "", [arg0], {
  582. (core.dtype("int64"), ): ("__nv_ll2double_rn", core.dtype("fp64")),
  583. }, is_pure=True, _semantic=_semantic)
  584. @core.extern
  585. def ll2double_rz(arg0, _semantic=None):
  586. return core.extern_elementwise("", "", [arg0], {
  587. (core.dtype("int64"), ): ("__nv_ll2double_rz", core.dtype("fp64")),
  588. }, is_pure=True, _semantic=_semantic)
  589. @core.extern
  590. def ll2double_rd(arg0, _semantic=None):
  591. return core.extern_elementwise("", "", [arg0], {
  592. (core.dtype("int64"), ): ("__nv_ll2double_rd", core.dtype("fp64")),
  593. }, is_pure=True, _semantic=_semantic)
  594. @core.extern
  595. def ll2double_ru(arg0, _semantic=None):
  596. return core.extern_elementwise("", "", [arg0], {
  597. (core.dtype("int64"), ): ("__nv_ll2double_ru", core.dtype("fp64")),
  598. }, is_pure=True, _semantic=_semantic)
  599. @core.extern
  600. def ull2double_rn(arg0, _semantic=None):
  601. return core.extern_elementwise("", "", [arg0], {
  602. (core.dtype("uint64"), ): ("__nv_ull2double_rn", core.dtype("fp64")),
  603. }, is_pure=True, _semantic=_semantic)
  604. @core.extern
  605. def ull2double_rz(arg0, _semantic=None):
  606. return core.extern_elementwise("", "", [arg0], {
  607. (core.dtype("uint64"), ): ("__nv_ull2double_rz", core.dtype("fp64")),
  608. }, is_pure=True, _semantic=_semantic)
  609. @core.extern
  610. def ull2double_rd(arg0, _semantic=None):
  611. return core.extern_elementwise("", "", [arg0], {
  612. (core.dtype("uint64"), ): ("__nv_ull2double_rd", core.dtype("fp64")),
  613. }, is_pure=True, _semantic=_semantic)
  614. @core.extern
  615. def ull2double_ru(arg0, _semantic=None):
  616. return core.extern_elementwise("", "", [arg0], {
  617. (core.dtype("uint64"), ): ("__nv_ull2double_ru", core.dtype("fp64")),
  618. }, is_pure=True, _semantic=_semantic)
  619. @core.extern
  620. def int_as_float(arg0, _semantic=None):
  621. return core.extern_elementwise("", "", [arg0], {
  622. (core.dtype("int32"), ): ("__nv_int_as_float", core.dtype("fp32")),
  623. }, is_pure=True, _semantic=_semantic)
  624. @core.extern
  625. def float_as_int(arg0, _semantic=None):
  626. return core.extern_elementwise("", "", [arg0], {
  627. (core.dtype("fp32"), ): ("__nv_float_as_int", core.dtype("int32")),
  628. }, is_pure=True, _semantic=_semantic)
  629. @core.extern
  630. def uint_as_float(arg0, _semantic=None):
  631. return core.extern_elementwise("", "", [arg0], {
  632. (core.dtype("uint32"), ): ("__nv_uint_as_float", core.dtype("fp32")),
  633. }, is_pure=True, _semantic=_semantic)
  634. @core.extern
  635. def float_as_uint(arg0, _semantic=None):
  636. return core.extern_elementwise("", "", [arg0], {
  637. (core.dtype("fp32"), ): ("__nv_float_as_uint", core.dtype("int32")),
  638. }, is_pure=True, _semantic=_semantic)
  639. @core.extern
  640. def longlong_as_double(arg0, _semantic=None):
  641. return core.extern_elementwise("", "", [arg0], {
  642. (core.dtype("int64"), ): ("__nv_longlong_as_double", core.dtype("fp64")),
  643. }, is_pure=True, _semantic=_semantic)
  644. @core.extern
  645. def double_as_longlong(arg0, _semantic=None):
  646. return core.extern_elementwise("", "", [arg0], {
  647. (core.dtype("fp64"), ): ("__nv_double_as_longlong", core.dtype("int64")),
  648. }, is_pure=True, _semantic=_semantic)
  649. @core.extern
  650. def fast_sinf(arg0, _semantic=None):
  651. return core.extern_elementwise("", "", [arg0], {
  652. (core.dtype("fp32"), ): ("__nv_fast_sinf", core.dtype("fp32")),
  653. }, is_pure=True, _semantic=_semantic)
  654. @core.extern
  655. def fast_cosf(arg0, _semantic=None):
  656. return core.extern_elementwise("", "", [arg0], {
  657. (core.dtype("fp32"), ): ("__nv_fast_cosf", core.dtype("fp32")),
  658. }, is_pure=True, _semantic=_semantic)
  659. @core.extern
  660. def fast_log2f(arg0, _semantic=None):
  661. return core.extern_elementwise("", "", [arg0], {
  662. (core.dtype("fp32"), ): ("__nv_fast_log2f", core.dtype("fp32")),
  663. }, is_pure=True, _semantic=_semantic)
  664. @core.extern
  665. def fast_logf(arg0, _semantic=None):
  666. return core.extern_elementwise("", "", [arg0], {
  667. (core.dtype("fp32"), ): ("__nv_fast_logf", core.dtype("fp32")),
  668. }, is_pure=True, _semantic=_semantic)
  669. @core.extern
  670. def fast_expf(arg0, _semantic=None):
  671. return core.extern_elementwise("", "", [arg0], {
  672. (core.dtype("fp32"), ): ("__nv_fast_expf", core.dtype("fp32")),
  673. }, is_pure=True, _semantic=_semantic)
  674. @core.extern
  675. def fast_tanf(arg0, _semantic=None):
  676. return core.extern_elementwise("", "", [arg0], {
  677. (core.dtype("fp32"), ): ("__nv_fast_tanf", core.dtype("fp32")),
  678. }, is_pure=True, _semantic=_semantic)
  679. @core.extern
  680. def fast_exp10f(arg0, _semantic=None):
  681. return core.extern_elementwise("", "", [arg0], {
  682. (core.dtype("fp32"), ): ("__nv_fast_exp10f", core.dtype("fp32")),
  683. }, is_pure=True, _semantic=_semantic)
  684. @core.extern
  685. def fast_log10f(arg0, _semantic=None):
  686. return core.extern_elementwise("", "", [arg0], {
  687. (core.dtype("fp32"), ): ("__nv_fast_log10f", core.dtype("fp32")),
  688. }, is_pure=True, _semantic=_semantic)
  689. @core.extern
  690. def fast_powf(arg0, arg1, _semantic=None):
  691. return core.extern_elementwise("", "", [arg0, arg1], {
  692. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fast_powf", core.dtype("fp32")),
  693. }, is_pure=True, _semantic=_semantic)
  694. @core.extern
  695. def hadd(arg0, arg1, _semantic=None):
  696. return core.extern_elementwise(
  697. "", "", [arg0, arg1], {
  698. (core.dtype("int32"), core.dtype("int32")): ("__nv_hadd", core.dtype("int32")),
  699. (core.dtype("uint32"), core.dtype("uint32")): ("__nv_uhadd", core.dtype("uint32")),
  700. }, is_pure=True, _semantic=_semantic)
  701. @core.extern
  702. def rhadd(arg0, arg1, _semantic=None):
  703. return core.extern_elementwise(
  704. "", "", [arg0, arg1], {
  705. (core.dtype("int32"), core.dtype("int32")): ("__nv_rhadd", core.dtype("int32")),
  706. (core.dtype("uint32"), core.dtype("uint32")): ("__nv_urhadd", core.dtype("uint32")),
  707. }, is_pure=True, _semantic=_semantic)
  708. @core.extern
  709. def sub_rn(arg0, arg1, _semantic=None):
  710. return core.extern_elementwise(
  711. "", "", [arg0, arg1], {
  712. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_rn", core.dtype("fp32")),
  713. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_rn", core.dtype("fp64")),
  714. }, is_pure=True, _semantic=_semantic)
  715. @core.extern
  716. def sub_rz(arg0, arg1, _semantic=None):
  717. return core.extern_elementwise(
  718. "", "", [arg0, arg1], {
  719. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_rz", core.dtype("fp32")),
  720. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_rz", core.dtype("fp64")),
  721. }, is_pure=True, _semantic=_semantic)
  722. @core.extern
  723. def sub_rd(arg0, arg1, _semantic=None):
  724. return core.extern_elementwise(
  725. "", "", [arg0, arg1], {
  726. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_rd", core.dtype("fp32")),
  727. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_rd", core.dtype("fp64")),
  728. }, is_pure=True, _semantic=_semantic)
  729. @core.extern
  730. def sub_ru(arg0, arg1, _semantic=None):
  731. return core.extern_elementwise(
  732. "", "", [arg0, arg1], {
  733. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_ru", core.dtype("fp32")),
  734. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_ru", core.dtype("fp64")),
  735. }, is_pure=True, _semantic=_semantic)
  736. @core.extern
  737. def rsqrt_rn(arg0, _semantic=None):
  738. return core.extern_elementwise("", "", [
  739. arg0,
  740. ], {
  741. (core.dtype("fp32"), ): ("__nv_frsqrt_rn", core.dtype("fp32")),
  742. }, is_pure=True, _semantic=_semantic)
  743. @core.extern
  744. def ffs(arg0, _semantic=None):
  745. return core.extern_elementwise(
  746. "", "", [
  747. arg0,
  748. ], {
  749. (core.dtype("int32"), ): ("__nv_ffs", core.dtype("int32")),
  750. (core.dtype("int64"), ): ("__nv_ffsll", core.dtype("int32")),
  751. }, is_pure=True, _semantic=_semantic)
  752. @core.extern
  753. def rint(arg0, _semantic=None):
  754. return core.extern_elementwise(
  755. "", "", [
  756. arg0,
  757. ], {
  758. (core.dtype("fp32"), ): ("__nv_rintf", core.dtype("fp32")),
  759. (core.dtype("fp64"), ): ("__nv_rint", core.dtype("fp64")),
  760. }, is_pure=True, _semantic=_semantic)
  761. @core.extern
  762. def llrint(arg0, _semantic=None):
  763. return core.extern_elementwise(
  764. "", "", [
  765. arg0,
  766. ], {
  767. (core.dtype("fp32"), ): ("__nv_llrintf", core.dtype("int64")),
  768. (core.dtype("fp64"), ): ("__nv_llrint", core.dtype("int64")),
  769. }, is_pure=True, _semantic=_semantic)
  770. @core.extern
  771. def nearbyint(arg0, _semantic=None):
  772. return core.extern_elementwise(
  773. "", "", [
  774. arg0,
  775. ], {
  776. (core.dtype("fp32"), ): ("__nv_nearbyintf", core.dtype("fp32")),
  777. (core.dtype("fp64"), ): ("__nv_nearbyint", core.dtype("fp64")),
  778. }, is_pure=True, _semantic=_semantic)
  779. @core.extern
  780. def isnan(arg0, _semantic=None):
  781. return core.extern_elementwise(
  782. "", "", [
  783. arg0,
  784. ], {
  785. (core.dtype("fp32"), ): ("__nv_isnanf", core.dtype("int32")),
  786. (core.dtype("fp64"), ): ("__nv_isnand", core.dtype("int32")),
  787. }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
  788. @core.extern
  789. def signbit(arg0, _semantic=None):
  790. return core.extern_elementwise(
  791. "", "", [
  792. arg0,
  793. ], {
  794. (core.dtype("fp32"), ): ("__nv_signbitf", core.dtype("int32")),
  795. (core.dtype("fp64"), ): ("__nv_signbitd", core.dtype("int32")),
  796. }, is_pure=True, _semantic=_semantic)
  797. @core.extern
  798. def copysign(arg0, arg1, _semantic=None):
  799. return core.extern_elementwise(
  800. "", "", [arg0, arg1], {
  801. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_copysignf", core.dtype("fp32")),
  802. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_copysign", core.dtype("fp64")),
  803. }, is_pure=True, _semantic=_semantic)
  804. @core.extern
  805. def finitef(arg0, _semantic=None):
  806. return core.extern_elementwise("", "", [arg0], {
  807. (core.dtype("fp32"), ): ("__nv_finitef", core.dtype("int32")),
  808. }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
  809. @core.extern
  810. def isinf(arg0, _semantic=None):
  811. return core.extern_elementwise(
  812. "", "", [arg0], {
  813. (core.dtype("fp32"), ): ("__nv_isinff", core.dtype("int32")),
  814. (core.dtype("fp64"), ): ("__nv_isinfd", core.dtype("int32")),
  815. }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
  816. @core.extern
  817. def nextafter(arg0, arg1, _semantic=None):
  818. return core.extern_elementwise(
  819. "", "", [arg0, arg1], {
  820. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_nextafterf", core.dtype("fp32")),
  821. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_nextafter", core.dtype("fp64")),
  822. }, is_pure=True, _semantic=_semantic)
  823. @core.extern
  824. def sin(arg0, _semantic=None):
  825. return core.extern_elementwise(
  826. "", "", [arg0], {
  827. (core.dtype("fp32"), ): ("__nv_sinf", core.dtype("fp32")),
  828. (core.dtype("fp64"), ): ("__nv_sin", core.dtype("fp64")),
  829. }, is_pure=True, _semantic=_semantic)
  830. @core.extern
  831. def cos(arg0, _semantic=None):
  832. return core.extern_elementwise(
  833. "", "", [arg0], {
  834. (core.dtype("fp32"), ): ("__nv_cosf", core.dtype("fp32")),
  835. (core.dtype("fp64"), ): ("__nv_cos", core.dtype("fp64")),
  836. }, is_pure=True, _semantic=_semantic)
  837. @core.extern
  838. def sinpi(arg0, _semantic=None):
  839. return core.extern_elementwise(
  840. "", "", [arg0], {
  841. (core.dtype("fp32"), ): ("__nv_sinpif", core.dtype("fp32")),
  842. (core.dtype("fp64"), ): ("__nv_sinpi", core.dtype("fp64")),
  843. }, is_pure=True, _semantic=_semantic)
  844. @core.extern
  845. def cospi(arg0, _semantic=None):
  846. return core.extern_elementwise(
  847. "", "", [arg0], {
  848. (core.dtype("fp32"), ): ("__nv_cospif", core.dtype("fp32")),
  849. (core.dtype("fp64"), ): ("__nv_cospi", core.dtype("fp64")),
  850. }, is_pure=True, _semantic=_semantic)
  851. @core.extern
  852. def tan(arg0, _semantic=None):
  853. return core.extern_elementwise(
  854. "", "", [arg0], {
  855. (core.dtype("fp32"), ): ("__nv_tanf", core.dtype("fp32")),
  856. (core.dtype("fp64"), ): ("__nv_tan", core.dtype("fp64")),
  857. }, is_pure=True, _semantic=_semantic)
  858. @core.extern
  859. def log2(arg0, _semantic=None):
  860. return core.extern_elementwise(
  861. "", "", [arg0], {
  862. (core.dtype("fp32"), ): ("__nv_log2f", core.dtype("fp32")),
  863. (core.dtype("fp64"), ): ("__nv_log2", core.dtype("fp64")),
  864. }, is_pure=True, _semantic=_semantic)
  865. @core.extern
  866. def exp(arg0, _semantic=None):
  867. return core.extern_elementwise(
  868. "", "", [arg0], {
  869. (core.dtype("fp32"), ): ("__nv_expf", core.dtype("fp32")),
  870. (core.dtype("fp64"), ): ("__nv_exp", core.dtype("fp64")),
  871. }, is_pure=True, _semantic=_semantic)
  872. @core.extern
  873. def exp10(arg0, _semantic=None):
  874. return core.extern_elementwise(
  875. "", "", [arg0], {
  876. (core.dtype("fp32"), ): ("__nv_exp10f", core.dtype("fp32")),
  877. (core.dtype("fp64"), ): ("__nv_exp10", core.dtype("fp64")),
  878. }, is_pure=True, _semantic=_semantic)
  879. @core.extern
  880. def cosh(arg0, _semantic=None):
  881. return core.extern_elementwise(
  882. "", "", [arg0], {
  883. (core.dtype("fp32"), ): ("__nv_coshf", core.dtype("fp32")),
  884. (core.dtype("fp64"), ): ("__nv_cosh", core.dtype("fp64")),
  885. }, is_pure=True, _semantic=_semantic)
  886. @core.extern
  887. def sinh(arg0, _semantic=None):
  888. return core.extern_elementwise(
  889. "", "", [arg0], {
  890. (core.dtype("fp32"), ): ("__nv_sinhf", core.dtype("fp32")),
  891. (core.dtype("fp64"), ): ("__nv_sinh", core.dtype("fp64")),
  892. }, is_pure=True, _semantic=_semantic)
  893. @core.extern
  894. def tanh(arg0, _semantic=None):
  895. return core.extern_elementwise(
  896. "", "", [arg0], {
  897. (core.dtype("fp32"), ): ("__nv_tanhf", core.dtype("fp32")),
  898. (core.dtype("fp64"), ): ("__nv_tanh", core.dtype("fp64")),
  899. }, is_pure=True, _semantic=_semantic)
  900. @core.extern
  901. def atan2(arg0, arg1, _semantic=None):
  902. return core.extern_elementwise(
  903. "", "", [arg0, arg1], {
  904. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_atan2f", core.dtype("fp32")),
  905. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_atan2", core.dtype("fp64")),
  906. }, is_pure=True, _semantic=_semantic)
  907. @core.extern
  908. def atan(arg0, _semantic=None):
  909. return core.extern_elementwise(
  910. "", "", [arg0], {
  911. (core.dtype("fp32"), ): ("__nv_atanf", core.dtype("fp32")),
  912. (core.dtype("fp64"), ): ("__nv_atan", core.dtype("fp64")),
  913. }, is_pure=True, _semantic=_semantic)
  914. @core.extern
  915. def asin(arg0, _semantic=None):
  916. return core.extern_elementwise(
  917. "", "", [arg0], {
  918. (core.dtype("fp32"), ): ("__nv_asinf", core.dtype("fp32")),
  919. (core.dtype("fp64"), ): ("__nv_asin", core.dtype("fp64")),
  920. }, is_pure=True, _semantic=_semantic)
  921. @core.extern
  922. def acos(arg0, _semantic=None):
  923. return core.extern_elementwise(
  924. "", "", [arg0], {
  925. (core.dtype("fp32"), ): ("__nv_acosf", core.dtype("fp32")),
  926. (core.dtype("fp64"), ): ("__nv_acos", core.dtype("fp64")),
  927. }, is_pure=True, _semantic=_semantic)
  928. @core.extern
  929. def log(arg0, _semantic=None):
  930. return core.extern_elementwise(
  931. "", "", [arg0], {
  932. (core.dtype("fp32"), ): ("__nv_logf", core.dtype("fp32")),
  933. (core.dtype("fp64"), ): ("__nv_log", core.dtype("fp64")),
  934. }, is_pure=True, _semantic=_semantic)
  935. @core.extern
  936. def log10(arg0, _semantic=None):
  937. return core.extern_elementwise(
  938. "", "", [arg0], {
  939. (core.dtype("fp32"), ): ("__nv_log10f", core.dtype("fp32")),
  940. (core.dtype("fp64"), ): ("__nv_log10", core.dtype("fp64")),
  941. }, is_pure=True, _semantic=_semantic)
  942. @core.extern
  943. def log1p(arg0, _semantic=None):
  944. return core.extern_elementwise(
  945. "", "", [arg0], {
  946. (core.dtype("fp32"), ): ("__nv_log1pf", core.dtype("fp32")),
  947. (core.dtype("fp64"), ): ("__nv_log1p", core.dtype("fp64")),
  948. }, is_pure=True, _semantic=_semantic)
  949. @core.extern
  950. def acosh(arg0, _semantic=None):
  951. return core.extern_elementwise(
  952. "", "", [arg0], {
  953. (core.dtype("fp32"), ): ("__nv_acoshf", core.dtype("fp32")),
  954. (core.dtype("fp64"), ): ("__nv_acosh", core.dtype("fp64")),
  955. }, is_pure=True, _semantic=_semantic)
  956. @core.extern
  957. def asinh(arg0, _semantic=None):
  958. return core.extern_elementwise(
  959. "", "", [arg0], {
  960. (core.dtype("fp32"), ): ("__nv_asinhf", core.dtype("fp32")),
  961. (core.dtype("fp64"), ): ("__nv_asinh", core.dtype("fp64")),
  962. }, is_pure=True, _semantic=_semantic)
  963. @core.extern
  964. def atanh(arg0, _semantic=None):
  965. return core.extern_elementwise(
  966. "", "", [arg0], {
  967. (core.dtype("fp32"), ): ("__nv_atanhf", core.dtype("fp32")),
  968. (core.dtype("fp64"), ): ("__nv_atanh", core.dtype("fp64")),
  969. }, is_pure=True, _semantic=_semantic)
  970. @core.extern
  971. def expm1(arg0, _semantic=None):
  972. return core.extern_elementwise(
  973. "", "", [arg0], {
  974. (core.dtype("fp32"), ): ("__nv_expm1f", core.dtype("fp32")),
  975. (core.dtype("fp64"), ): ("__nv_expm1", core.dtype("fp64")),
  976. }, is_pure=True, _semantic=_semantic)
  977. @core.extern
  978. def hypot(arg0, arg1, _semantic=None):
  979. return core.extern_elementwise(
  980. "", "", [arg0, arg1], {
  981. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_hypotf", core.dtype("fp32")),
  982. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_hypot", core.dtype("fp64")),
  983. }, is_pure=True, _semantic=_semantic)
  984. @core.extern
  985. def rhypot(arg0, arg1, _semantic=None):
  986. return core.extern_elementwise(
  987. "", "", [arg0, arg1], {
  988. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_rhypotf", core.dtype("fp32")),
  989. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_rhypot", core.dtype("fp64")),
  990. }, is_pure=True, _semantic=_semantic)
  991. @core.extern
  992. def norm3d(arg0, arg1, arg2, _semantic=None):
  993. return core.extern_elementwise(
  994. "", "", [arg0, arg1, arg2], {
  995. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_norm3df", core.dtype("fp32")),
  996. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_norm3d", core.dtype("fp64")),
  997. }, is_pure=True, _semantic=_semantic)
  998. @core.extern
  999. def rnorm3d(arg0, arg1, arg2, _semantic=None):
  1000. return core.extern_elementwise(
  1001. "", "", [arg0, arg1, arg2], {
  1002. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_rnorm3df", core.dtype("fp32")),
  1003. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_rnorm3d", core.dtype("fp64")),
  1004. }, is_pure=True, _semantic=_semantic)
  1005. @core.extern
  1006. def norm4d(arg0, arg1, arg2, arg3, _semantic=None):
  1007. return core.extern_elementwise(
  1008. "", "", [arg0, arg1, arg2, arg3], {
  1009. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")):
  1010. ("__nv_norm4df", core.dtype("fp32")),
  1011. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")):
  1012. ("__nv_norm4d", core.dtype("fp64")),
  1013. }, is_pure=True, _semantic=_semantic)
  1014. @core.extern
  1015. def rnorm4d(arg0, arg1, arg2, arg3, _semantic=None):
  1016. return core.extern_elementwise(
  1017. "", "", [arg0, arg1, arg2, arg3], {
  1018. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")):
  1019. ("__nv_rnorm4df", core.dtype("fp32")),
  1020. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")):
  1021. ("__nv_rnorm4d", core.dtype("fp64")),
  1022. }, is_pure=True, _semantic=_semantic)
  1023. @core.extern
  1024. def cbrt(arg0, _semantic=None):
  1025. return core.extern_elementwise(
  1026. "", "", [arg0], {
  1027. (core.dtype("fp32"), ): ("__nv_cbrtf", core.dtype("fp32")),
  1028. (core.dtype("fp64"), ): ("__nv_cbrt", core.dtype("fp64")),
  1029. }, is_pure=True, _semantic=_semantic)
  1030. @core.extern
  1031. def rcbrt(arg0, _semantic=None):
  1032. return core.extern_elementwise(
  1033. "", "", [arg0], {
  1034. (core.dtype("fp32"), ): ("__nv_rcbrtf", core.dtype("fp32")),
  1035. (core.dtype("fp64"), ): ("__nv_rcbrt", core.dtype("fp64")),
  1036. }, is_pure=True, _semantic=_semantic)
  1037. @core.extern
  1038. def j0(arg0, _semantic=None):
  1039. return core.extern_elementwise("", "", [arg0], {
  1040. (core.dtype("fp32"), ): ("__nv_j0f", core.dtype("fp32")),
  1041. (core.dtype("fp64"), ): ("__nv_j0", core.dtype("fp64")),
  1042. }, is_pure=True, _semantic=_semantic)
  1043. @core.extern
  1044. def j1(arg0, _semantic=None):
  1045. return core.extern_elementwise("", "", [arg0], {
  1046. (core.dtype("fp32"), ): ("__nv_j1f", core.dtype("fp32")),
  1047. (core.dtype("fp64"), ): ("__nv_j1", core.dtype("fp64")),
  1048. }, is_pure=True, _semantic=_semantic)
  1049. @core.extern
  1050. def y0(arg0, _semantic=None):
  1051. return core.extern_elementwise("", "", [arg0], {
  1052. (core.dtype("fp32"), ): ("__nv_y0f", core.dtype("fp32")),
  1053. (core.dtype("fp64"), ): ("__nv_y0", core.dtype("fp64")),
  1054. }, is_pure=True, _semantic=_semantic)
  1055. @core.extern
  1056. def y1(arg0, _semantic=None):
  1057. return core.extern_elementwise("", "", [arg0], {
  1058. (core.dtype("fp32"), ): ("__nv_y1f", core.dtype("fp32")),
  1059. (core.dtype("fp64"), ): ("__nv_y1", core.dtype("fp64")),
  1060. }, is_pure=True, _semantic=_semantic)
  1061. @core.extern
  1062. def yn(arg0, arg1, _semantic=None):
  1063. return core.extern_elementwise(
  1064. "", "", [arg0, arg1], {
  1065. (core.dtype("int32"), core.dtype("fp32")): ("__nv_ynf", core.dtype("fp32")),
  1066. (core.dtype("int32"), core.dtype("fp64")): ("__nv_yn", core.dtype("fp64")),
  1067. }, is_pure=True, _semantic=_semantic)
  1068. @core.extern
  1069. def jn(arg0, arg1, _semantic=None):
  1070. return core.extern_elementwise(
  1071. "", "", [arg0, arg1], {
  1072. (core.dtype("int32"), core.dtype("fp32")): ("__nv_jnf", core.dtype("fp32")),
  1073. (core.dtype("int32"), core.dtype("fp64")): ("__nv_jn", core.dtype("fp64")),
  1074. }, is_pure=True, _semantic=_semantic)
  1075. @core.extern
  1076. def cyl_bessel_i0(arg0, _semantic=None):
  1077. return core.extern_elementwise(
  1078. "", "", [arg0], {
  1079. (core.dtype("fp32"), ): ("__nv_cyl_bessel_i0f", core.dtype("fp32")),
  1080. (core.dtype("fp64"), ): ("__nv_cyl_bessel_i0", core.dtype("fp64")),
  1081. }, is_pure=True, _semantic=_semantic)
  1082. @core.extern
  1083. def cyl_bessel_i1(arg0, _semantic=None):
  1084. return core.extern_elementwise(
  1085. "", "", [arg0], {
  1086. (core.dtype("fp32"), ): ("__nv_cyl_bessel_i1f", core.dtype("fp32")),
  1087. (core.dtype("fp64"), ): ("__nv_cyl_bessel_i1", core.dtype("fp64")),
  1088. }, is_pure=True, _semantic=_semantic)
  1089. @core.extern
  1090. def erf(arg0, _semantic=None):
  1091. return core.extern_elementwise(
  1092. "", "", [arg0], {
  1093. (core.dtype("fp32"), ): ("__nv_erff", core.dtype("fp32")),
  1094. (core.dtype("fp64"), ): ("__nv_erf", core.dtype("fp64")),
  1095. }, is_pure=True, _semantic=_semantic)
  1096. @core.extern
  1097. def erfinv(arg0, _semantic=None):
  1098. return core.extern_elementwise(
  1099. "", "", [arg0], {
  1100. (core.dtype("fp32"), ): ("__nv_erfinvf", core.dtype("fp32")),
  1101. (core.dtype("fp64"), ): ("__nv_erfinv", core.dtype("fp64")),
  1102. }, is_pure=True, _semantic=_semantic)
  1103. @core.extern
  1104. def erfc(arg0, _semantic=None):
  1105. return core.extern_elementwise(
  1106. "", "", [arg0], {
  1107. (core.dtype("fp32"), ): ("__nv_erfcf", core.dtype("fp32")),
  1108. (core.dtype("fp64"), ): ("__nv_erfc", core.dtype("fp64")),
  1109. }, is_pure=True, _semantic=_semantic)
  1110. @core.extern
  1111. def erfcx(arg0, _semantic=None):
  1112. return core.extern_elementwise(
  1113. "", "", [arg0], {
  1114. (core.dtype("fp32"), ): ("__nv_erfcxf", core.dtype("fp32")),
  1115. (core.dtype("fp64"), ): ("__nv_erfcx", core.dtype("fp64")),
  1116. }, is_pure=True, _semantic=_semantic)
  1117. @core.extern
  1118. def erfcinv(arg0, _semantic=None):
  1119. return core.extern_elementwise(
  1120. "", "", [arg0], {
  1121. (core.dtype("fp32"), ): ("__nv_erfcinvf", core.dtype("fp32")),
  1122. (core.dtype("fp64"), ): ("__nv_erfcinv", core.dtype("fp64")),
  1123. }, is_pure=True, _semantic=_semantic)
  1124. @core.extern
  1125. def normcdfinv(arg0, _semantic=None):
  1126. return core.extern_elementwise(
  1127. "", "", [arg0], {
  1128. (core.dtype("fp32"), ): ("__nv_normcdfinvf", core.dtype("fp32")),
  1129. (core.dtype("fp64"), ): ("__nv_normcdfinv", core.dtype("fp64")),
  1130. }, is_pure=True, _semantic=_semantic)
  1131. @core.extern
  1132. def normcdf(arg0, _semantic=None):
  1133. return core.extern_elementwise(
  1134. "", "", [arg0], {
  1135. (core.dtype("fp32"), ): ("__nv_normcdff", core.dtype("fp32")),
  1136. (core.dtype("fp64"), ): ("__nv_normcdf", core.dtype("fp64")),
  1137. }, is_pure=True, _semantic=_semantic)
  1138. @core.extern
  1139. def lgamma(arg0, _semantic=None):
  1140. return core.extern_elementwise(
  1141. "", "", [arg0], {
  1142. (core.dtype("fp32"), ): ("__nv_lgammaf", core.dtype("fp32")),
  1143. (core.dtype("fp64"), ): ("__nv_lgamma", core.dtype("fp64")),
  1144. }, is_pure=True, _semantic=_semantic)
  1145. @core.extern
  1146. def ldexp(arg0, arg1, _semantic=None):
  1147. return core.extern_elementwise(
  1148. "", "", [arg0, arg1], {
  1149. (core.dtype("fp32"), core.dtype("int32")): ("__nv_ldexpf", core.dtype("fp32")),
  1150. (core.dtype("fp64"), core.dtype("int32")): ("__nv_ldexp", core.dtype("fp64")),
  1151. }, is_pure=True, _semantic=_semantic)
  1152. @core.extern
  1153. def scalbn(arg0, arg1, _semantic=None):
  1154. return core.extern_elementwise(
  1155. "", "", [arg0, arg1], {
  1156. (core.dtype("fp32"), core.dtype("int32")): ("__nv_scalbnf", core.dtype("fp32")),
  1157. (core.dtype("fp64"), core.dtype("int32")): ("__nv_scalbn", core.dtype("fp64")),
  1158. }, is_pure=True, _semantic=_semantic)
  1159. @core.extern
  1160. def fmod(arg0, arg1, _semantic=None):
  1161. return core.extern_elementwise(
  1162. "", "", [arg0, arg1], {
  1163. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmodf", core.dtype("fp32")),
  1164. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_fmod", core.dtype("fp64")),
  1165. }, is_pure=True, _semantic=_semantic)
  1166. @core.extern
  1167. def remainder(arg0, arg1, _semantic=None):
  1168. return core.extern_elementwise(
  1169. "", "", [arg0, arg1], {
  1170. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_remainderf", core.dtype("fp32")),
  1171. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_remainder", core.dtype("fp64")),
  1172. }, is_pure=True, _semantic=_semantic)
  1173. @core.extern
  1174. def fma(arg0, arg1, arg2, _semantic=None):
  1175. return core.extern_elementwise(
  1176. "", "", [arg0, arg1, arg2], {
  1177. (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf", core.dtype("fp32")),
  1178. (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma", core.dtype("fp64")),
  1179. }, is_pure=True, _semantic=_semantic)
  1180. @core.extern
  1181. def pow(arg0, arg1, _semantic=None):
  1182. return core.extern_elementwise(
  1183. "", "", [arg0, arg1], {
  1184. (core.dtype("fp32"), core.dtype("int32")): ("__nv_powif", core.dtype("fp32")),
  1185. (core.dtype("fp64"), core.dtype("int32")): ("__nv_powi", core.dtype("fp64")),
  1186. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_powf", core.dtype("fp32")),
  1187. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_pow", core.dtype("fp64")),
  1188. }, is_pure=True, _semantic=_semantic)
  1189. @core.extern
  1190. def tgamma(arg0, _semantic=None):
  1191. return core.extern_elementwise(
  1192. "", "", [arg0], {
  1193. (core.dtype("fp32"), ): ("__nv_tgammaf", core.dtype("fp32")),
  1194. (core.dtype("fp64"), ): ("__nv_tgamma", core.dtype("fp64")),
  1195. }, is_pure=True, _semantic=_semantic)
  1196. @core.extern
  1197. def round(arg0, _semantic=None):
  1198. return core.extern_elementwise(
  1199. "", "", [arg0], {
  1200. (core.dtype("fp32"), ): ("__nv_roundf", core.dtype("fp32")),
  1201. (core.dtype("fp64"), ): ("__nv_round", core.dtype("fp64")),
  1202. }, is_pure=True, _semantic=_semantic)
  1203. @core.extern
  1204. def llround(arg0, _semantic=None):
  1205. return core.extern_elementwise(
  1206. "", "", [arg0], {
  1207. (core.dtype("fp32"), ): ("__nv_llroundf", core.dtype("int64")),
  1208. (core.dtype("fp64"), ): ("__nv_llround", core.dtype("int64")),
  1209. }, is_pure=True, _semantic=_semantic)
  1210. @core.extern
  1211. def fdim(arg0, arg1, _semantic=None):
  1212. return core.extern_elementwise(
  1213. "", "", [arg0, arg1], {
  1214. (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdimf", core.dtype("fp32")),
  1215. (core.dtype("fp64"), core.dtype("fp64")): ("__nv_fdim", core.dtype("fp64")),
  1216. }, is_pure=True, _semantic=_semantic)
  1217. @core.extern
  1218. def ilogb(arg0, _semantic=None):
  1219. return core.extern_elementwise(
  1220. "", "", [arg0], {
  1221. (core.dtype("fp32"), ): ("__nv_ilogbf", core.dtype("int32")),
  1222. (core.dtype("fp64"), ): ("__nv_ilogb", core.dtype("int32")),
  1223. }, is_pure=True, _semantic=_semantic)
  1224. @core.extern
  1225. def logb(arg0, _semantic=None):
  1226. return core.extern_elementwise(
  1227. "", "", [arg0], {
  1228. (core.dtype("fp32"), ): ("__nv_logbf", core.dtype("fp32")),
  1229. (core.dtype("fp64"), ): ("__nv_logb", core.dtype("fp64")),
  1230. }, is_pure=True, _semantic=_semantic)
  1231. @core.extern
  1232. def isfinited(arg0, _semantic=None):
  1233. return core.extern_elementwise("", "", [arg0], {
  1234. (core.dtype("fp64"), ): ("__nv_isfinited", core.dtype("int32")),
  1235. }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)