10b57cec5SDimitry Andric//===--- BuiltinsPTX.def - PTX Builtin function database ----*- C++ -*-===// 20b57cec5SDimitry Andric// 30b57cec5SDimitry Andric// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 40b57cec5SDimitry Andric// See https://llvm.org/LICENSE.txt for license information. 50b57cec5SDimitry Andric// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 60b57cec5SDimitry Andric// 70b57cec5SDimitry Andric//===----------------------------------------------------------------------===// 80b57cec5SDimitry Andric// 90b57cec5SDimitry Andric// This file defines the PTX-specific builtin function database. Users of 100b57cec5SDimitry Andric// this file must define the BUILTIN macro to make use of this information. 110b57cec5SDimitry Andric// 120b57cec5SDimitry Andric//===----------------------------------------------------------------------===// 130b57cec5SDimitry Andric 140b57cec5SDimitry Andric// The format of this database matches clang/Basic/Builtins.def. 150b57cec5SDimitry Andric 160b57cec5SDimitry Andric#if defined(BUILTIN) && !defined(TARGET_BUILTIN) 170b57cec5SDimitry Andric# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS) 180b57cec5SDimitry Andric#endif 190b57cec5SDimitry Andric 2081ad6265SDimitry Andric#pragma push_macro("SM_53") 210b57cec5SDimitry Andric#pragma push_macro("SM_70") 220b57cec5SDimitry Andric#pragma push_macro("SM_72") 230b57cec5SDimitry Andric#pragma push_macro("SM_75") 245ffd83dbSDimitry Andric#pragma push_macro("SM_80") 25fe6060f1SDimitry Andric#pragma push_macro("SM_86") 26bdd1243dSDimitry Andric#pragma push_macro("SM_87") 27bdd1243dSDimitry Andric#pragma push_macro("SM_89") 28bdd1243dSDimitry Andric#pragma push_macro("SM_90") 295f757f3fSDimitry Andric#pragma push_macro("SM_90a") 305f757f3fSDimitry Andric#define SM_90a "sm_90a" 315f757f3fSDimitry Andric#define SM_90 "sm_90|" SM_90a 32bdd1243dSDimitry Andric#define SM_89 "sm_89|" SM_90 33bdd1243dSDimitry Andric#define SM_87 "sm_87|" SM_89 34bdd1243dSDimitry Andric#define SM_86 "sm_86|" SM_87 35fe6060f1SDimitry Andric#define SM_80 "sm_80|" SM_86 365ffd83dbSDimitry Andric#define SM_75 "sm_75|" SM_80 370b57cec5SDimitry Andric#define SM_72 "sm_72|" SM_75 380b57cec5SDimitry Andric#define SM_70 "sm_70|" SM_72 390b57cec5SDimitry Andric 400b57cec5SDimitry Andric#pragma push_macro("SM_60") 410b57cec5SDimitry Andric#define SM_60 "sm_60|sm_61|sm_62|" SM_70 4281ad6265SDimitry Andric#define SM_53 "sm_53|" SM_60 430b57cec5SDimitry Andric 4481ad6265SDimitry Andric#pragma push_macro("PTX42") 450b57cec5SDimitry Andric#pragma push_macro("PTX60") 460b57cec5SDimitry Andric#pragma push_macro("PTX61") 470b57cec5SDimitry Andric#pragma push_macro("PTX63") 480b57cec5SDimitry Andric#pragma push_macro("PTX64") 495ffd83dbSDimitry Andric#pragma push_macro("PTX65") 505ffd83dbSDimitry Andric#pragma push_macro("PTX70") 51fe6060f1SDimitry Andric#pragma push_macro("PTX71") 52fe6060f1SDimitry Andric#pragma push_macro("PTX72") 53349cc55cSDimitry Andric#pragma push_macro("PTX73") 54349cc55cSDimitry Andric#pragma push_macro("PTX74") 55349cc55cSDimitry Andric#pragma push_macro("PTX75") 56bdd1243dSDimitry Andric#pragma push_macro("PTX76") 57bdd1243dSDimitry Andric#pragma push_macro("PTX77") 58bdd1243dSDimitry Andric#pragma push_macro("PTX78") 5906c3fb27SDimitry Andric#pragma push_macro("PTX80") 6006c3fb27SDimitry Andric#pragma push_macro("PTX81") 615f757f3fSDimitry Andric#pragma push_macro("PTX82") 625f757f3fSDimitry Andric#pragma push_macro("PTX83") 635f757f3fSDimitry Andric#define PTX83 "ptx83" 645f757f3fSDimitry Andric#define PTX82 "ptx82|" PTX83 655f757f3fSDimitry Andric#define PTX81 "ptx81|" PTX82 6606c3fb27SDimitry Andric#define PTX80 "ptx80|" PTX81 6706c3fb27SDimitry Andric#define PTX78 "ptx78|" PTX80 68bdd1243dSDimitry Andric#define PTX77 "ptx77|" PTX78 69bdd1243dSDimitry Andric#define PTX76 "ptx76|" PTX77 70bdd1243dSDimitry Andric#define PTX75 "ptx75|" PTX76 71349cc55cSDimitry Andric#define PTX74 "ptx74|" PTX75 72349cc55cSDimitry Andric#define PTX73 "ptx73|" PTX74 73349cc55cSDimitry Andric#define PTX72 "ptx72|" PTX73 74fe6060f1SDimitry Andric#define PTX71 "ptx71|" PTX72 75fe6060f1SDimitry Andric#define PTX70 "ptx70|" PTX71 765ffd83dbSDimitry Andric#define PTX65 "ptx65|" PTX70 775ffd83dbSDimitry Andric#define PTX64 "ptx64|" PTX65 780b57cec5SDimitry Andric#define PTX63 "ptx63|" PTX64 790b57cec5SDimitry Andric#define PTX61 "ptx61|" PTX63 800b57cec5SDimitry Andric#define PTX60 "ptx60|" PTX61 8181ad6265SDimitry Andric#define PTX42 "ptx42|" PTX60 820b57cec5SDimitry Andric 830b57cec5SDimitry Andric#pragma push_macro("AND") 84e8d8bef9SDimitry Andric#define AND(a, b) "(" a "),(" b ")" 850b57cec5SDimitry Andric 860b57cec5SDimitry Andric// Special Registers 870b57cec5SDimitry Andric 880b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_tid_x, "i", "nc") 890b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_tid_y, "i", "nc") 900b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_tid_z, "i", "nc") 910b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_tid_w, "i", "nc") 920b57cec5SDimitry Andric 930b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ntid_x, "i", "nc") 940b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ntid_y, "i", "nc") 950b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ntid_z, "i", "nc") 960b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ntid_w, "i", "nc") 970b57cec5SDimitry Andric 980b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ctaid_x, "i", "nc") 990b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ctaid_y, "i", "nc") 1000b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ctaid_z, "i", "nc") 1010b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_ctaid_w, "i", "nc") 1020b57cec5SDimitry Andric 1030b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_nctaid_x, "i", "nc") 1040b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_nctaid_y, "i", "nc") 1050b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_nctaid_z, "i", "nc") 1060b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_nctaid_w, "i", "nc") 1070b57cec5SDimitry Andric 10806c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_x, "i", "nc", AND(SM_90, PTX78)) 10906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_y, "i", "nc", AND(SM_90, PTX78)) 11006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_z, "i", "nc", AND(SM_90, PTX78)) 11106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_w, "i", "nc", AND(SM_90, PTX78)) 11206c3fb27SDimitry Andric 11306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_x, "i", "nc", AND(SM_90, PTX78)) 11406c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_y, "i", "nc", AND(SM_90, PTX78)) 11506c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_z, "i", "nc", AND(SM_90, PTX78)) 11606c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_w, "i", "nc", AND(SM_90, PTX78)) 11706c3fb27SDimitry Andric 11806c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_x, "i", "nc", AND(SM_90, PTX78)) 11906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_y, "i", "nc", AND(SM_90, PTX78)) 12006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_z, "i", "nc", AND(SM_90, PTX78)) 12106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_w, "i", "nc", AND(SM_90, PTX78)) 12206c3fb27SDimitry Andric 12306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_x, "i", "nc", AND(SM_90, PTX78)) 12406c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_y, "i", "nc", AND(SM_90, PTX78)) 12506c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_z, "i", "nc", AND(SM_90, PTX78)) 12606c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_w, "i", "nc", AND(SM_90, PTX78)) 12706c3fb27SDimitry Andric 12806c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctarank, "i", "nc", AND(SM_90, PTX78)) 12906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctarank, "i", "nc", AND(SM_90, PTX78)) 13006c3fb27SDimitry Andric 13106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_is_explicit_cluster, "b", "nc", AND(SM_90, PTX78)) 13206c3fb27SDimitry Andric 1330b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_laneid, "i", "nc") 1340b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_warpid, "i", "nc") 1350b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_nwarpid, "i", "nc") 1360b57cec5SDimitry Andric 1370b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_smid, "i", "nc") 1380b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_nsmid, "i", "nc") 1390b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_gridid, "i", "nc") 1400b57cec5SDimitry Andric 1410b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_lanemask_eq, "i", "nc") 1420b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_lanemask_le, "i", "nc") 1430b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_lanemask_lt, "i", "nc") 1440b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_lanemask_ge, "i", "nc") 1450b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_lanemask_gt, "i", "nc") 1460b57cec5SDimitry Andric 1470b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_clock, "i", "n") 1480b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_clock64, "LLi", "n") 1490b57cec5SDimitry Andric 1500b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_pm0, "i", "n") 1510b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_pm1, "i", "n") 1520b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_pm2, "i", "n") 1530b57cec5SDimitry AndricBUILTIN(__nvvm_read_ptx_sreg_pm3, "i", "n") 1540b57cec5SDimitry Andric 1550b57cec5SDimitry Andric// MISC 1560b57cec5SDimitry Andric 1570b57cec5SDimitry AndricBUILTIN(__nvvm_prmt, "UiUiUiUi", "") 1580b57cec5SDimitry Andric 1590b57cec5SDimitry Andric// Min Max 1600b57cec5SDimitry Andric 16181ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_f16, "hhh", "", AND(SM_80, PTX70)) 16281ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_f16, "hhh", "", AND(SM_80, PTX70)) 16381ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_f16, "hhh", "", AND(SM_80, PTX70)) 16481ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_f16, "hhh", "", AND(SM_80, PTX70)) 16581ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_xorsign_abs_f16, "hhh", "", AND(SM_86, PTX72)) 16681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_xorsign_abs_f16, "hhh", "", AND(SM_86, PTX72)) 16781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_f16, "hhh", "", AND(SM_86, PTX72)) 16881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_xorsign_abs_f16, "hhh", "", 16981ad6265SDimitry Andric AND(SM_86, PTX72)) 17081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 17181ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 17281ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 17381ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 17481ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_xorsign_abs_f16x2, "V2hV2hV2h", "", 17581ad6265SDimitry Andric AND(SM_86, PTX72)) 17681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_xorsign_abs_f16x2, "V2hV2hV2h", "", 17781ad6265SDimitry Andric AND(SM_86, PTX72)) 17881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_f16x2, "V2hV2hV2h", "", 17981ad6265SDimitry Andric AND(SM_86, PTX72)) 18081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_xorsign_abs_f16x2, "V2hV2hV2h", "", 18181ad6265SDimitry Andric AND(SM_86, PTX72)) 18206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_bf16, "yyy", "", AND(SM_80, PTX70)) 18306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_bf16, "yyy", "", AND(SM_80, PTX70)) 18406c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_bf16, "yyy", "", AND(SM_80, PTX70)) 18506c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_bf16, "yyy", "", AND(SM_80, PTX70)) 18606c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_xorsign_abs_bf16, "yyy", "", AND(SM_86, PTX72)) 18706c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_bf16, "yyy", "", 18881ad6265SDimitry Andric AND(SM_86, PTX72)) 18906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 19006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 19106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 19206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 19306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_xorsign_abs_bf16x2, "V2yV2yV2y", "", 19481ad6265SDimitry Andric AND(SM_86, PTX72)) 19506c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_bf16x2, "V2yV2yV2y", "", 19681ad6265SDimitry Andric AND(SM_86, PTX72)) 1970b57cec5SDimitry AndricBUILTIN(__nvvm_fmin_f, "fff", "") 19881ad6265SDimitry AndricBUILTIN(__nvvm_fmin_ftz_f, "fff", "") 19981ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_f, "fff", "", AND(SM_80, PTX70)) 20081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_f, "fff", "", AND(SM_80, PTX70)) 20181ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 20281ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 20381ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 20481ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmin_ftz_nan_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 2050b57cec5SDimitry AndricBUILTIN(__nvvm_fmin_d, "ddd", "") 2060b57cec5SDimitry Andric 20781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_f16, "hhh", "", AND(SM_80, PTX70)) 20881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_f16, "hhh", "", AND(SM_80, PTX70)) 20981ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_f16, "hhh", "", AND(SM_80, PTX70)) 21081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_f16, "hhh", "", AND(SM_80, PTX70)) 21181ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_xorsign_abs_f16, "hhh", "", AND(SM_86, PTX72)) 21281ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_xorsign_abs_f16, "hhh", "", AND(SM_86, PTX72)) 21381ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_f16, "hhh", "", AND(SM_86, PTX72)) 21481ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_xorsign_abs_f16, "hhh", "", 21581ad6265SDimitry Andric AND(SM_86, PTX72)) 21681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 21781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 21881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 21981ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_f16x2, "V2hV2hV2h", "", AND(SM_80, PTX70)) 22081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_xorsign_abs_f16x2, "V2hV2hV2h", "", 22181ad6265SDimitry Andric AND(SM_86, PTX72)) 22281ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_xorsign_abs_f16x2, "V2hV2hV2h", "", 22381ad6265SDimitry Andric AND(SM_86, PTX72)) 22481ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_f16x2, "V2hV2hV2h", "", 22581ad6265SDimitry Andric AND(SM_86, PTX72)) 22681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_xorsign_abs_f16x2, "V2hV2hV2h", "", 22781ad6265SDimitry Andric AND(SM_86, PTX72)) 22806c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_bf16, "yyy", "", AND(SM_80, PTX70)) 22906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_bf16, "yyy", "", AND(SM_80, PTX70)) 23006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_bf16, "yyy", "", AND(SM_80, PTX70)) 23106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_bf16, "yyy", "", AND(SM_80, PTX70)) 23206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_xorsign_abs_bf16, "yyy", "", AND(SM_86, PTX72)) 23306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_bf16, "yyy", "", 23481ad6265SDimitry Andric AND(SM_86, PTX72)) 23506c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 23606c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 23706c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 23806c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_bf16x2, "V2yV2yV2y", "", AND(SM_80, PTX70)) 23906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_xorsign_abs_bf16x2, "V2yV2yV2y", "", 24081ad6265SDimitry Andric AND(SM_86, PTX72)) 24106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_bf16x2, "V2yV2yV2y", "", 24281ad6265SDimitry Andric AND(SM_86, PTX72)) 24381ad6265SDimitry AndricBUILTIN(__nvvm_fmax_f, "fff", "") 24481ad6265SDimitry AndricBUILTIN(__nvvm_fmax_ftz_f, "fff", "") 24581ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_f, "fff", "", AND(SM_80, PTX70)) 24681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_f, "fff", "", AND(SM_80, PTX70)) 24781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 24881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 24981ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 25081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fmax_ftz_nan_xorsign_abs_f, "fff", "", AND(SM_86, PTX72)) 25181ad6265SDimitry AndricBUILTIN(__nvvm_fmax_d, "ddd", "") 25281ad6265SDimitry Andric 2530b57cec5SDimitry Andric// Multiplication 2540b57cec5SDimitry Andric 2550b57cec5SDimitry AndricBUILTIN(__nvvm_mulhi_i, "iii", "") 2560b57cec5SDimitry AndricBUILTIN(__nvvm_mulhi_ui, "UiUiUi", "") 2570b57cec5SDimitry AndricBUILTIN(__nvvm_mulhi_ll, "LLiLLiLLi", "") 2580b57cec5SDimitry AndricBUILTIN(__nvvm_mulhi_ull, "ULLiULLiULLi", "") 2590b57cec5SDimitry Andric 2600b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rn_ftz_f, "fff", "") 2610b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rn_f, "fff", "") 2620b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rz_ftz_f, "fff", "") 2630b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rz_f, "fff", "") 2640b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rm_ftz_f, "fff", "") 2650b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rm_f, "fff", "") 2660b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rp_ftz_f, "fff", "") 2670b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rp_f, "fff", "") 2680b57cec5SDimitry Andric 2690b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rn_d, "ddd", "") 2700b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rz_d, "ddd", "") 2710b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rm_d, "ddd", "") 2720b57cec5SDimitry AndricBUILTIN(__nvvm_mul_rp_d, "ddd", "") 2730b57cec5SDimitry Andric 2740b57cec5SDimitry AndricBUILTIN(__nvvm_mul24_i, "iii", "") 2750b57cec5SDimitry AndricBUILTIN(__nvvm_mul24_ui, "UiUiUi", "") 2760b57cec5SDimitry Andric 2770b57cec5SDimitry Andric// Div 2780b57cec5SDimitry Andric 2790b57cec5SDimitry AndricBUILTIN(__nvvm_div_approx_ftz_f, "fff", "") 2800b57cec5SDimitry AndricBUILTIN(__nvvm_div_approx_f, "fff", "") 2810b57cec5SDimitry Andric 2820b57cec5SDimitry AndricBUILTIN(__nvvm_div_rn_ftz_f, "fff", "") 2830b57cec5SDimitry AndricBUILTIN(__nvvm_div_rn_f, "fff", "") 2840b57cec5SDimitry AndricBUILTIN(__nvvm_div_rz_ftz_f, "fff", "") 2850b57cec5SDimitry AndricBUILTIN(__nvvm_div_rz_f, "fff", "") 2860b57cec5SDimitry AndricBUILTIN(__nvvm_div_rm_ftz_f, "fff", "") 2870b57cec5SDimitry AndricBUILTIN(__nvvm_div_rm_f, "fff", "") 2880b57cec5SDimitry AndricBUILTIN(__nvvm_div_rp_ftz_f, "fff", "") 2890b57cec5SDimitry AndricBUILTIN(__nvvm_div_rp_f, "fff", "") 2900b57cec5SDimitry Andric 2910b57cec5SDimitry AndricBUILTIN(__nvvm_div_rn_d, "ddd", "") 2920b57cec5SDimitry AndricBUILTIN(__nvvm_div_rz_d, "ddd", "") 2930b57cec5SDimitry AndricBUILTIN(__nvvm_div_rm_d, "ddd", "") 2940b57cec5SDimitry AndricBUILTIN(__nvvm_div_rp_d, "ddd", "") 2950b57cec5SDimitry Andric 2960b57cec5SDimitry Andric// Sad 2970b57cec5SDimitry Andric 2980b57cec5SDimitry AndricBUILTIN(__nvvm_sad_i, "iiii", "") 2990b57cec5SDimitry AndricBUILTIN(__nvvm_sad_ui, "UiUiUiUi", "") 3000b57cec5SDimitry Andric 3010b57cec5SDimitry Andric// Floor, Ceil 3020b57cec5SDimitry Andric 3030b57cec5SDimitry AndricBUILTIN(__nvvm_floor_ftz_f, "ff", "") 3040b57cec5SDimitry AndricBUILTIN(__nvvm_floor_f, "ff", "") 3050b57cec5SDimitry AndricBUILTIN(__nvvm_floor_d, "dd", "") 3060b57cec5SDimitry Andric 3070b57cec5SDimitry AndricBUILTIN(__nvvm_ceil_ftz_f, "ff", "") 3080b57cec5SDimitry AndricBUILTIN(__nvvm_ceil_f, "ff", "") 3090b57cec5SDimitry AndricBUILTIN(__nvvm_ceil_d, "dd", "") 3100b57cec5SDimitry Andric 3110b57cec5SDimitry Andric// Abs 3120b57cec5SDimitry Andric 3130b57cec5SDimitry AndricBUILTIN(__nvvm_fabs_ftz_f, "ff", "") 3140b57cec5SDimitry AndricBUILTIN(__nvvm_fabs_f, "ff", "") 3150b57cec5SDimitry AndricBUILTIN(__nvvm_fabs_d, "dd", "") 3160b57cec5SDimitry Andric 3170b57cec5SDimitry Andric// Round 3180b57cec5SDimitry Andric 3190b57cec5SDimitry AndricBUILTIN(__nvvm_round_ftz_f, "ff", "") 3200b57cec5SDimitry AndricBUILTIN(__nvvm_round_f, "ff", "") 3210b57cec5SDimitry AndricBUILTIN(__nvvm_round_d, "dd", "") 3220b57cec5SDimitry Andric 3230b57cec5SDimitry Andric// Trunc 3240b57cec5SDimitry Andric 3250b57cec5SDimitry AndricBUILTIN(__nvvm_trunc_ftz_f, "ff", "") 3260b57cec5SDimitry AndricBUILTIN(__nvvm_trunc_f, "ff", "") 3270b57cec5SDimitry AndricBUILTIN(__nvvm_trunc_d, "dd", "") 3280b57cec5SDimitry Andric 3290b57cec5SDimitry Andric// Saturate 3300b57cec5SDimitry Andric 3310b57cec5SDimitry AndricBUILTIN(__nvvm_saturate_ftz_f, "ff", "") 3320b57cec5SDimitry AndricBUILTIN(__nvvm_saturate_f, "ff", "") 3330b57cec5SDimitry AndricBUILTIN(__nvvm_saturate_d, "dd", "") 3340b57cec5SDimitry Andric 3350b57cec5SDimitry Andric// Exp2, Log2 3360b57cec5SDimitry Andric 3370b57cec5SDimitry AndricBUILTIN(__nvvm_ex2_approx_ftz_f, "ff", "") 3380b57cec5SDimitry AndricBUILTIN(__nvvm_ex2_approx_f, "ff", "") 3390b57cec5SDimitry AndricBUILTIN(__nvvm_ex2_approx_d, "dd", "") 34081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_ex2_approx_f16, "hh", "", AND(SM_75, PTX70)) 34181ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_ex2_approx_f16x2, "V2hV2h", "", AND(SM_75, PTX70)) 3420b57cec5SDimitry Andric 3430b57cec5SDimitry AndricBUILTIN(__nvvm_lg2_approx_ftz_f, "ff", "") 3440b57cec5SDimitry AndricBUILTIN(__nvvm_lg2_approx_f, "ff", "") 3450b57cec5SDimitry AndricBUILTIN(__nvvm_lg2_approx_d, "dd", "") 3460b57cec5SDimitry Andric 3470b57cec5SDimitry Andric// Sin, Cos 3480b57cec5SDimitry Andric 3490b57cec5SDimitry AndricBUILTIN(__nvvm_sin_approx_ftz_f, "ff", "") 3500b57cec5SDimitry AndricBUILTIN(__nvvm_sin_approx_f, "ff", "") 3510b57cec5SDimitry Andric 3520b57cec5SDimitry AndricBUILTIN(__nvvm_cos_approx_ftz_f, "ff", "") 3530b57cec5SDimitry AndricBUILTIN(__nvvm_cos_approx_f, "ff", "") 3540b57cec5SDimitry Andric 3550b57cec5SDimitry Andric// Fma 3560b57cec5SDimitry Andric 35781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_f16, "hhhh", "", AND(SM_53, PTX42)) 35881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_ftz_f16, "hhhh", "", AND(SM_53, PTX42)) 35981ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_sat_f16, "hhhh", "", AND(SM_53, PTX42)) 36081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_ftz_sat_f16, "hhhh", "", AND(SM_53, PTX42)) 36181ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_relu_f16, "hhhh", "", AND(SM_80, PTX70)) 36281ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_ftz_relu_f16, "hhhh", "", AND(SM_80, PTX70)) 36381ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42)) 36481ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_ftz_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42)) 36581ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_sat_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42)) 36681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_ftz_sat_f16x2, "V2hV2hV2hV2h", "", AND(SM_53, PTX42)) 36781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_relu_f16x2, "V2hV2hV2hV2h", "", AND(SM_80, PTX70)) 36881ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_ftz_relu_f16x2, "V2hV2hV2hV2h", "", AND(SM_80, PTX70)) 36906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_bf16, "yyyy", "", AND(SM_80, PTX70)) 37006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_relu_bf16, "yyyy", "", AND(SM_80, PTX70)) 37106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_bf16x2, "V2yV2yV2yV2y", "", AND(SM_80, PTX70)) 37206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fma_rn_relu_bf16x2, "V2yV2yV2yV2y", "", AND(SM_80, PTX70)) 3730b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rn_ftz_f, "ffff", "") 3740b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rn_f, "ffff", "") 3750b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rz_ftz_f, "ffff", "") 3760b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rz_f, "ffff", "") 3770b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rm_ftz_f, "ffff", "") 3780b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rm_f, "ffff", "") 3790b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rp_ftz_f, "ffff", "") 3800b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rp_f, "ffff", "") 3810b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rn_d, "dddd", "") 3820b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rz_d, "dddd", "") 3830b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rm_d, "dddd", "") 3840b57cec5SDimitry AndricBUILTIN(__nvvm_fma_rp_d, "dddd", "") 3850b57cec5SDimitry Andric 3860b57cec5SDimitry Andric// Rcp 3870b57cec5SDimitry Andric 3880b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rn_ftz_f, "ff", "") 3890b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rn_f, "ff", "") 3900b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rz_ftz_f, "ff", "") 3910b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rz_f, "ff", "") 3920b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rm_ftz_f, "ff", "") 3930b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rm_f, "ff", "") 3940b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rp_ftz_f, "ff", "") 3950b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rp_f, "ff", "") 3960b57cec5SDimitry Andric 3970b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rn_d, "dd", "") 3980b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rz_d, "dd", "") 3990b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rm_d, "dd", "") 4000b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_rp_d, "dd", "") 40181ad6265SDimitry Andric 40281ad6265SDimitry AndricBUILTIN(__nvvm_rcp_approx_ftz_f, "ff", "") 4030b57cec5SDimitry AndricBUILTIN(__nvvm_rcp_approx_ftz_d, "dd", "") 4040b57cec5SDimitry Andric 4050b57cec5SDimitry Andric// Sqrt 4060b57cec5SDimitry Andric 4070b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rn_ftz_f, "ff", "") 4080b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rn_f, "ff", "") 4090b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rz_ftz_f, "ff", "") 4100b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rz_f, "ff", "") 4110b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rm_ftz_f, "ff", "") 4120b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rm_f, "ff", "") 4130b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rp_ftz_f, "ff", "") 4140b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rp_f, "ff", "") 4150b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_approx_ftz_f, "ff", "") 4160b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_approx_f, "ff", "") 4170b57cec5SDimitry Andric 4180b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rn_d, "dd", "") 4190b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rz_d, "dd", "") 4200b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rm_d, "dd", "") 4210b57cec5SDimitry AndricBUILTIN(__nvvm_sqrt_rp_d, "dd", "") 4220b57cec5SDimitry Andric 4230b57cec5SDimitry Andric// Rsqrt 4240b57cec5SDimitry Andric 4250b57cec5SDimitry AndricBUILTIN(__nvvm_rsqrt_approx_ftz_f, "ff", "") 4260b57cec5SDimitry AndricBUILTIN(__nvvm_rsqrt_approx_f, "ff", "") 4270b57cec5SDimitry AndricBUILTIN(__nvvm_rsqrt_approx_d, "dd", "") 4280b57cec5SDimitry Andric 4290b57cec5SDimitry Andric// Add 4300b57cec5SDimitry Andric 4310b57cec5SDimitry AndricBUILTIN(__nvvm_add_rn_ftz_f, "fff", "") 4320b57cec5SDimitry AndricBUILTIN(__nvvm_add_rn_f, "fff", "") 4330b57cec5SDimitry AndricBUILTIN(__nvvm_add_rz_ftz_f, "fff", "") 4340b57cec5SDimitry AndricBUILTIN(__nvvm_add_rz_f, "fff", "") 4350b57cec5SDimitry AndricBUILTIN(__nvvm_add_rm_ftz_f, "fff", "") 4360b57cec5SDimitry AndricBUILTIN(__nvvm_add_rm_f, "fff", "") 4370b57cec5SDimitry AndricBUILTIN(__nvvm_add_rp_ftz_f, "fff", "") 4380b57cec5SDimitry AndricBUILTIN(__nvvm_add_rp_f, "fff", "") 4390b57cec5SDimitry Andric 4400b57cec5SDimitry AndricBUILTIN(__nvvm_add_rn_d, "ddd", "") 4410b57cec5SDimitry AndricBUILTIN(__nvvm_add_rz_d, "ddd", "") 4420b57cec5SDimitry AndricBUILTIN(__nvvm_add_rm_d, "ddd", "") 4430b57cec5SDimitry AndricBUILTIN(__nvvm_add_rp_d, "ddd", "") 4440b57cec5SDimitry Andric 4450b57cec5SDimitry Andric// Convert 4460b57cec5SDimitry Andric 4470b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rn_ftz, "fd", "") 4480b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rn, "fd", "") 4490b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rz_ftz, "fd", "") 4500b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rz, "fd", "") 4510b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rm_ftz, "fd", "") 4520b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rm, "fd", "") 4530b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rp_ftz, "fd", "") 4540b57cec5SDimitry AndricBUILTIN(__nvvm_d2f_rp, "fd", "") 4550b57cec5SDimitry Andric 4560b57cec5SDimitry AndricBUILTIN(__nvvm_d2i_rn, "id", "") 4570b57cec5SDimitry AndricBUILTIN(__nvvm_d2i_rz, "id", "") 4580b57cec5SDimitry AndricBUILTIN(__nvvm_d2i_rm, "id", "") 4590b57cec5SDimitry AndricBUILTIN(__nvvm_d2i_rp, "id", "") 4600b57cec5SDimitry Andric 4610b57cec5SDimitry AndricBUILTIN(__nvvm_d2ui_rn, "Uid", "") 4620b57cec5SDimitry AndricBUILTIN(__nvvm_d2ui_rz, "Uid", "") 4630b57cec5SDimitry AndricBUILTIN(__nvvm_d2ui_rm, "Uid", "") 4640b57cec5SDimitry AndricBUILTIN(__nvvm_d2ui_rp, "Uid", "") 4650b57cec5SDimitry Andric 4660b57cec5SDimitry AndricBUILTIN(__nvvm_i2d_rn, "di", "") 4670b57cec5SDimitry AndricBUILTIN(__nvvm_i2d_rz, "di", "") 4680b57cec5SDimitry AndricBUILTIN(__nvvm_i2d_rm, "di", "") 4690b57cec5SDimitry AndricBUILTIN(__nvvm_i2d_rp, "di", "") 4700b57cec5SDimitry Andric 4710b57cec5SDimitry AndricBUILTIN(__nvvm_ui2d_rn, "dUi", "") 4720b57cec5SDimitry AndricBUILTIN(__nvvm_ui2d_rz, "dUi", "") 4730b57cec5SDimitry AndricBUILTIN(__nvvm_ui2d_rm, "dUi", "") 4740b57cec5SDimitry AndricBUILTIN(__nvvm_ui2d_rp, "dUi", "") 4750b57cec5SDimitry Andric 4760b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rn_ftz, "if", "") 4770b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rn, "if", "") 4780b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rz_ftz, "if", "") 4790b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rz, "if", "") 4800b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rm_ftz, "if", "") 4810b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rm, "if", "") 4820b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rp_ftz, "if", "") 4830b57cec5SDimitry AndricBUILTIN(__nvvm_f2i_rp, "if", "") 4840b57cec5SDimitry Andric 4850b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rn_ftz, "Uif", "") 4860b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rn, "Uif", "") 4870b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rz_ftz, "Uif", "") 4880b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rz, "Uif", "") 4890b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rm_ftz, "Uif", "") 4900b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rm, "Uif", "") 4910b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rp_ftz, "Uif", "") 4920b57cec5SDimitry AndricBUILTIN(__nvvm_f2ui_rp, "Uif", "") 4930b57cec5SDimitry Andric 4940b57cec5SDimitry AndricBUILTIN(__nvvm_i2f_rn, "fi", "") 4950b57cec5SDimitry AndricBUILTIN(__nvvm_i2f_rz, "fi", "") 4960b57cec5SDimitry AndricBUILTIN(__nvvm_i2f_rm, "fi", "") 4970b57cec5SDimitry AndricBUILTIN(__nvvm_i2f_rp, "fi", "") 4980b57cec5SDimitry Andric 4990b57cec5SDimitry AndricBUILTIN(__nvvm_ui2f_rn, "fUi", "") 5000b57cec5SDimitry AndricBUILTIN(__nvvm_ui2f_rz, "fUi", "") 5010b57cec5SDimitry AndricBUILTIN(__nvvm_ui2f_rm, "fUi", "") 5020b57cec5SDimitry AndricBUILTIN(__nvvm_ui2f_rp, "fUi", "") 5030b57cec5SDimitry Andric 5040b57cec5SDimitry AndricBUILTIN(__nvvm_lohi_i2d, "dii", "") 5050b57cec5SDimitry Andric 5060b57cec5SDimitry AndricBUILTIN(__nvvm_d2i_lo, "id", "") 5070b57cec5SDimitry AndricBUILTIN(__nvvm_d2i_hi, "id", "") 5080b57cec5SDimitry Andric 5090b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rn_ftz, "LLif", "") 5100b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rn, "LLif", "") 5110b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rz_ftz, "LLif", "") 5120b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rz, "LLif", "") 5130b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rm_ftz, "LLif", "") 5140b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rm, "LLif", "") 5150b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rp_ftz, "LLif", "") 5160b57cec5SDimitry AndricBUILTIN(__nvvm_f2ll_rp, "LLif", "") 5170b57cec5SDimitry Andric 5180b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rn_ftz, "ULLif", "") 5190b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rn, "ULLif", "") 5200b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rz_ftz, "ULLif", "") 5210b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rz, "ULLif", "") 5220b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rm_ftz, "ULLif", "") 5230b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rm, "ULLif", "") 5240b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rp_ftz, "ULLif", "") 5250b57cec5SDimitry AndricBUILTIN(__nvvm_f2ull_rp, "ULLif", "") 5260b57cec5SDimitry Andric 5270b57cec5SDimitry AndricBUILTIN(__nvvm_d2ll_rn, "LLid", "") 5280b57cec5SDimitry AndricBUILTIN(__nvvm_d2ll_rz, "LLid", "") 5290b57cec5SDimitry AndricBUILTIN(__nvvm_d2ll_rm, "LLid", "") 5300b57cec5SDimitry AndricBUILTIN(__nvvm_d2ll_rp, "LLid", "") 5310b57cec5SDimitry Andric 5320b57cec5SDimitry AndricBUILTIN(__nvvm_d2ull_rn, "ULLid", "") 5330b57cec5SDimitry AndricBUILTIN(__nvvm_d2ull_rz, "ULLid", "") 5340b57cec5SDimitry AndricBUILTIN(__nvvm_d2ull_rm, "ULLid", "") 5350b57cec5SDimitry AndricBUILTIN(__nvvm_d2ull_rp, "ULLid", "") 5360b57cec5SDimitry Andric 5370b57cec5SDimitry AndricBUILTIN(__nvvm_ll2f_rn, "fLLi", "") 5380b57cec5SDimitry AndricBUILTIN(__nvvm_ll2f_rz, "fLLi", "") 5390b57cec5SDimitry AndricBUILTIN(__nvvm_ll2f_rm, "fLLi", "") 5400b57cec5SDimitry AndricBUILTIN(__nvvm_ll2f_rp, "fLLi", "") 5410b57cec5SDimitry Andric 5420b57cec5SDimitry AndricBUILTIN(__nvvm_ull2f_rn, "fULLi", "") 5430b57cec5SDimitry AndricBUILTIN(__nvvm_ull2f_rz, "fULLi", "") 5440b57cec5SDimitry AndricBUILTIN(__nvvm_ull2f_rm, "fULLi", "") 5450b57cec5SDimitry AndricBUILTIN(__nvvm_ull2f_rp, "fULLi", "") 5460b57cec5SDimitry Andric 5470b57cec5SDimitry AndricBUILTIN(__nvvm_ll2d_rn, "dLLi", "") 5480b57cec5SDimitry AndricBUILTIN(__nvvm_ll2d_rz, "dLLi", "") 5490b57cec5SDimitry AndricBUILTIN(__nvvm_ll2d_rm, "dLLi", "") 5500b57cec5SDimitry AndricBUILTIN(__nvvm_ll2d_rp, "dLLi", "") 5510b57cec5SDimitry Andric 5520b57cec5SDimitry AndricBUILTIN(__nvvm_ull2d_rn, "dULLi", "") 5530b57cec5SDimitry AndricBUILTIN(__nvvm_ull2d_rz, "dULLi", "") 5540b57cec5SDimitry AndricBUILTIN(__nvvm_ull2d_rm, "dULLi", "") 5550b57cec5SDimitry AndricBUILTIN(__nvvm_ull2d_rp, "dULLi", "") 5560b57cec5SDimitry Andric 5570b57cec5SDimitry AndricBUILTIN(__nvvm_f2h_rn_ftz, "Usf", "") 5580b57cec5SDimitry AndricBUILTIN(__nvvm_f2h_rn, "Usf", "") 5590b57cec5SDimitry Andric 56006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_ff2bf16x2_rn, "V2yff", "", AND(SM_80,PTX70)) 56106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_ff2bf16x2_rn_relu, "V2yff", "", AND(SM_80,PTX70)) 56206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_ff2bf16x2_rz, "V2yff", "", AND(SM_80,PTX70)) 56306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_ff2bf16x2_rz_relu, "V2yff", "", AND(SM_80,PTX70)) 56404eeddc0SDimitry Andric 56504eeddc0SDimitry AndricTARGET_BUILTIN(__nvvm_ff2f16x2_rn, "V2hff", "", AND(SM_80,PTX70)) 56604eeddc0SDimitry AndricTARGET_BUILTIN(__nvvm_ff2f16x2_rn_relu, "V2hff", "", AND(SM_80,PTX70)) 56704eeddc0SDimitry AndricTARGET_BUILTIN(__nvvm_ff2f16x2_rz, "V2hff", "", AND(SM_80,PTX70)) 56804eeddc0SDimitry AndricTARGET_BUILTIN(__nvvm_ff2f16x2_rz_relu, "V2hff", "", AND(SM_80,PTX70)) 56904eeddc0SDimitry Andric 57006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_f2bf16_rn, "yf", "", AND(SM_80,PTX70)) 57106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_f2bf16_rn_relu, "yf", "", AND(SM_80,PTX70)) 57206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_f2bf16_rz, "yf", "", AND(SM_80,PTX70)) 57306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_f2bf16_rz_relu, "yf", "", AND(SM_80,PTX70)) 57404eeddc0SDimitry Andric 57504eeddc0SDimitry AndricTARGET_BUILTIN(__nvvm_f2tf32_rna, "ZUif", "", AND(SM_80,PTX70)) 57604eeddc0SDimitry Andric 5770b57cec5SDimitry Andric// Bitcast 5780b57cec5SDimitry Andric 5790b57cec5SDimitry AndricBUILTIN(__nvvm_bitcast_f2i, "if", "") 5800b57cec5SDimitry AndricBUILTIN(__nvvm_bitcast_i2f, "fi", "") 5810b57cec5SDimitry Andric 5820b57cec5SDimitry AndricBUILTIN(__nvvm_bitcast_ll2d, "dLLi", "") 5830b57cec5SDimitry AndricBUILTIN(__nvvm_bitcast_d2ll, "LLid", "") 5840b57cec5SDimitry Andric 5850b57cec5SDimitry Andric// FNS 5860b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_fns, "UiUiUii", "n", PTX60) 5870b57cec5SDimitry Andric 5880b57cec5SDimitry Andric// Sync 5890b57cec5SDimitry Andric 5900b57cec5SDimitry AndricBUILTIN(__syncthreads, "v", "") 5910b57cec5SDimitry AndricBUILTIN(__nvvm_bar0_popc, "ii", "") 5920b57cec5SDimitry AndricBUILTIN(__nvvm_bar0_and, "ii", "") 5930b57cec5SDimitry AndricBUILTIN(__nvvm_bar0_or, "ii", "") 5940b57cec5SDimitry AndricBUILTIN(__nvvm_bar_sync, "vi", "n") 5950b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi", "n", PTX60) 5960b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_barrier_sync, "vUi", "n", PTX60) 5970b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi", "n", PTX60) 5980b57cec5SDimitry Andric 59906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_barrier_cluster_arrive, "v", "n", AND(SM_90,PTX78)) 60006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_barrier_cluster_arrive_relaxed, "v", "n", AND(SM_90,PTX80)) 60106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_barrier_cluster_wait, "v", "n", AND(SM_90,PTX78)) 60206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_fence_sc_cluster, "v", "n", AND(SM_90,PTX78)) 60306c3fb27SDimitry Andric 6040b57cec5SDimitry Andric// Shuffle 6050b57cec5SDimitry Andric 6060b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_down_i32, "iiii", "") 6070b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_down_f32, "ffii", "") 6080b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_up_i32, "iiii", "") 6090b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_up_f32, "ffii", "") 6100b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_bfly_i32, "iiii", "") 6110b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_bfly_f32, "ffii", "") 6120b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_idx_i32, "iiii", "") 6130b57cec5SDimitry AndricBUILTIN(__nvvm_shfl_idx_f32, "ffii", "") 6140b57cec5SDimitry Andric 6150b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_down_i32, "iUiiii", "", PTX60) 6160b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_down_f32, "fUifii", "", PTX60) 6170b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_up_i32, "iUiiii", "", PTX60) 6180b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_up_f32, "fUifii", "", PTX60) 6190b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_bfly_i32, "iUiiii", "", PTX60) 6200b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_bfly_f32, "fUifii", "", PTX60) 6210b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_idx_i32, "iUiiii", "", PTX60) 6220b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_shfl_sync_idx_f32, "fUifii", "", PTX60) 6230b57cec5SDimitry Andric 6240b57cec5SDimitry Andric// Vote 6250b57cec5SDimitry AndricBUILTIN(__nvvm_vote_all, "bb", "") 6260b57cec5SDimitry AndricBUILTIN(__nvvm_vote_any, "bb", "") 6270b57cec5SDimitry AndricBUILTIN(__nvvm_vote_uni, "bb", "") 6280b57cec5SDimitry AndricBUILTIN(__nvvm_vote_ballot, "Uib", "") 6290b57cec5SDimitry Andric 6300b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_vote_all_sync, "bUib", "", PTX60) 6310b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_vote_any_sync, "bUib", "", PTX60) 6320b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib", "", PTX60) 6330b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib", "", PTX60) 6340b57cec5SDimitry Andric 6350b57cec5SDimitry Andric// Match 63681ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi", "", AND(SM_70,PTX60)) 63781ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_match_any_sync_i64, "UiUiWi", "", AND(SM_70,PTX60)) 6380b57cec5SDimitry Andric// These return a pair {value, predicate}, which requires custom lowering. 63981ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*", "", AND(SM_70,PTX60)) 64081ad6265SDimitry AndricTARGET_BUILTIN(__nvvm_match_all_sync_i64p, "UiUiWii*", "", AND(SM_70,PTX60)) 6410b57cec5SDimitry Andric 642fe6060f1SDimitry Andric// Redux 643fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_add, "iii", "", AND(SM_80,PTX70)) 644fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_min, "iii", "", AND(SM_80,PTX70)) 645fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_max, "iii", "", AND(SM_80,PTX70)) 646fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_umin, "UiUii", "", AND(SM_80,PTX70)) 647fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_umax, "UiUii", "", AND(SM_80,PTX70)) 648fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_and, "iii", "", AND(SM_80,PTX70)) 649fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_xor, "iii", "", AND(SM_80,PTX70)) 650fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_redux_sync_or, "iii", "", AND(SM_80,PTX70)) 651fe6060f1SDimitry Andric 6520b57cec5SDimitry Andric// Membar 6530b57cec5SDimitry Andric 6540b57cec5SDimitry AndricBUILTIN(__nvvm_membar_cta, "v", "") 6550b57cec5SDimitry AndricBUILTIN(__nvvm_membar_gl, "v", "") 6560b57cec5SDimitry AndricBUILTIN(__nvvm_membar_sys, "v", "") 6570b57cec5SDimitry Andric 658fe6060f1SDimitry Andric// mbarrier 659fe6060f1SDimitry Andric 660fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_init, "vWi*i", "", AND(SM_80,PTX70)) 661fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_init_shared, "vWi*3i", "", AND(SM_80,PTX70)) 662fe6060f1SDimitry Andric 663fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_inval, "vWi*", "", AND(SM_80,PTX70)) 664fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_inval_shared, "vWi*3", "", AND(SM_80,PTX70)) 665fe6060f1SDimitry Andric 666fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive, "WiWi*", "", AND(SM_80,PTX70)) 667fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_shared, "WiWi*3", "", AND(SM_80,PTX70)) 668fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_noComplete, "WiWi*i", "", AND(SM_80,PTX70)) 669fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_noComplete_shared, "WiWi*3i", "", AND(SM_80,PTX70)) 670fe6060f1SDimitry Andric 671fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_drop, "WiWi*", "", AND(SM_80,PTX70)) 672fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_shared, "WiWi*3", "", AND(SM_80,PTX70)) 673fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_noComplete, "WiWi*i", "", AND(SM_80,PTX70)) 674fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_noComplete_shared, "WiWi*3i", "", AND(SM_80,PTX70)) 675fe6060f1SDimitry Andric 676fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_test_wait, "bWi*Wi", "", AND(SM_80,PTX70)) 677fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_test_wait_shared, "bWi*3Wi", "", AND(SM_80,PTX70)) 678fe6060f1SDimitry Andric 679fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_mbarrier_pending_count, "iWi", "", AND(SM_80,PTX70)) 680fe6060f1SDimitry Andric 6810b57cec5SDimitry Andric// Memcpy, Memset 6820b57cec5SDimitry Andric 6830b57cec5SDimitry AndricBUILTIN(__nvvm_memcpy, "vUc*Uc*zi","") 6840b57cec5SDimitry AndricBUILTIN(__nvvm_memset, "vUc*Uczi","") 6850b57cec5SDimitry Andric 6860b57cec5SDimitry Andric// Image 6870b57cec5SDimitry Andric 6880b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image2Dfi_, "V4fiiii", "") 6890b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image2Dff_, "V4fiiff", "") 6900b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image2Dii_, "V4iiiii", "") 6910b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image2Dif_, "V4iiiff", "") 6920b57cec5SDimitry Andric 6930b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image3Dfi_, "V4fiiiiii", "") 6940b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image3Dff_, "V4fiiffff", "") 6950b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image3Dii_, "V4iiiiiii", "") 6960b57cec5SDimitry AndricBUILTIN(__builtin_ptx_read_image3Dif_, "V4iiiffff", "") 6970b57cec5SDimitry Andric 6980b57cec5SDimitry AndricBUILTIN(__builtin_ptx_write_image2Df_, "viiiffff", "") 6990b57cec5SDimitry AndricBUILTIN(__builtin_ptx_write_image2Di_, "viiiiiii", "") 7000b57cec5SDimitry AndricBUILTIN(__builtin_ptx_write_image2Dui_, "viiiUiUiUiUi", "") 7010b57cec5SDimitry AndricBUILTIN(__builtin_ptx_get_image_depthi_, "ii", "") 7020b57cec5SDimitry AndricBUILTIN(__builtin_ptx_get_image_heighti_, "ii", "") 7030b57cec5SDimitry AndricBUILTIN(__builtin_ptx_get_image_widthi_, "ii", "") 7040b57cec5SDimitry AndricBUILTIN(__builtin_ptx_get_image_channel_data_typei_, "ii", "") 7050b57cec5SDimitry AndricBUILTIN(__builtin_ptx_get_image_channel_orderi_, "ii", "") 7060b57cec5SDimitry Andric 7070b57cec5SDimitry Andric// Atomic 7080b57cec5SDimitry Andric// 7090b57cec5SDimitry Andric// We need the atom intrinsics because 7100b57cec5SDimitry Andric// - they are used in converging analysis 7110b57cec5SDimitry Andric// - they are used in address space analysis and optimization 7120b57cec5SDimitry Andric// So it does not hurt to expose them as builtins. 7130b57cec5SDimitry Andric// 7140b57cec5SDimitry AndricBUILTIN(__nvvm_atom_add_gen_i, "iiD*i", "n") 7150b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_add_gen_i, "iiD*i", "n", SM_60) 7160b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_add_gen_i, "iiD*i", "n", SM_60) 7170b57cec5SDimitry AndricBUILTIN(__nvvm_atom_add_gen_l, "LiLiD*Li", "n") 7180b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_add_gen_l, "LiLiD*Li", "n", SM_60) 7190b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_add_gen_l, "LiLiD*Li", "n", SM_60) 7200b57cec5SDimitry AndricBUILTIN(__nvvm_atom_add_gen_ll, "LLiLLiD*LLi", "n") 7210b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_add_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7220b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_add_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7230b57cec5SDimitry AndricBUILTIN(__nvvm_atom_add_gen_f, "ffD*f", "n") 7240b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_add_gen_f, "ffD*f", "n", SM_60) 7250b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_add_gen_f, "ffD*f", "n", SM_60) 7260b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_add_gen_d, "ddD*d", "n", SM_60) 7270b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d", "n", SM_60) 7280b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d", "n", SM_60) 7290b57cec5SDimitry Andric 7300b57cec5SDimitry AndricBUILTIN(__nvvm_atom_sub_gen_i, "iiD*i", "n") 7310b57cec5SDimitry AndricBUILTIN(__nvvm_atom_sub_gen_l, "LiLiD*Li", "n") 7320b57cec5SDimitry AndricBUILTIN(__nvvm_atom_sub_gen_ll, "LLiLLiD*LLi", "n") 7330b57cec5SDimitry Andric 7340b57cec5SDimitry AndricBUILTIN(__nvvm_atom_xchg_gen_i, "iiD*i", "n") 7350b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_i, "iiD*i", "n", SM_60) 7360b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_i, "iiD*i", "n", SM_60) 7370b57cec5SDimitry AndricBUILTIN(__nvvm_atom_xchg_gen_l, "LiLiD*Li", "n") 7380b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_l, "LiLiD*Li", "n", SM_60) 7390b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li", "n", SM_60) 7400b57cec5SDimitry AndricBUILTIN(__nvvm_atom_xchg_gen_ll, "LLiLLiD*LLi", "n") 7410b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7420b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7430b57cec5SDimitry Andric 7440b57cec5SDimitry AndricBUILTIN(__nvvm_atom_max_gen_i, "iiD*i", "n") 7450b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i", "n", SM_60) 7460b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_max_gen_i, "iiD*i", "n", SM_60) 7470b57cec5SDimitry AndricBUILTIN(__nvvm_atom_max_gen_ui, "UiUiD*Ui", "n") 7480b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_max_gen_ui, "UiUiD*Ui", "n", SM_60) 7490b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_max_gen_ui, "UiUiD*Ui", "n", SM_60) 7500b57cec5SDimitry AndricBUILTIN(__nvvm_atom_max_gen_l, "LiLiD*Li", "n") 7510b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_max_gen_l, "LiLiD*Li", "n", SM_60) 7520b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_max_gen_l, "LiLiD*Li", "n", SM_60) 7530b57cec5SDimitry AndricBUILTIN(__nvvm_atom_max_gen_ul, "ULiULiD*ULi", "n") 7540b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_max_gen_ul, "ULiULiD*ULi", "n", SM_60) 7550b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_max_gen_ul, "ULiULiD*ULi", "n", SM_60) 7560b57cec5SDimitry AndricBUILTIN(__nvvm_atom_max_gen_ll, "LLiLLiD*LLi", "n") 7570b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_max_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7580b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_max_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7590b57cec5SDimitry AndricBUILTIN(__nvvm_atom_max_gen_ull, "ULLiULLiD*ULLi", "n") 7600b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_max_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) 7610b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_max_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) 7620b57cec5SDimitry Andric 7630b57cec5SDimitry AndricBUILTIN(__nvvm_atom_min_gen_i, "iiD*i", "n") 7640b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_min_gen_i, "iiD*i", "n", SM_60) 7650b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_min_gen_i, "iiD*i", "n", SM_60) 7660b57cec5SDimitry AndricBUILTIN(__nvvm_atom_min_gen_ui, "UiUiD*Ui", "n") 7670b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_min_gen_ui, "UiUiD*Ui", "n", SM_60) 7680b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_min_gen_ui, "UiUiD*Ui", "n", SM_60) 7690b57cec5SDimitry AndricBUILTIN(__nvvm_atom_min_gen_l, "LiLiD*Li", "n") 7700b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_min_gen_l, "LiLiD*Li", "n", SM_60) 7710b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_min_gen_l, "LiLiD*Li", "n", SM_60) 7720b57cec5SDimitry AndricBUILTIN(__nvvm_atom_min_gen_ul, "ULiULiD*ULi", "n") 7730b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_min_gen_ul, "ULiULiD*ULi", "n", SM_60) 7740b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_min_gen_ul, "ULiULiD*ULi", "n", SM_60) 7750b57cec5SDimitry AndricBUILTIN(__nvvm_atom_min_gen_ll, "LLiLLiD*LLi", "n") 7760b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_min_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7770b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_min_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7780b57cec5SDimitry AndricBUILTIN(__nvvm_atom_min_gen_ull, "ULLiULLiD*ULLi", "n") 7790b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_min_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) 7800b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_min_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) 7810b57cec5SDimitry Andric 7820b57cec5SDimitry AndricBUILTIN(__nvvm_atom_inc_gen_ui, "UiUiD*Ui", "n") 7830b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_inc_gen_ui, "UiUiD*Ui", "n", SM_60) 7840b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_inc_gen_ui, "UiUiD*Ui", "n", SM_60) 7850b57cec5SDimitry AndricBUILTIN(__nvvm_atom_dec_gen_ui, "UiUiD*Ui", "n") 7860b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_dec_gen_ui, "UiUiD*Ui", "n", SM_60) 7870b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_dec_gen_ui, "UiUiD*Ui", "n", SM_60) 7880b57cec5SDimitry Andric 7890b57cec5SDimitry AndricBUILTIN(__nvvm_atom_and_gen_i, "iiD*i", "n") 7900b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_and_gen_i, "iiD*i", "n", SM_60) 7910b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_and_gen_i, "iiD*i", "n", SM_60) 7920b57cec5SDimitry AndricBUILTIN(__nvvm_atom_and_gen_l, "LiLiD*Li", "n") 7930b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_and_gen_l, "LiLiD*Li", "n", SM_60) 7940b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_and_gen_l, "LiLiD*Li", "n", SM_60) 7950b57cec5SDimitry AndricBUILTIN(__nvvm_atom_and_gen_ll, "LLiLLiD*LLi", "n") 7960b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_and_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7970b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_and_gen_ll, "LLiLLiD*LLi", "n", SM_60) 7980b57cec5SDimitry Andric 7990b57cec5SDimitry AndricBUILTIN(__nvvm_atom_or_gen_i, "iiD*i", "n") 8000b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_or_gen_i, "iiD*i", "n", SM_60) 8010b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_or_gen_i, "iiD*i", "n", SM_60) 8020b57cec5SDimitry AndricBUILTIN(__nvvm_atom_or_gen_l, "LiLiD*Li", "n") 8030b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_or_gen_l, "LiLiD*Li", "n", SM_60) 8040b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_or_gen_l, "LiLiD*Li", "n", SM_60) 8050b57cec5SDimitry AndricBUILTIN(__nvvm_atom_or_gen_ll, "LLiLLiD*LLi", "n") 8060b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_or_gen_ll, "LLiLLiD*LLi", "n", SM_60) 8070b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_or_gen_ll, "LLiLLiD*LLi", "n", SM_60) 8080b57cec5SDimitry Andric 8090b57cec5SDimitry AndricBUILTIN(__nvvm_atom_xor_gen_i, "iiD*i", "n") 8100b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_xor_gen_i, "iiD*i", "n", SM_60) 8110b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_xor_gen_i, "iiD*i", "n", SM_60) 8120b57cec5SDimitry AndricBUILTIN(__nvvm_atom_xor_gen_l, "LiLiD*Li", "n") 8130b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_xor_gen_l, "LiLiD*Li", "n", SM_60) 8140b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_xor_gen_l, "LiLiD*Li", "n", SM_60) 8150b57cec5SDimitry AndricBUILTIN(__nvvm_atom_xor_gen_ll, "LLiLLiD*LLi", "n") 8160b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_xor_gen_ll, "LLiLLiD*LLi", "n", SM_60) 8170b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_xor_gen_ll, "LLiLLiD*LLi", "n", SM_60) 8180b57cec5SDimitry Andric 8190b57cec5SDimitry AndricBUILTIN(__nvvm_atom_cas_gen_i, "iiD*ii", "n") 8200b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_cas_gen_i, "iiD*ii", "n", SM_60) 8210b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_cas_gen_i, "iiD*ii", "n", SM_60) 8220b57cec5SDimitry AndricBUILTIN(__nvvm_atom_cas_gen_l, "LiLiD*LiLi", "n") 8230b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_cas_gen_l, "LiLiD*LiLi", "n", SM_60) 8240b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_60) 8250b57cec5SDimitry AndricBUILTIN(__nvvm_atom_cas_gen_ll, "LLiLLiD*LLiLLi", "n") 8260b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60) 8270b57cec5SDimitry AndricTARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60) 8280b57cec5SDimitry Andric 8290b57cec5SDimitry Andric// Compiler Error Warn 8300b57cec5SDimitry AndricBUILTIN(__nvvm_compiler_error, "vcC*4", "n") 8310b57cec5SDimitry AndricBUILTIN(__nvvm_compiler_warn, "vcC*4", "n") 8320b57cec5SDimitry Andric 83306c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_c, "ccC*", "") 83406c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_sc, "ScScC*", "") 83506c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_s, "ssC*", "") 83606c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_i, "iiC*", "") 83706c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_l, "LiLiC*", "") 83806c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ll, "LLiLLiC*", "") 83906c3fb27SDimitry Andric 84006c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_uc, "UcUcC*", "") 84106c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_us, "UsUsC*", "") 84206c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ui, "UiUiC*", "") 84306c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ul, "ULiULiC*", "") 84406c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ull, "ULLiULLiC*", "") 84506c3fb27SDimitry Andric 84606c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_h, "hhC*", "") 84706c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_f, "ffC*", "") 84806c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_d, "ddC*", "") 84906c3fb27SDimitry Andric 85006c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_c2, "E2cE2cC*", "") 85106c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_sc2, "E2ScE2ScC*", "") 85206c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_c4, "E4cE4cC*", "") 85306c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_sc4, "E4ScE4ScC*", "") 85406c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_s2, "E2sE2sC*", "") 85506c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_s4, "E4sE4sC*", "") 85606c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_i2, "E2iE2iC*", "") 85706c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_i4, "E4iE4iC*", "") 85806c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_l2, "E2LiE2LiC*", "") 85906c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ll2, "E2LLiE2LLiC*", "") 86006c3fb27SDimitry Andric 86106c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_uc2, "E2UcE2UcC*", "") 86206c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_uc4, "E4UcE4UcC*", "") 86306c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_us2, "E2UsE2UsC*", "") 86406c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_us4, "E4UsE4UsC*", "") 86506c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ui2, "E2UiE2UiC*", "") 86606c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ui4, "E4UiE4UiC*", "") 86706c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ul2, "E2ULiE2ULiC*", "") 86806c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_ull2, "E2ULLiE2ULLiC*", "") 86906c3fb27SDimitry Andric 87006c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_h2, "E2hE2hC*", "") 87106c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_f2, "E2fE2fC*", "") 87206c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_f4, "E4fE4fC*", "") 87306c3fb27SDimitry AndricBUILTIN(__nvvm_ldu_d2, "E2dE2dC*", "") 87406c3fb27SDimitry Andric 8750b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_c, "ccC*", "") 87606c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_sc, "ScScC*", "") 8770b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_s, "ssC*", "") 8780b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_i, "iiC*", "") 8790b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_l, "LiLiC*", "") 8800b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ll, "LLiLLiC*", "") 8810b57cec5SDimitry Andric 8820b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_uc, "UcUcC*", "") 8830b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_us, "UsUsC*", "") 8840b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ui, "UiUiC*", "") 8850b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ul, "ULiULiC*", "") 8860b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ull, "ULLiULLiC*", "") 8870b57cec5SDimitry Andric 88806c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_h, "hhC*", "") 8890b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_f, "ffC*", "") 8900b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_d, "ddC*", "") 8910b57cec5SDimitry Andric 8920b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_c2, "E2cE2cC*", "") 89306c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_sc2, "E2ScE2ScC*", "") 8940b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_c4, "E4cE4cC*", "") 89506c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_sc4, "E4ScE4ScC*", "") 8960b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_s2, "E2sE2sC*", "") 8970b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_s4, "E4sE4sC*", "") 8980b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_i2, "E2iE2iC*", "") 8990b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_i4, "E4iE4iC*", "") 90006c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_l2, "E2LiE2LiC*", "") 9010b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ll2, "E2LLiE2LLiC*", "") 9020b57cec5SDimitry Andric 9030b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_uc2, "E2UcE2UcC*", "") 9040b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_uc4, "E4UcE4UcC*", "") 9050b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_us2, "E2UsE2UsC*", "") 9060b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_us4, "E4UsE4UsC*", "") 9070b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ui2, "E2UiE2UiC*", "") 9080b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ui4, "E4UiE4UiC*", "") 90906c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_ul2, "E2ULiE2ULiC*", "") 9100b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_ull2, "E2ULLiE2ULLiC*", "") 9110b57cec5SDimitry Andric 91206c3fb27SDimitry AndricBUILTIN(__nvvm_ldg_h2, "E2hE2hC*", "") 9130b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_f2, "E2fE2fC*", "") 9140b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_f4, "E4fE4fC*", "") 9150b57cec5SDimitry AndricBUILTIN(__nvvm_ldg_d2, "E2dE2dC*", "") 9160b57cec5SDimitry Andric 917349cc55cSDimitry Andric// Address space predicates. 918349cc55cSDimitry AndricBUILTIN(__nvvm_isspacep_const, "bvC*", "nc") 919349cc55cSDimitry AndricBUILTIN(__nvvm_isspacep_global, "bvC*", "nc") 920349cc55cSDimitry AndricBUILTIN(__nvvm_isspacep_local, "bvC*", "nc") 921349cc55cSDimitry AndricBUILTIN(__nvvm_isspacep_shared, "bvC*", "nc") 92206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_isspacep_shared_cluster,"bvC*", "nc", AND(SM_90,PTX78)) 923349cc55cSDimitry Andric 9240b57cec5SDimitry Andric// Builtins to support WMMA instructions on sm_70 9250b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_ld_a, "vi*iC*UiIi", "", AND(SM_70,PTX60)) 9260b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_ld_b, "vi*iC*UiIi", "", AND(SM_70,PTX60)) 9270b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_ld_c_f16, "vi*iC*UiIi", "", AND(SM_70,PTX60)) 9280b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_ld_c_f32, "vf*fC*UiIi", "", AND(SM_70,PTX60)) 9295f757f3fSDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_st_c_f16, "vi*iC*UiIi", "", AND(SM_70,PTX60)) 9305f757f3fSDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_st_c_f32, "vf*fC*UiIi", "", AND(SM_70,PTX60)) 9310b57cec5SDimitry Andric 9320b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_ld_a, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9330b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_ld_b, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9340b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_ld_c_f16, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9350b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_ld_c_f32, "vf*fC*UiIi", "", AND(SM_70,PTX61)) 9365f757f3fSDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_st_c_f16, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9375f757f3fSDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_st_c_f32, "vf*fC*UiIi", "", AND(SM_70,PTX61)) 9380b57cec5SDimitry Andric 9390b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_ld_a, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9400b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_ld_b, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9410b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_ld_c_f16, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9420b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_ld_c_f32, "vf*fC*UiIi", "", AND(SM_70,PTX61)) 9435f757f3fSDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_st_c_f16, "vi*iC*UiIi", "", AND(SM_70,PTX61)) 9445f757f3fSDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_st_c_f32, "vf*fC*UiIi", "", AND(SM_70,PTX61)) 9450b57cec5SDimitry Andric 9460b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", AND(SM_70,PTX60)) 9470b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", AND(SM_70,PTX60)) 9480b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", AND(SM_70,PTX60)) 9490b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m16n16k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", AND(SM_70,PTX60)) 9500b57cec5SDimitry Andric 9510b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", AND(SM_70,PTX61)) 9520b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", AND(SM_70,PTX61)) 9530b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", AND(SM_70,PTX61)) 9540b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m32n8k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", AND(SM_70,PTX61)) 9550b57cec5SDimitry Andric 9560b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", AND(SM_70,PTX61)) 9570b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", AND(SM_70,PTX61)) 9580b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", AND(SM_70,PTX61)) 9590b57cec5SDimitry AndricTARGET_BUILTIN(__hmma_m8n32k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", AND(SM_70,PTX61)) 9600b57cec5SDimitry Andric 9610b57cec5SDimitry Andric// Builtins to support integer and sub-integer WMMA instructions on sm_72/sm_75 9620b57cec5SDimitry AndricTARGET_BUILTIN(__bmma_m8n8k128_ld_a_b1, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9630b57cec5SDimitry AndricTARGET_BUILTIN(__bmma_m8n8k128_ld_b_b1, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9640b57cec5SDimitry AndricTARGET_BUILTIN(__bmma_m8n8k128_ld_c, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 965bdd1243dSDimitry AndricTARGET_BUILTIN(__bmma_m8n8k128_mma_and_popc_b1, "vi*iC*iC*iC*Ii", "", AND(SM_80,PTX71)) 9660b57cec5SDimitry AndricTARGET_BUILTIN(__bmma_m8n8k128_mma_xor_popc_b1, "vi*iC*iC*iC*Ii", "", AND(SM_75,PTX63)) 9670b57cec5SDimitry AndricTARGET_BUILTIN(__bmma_m8n8k128_st_c_i32, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9680b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_ld_a_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9690b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_ld_a_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9700b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_ld_b_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9710b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_ld_b_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9720b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_ld_c, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9730b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_mma_s8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63)) 9740b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_mma_u8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63)) 9750b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m16n16k16_st_c_i32, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9760b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_ld_a_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9770b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_ld_a_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9780b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_ld_b_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9790b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_ld_b_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9800b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_ld_c, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9810b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_mma_s8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63)) 9820b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_mma_u8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63)) 9830b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m32n8k16_st_c_i32, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9840b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_ld_a_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9850b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_ld_a_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9860b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_ld_b_s8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9870b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_ld_b_u8, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9880b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_ld_c, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9890b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_mma_s8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63)) 9900b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_mma_u8, "vi*iC*iC*iC*IiIi", "", AND(SM_72,PTX63)) 9910b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n32k16_st_c_i32, "vi*iC*UiIi", "", AND(SM_72,PTX63)) 9920b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_ld_a_s4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9930b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_ld_a_u4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9940b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_ld_b_s4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9950b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_ld_b_u4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9960b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_ld_c, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 9970b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_mma_s4, "vi*iC*iC*iC*IiIi", "", AND(SM_75,PTX63)) 9980b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_mma_u4, "vi*iC*iC*iC*IiIi", "", AND(SM_75,PTX63)) 9990b57cec5SDimitry AndricTARGET_BUILTIN(__imma_m8n8k32_st_c_i32, "vi*iC*UiIi", "", AND(SM_75,PTX63)) 10000b57cec5SDimitry Andric 1001fe6060f1SDimitry Andric// Builtins to support double and alternate float WMMA instructions on sm_80 1002fe6060f1SDimitry AndricTARGET_BUILTIN(__dmma_m8n8k4_ld_a, "vd*dC*UiIi", "", AND(SM_80,PTX70)) 1003fe6060f1SDimitry AndricTARGET_BUILTIN(__dmma_m8n8k4_ld_b, "vd*dC*UiIi", "", AND(SM_80,PTX70)) 1004fe6060f1SDimitry AndricTARGET_BUILTIN(__dmma_m8n8k4_ld_c, "vd*dC*UiIi", "", AND(SM_80,PTX70)) 1005fe6060f1SDimitry AndricTARGET_BUILTIN(__dmma_m8n8k4_st_c_f64, "vd*dC*UiIi", "", AND(SM_80,PTX70)) 1006fe6060f1SDimitry AndricTARGET_BUILTIN(__dmma_m8n8k4_mma_f64, "vd*dC*dC*dC*IiIi", "", AND(SM_80,PTX70)) 1007fe6060f1SDimitry Andric 1008fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m16n16k16_ld_a, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1009fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m16n16k16_ld_b, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1010fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m16n16k16_mma_f32, "vf*iC*iC*fC*IiIi", "", AND(SM_80,PTX70)) 1011fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m8n32k16_ld_a, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1012fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m8n32k16_ld_b, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1013fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m8n32k16_mma_f32, "vf*iC*iC*fC*IiIi", "", AND(SM_80,PTX70)) 1014fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m32n8k16_ld_a, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1015fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m32n8k16_ld_b, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1016fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_bf16_m32n8k16_mma_f32, "vf*iC*iC*fC*IiIi", "", AND(SM_80,PTX70)) 1017fe6060f1SDimitry Andric 1018fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_tf32_m16n16k8_ld_a, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1019fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_tf32_m16n16k8_ld_b, "vi*iC*UiIi", "", AND(SM_80,PTX70)) 1020fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_tf32_m16n16k8_ld_c, "vf*fC*UiIi", "", AND(SM_80,PTX70)) 1021fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_m16n16k8_st_c_f32, "vf*fC*UiIi", "", AND(SM_80,PTX70)) 1022fe6060f1SDimitry AndricTARGET_BUILTIN(__mma_tf32_m16n16k8_mma_f32, "vf*iC*iC*fC*IiIi", "", AND(SM_80,PTX70)) 1023fe6060f1SDimitry Andric 1024fe6060f1SDimitry Andric// Async Copy 1025fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive, "vWi*", "", AND(SM_80,PTX70)) 1026fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_shared, "vWi*3", "", AND(SM_80,PTX70)) 1027fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc, "vWi*", "", AND(SM_80,PTX70)) 1028fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc_shared, "vWi*3", "", AND(SM_80,PTX70)) 1029fe6060f1SDimitry Andric 103006c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_4, "vv*3vC*1.", "", AND(SM_80,PTX70)) 103106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_8, "vv*3vC*1.", "", AND(SM_80,PTX70)) 103206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1.", "", AND(SM_80,PTX70)) 103306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_cg_shared_global_16, "vv*3vC*1.", "", AND(SM_80,PTX70)) 1034fe6060f1SDimitry Andric 1035fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_commit_group, "v", "", AND(SM_80,PTX70)) 1036fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_wait_group, "vIi", "", AND(SM_80,PTX70)) 1037fe6060f1SDimitry AndricTARGET_BUILTIN(__nvvm_cp_async_wait_all, "v", "", AND(SM_80,PTX70)) 1038fe6060f1SDimitry Andric 103981ad6265SDimitry Andric 104081ad6265SDimitry Andric// bf16, bf16x2 abs, neg 104106c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_abs_bf16, "yy", "", AND(SM_80,PTX70)) 104206c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_abs_bf16x2, "V2yV2y", "", AND(SM_80,PTX70)) 104306c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_neg_bf16, "yy", "", AND(SM_80,PTX70)) 104406c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_neg_bf16x2, "V2yV2y", "", AND(SM_80,PTX70)) 104506c3fb27SDimitry Andric 104606c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_mapa, "v*v*i", "", AND(SM_90, PTX78)) 104706c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_mapa_shared_cluster, "v*3v*3i", "", AND(SM_90, PTX78)) 104806c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_getctarank, "iv*", "", AND(SM_90, PTX78)) 104906c3fb27SDimitry AndricTARGET_BUILTIN(__nvvm_getctarank_shared_cluster, "iv*3", "", AND(SM_90,PTX78)) 105081ad6265SDimitry Andric 10510b57cec5SDimitry Andric#undef BUILTIN 10520b57cec5SDimitry Andric#undef TARGET_BUILTIN 10530b57cec5SDimitry Andric#pragma pop_macro("AND") 105481ad6265SDimitry Andric#pragma pop_macro("SM_53") 10550b57cec5SDimitry Andric#pragma pop_macro("SM_60") 10560b57cec5SDimitry Andric#pragma pop_macro("SM_70") 10570b57cec5SDimitry Andric#pragma pop_macro("SM_72") 10580b57cec5SDimitry Andric#pragma pop_macro("SM_75") 10595ffd83dbSDimitry Andric#pragma pop_macro("SM_80") 1060fe6060f1SDimitry Andric#pragma pop_macro("SM_86") 1061bdd1243dSDimitry Andric#pragma pop_macro("SM_87") 1062bdd1243dSDimitry Andric#pragma pop_macro("SM_89") 1063bdd1243dSDimitry Andric#pragma pop_macro("SM_90") 10645f757f3fSDimitry Andric#pragma pop_macro("SM_90a") 106581ad6265SDimitry Andric#pragma pop_macro("PTX42") 10660b57cec5SDimitry Andric#pragma pop_macro("PTX60") 10670b57cec5SDimitry Andric#pragma pop_macro("PTX61") 10680b57cec5SDimitry Andric#pragma pop_macro("PTX63") 10690b57cec5SDimitry Andric#pragma pop_macro("PTX64") 10705ffd83dbSDimitry Andric#pragma pop_macro("PTX65") 10715ffd83dbSDimitry Andric#pragma pop_macro("PTX70") 1072fe6060f1SDimitry Andric#pragma pop_macro("PTX71") 1073fe6060f1SDimitry Andric#pragma pop_macro("PTX72") 1074349cc55cSDimitry Andric#pragma pop_macro("PTX73") 1075349cc55cSDimitry Andric#pragma pop_macro("PTX74") 1076349cc55cSDimitry Andric#pragma pop_macro("PTX75") 1077bdd1243dSDimitry Andric#pragma pop_macro("PTX76") 1078bdd1243dSDimitry Andric#pragma pop_macro("PTX77") 1079bdd1243dSDimitry Andric#pragma pop_macro("PTX78") 108006c3fb27SDimitry Andric#pragma pop_macro("PTX80") 108106c3fb27SDimitry Andric#pragma pop_macro("PTX81") 10825f757f3fSDimitry Andric#pragma pop_macro("PTX82") 10835f757f3fSDimitry Andric#pragma pop_macro("PTX83") 1084