[NVPTX, CUDA] Improved feature constraints on NVPTX target builtins.

[tra]

Summary

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.

[tra]

Run the tests with -target-cpu sm_61 to make sure intrinsics that require sm_60 are still enabled.

[jlebar]

clang/include/clang/Basic/Cuda.h

55 ↗

(On Diff #140334)

Why 'static'?

59 ↗

(On Diff #140334)

Could we instead rely on the fact that the CudaArch enum is dense and goes from UNKNOWN to LAST? One less bug to make...

clang/lib/Basic/Targets/NVPTX.cpp

182 ↗

(On Diff #140334)

Nit, if you spell it /*FeatureMap*/ here, it'll be clearer that you're intentionally ignoring that arg.

[tra]

clang/include/clang/Basic/Cuda.h

55 ↗

(On Diff #140334)

Old habits. Just 'inline' should do.

59 ↗

(On Diff #140334)

Actually this function is not needed. It's a leftover from an old iteration of the code. I'll remove it.
As for why array -- that's what I wanted to do. Turns out that iterating over enums is a pain. I can't increment/decrement it and had to cast to an int for that. And then back, in order to use it.

[tra]

Removed unneeded function.

[tra]

Commented out unused argument.

[echristo]

The llvm change and corresponding switch from satom->sm_60 in the front end is fine.

Let's talk about the rest of it more. I'm not sure I'm seeing the need here rather than the annotations that are already here. Can you elaborate more here on why we need an additional method when you've already got subtarget features for each of the ptx versions anyhow?

Thanks.

-eric

[tra]

In D45061#1053795, @echristo wrote:

Let's talk about the rest of it more. I'm not sure I'm seeing the need here rather than the annotations that are already here. Can you elaborate more here on why we need an additional method when you've already got subtarget features for each of the ptx versions anyhow?

The patch intends to address two issues:

• mismatch on constraints between llvm and clang. E.g. hasPTX60() on LLVM side means "ptx60 or newer". "ptx60" in TARGET_BUILTIN on clang side means "ptx60" *only*. It is possible to address this within existing implementation by enumerating all PTX versions that are newer. It works OK for ptx60 as we only need to write "ptx60|ptx61". It gets more interesting for older GPUs E.g "ptx31+" would have to become "ptx31|ptx32|ptx40|ptx41|ptx42|ptx43|ptx50|ptx60|ptx61". Similar enumeration will need to happen for GPU version, which brings us to the next point
• NVIDIA keeps growing PTX versions and GPU variants. Recently they've changed CUDA release frequency to ~1/quarter and they tend to add minor variants of PTX and GPU versions fairly frequently. It's going to be an unnecessary maintenance headache as after every new introduced variant I'll need to go and update all NVPTX builtins that have nothing to do with the CUDA changes. Granted, it's not a showstopper, but it is an annoyance that is guaranteed to stay.

With this patch, TARGET_BUILTIN constraints become semantically identical to LLVM and we no longer need to chase every bump in PTX version or GPU variant.

[tra]

@echristo convinced me that this functionality can be implemented without growing a target-specific hook for custom interpretation of constraints used in TARGET_BUILTIN. Instead, we can hide unwieldy feature lists behind a macro.

[echristo]

Guessing that SM_60 (etc) are defines?

Anyhow LGTM. I'm not sure you can split up the satom part of the patch, but if you can that'd be great.

-eric

[tra]

In D45061#1063791, @echristo wrote:

Guessing that SM_60 (etc) are defines?

Anyhow LGTM. I'm not sure you can split up the satom part of the patch, but if you can that'd be great.

-eric

I'll split out removal of satom feature into a separate commit.

include/clang/Basic/BuiltinsNVPTX.def

 //===----------------------------------------------------------------------===//
 
 // The format of this database matches clang/Basic/Builtins.def.
 
 #if defined(BUILTIN) && !defined(TARGET_BUILTIN)
 #   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")
 BUILTIN(__nvvm_read_ptx_sreg_tid_y, "i", "nc")
 BUILTIN(__nvvm_read_ptx_sreg_tid_z, "i", "nc")
 BUILTIN(__nvvm_read_ptx_sreg_tid_w, "i", "nc")
 
 BUILTIN(__nvvm_read_ptx_sreg_ntid_x, "i", "nc")

 
 BUILTIN(__nvvm_bitcast_f2i, "if", "")
 BUILTIN(__nvvm_bitcast_i2f, "fi", "")
 
 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
 
 BUILTIN(__syncthreads, "v", "")
 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
 
 BUILTIN(__nvvm_shfl_down_i32, "iiii", "")
 BUILTIN(__nvvm_shfl_down_f32, "ffii", "")
 BUILTIN(__nvvm_shfl_up_i32, "iiii", "")
 BUILTIN(__nvvm_shfl_up_f32, "ffii", "")
 BUILTIN(__nvvm_shfl_bfly_i32, "iiii", "")
 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", "")
 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
 
 BUILTIN(__nvvm_membar_cta, "v", "")
 BUILTIN(__nvvm_membar_gl, "v", "")
 BUILTIN(__nvvm_membar_sys, "v", "")
 
 // Memcpy, Memset

 // We need the atom intrinsics because
 // - they are used in converging analysis
 // - they are used in address space analysis and optimization
 // So it does not hurt to expose them as builtins.
 //
 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")
 BUILTIN(__nvvm_atom_sub_gen_i, "iiD*i", "n")
 BUILTIN(__nvvm_atom_sub_g_l, "LiLiD*1Li", "n")
 BUILTIN(__nvvm_atom_sub_s_l, "LiLiD*3Li", "n")
 BUILTIN(__nvvm_atom_sub_gen_l, "LiLiD*Li", "n")
 BUILTIN(__nvvm_atom_sub_g_ll, "LLiLLiD*1LLi", "n")
 BUILTIN(__nvvm_atom_sub_s_ll, "LLiLLiD*3LLi", "n")
 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")
 BUILTIN(__nvvm_compiler_warn, "vcC*4", "n")
 
 // __ldg.  This is not implemented as a builtin by nvcc.
 BUILTIN(__nvvm_ldg_c, "ccC*", "")
 BUILTIN(__nvvm_ldg_s, "ssC*", "")

 BUILTIN(__nvvm_ldg_ui4, "E4UiE4UiC*", "")
 BUILTIN(__nvvm_ldg_ull2, "E2ULLiE2ULLiC*", "")
 
 BUILTIN(__nvvm_ldg_f2, "E2fE2fC*", "")
 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")

lib/Basic/Targets/NVPTX.h

static const unsigned NVPTXAddrSpaceMap[] = {
     4, // cuda_constant
     3, // cuda_shared
 };
 
 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<TargetInfo> HostTarget;
 
 public:
   NVPTXTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts,
                   unsigned TargetPointerWidth);
 
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override;
 
   bool
   initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
                  StringRef CPU,
                  const std::vector<std::string> &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);
   }
 
   bool hasFeature(StringRef Feature) const override;
 
   ArrayRef<const char *> getGCCRegNames() const override;
 
   ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {

lib/Basic/Targets/NVPTX.cpp

 
 NVPTXTargetInfo::NVPTXTargetInfo(const llvm::Triple &Triple,
                                  const TargetOptions &Opts,
                                  unsigned TargetPointerWidth)
     : TargetInfo(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<unsigned>(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;
   UseAddrSpaceMapMangling = true;
 
   // Define available target features
   // These must be defined in sorted order!
   NoAsmVariants = true;

test/CodeGen/builtins-nvptx-ptx50.cu

 // We have to keep all builtins that depend on particular target feature in the
 // same function, because the codegen will stop after the very first function
 // that encounters an error, so -verify will not be able to find errors in
 // subsequent functions.
 
 // 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);
 }

test/CodeGen/builtins-nvptx.c

 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \
 // RUN:            -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
 // RUN:   | FileCheck -check-prefix=CHECK -check-prefix=LP32 %s
 // 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
 
 #define __device__ __attribute__((device))
 #define __global__ __attribute__((global))
 #define __shared__ __attribute__((shared))
 #define __constant__ __attribute__((constant))
 
__device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip,
 
 
   //////////////////////////////////////////////////////////////////
   // Atomics with scope (only supported on sm_60+).
 
 #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
 
   // CHECK: ret
 }
 
 // CHECK-LABEL: nvvm_ldg
 __device__ void nvvm_ldg(const void *p) {