[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

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.

clang/include/clang/Basic/BuiltinsNVPTX.def

 // 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*", "")

clang/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["satom"] = GPU >= CudaArch::SM_60;
-    return TargetInfo::initFeatureMap(Features, Diags, CPU, FeaturesVec);
-  }
-
   bool hasFeature(StringRef Feature) const override;
+  Optional<bool> hasRequiredFeature(const llvm::StringMap<bool> FeatureMap,
+                                    const StringRef ReqFeature) const override;
 
   ArrayRef<const char *> getGCCRegNames() const override;
 
   ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
     // No aliases.
     return None;
   }
 

clang/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;
NVPTXTargetInfo::NVPTXTargetInfo(const llvm::Triple &Triple,
   //   TODO: nvcc emits a warning when using long double on device; we should
   //   do the same.
 }
 
 ArrayRef<const char *> NVPTXTargetInfo::getGCCRegNames() const {
   return llvm::makeArrayRef(GCCRegNames);
 }
 
-bool NVPTXTargetInfo::hasFeature(StringRef Feature) const {
-  return llvm::StringSwitch<bool>(Feature)
+static Optional<bool> hasFeatureEnabled(CudaArch GPU, uint32_t PTXVersion,
+                                       const StringRef Feature) {
+  // Handle known GPU names.
+  CudaArch Arch = StringToCudaArch(Feature);
+  if (Arch != CudaArch::UNKNOWN)
+    return GPU >= Arch;
+
+  return llvm::StringSwitch<Optional<bool>>(Feature)
       .Cases("ptx", "nvptx", true)
-      .Case("satom", GPU >= CudaArch::SM_60) // Atomics w/ scope.
-      .Default(false);
+      .Case("ptx61", PTXVersion >= 61)
+      .Case("ptx60", PTXVersion >= 60)
+      .Case("ptx50", PTXVersion >= 50)
+      .Case("ptx43", PTXVersion >= 43)
+      .Case("ptx42", PTXVersion >= 42)
+      .Case("ptx41", PTXVersion >= 41)
+      .Case("ptx40", PTXVersion >= 40)
+      .Case("ptx32", PTXVersion >= 32)
+      .Default({});
+}
+
+Optional<bool>
+NVPTXTargetInfo::hasRequiredFeature(const llvm::StringMap<bool> FeatureMap,
+                                    const StringRef ReqFeature) const {
+  return hasFeatureEnabled(GPU, PTXVersion, ReqFeature);
+}
+
+bool NVPTXTargetInfo::hasFeature(StringRef Feature) const {
+  if (Optional<bool> F = hasFeatureEnabled(GPU, PTXVersion, Feature))
+    return F.getValue();
+  return false;
 }
 
 void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
                                        MacroBuilder &Builder) const {
   Builder.defineMacro("__PTX__");
   Builder.defineMacro("__NVPTX__");
   if (Opts.CUDAIsDevice) {
     // Set __CUDA_ARCH__ for the GPU specified.

clang/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);
 }

clang/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) {

llvm/lib/Target/NVPTX/NVPTX.td

def SM60 : SubtargetFeature<"sm_60", "SmVersion", "60",
                              "Target SM 6.0">;
 def SM61 : SubtargetFeature<"sm_61", "SmVersion", "61",
                              "Target SM 6.1">;
 def SM62 : SubtargetFeature<"sm_62", "SmVersion", "62",
                              "Target SM 6.2">;
 def SM70 : SubtargetFeature<"sm_70", "SmVersion", "70",
                              "Target SM 7.0">;
 
-def SATOM : SubtargetFeature<"satom", "HasAtomScope", "true",
-                             "Atomic operations with scope">;
-
 // PTX Versions
 def PTX32 : SubtargetFeature<"ptx32", "PTXVersion", "32",
                              "Use PTX version 3.2">;
 def PTX40 : SubtargetFeature<"ptx40", "PTXVersion", "40",
                              "Use PTX version 4.0">;
 def PTX41 : SubtargetFeature<"ptx41", "PTXVersion", "41",
                              "Use PTX version 4.1">;
 def PTX42 : SubtargetFeature<"ptx42", "PTXVersion", "42",

 def : Proc<"sm_21", [SM21]>;
 def : Proc<"sm_30", [SM30]>;
 def : Proc<"sm_32", [SM32, PTX40]>;
 def : Proc<"sm_35", [SM35]>;
 def : Proc<"sm_37", [SM37, PTX41]>;
 def : Proc<"sm_50", [SM50, PTX40]>;
 def : Proc<"sm_52", [SM52, PTX41]>;
 def : Proc<"sm_53", [SM53, PTX42]>;
-def : Proc<"sm_60", [SM60, PTX50, SATOM]>;
-def : Proc<"sm_61", [SM61, PTX50, SATOM]>;
-def : Proc<"sm_62", [SM62, PTX50, SATOM]>;
-def : Proc<"sm_70", [SM70, PTX60, SATOM]>;
+def : Proc<"sm_60", [SM60, PTX50]>;
+def : Proc<"sm_61", [SM61, PTX50]>;
+def : Proc<"sm_62", [SM62, PTX50]>;
+def : Proc<"sm_70", [SM70, PTX60]>;
 
 def NVPTXInstrInfo : InstrInfo {
 }
 
 def NVPTX : Target {
   let InstructionSet = NVPTXInstrInfo;
 }

llvm/lib/Target/NVPTX/NVPTXSubtarget.h

class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
   NVPTXInstrInfo InstrInfo;
   NVPTXTargetLowering TLInfo;
   SelectionDAGTargetInfo TSInfo;
 
   // NVPTX does not have any call stack frame, but need a NVPTX specific
   // FrameLowering class because TargetFrameLowering is abstract.
   NVPTXFrameLowering FrameLowering;
 
-protected:
-  // Processor supports scoped atomic operations.
-  bool HasAtomScope;
-
 public:
   /// This constructor initializes the data members to match that
   /// of the specified module.
   ///
   NVPTXSubtarget(const Triple &TT, const std::string &CPU,
                  const std::string &FS, const NVPTXTargetMachine &TM);
 
   const TargetFrameLowering *getFrameLowering() const override {
     return &FrameLowering;
   }
   const NVPTXInstrInfo *getInstrInfo() const override { return &InstrInfo; }
   const NVPTXRegisterInfo *getRegisterInfo() const override {
     return &InstrInfo.getRegisterInfo();
   }
   const NVPTXTargetLowering *getTargetLowering() const override {
     return &TLInfo;
   }
   const SelectionDAGTargetInfo *getSelectionDAGInfo() const override {
     return &TSInfo;
   }
 
   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; }
   bool hasLDU() const { return ((SmVersion >= 20) && (SmVersion < 30)); }
   inline bool hasHWROT32() const { return SmVersion >= 32; }
   bool hasImageHandles() const;
   bool hasFP16Math() const { return SmVersion >= 53; }
   bool allowFP16Math() const;