From: Artem Belevich Date: Wed, 11 Apr 2018 17:51:19 +0000 (+0000) Subject: [NVPTX, CUDA] Improved feature constraints on NVPTX target builtins. X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=24e8a680e545dbdef3483bbf4ab47a05b14ab529;p=platform%2Fupstream%2Fllvm.git [NVPTX, CUDA] Improved feature constraints on NVPTX target builtins. When NVPTX TARGET_BUILTIN specifies sm_XX or ptxYY as required feature, consider those features available if we're compiling for GPU >= sm_XX or have enabled PTX version >= ptxYY. Differential Revision: https://reviews.llvm.org/D45061 llvm-svn: 329829 --- diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def index 7bab73a..ae33315 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.def +++ b/clang/include/clang/Basic/BuiltinsNVPTX.def @@ -18,6 +18,12 @@ # define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS) #endif +#pragma push_macro("SM_60") +#define SM_60 "sm_60|sm_61|sm_62|sm_70|sm_71" + +#pragma push_macro("PTX60") +#define PTX60 "ptx60|ptx61" + // Special Registers BUILTIN(__nvvm_read_ptx_sreg_tid_x, "i", "nc") @@ -372,7 +378,7 @@ BUILTIN(__nvvm_bitcast_ll2d, "dLLi", "") BUILTIN(__nvvm_bitcast_d2ll, "LLid", "") // FNS -TARGET_BUILTIN(__nvvm_fns, "UiUiUii", "n", "ptx60") +TARGET_BUILTIN(__nvvm_fns, "UiUiUii", "n", PTX60) // Sync @@ -381,9 +387,9 @@ BUILTIN(__nvvm_bar0_popc, "ii", "") BUILTIN(__nvvm_bar0_and, "ii", "") BUILTIN(__nvvm_bar0_or, "ii", "") BUILTIN(__nvvm_bar_sync, "vi", "n") -TARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi", "n", "ptx60") -TARGET_BUILTIN(__nvvm_barrier_sync, "vUi", "n", "ptx60") -TARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi", "n", "ptx60") +TARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi", "n", PTX60) +TARGET_BUILTIN(__nvvm_barrier_sync, "vUi", "n", PTX60) +TARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi", "n", PTX60) // Shuffle @@ -396,14 +402,14 @@ BUILTIN(__nvvm_shfl_bfly_f32, "ffii", "") BUILTIN(__nvvm_shfl_idx_i32, "iiii", "") BUILTIN(__nvvm_shfl_idx_f32, "ffii", "") -TARGET_BUILTIN(__nvvm_shfl_sync_down_i32, "iUiiii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_down_f32, "fUifii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_up_i32, "iUiiii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_up_f32, "fUifii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_bfly_i32, "iUiiii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_bfly_f32, "fUifii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_idx_i32, "iUiiii", "", "ptx60") -TARGET_BUILTIN(__nvvm_shfl_sync_idx_f32, "fUifii", "", "ptx60") +TARGET_BUILTIN(__nvvm_shfl_sync_down_i32, "iUiiii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_down_f32, "fUifii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_up_i32, "iUiiii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_up_f32, "fUifii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_bfly_i32, "iUiiii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_bfly_f32, "fUifii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_idx_i32, "iUiiii", "", PTX60) +TARGET_BUILTIN(__nvvm_shfl_sync_idx_f32, "fUifii", "", PTX60) // Vote BUILTIN(__nvvm_vote_all, "bb", "") @@ -411,17 +417,17 @@ BUILTIN(__nvvm_vote_any, "bb", "") BUILTIN(__nvvm_vote_uni, "bb", "") BUILTIN(__nvvm_vote_ballot, "Uib", "") -TARGET_BUILTIN(__nvvm_vote_all_sync, "bUib", "", "ptx60") -TARGET_BUILTIN(__nvvm_vote_any_sync, "bUib", "", "ptx60") -TARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib", "", "ptx60") -TARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib", "", "ptx60") +TARGET_BUILTIN(__nvvm_vote_all_sync, "bUib", "", PTX60) +TARGET_BUILTIN(__nvvm_vote_any_sync, "bUib", "", PTX60) +TARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib", "", PTX60) +TARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib", "", PTX60) // Match -TARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi", "", "ptx60") -TARGET_BUILTIN(__nvvm_match_any_sync_i64, "WiUiWi", "", "ptx60") +TARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi", "", PTX60) +TARGET_BUILTIN(__nvvm_match_any_sync_i64, "WiUiWi", "", PTX60) // These return a pair {value, predicate}, which requires custom lowering. -TARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*", "", "ptx60") -TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", "ptx60") +TARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*", "", PTX60) +TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", PTX60) // Membar @@ -465,28 +471,28 @@ BUILTIN(__builtin_ptx_get_image_channel_orderi_, "ii", "") BUILTIN(__nvvm_atom_add_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_add_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_add_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_add_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_add_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_add_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_add_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_add_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_add_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_add_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_add_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_add_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_add_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_add_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_add_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_add_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_add_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_add_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_add_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_add_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_add_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_add_g_f, "ffD*1f", "n") BUILTIN(__nvvm_atom_add_s_f, "ffD*3f", "n") BUILTIN(__nvvm_atom_add_gen_f, "ffD*f", "n") -TARGET_BUILTIN(__nvvm_atom_cta_add_gen_f, "ffD*f", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_add_gen_f, "ffD*f", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_add_gen_f, "ffD*f", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_add_gen_f, "ffD*f", "n", SM_60) BUILTIN(__nvvm_atom_add_g_d, "ddD*1d", "n") BUILTIN(__nvvm_atom_add_s_d, "ddD*3d", "n") -TARGET_BUILTIN(__nvvm_atom_add_gen_d, "ddD*d", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_add_gen_d, "ddD*d", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d", "n", SM_60) BUILTIN(__nvvm_atom_sub_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_sub_s_i, "iiD*3i", "n") @@ -501,155 +507,155 @@ BUILTIN(__nvvm_atom_sub_gen_ll, "LLiLLiD*LLi", "n") BUILTIN(__nvvm_atom_xchg_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_xchg_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_xchg_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_xchg_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_xchg_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_xchg_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_xchg_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_xchg_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_xchg_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_max_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_max_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_max_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_max_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_max_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_max_g_ui, "UiUiD*1Ui", "n") BUILTIN(__nvvm_atom_max_s_ui, "UiUiD*3Ui", "n") BUILTIN(__nvvm_atom_max_gen_ui, "UiUiD*Ui", "n") -TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ui, "UiUiD*Ui", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ui, "UiUiD*Ui", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ui, "UiUiD*Ui", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ui, "UiUiD*Ui", "n", SM_60) BUILTIN(__nvvm_atom_max_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_max_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_max_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_max_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_max_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_max_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_max_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_max_g_ul, "ULiULiD*1ULi", "n") BUILTIN(__nvvm_atom_max_s_ul, "ULiULiD*3ULi", "n") BUILTIN(__nvvm_atom_max_gen_ul, "ULiULiD*ULi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ul, "ULiULiD*ULi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ul, "ULiULiD*ULi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ul, "ULiULiD*ULi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ul, "ULiULiD*ULi", "n", SM_60) BUILTIN(__nvvm_atom_max_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_max_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_max_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_max_g_ull, "ULLiULLiD*1ULLi", "n") BUILTIN(__nvvm_atom_max_s_ull, "ULLiULLiD*3ULLi", "n") BUILTIN(__nvvm_atom_max_gen_ull, "ULLiULLiD*ULLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ull, "ULLiULLiD*ULLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ull, "ULLiULLiD*ULLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) BUILTIN(__nvvm_atom_min_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_min_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_min_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_min_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_min_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_min_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_min_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_min_g_ui, "UiUiD*1Ui", "n") BUILTIN(__nvvm_atom_min_s_ui, "UiUiD*3Ui", "n") BUILTIN(__nvvm_atom_min_gen_ui, "UiUiD*Ui", "n") -TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ui, "UiUiD*Ui", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ui, "UiUiD*Ui", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ui, "UiUiD*Ui", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ui, "UiUiD*Ui", "n", SM_60) BUILTIN(__nvvm_atom_min_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_min_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_min_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_min_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_min_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_min_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_min_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_min_g_ul, "ULiULiD*1ULi", "n") BUILTIN(__nvvm_atom_min_s_ul, "ULiULiD*3ULi", "n") BUILTIN(__nvvm_atom_min_gen_ul, "ULiULiD*ULi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ul, "ULiULiD*ULi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ul, "ULiULiD*ULi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ul, "ULiULiD*ULi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ul, "ULiULiD*ULi", "n", SM_60) BUILTIN(__nvvm_atom_min_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_min_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_min_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_min_g_ull, "ULLiULLiD*1ULLi", "n") BUILTIN(__nvvm_atom_min_s_ull, "ULLiULLiD*3ULLi", "n") BUILTIN(__nvvm_atom_min_gen_ull, "ULLiULLiD*ULLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ull, "ULLiULLiD*ULLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ull, "ULLiULLiD*ULLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ull, "ULLiULLiD*ULLi", "n", SM_60) BUILTIN(__nvvm_atom_inc_g_ui, "UiUiD*1Ui", "n") BUILTIN(__nvvm_atom_inc_s_ui, "UiUiD*3Ui", "n") BUILTIN(__nvvm_atom_inc_gen_ui, "UiUiD*Ui", "n") -TARGET_BUILTIN(__nvvm_atom_cta_inc_gen_ui, "UiUiD*Ui", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_inc_gen_ui, "UiUiD*Ui", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_inc_gen_ui, "UiUiD*Ui", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_inc_gen_ui, "UiUiD*Ui", "n", SM_60) BUILTIN(__nvvm_atom_dec_g_ui, "UiUiD*1Ui", "n") BUILTIN(__nvvm_atom_dec_s_ui, "UiUiD*3Ui", "n") BUILTIN(__nvvm_atom_dec_gen_ui, "UiUiD*Ui", "n") -TARGET_BUILTIN(__nvvm_atom_cta_dec_gen_ui, "UiUiD*Ui", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_dec_gen_ui, "UiUiD*Ui", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_dec_gen_ui, "UiUiD*Ui", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_dec_gen_ui, "UiUiD*Ui", "n", SM_60) BUILTIN(__nvvm_atom_and_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_and_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_and_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_and_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_and_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_and_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_and_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_and_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_and_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_and_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_and_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_and_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_and_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_and_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_and_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_and_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_and_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_and_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_and_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_and_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_and_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_or_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_or_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_or_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_or_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_or_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_or_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_or_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_or_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_or_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_or_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_or_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_or_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_or_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_or_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_or_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_or_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_or_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_or_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_or_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_or_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_or_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_xor_g_i, "iiD*1i", "n") BUILTIN(__nvvm_atom_xor_s_i, "iiD*3i", "n") BUILTIN(__nvvm_atom_xor_gen_i, "iiD*i", "n") -TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_i, "iiD*i", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_i, "iiD*i", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_i, "iiD*i", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_i, "iiD*i", "n", SM_60) BUILTIN(__nvvm_atom_xor_g_l, "LiLiD*1Li", "n") BUILTIN(__nvvm_atom_xor_s_l, "LiLiD*3Li", "n") BUILTIN(__nvvm_atom_xor_gen_l, "LiLiD*Li", "n") -TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_l, "LiLiD*Li", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_l, "LiLiD*Li", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_l, "LiLiD*Li", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_xor_g_ll, "LLiLLiD*1LLi", "n") BUILTIN(__nvvm_atom_xor_s_ll, "LLiLLiD*3LLi", "n") BUILTIN(__nvvm_atom_xor_gen_ll, "LLiLLiD*LLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_ll, "LLiLLiD*LLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_ll, "LLiLLiD*LLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_ll, "LLiLLiD*LLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_ll, "LLiLLiD*LLi", "n", SM_60) BUILTIN(__nvvm_atom_cas_g_i, "iiD*1ii", "n") BUILTIN(__nvvm_atom_cas_s_i, "iiD*3ii", "n") BUILTIN(__nvvm_atom_cas_gen_i, "iiD*ii", "n") -TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_i, "iiD*ii", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_i, "iiD*ii", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_i, "iiD*ii", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_i, "iiD*ii", "n", SM_60) BUILTIN(__nvvm_atom_cas_g_l, "LiLiD*1LiLi", "n") BUILTIN(__nvvm_atom_cas_s_l, "LiLiD*3LiLi", "n") BUILTIN(__nvvm_atom_cas_gen_l, "LiLiD*LiLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_l, "LiLiD*LiLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_l, "LiLiD*LiLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_60) BUILTIN(__nvvm_atom_cas_g_ll, "LLiLLiD*1LLiLLi", "n") BUILTIN(__nvvm_atom_cas_s_ll, "LLiLLiD*3LLiLLi", "n") BUILTIN(__nvvm_atom_cas_gen_ll, "LLiLLiD*LLiLLi", "n") -TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", "satom") -TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", "satom") +TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60) // Compiler Error Warn BUILTIN(__nvvm_compiler_error, "vcC*4", "n") @@ -692,17 +698,19 @@ BUILTIN(__nvvm_ldg_f4, "E4fE4fC*", "") BUILTIN(__nvvm_ldg_d2, "E2dE2dC*", "") // Builtins to support WMMA instructions on sm_70 -TARGET_BUILTIN(__hmma_m16n16k16_ld_a, "vi*iC*UiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_ld_b, "vi*iC*UiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f16, "vi*iC*UiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f32, "vf*fC*UiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_st_c_f16, "vi*i*UiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_st_c_f32, "vf*f*UiIi", "", "ptx60") - -TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", "ptx60") -TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", "ptx60") +TARGET_BUILTIN(__hmma_m16n16k16_ld_a, "vi*iC*UiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_ld_b, "vi*iC*UiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f16, "vi*iC*UiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f32, "vf*fC*UiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_st_c_f16, "vi*i*UiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_st_c_f32, "vf*f*UiIi", "", PTX60) + +TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f16, "vi*iC*iC*iC*IiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f16, "vf*iC*iC*iC*IiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f32, "vf*iC*iC*fC*IiIi", "", PTX60) +TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f32, "vi*iC*iC*fC*IiIi", "", PTX60) #undef BUILTIN #undef TARGET_BUILTIN +#pragma pop_macro("SM_60") +#pragma pop_macro("PTX60") diff --git a/clang/lib/Basic/Targets/NVPTX.cpp b/clang/lib/Basic/Targets/NVPTX.cpp index c497e2e..3afb32f 100644 --- a/clang/lib/Basic/Targets/NVPTX.cpp +++ b/clang/lib/Basic/Targets/NVPTX.cpp @@ -40,6 +40,22 @@ NVPTXTargetInfo::NVPTXTargetInfo(const llvm::Triple &Triple, assert((TargetPointerWidth == 32 || TargetPointerWidth == 64) && "NVPTX only supports 32- and 64-bit modes."); + PTXVersion = 32; + for (const StringRef Feature : Opts.FeaturesAsWritten) { + if (!Feature.startswith("+ptx")) + continue; + PTXVersion = llvm::StringSwitch(Feature) + .Case("+ptx61", 61) + .Case("+ptx60", 60) + .Case("+ptx50", 50) + .Case("+ptx43", 43) + .Case("+ptx42", 42) + .Case("+ptx41", 41) + .Case("+ptx40", 40) + .Case("+ptx32", 32) + .Default(32); + } + TLSSupported = false; VLASupported = false; AddrSpaceMap = &NVPTXAddrSpaceMap; diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h index 83d7dfb..26d4c9d 100644 --- a/clang/lib/Basic/Targets/NVPTX.h +++ b/clang/lib/Basic/Targets/NVPTX.h @@ -40,6 +40,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo { static const char *const GCCRegNames[]; static const Builtin::Info BuiltinInfo[]; CudaArch GPU; + uint32_t PTXVersion; std::unique_ptr HostTarget; public: @@ -55,7 +56,9 @@ public: initFeatureMap(llvm::StringMap &Features, DiagnosticsEngine &Diags, StringRef CPU, const std::vector &FeaturesVec) const override { + Features[CudaArchToString(GPU)] = true; Features["satom"] = GPU >= CudaArch::SM_60; + Features["ptx" + std::to_string(PTXVersion)] = true; return TargetInfo::initFeatureMap(Features, Diags, CPU, FeaturesVec); } diff --git a/clang/test/CodeGen/builtins-nvptx-ptx50.cu b/clang/test/CodeGen/builtins-nvptx-ptx50.cu index e85be44..72e1aec 100644 --- a/clang/test/CodeGen/builtins-nvptx-ptx50.cu +++ b/clang/test/CodeGen/builtins-nvptx-ptx50.cu @@ -18,6 +18,6 @@ // CHECK-LABEL: test_fn __device__ void test_fn(double d, double* double_ptr) { // CHECK: call double @llvm.nvvm.atomic.load.add.f64.p0f64 - // expected-error@+1 {{'__nvvm_atom_add_gen_d' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_add_gen_d' needs target feature sm_60}} __nvvm_atom_add_gen_d(double_ptr, d); } diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index 89a9823..16f41ba 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -5,6 +5,9 @@ // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_60 \ // RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ // RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s +// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_61 \ +// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ +// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s // RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_53 \ // RUN: -DERROR_CHECK -fcuda-is-device -S -o /dev/null -x cuda -verify %s @@ -292,245 +295,245 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, #if ERROR_CHECK || __CUDA_ARCH__ >= 600 // CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_add_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_i' needs target feature sm_60}} __nvvm_atom_cta_add_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_add_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_l' needs target feature sm_60}} __nvvm_atom_cta_add_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_add_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_add_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_add_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_i' needs target feature sm_60}} __nvvm_atom_sys_add_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_add_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_l' needs target feature sm_60}} __nvvm_atom_sys_add_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_add_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_add_gen_ll(&sll, ll); // CHECK: call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0f32 - // expected-error@+1 {{'__nvvm_atom_cta_add_gen_f' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_f' needs target feature sm_60}} __nvvm_atom_cta_add_gen_f(fp, f); // CHECK: call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0f64 - // expected-error@+1 {{'__nvvm_atom_cta_add_gen_d' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_d' needs target feature sm_60}} __nvvm_atom_cta_add_gen_d(dfp, df); // CHECK: call float @llvm.nvvm.atomic.add.gen.f.sys.f32.p0f32 - // expected-error@+1 {{'__nvvm_atom_sys_add_gen_f' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_f' needs target feature sm_60}} __nvvm_atom_sys_add_gen_f(fp, f); // CHECK: call double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0f64 - // expected-error@+1 {{'__nvvm_atom_sys_add_gen_d' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_d' needs target feature sm_60}} __nvvm_atom_sys_add_gen_d(dfp, df); // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_i' needs target feature sm_60}} __nvvm_atom_cta_xchg_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_l' needs target feature sm_60}} __nvvm_atom_cta_xchg_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_xchg_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature sm_60}} __nvvm_atom_sys_xchg_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_l' needs target feature sm_60}} __nvvm_atom_sys_xchg_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_xchg_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature sm_60}} __nvvm_atom_cta_max_gen_i(ip, i); // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ui' needs target feature sm_60}} __nvvm_atom_cta_max_gen_ui((unsigned int *)ip, i); // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_max_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_l' needs target feature sm_60}} __nvvm_atom_cta_max_gen_l(&dl, l); // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ul' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ul' needs target feature sm_60}} __nvvm_atom_cta_max_gen_ul((unsigned long *)lp, l); // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_max_gen_ll(&sll, ll); // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ull' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ull' needs target feature sm_60}} __nvvm_atom_cta_max_gen_ull((unsigned long long *)llp, ll); // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_max_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_i' needs target feature sm_60}} __nvvm_atom_sys_max_gen_i(ip, i); // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ui' needs target feature sm_60}} __nvvm_atom_sys_max_gen_ui((unsigned int *)ip, i); // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_max_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_l' needs target feature sm_60}} __nvvm_atom_sys_max_gen_l(&dl, l); // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ul' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ul' needs target feature sm_60}} __nvvm_atom_sys_max_gen_ul((unsigned long *)lp, l); // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_max_gen_ll(&sll, ll); // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ull' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ull' needs target feature sm_60}} __nvvm_atom_sys_max_gen_ull((unsigned long long *)llp, ll); // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_min_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_i' needs target feature sm_60}} __nvvm_atom_cta_min_gen_i(ip, i); // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ui' needs target feature sm_60}} __nvvm_atom_cta_min_gen_ui((unsigned int *)ip, i); // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_min_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_l' needs target feature sm_60}} __nvvm_atom_cta_min_gen_l(&dl, l); // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ul' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ul' needs target feature sm_60}} __nvvm_atom_cta_min_gen_ul((unsigned long *)lp, l); // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_min_gen_ll(&sll, ll); // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ull' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ull' needs target feature sm_60}} __nvvm_atom_cta_min_gen_ull((unsigned long long *)llp, ll); // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_min_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_i' needs target feature sm_60}} __nvvm_atom_sys_min_gen_i(ip, i); // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ui' needs target feature sm_60}} __nvvm_atom_sys_min_gen_ui((unsigned int *)ip, i); // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_min_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_l' needs target feature sm_60}} __nvvm_atom_sys_min_gen_l(&dl, l); // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ul' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ul' needs target feature sm_60}} __nvvm_atom_sys_min_gen_ul((unsigned long *)lp, l); // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_min_gen_ll(&sll, ll); // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ull' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ull' needs target feature sm_60}} __nvvm_atom_sys_min_gen_ull((unsigned long long *)llp, ll); // CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_inc_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_inc_gen_ui' needs target feature sm_60}} __nvvm_atom_cta_inc_gen_ui((unsigned int *)ip, i); // CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_inc_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_inc_gen_ui' needs target feature sm_60}} __nvvm_atom_sys_inc_gen_ui((unsigned int *)ip, i); // CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_dec_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_dec_gen_ui' needs target feature sm_60}} __nvvm_atom_cta_dec_gen_ui((unsigned int *)ip, i); // CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_dec_gen_ui' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_dec_gen_ui' needs target feature sm_60}} __nvvm_atom_sys_dec_gen_ui((unsigned int *)ip, i); // CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_and_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_and_gen_i' needs target feature sm_60}} __nvvm_atom_cta_and_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_and_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_and_gen_l' needs target feature sm_60}} __nvvm_atom_cta_and_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_and_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_and_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_and_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_and_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_and_gen_i' needs target feature sm_60}} __nvvm_atom_sys_and_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_and_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_and_gen_l' needs target feature sm_60}} __nvvm_atom_sys_and_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_and_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_and_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_and_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_or_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_or_gen_i' needs target feature sm_60}} __nvvm_atom_cta_or_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_or_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_or_gen_l' needs target feature sm_60}} __nvvm_atom_cta_or_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_or_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_or_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_or_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_or_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_or_gen_i' needs target feature sm_60}} __nvvm_atom_sys_or_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_or_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_or_gen_l' needs target feature sm_60}} __nvvm_atom_sys_or_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_or_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_or_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_or_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_i' needs target feature sm_60}} __nvvm_atom_cta_xor_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_l' needs target feature sm_60}} __nvvm_atom_cta_xor_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_xor_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_i' needs target feature sm_60}} __nvvm_atom_sys_xor_gen_i(ip, i); // LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_l' needs target feature sm_60}} __nvvm_atom_sys_xor_gen_l(&dl, l); // CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_xor_gen_ll(&sll, ll); // CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_i' needs target feature sm_60}} __nvvm_atom_cta_cas_gen_i(ip, i, 0); // LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_l' needs target feature sm_60}} __nvvm_atom_cta_cas_gen_l(&dl, l, 0); // CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_cas_gen_ll(&sll, ll, 0); // CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0i32 - // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_i' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_i' needs target feature sm_60}} __nvvm_atom_sys_cas_gen_i(ip, i, 0); // LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0i32 // LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_l' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_l' needs target feature sm_60}} __nvvm_atom_sys_cas_gen_l(&dl, l, 0); // CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0i64 - // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_ll' needs target feature satom}} + // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_cas_gen_ll(&sll, ll, 0); #endif diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h index b89bd41..aa57e3e 100644 --- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.h +++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.h @@ -74,7 +74,7 @@ public: } bool hasAtomAddF64() const { return SmVersion >= 60; } - bool hasAtomScope() const { return HasAtomScope; } + bool hasAtomScope() const { return SmVersion >= 60; } bool hasAtomBitwise64() const { return SmVersion >= 32; } bool hasAtomMinMax64() const { return SmVersion >= 32; } bool hasLDG() const { return SmVersion >= 32; }