| 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735736737738739740741742743744745746747748749750751752753754755756757758759760761762763764765766767768769770771772773774775776777778779780781782783784785786787788789790791792793794795796797798799800801802803804805806807808809810811812813814815816817818819820821822823824825826827828829830831832833834835836837838839840841842843844845846847848849850851852853854855856857858859860861862863864865866867868869870871872873874875876877878879880881882883884885886887888889890891892893894895896897898899900901902903904905906907908909910911912913914915916917918919920921922923924925926927928929930931932933934935936937938939940941942943944945946947948949950951952953954955956957958959960961962963964965966967968969970971972973974975976977978979980981982983984985986987988989990991992993994995996997998999100010011002100310041005100610071008100910101011101210131014101510161017101810191020102110221023102410251026102710281029103010311032103310341035103610371038103910401041104210431044104510461047104810491050105110521053105410551056105710581059106010611062106310641065106610671068106910701071107210731074107510761077107810791080108110821083108410851086108710881089109010911092109310941095109610971098109911001101110211031104110511061107110811091110111111121113111411151116111711181119112011211122112311241125112611271128112911301131113211331134113511361137113811391140114111421143114411451146114711481149115011511152115311541155115611571158115911601161116211631164116511661167116811691170117111721173117411751176117711781179118011811182118311841185118611871188118911901191119211931194119511961197119811991200120112021203120412051206120712081209121012111212121312141215121612171218121912201221122212231224122512261227122812291230123112321233123412351236123712381239124012411242124312441245124612471248124912501251125212531254125512561257125812591260126112621263126412651266126712681269127012711272127312741275127612771278127912801281128212831284128512861287128812891290129112921293129412951296129712981299130013011302130313041305130613071308130913101311131213131314131513161317131813191320132113221323132413251326132713281329133013311332133313341335133613371338133913401341134213431344134513461347134813491350135113521353135413551356135713581359136013611362136313641365136613671368136913701371137213731374137513761377137813791380138113821383138413851386138713881389139013911392139313941395139613971398139914001401140214031404140514061407140814091410141114121413141414151416141714181419142014211422142314241425142614271428142914301431143214331434143514361437143814391440144114421443144414451446144714481449145014511452145314541455145614571458145914601461146214631464146514661467146814691470147114721473147414751476147714781479148014811482148314841485148614871488148914901491149214931494149514961497149814991500150115021503150415051506150715081509151015111512151315141515151615171518151915201521152215231524152515261527152815291530153115321533153415351536153715381539154015411542154315441545154615471548154915501551155215531554155515561557155815591560156115621563156415651566156715681569157015711572157315741575157615771578157915801581158215831584158515861587158815891590159115921593159415951596159715981599160016011602160316041605160616071608160916101611161216131614161516161617161816191620162116221623162416251626162716281629 |
- from triton.language import core
- @core.extern
- def clz(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_clz", core.dtype("int32")),
- (core.dtype("int64"), ): ("__nv_clzll", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def popc(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_popc", core.dtype("int32")),
- (core.dtype("int64"), ): ("__nv_popcll", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def byte_perm(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise("", "", [arg0, arg1, arg2], {
- (core.dtype("int32"), core.dtype("int32"), core.dtype("int32")): ("__nv_byte_perm", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def mulhi(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("int32")): ("__nv_mulhi", core.dtype("int32")),
- (core.dtype("uint32"), core.dtype("uint32")): ("__nv_umulhi", core.dtype("uint32")),
- (core.dtype("int64"), core.dtype("int64")): ("__nv_mul64hi", core.dtype("int64")),
- (core.dtype("uint64"), core.dtype("uint64")): ("__nv_umul64hi", core.dtype("uint64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def mul24(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("int32")): ("__nv_mul24", core.dtype("int32")),
- (core.dtype("uint32"), core.dtype("uint32")): ("__nv_umul24", core.dtype("uint32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def brev(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_brev", core.dtype("int32")),
- (core.dtype("int64"), ): ("__nv_brevll", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sad(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("int32"), core.dtype("int32"), core.dtype("uint32")): ("__nv_sad", core.dtype("int32")),
- (core.dtype("uint32"), core.dtype("uint32"), core.dtype("uint32")): ("__nv_usad", core.dtype("uint32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def abs(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_abs", core.dtype("int32")),
- (core.dtype("int64"), ): ("__nv_llabs", core.dtype("int64")),
- (core.dtype("fp32"), ): ("__nv_fabsf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_fabs", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def floor(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_floorf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_floor", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rcp64h(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_rcp64h", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rsqrt(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_rsqrtf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_rsqrt", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ceil(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_ceil", core.dtype("fp64")),
- (core.dtype("fp32"), ): ("__nv_ceilf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def trunc(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_trunc", core.dtype("fp64")),
- (core.dtype("fp32"), ): ("__nv_truncf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def exp2(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_exp2f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_exp2", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def saturatef(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_saturatef", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fma_rn(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_rn", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fma_rz(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_rz", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fma_rd(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_rd", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fma_ru(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf_ru", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_dividef(arg0, arg1, _semantic=None):
- return core.extern_elementwise("", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fast_fdividef", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def div_rn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_rn", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def div_rz(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_rz", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def div_rd(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_rd", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def div_ru(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdiv_ru", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_ddiv_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rcp_rn(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_frcp_rn", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_drcp_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rcp_rz(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_frcp_rz", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_drcp_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rcp_rd(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_frcp_rd", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_drcp_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rcp_ru(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_frcp_ru", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_drcp_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sqrt_rn(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fsqrt_rn", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_dsqrt_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sqrt_rz(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fsqrt_rz", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_dsqrt_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sqrt_rd(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fsqrt_rd", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_dsqrt_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sqrt_ru(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fsqrt_ru", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_dsqrt_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sqrt(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_sqrtf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_sqrt", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def add_rn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_rn", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def add_rz(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_rz", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def add_rd(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_rd", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def add_ru(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dadd_ru", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fadd_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def mul_rn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dmul_rn", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmul_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def mul_rz(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dmul_rz", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmul_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def mul_rd(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dmul_rd", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmul_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def mul_ru(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- arg1,
- ], {
- (
- core.dtype("fp64"),
- core.dtype("fp64"),
- ): ("__nv_dmul_ru", core.dtype("fp64")),
- (
- core.dtype("fp32"),
- core.dtype("fp32"),
- ): ("__nv_fmul_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2float_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2float_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2float_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2float_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2float_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2float_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2float_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2float_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2int_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2int_rn", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2int_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2int_rz", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2int_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2int_rd", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2int_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2int_ru", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2uint_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2uint_rn", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2uint_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2uint_rz", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2uint_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2uint_rd", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2uint_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2uint_ru", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def int2double_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_int2double_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def uint2double_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint32"), ): ("__nv_uint2double_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2int_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2int_rn", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2int_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2int_rz", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2int_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2int_rd", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2int_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2int_ru", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2uint_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2uint_rn", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2uint_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2uint_rz", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2uint_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2uint_rd", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2uint_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2uint_ru", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def int2float_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_int2float_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def int2float_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_int2float_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def int2float_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_int2float_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def int2float_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_int2float_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def uint2float_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint32"), ): ("__nv_uint2float_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def uint2float_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint32"), ): ("__nv_uint2float_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def uint2float_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint32"), ): ("__nv_uint2float_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def uint2float_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint32"), ): ("__nv_uint2float_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def hiloint2double(arg0, arg1, _semantic=None):
- return core.extern_elementwise("", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("int32")): ("__nv_hiloint2double", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2loint(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2loint", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2hiint(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2hiint", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ll_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ll_rn", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ll_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ll_rz", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ll_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ll_rd", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ll_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ll_ru", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ull_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ull_rn", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ull_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ull_rz", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ull_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ull_rd", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float2ull_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float2ull_ru", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ll_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ll_rn", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ll_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ll_rz", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ll_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ll_rd", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ll_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ll_ru", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ull_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ull_rn", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ull_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ull_rz", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ull_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ull_rd", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double2ull_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double2ull_ru", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2float_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2float_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2float_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2float_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2float_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2float_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2float_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2float_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2float_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2float_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2float_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2float_rz", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2float_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2float_rd", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2float_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2float_ru", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2double_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2double_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2double_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2double_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2double_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2double_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ll2double_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_ll2double_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2double_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2double_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2double_rz(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2double_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2double_rd(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2double_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ull2double_ru(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint64"), ): ("__nv_ull2double_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def int_as_float(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int32"), ): ("__nv_int_as_float", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float_as_int(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float_as_int", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def uint_as_float(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("uint32"), ): ("__nv_uint_as_float", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def float_as_uint(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_float_as_uint", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def longlong_as_double(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("int64"), ): ("__nv_longlong_as_double", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def double_as_longlong(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_double_as_longlong", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_sinf(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_sinf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_cosf(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_cosf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_log2f(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_log2f", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_logf(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_logf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_expf(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_expf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_tanf(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_tanf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_exp10f(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_exp10f", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_log10f(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_fast_log10f", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fast_powf(arg0, arg1, _semantic=None):
- return core.extern_elementwise("", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fast_powf", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def hadd(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("int32")): ("__nv_hadd", core.dtype("int32")),
- (core.dtype("uint32"), core.dtype("uint32")): ("__nv_uhadd", core.dtype("uint32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rhadd(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("int32")): ("__nv_rhadd", core.dtype("int32")),
- (core.dtype("uint32"), core.dtype("uint32")): ("__nv_urhadd", core.dtype("uint32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sub_rn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_rn", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_rn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sub_rz(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_rz", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_rz", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sub_rd(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_rd", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_rd", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sub_ru(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fsub_ru", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_dsub_ru", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rsqrt_rn(arg0, _semantic=None):
- return core.extern_elementwise("", "", [
- arg0,
- ], {
- (core.dtype("fp32"), ): ("__nv_frsqrt_rn", core.dtype("fp32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ffs(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- ], {
- (core.dtype("int32"), ): ("__nv_ffs", core.dtype("int32")),
- (core.dtype("int64"), ): ("__nv_ffsll", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rint(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- ], {
- (core.dtype("fp32"), ): ("__nv_rintf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_rint", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def llrint(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- ], {
- (core.dtype("fp32"), ): ("__nv_llrintf", core.dtype("int64")),
- (core.dtype("fp64"), ): ("__nv_llrint", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def nearbyint(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- ], {
- (core.dtype("fp32"), ): ("__nv_nearbyintf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_nearbyint", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def isnan(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- ], {
- (core.dtype("fp32"), ): ("__nv_isnanf", core.dtype("int32")),
- (core.dtype("fp64"), ): ("__nv_isnand", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
- @core.extern
- def signbit(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [
- arg0,
- ], {
- (core.dtype("fp32"), ): ("__nv_signbitf", core.dtype("int32")),
- (core.dtype("fp64"), ): ("__nv_signbitd", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def copysign(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_copysignf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_copysign", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def finitef(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_finitef", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
- @core.extern
- def isinf(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_isinff", core.dtype("int32")),
- (core.dtype("fp64"), ): ("__nv_isinfd", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
- @core.extern
- def nextafter(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_nextafterf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_nextafter", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sin(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_sinf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_sin", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def cos(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_cosf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_cos", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sinpi(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_sinpif", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_sinpi", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def cospi(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_cospif", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_cospi", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def tan(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_tanf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_tan", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def log2(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_log2f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_log2", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def exp(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_expf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_exp", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def exp10(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_exp10f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_exp10", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def cosh(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_coshf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_cosh", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def sinh(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_sinhf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_sinh", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def tanh(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_tanhf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_tanh", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def atan2(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_atan2f", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_atan2", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def atan(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_atanf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_atan", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def asin(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_asinf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_asin", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def acos(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_acosf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_acos", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def log(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_logf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_log", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def log10(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_log10f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_log10", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def log1p(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_log1pf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_log1p", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def acosh(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_acoshf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_acosh", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def asinh(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_asinhf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_asinh", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def atanh(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_atanhf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_atanh", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def expm1(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_expm1f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_expm1", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def hypot(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_hypotf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_hypot", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rhypot(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_rhypotf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_rhypot", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def norm3d(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_norm3df", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_norm3d", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rnorm3d(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_rnorm3df", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_rnorm3d", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def norm4d(arg0, arg1, arg2, arg3, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2, arg3], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")):
- ("__nv_norm4df", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")):
- ("__nv_norm4d", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rnorm4d(arg0, arg1, arg2, arg3, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2, arg3], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")):
- ("__nv_rnorm4df", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")):
- ("__nv_rnorm4d", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def cbrt(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_cbrtf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_cbrt", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def rcbrt(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_rcbrtf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_rcbrt", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def j0(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_j0f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_j0", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def j1(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_j1f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_j1", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def y0(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_y0f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_y0", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def y1(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_y1f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_y1", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def yn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("fp32")): ("__nv_ynf", core.dtype("fp32")),
- (core.dtype("int32"), core.dtype("fp64")): ("__nv_yn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def jn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("int32"), core.dtype("fp32")): ("__nv_jnf", core.dtype("fp32")),
- (core.dtype("int32"), core.dtype("fp64")): ("__nv_jn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def cyl_bessel_i0(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_cyl_bessel_i0f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_cyl_bessel_i0", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def cyl_bessel_i1(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_cyl_bessel_i1f", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_cyl_bessel_i1", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def erf(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_erff", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_erf", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def erfinv(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_erfinvf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_erfinv", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def erfc(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_erfcf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_erfc", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def erfcx(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_erfcxf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_erfcx", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def erfcinv(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_erfcinvf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_erfcinv", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def normcdfinv(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_normcdfinvf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_normcdfinv", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def normcdf(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_normcdff", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_normcdf", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def lgamma(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_lgammaf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_lgamma", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ldexp(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("int32")): ("__nv_ldexpf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("int32")): ("__nv_ldexp", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def scalbn(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("int32")): ("__nv_scalbnf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("int32")): ("__nv_scalbn", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fmod(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmodf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_fmod", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def remainder(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_remainderf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_remainder", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fma(arg0, arg1, arg2, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1, arg2], {
- (core.dtype("fp32"), core.dtype("fp32"), core.dtype("fp32")): ("__nv_fmaf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64"), core.dtype("fp64")): ("__nv_fma", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def pow(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("int32")): ("__nv_powif", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("int32")): ("__nv_powi", core.dtype("fp64")),
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_powf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_pow", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def tgamma(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_tgammaf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_tgamma", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def round(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_roundf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_round", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def llround(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_llroundf", core.dtype("int64")),
- (core.dtype("fp64"), ): ("__nv_llround", core.dtype("int64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def fdim(arg0, arg1, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0, arg1], {
- (core.dtype("fp32"), core.dtype("fp32")): ("__nv_fdimf", core.dtype("fp32")),
- (core.dtype("fp64"), core.dtype("fp64")): ("__nv_fdim", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def ilogb(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_ilogbf", core.dtype("int32")),
- (core.dtype("fp64"), ): ("__nv_ilogb", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def logb(arg0, _semantic=None):
- return core.extern_elementwise(
- "", "", [arg0], {
- (core.dtype("fp32"), ): ("__nv_logbf", core.dtype("fp32")),
- (core.dtype("fp64"), ): ("__nv_logb", core.dtype("fp64")),
- }, is_pure=True, _semantic=_semantic)
- @core.extern
- def isfinited(arg0, _semantic=None):
- return core.extern_elementwise("", "", [arg0], {
- (core.dtype("fp64"), ): ("__nv_isfinited", core.dtype("int32")),
- }, is_pure=True, _semantic=_semantic).to(core.int1, _semantic=_semantic)
|