[CUDA] Added __nvvm_atom_{sys|cta}_* builtins for sm_60 GPUs.

include/clang/Basic/BuiltinsNVPTX.def

 //===--- BuiltinsPTX.def - PTX Builtin function database ----*- C++ -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the PTX-specific builtin function database.  Users of
 // this file must define the BUILTIN macro to make use of this information.
 //
 //===----------------------------------------------------------------------===//
 
 // 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
+
 // 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")

 // 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")
 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")
 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")
 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")
 BUILTIN(__nvvm_atom_add_g_d, "ddD*1d", "n")
 BUILTIN(__nvvm_atom_add_s_d, "ddD*3d", "n")
 BUILTIN(__nvvm_atom_add_gen_d, "ddD*d", "n")
+TARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d", "n", "satom")
+TARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d", "n", "satom")
 
 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")
 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")
 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")
 
 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")
 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")
 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")
 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")
 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")
 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")
 
 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")
 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")
 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")
 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")
 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")
 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")
 
 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")
 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")
 
 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")
 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")
 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")
 
 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")
 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")
 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")
 
 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")
 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")
 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")
 
 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")
 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")
 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")
 
 // 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*", "")
 
 #undef BUILTIN
+#undef TARGET_BUILTIN

lib/Basic/Targets.cpp

if (Opts.CUDAIsDevice) {
       }();
       Builder.defineMacro("__CUDA_ARCH__", CUDAArchCode);
     }
   }
   ArrayRef<Builtin::Info> getTargetBuiltins() const override {
     return llvm::makeArrayRef(BuiltinInfo,
                          clang::NVPTX::LastTSBuiltin - Builtin::FirstTSBuiltin);
   }
+  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 {
-    return Feature == "ptx" || Feature == "nvptx";
+    return llvm::StringSwitch<bool>(Feature)
+        .Cases("ptx", "nvptx", true)
+        .Case("satom", GPU >= CudaArch::SM_60)  // Atomics w/ scope.
+        .Default(false);
   }
 
   ArrayRef<const char *> getGCCRegNames() const override;
   ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
     // No aliases.
     return None;
   }
   bool validateAsmConstraint(const char *&Name,
public:
   }
 };
 
 const Builtin::Info NVPTXTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                                               \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
 #define LIBBUILTIN(ID, TYPE, ATTRS, HEADER)                                    \
   { #ID, TYPE, ATTRS, HEADER, ALL_LANGUAGES, nullptr },
+#define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE)                               \
+  { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, FEATURE },
 #include "clang/Basic/BuiltinsNVPTX.def"
 };
 
 const char *const NVPTXTargetInfo::GCCRegNames[] = {"r0"};
 
 ArrayRef<const char *> NVPTXTargetInfo::getGCCRegNames() const {
   return llvm::makeArrayRef(GCCRegNames);
 }

lib/CodeGen/CGBuiltin.cpp

auto MakeLdg = [&](unsigned IntrinsicID) {
     AlignmentSource AlignSource;
     clang::CharUnits Align =
         getNaturalPointeeTypeAlignment(E->getArg(0)->getType(), &AlignSource);
     return Builder.CreateCall(
         CGM.getIntrinsic(IntrinsicID, {Ptr->getType()->getPointerElementType(),
                                        Ptr->getType()}),
         {Ptr, ConstantInt::get(Builder.getInt32Ty(), Align.getQuantity())});
   };
-
+  auto MakeScopedAtomic = [&](unsigned IntrinsicID) {
+    Value *Ptr = EmitScalarExpr(E->getArg(0));
+    return Builder.CreateCall(
+        CGM.getIntrinsic(IntrinsicID, {Ptr->getType()->getPointerElementType(),
+                                       Ptr->getType()}),
+        {Ptr, EmitScalarExpr(E->getArg(1))});
+  };
   switch (BuiltinID) {
   case NVPTX::BI__nvvm_atom_add_gen_i:
   case NVPTX::BI__nvvm_atom_add_gen_l:
   case NVPTX::BI__nvvm_atom_add_gen_ll:
     return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Add, E);
 
   case NVPTX::BI__nvvm_atom_sub_gen_i:
   case NVPTX::BI__nvvm_atom_sub_gen_l:
case NVPTX::BI__nvvm_ldg_ull2:
     // of its member: n*alignof(t)."
     return MakeLdg(Intrinsic::nvvm_ldg_global_i);
   case NVPTX::BI__nvvm_ldg_f:
   case NVPTX::BI__nvvm_ldg_f2:
   case NVPTX::BI__nvvm_ldg_f4:
   case NVPTX::BI__nvvm_ldg_d:
   case NVPTX::BI__nvvm_ldg_d2:
     return MakeLdg(Intrinsic::nvvm_ldg_global_f);
+
+  case NVPTX::BI__nvvm_atom_cta_add_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_add_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_add_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_add_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_add_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_add_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_add_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_add_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_add_gen_f:
+  case NVPTX::BI__nvvm_atom_cta_add_gen_d:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_add_gen_f_cta);
+  case NVPTX::BI__nvvm_atom_sys_add_gen_f:
+  case NVPTX::BI__nvvm_atom_sys_add_gen_d:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_add_gen_f_sys);
+  case NVPTX::BI__nvvm_atom_cta_xchg_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_xchg_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_xchg_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_xchg_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_xchg_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_xchg_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_max_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_max_gen_ui:
+  case NVPTX::BI__nvvm_atom_cta_max_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_max_gen_ul:
+  case NVPTX::BI__nvvm_atom_cta_max_gen_ll:
+  case NVPTX::BI__nvvm_atom_cta_max_gen_ull:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_max_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_max_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_max_gen_ui:
+  case NVPTX::BI__nvvm_atom_sys_max_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_max_gen_ul:
+  case NVPTX::BI__nvvm_atom_sys_max_gen_ll:
+  case NVPTX::BI__nvvm_atom_sys_max_gen_ull:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_max_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_min_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_min_gen_ui:
+  case NVPTX::BI__nvvm_atom_cta_min_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_min_gen_ul:
+  case NVPTX::BI__nvvm_atom_cta_min_gen_ll:
+  case NVPTX::BI__nvvm_atom_cta_min_gen_ull:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_min_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_min_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_min_gen_ui:
+  case NVPTX::BI__nvvm_atom_sys_min_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_min_gen_ul:
+  case NVPTX::BI__nvvm_atom_sys_min_gen_ll:
+  case NVPTX::BI__nvvm_atom_sys_min_gen_ull:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_min_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_inc_gen_ui:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_inc_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_cta_dec_gen_ui:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_dec_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_inc_gen_ui:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_inc_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_sys_dec_gen_ui:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_dec_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_and_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_and_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_and_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_and_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_and_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_and_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_and_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_and_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_or_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_or_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_or_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_or_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_or_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_or_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_or_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_or_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_xor_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_xor_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_xor_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_xor_gen_i_cta);
+  case NVPTX::BI__nvvm_atom_sys_xor_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_xor_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_xor_gen_ll:
+    return MakeScopedAtomic(Intrinsic::nvvm_atomic_xor_gen_i_sys);
+  case NVPTX::BI__nvvm_atom_cta_cas_gen_i:
+  case NVPTX::BI__nvvm_atom_cta_cas_gen_l:
+  case NVPTX::BI__nvvm_atom_cta_cas_gen_ll: {
+    Value *Ptr = EmitScalarExpr(E->getArg(0));
+    return Builder.CreateCall(
+        CGM.getIntrinsic(
+            Intrinsic::nvvm_atomic_cas_gen_i_cta,
+            {Ptr->getType()->getPointerElementType(), Ptr->getType()}),
+        {Ptr, EmitScalarExpr(E->getArg(1)), EmitScalarExpr(E->getArg(2))});
+  }
+  case NVPTX::BI__nvvm_atom_sys_cas_gen_i:
+  case NVPTX::BI__nvvm_atom_sys_cas_gen_l:
+  case NVPTX::BI__nvvm_atom_sys_cas_gen_ll: {
+    Value *Ptr = EmitScalarExpr(E->getArg(0));
+    return Builder.CreateCall(
+        CGM.getIntrinsic(
+            Intrinsic::nvvm_atomic_cas_gen_i_sys,
+            {Ptr->getType()->getPointerElementType(), Ptr->getType()}),
+        {Ptr, EmitScalarExpr(E->getArg(1)), EmitScalarExpr(E->getArg(2))});
+  }
   default:
     return nullptr;
   }
 }
 
 Value *CodeGenFunction::EmitWebAssemblyBuiltinExpr(unsigned BuiltinID,
                                                    const CallExpr *E) {
   switch (BuiltinID) {

test/CodeGen/builtins-nvptx.c

 // REQUIRES: nvptx-registered-target
-// RUN: %clang_cc1 -triple nvptx-unknown-unknown -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 -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_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 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__ int read_tid() {
 

 __shared__ int si;
 __device__ long dl;
 __shared__ long sl;
 __device__ long long dll;
 __shared__ long long sll;
 
 // Check for atomic intrinsics
 // CHECK-LABEL: nvvm_atom
-__device__ void nvvm_atom(float *fp, float f, int *ip, int i, unsigned int *uip, unsigned ui, long *lp, long l,
-                          long long *llp, long long ll) {
+__device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip,
+                          int i, unsigned int *uip, unsigned ui, long *lp,
+                          long l, long long *llp, long long ll) {
   // CHECK: atomicrmw add
   __nvvm_atom_add_gen_i(ip, i);
   // CHECK: atomicrmw add
   __nvvm_atom_add_gen_l(&dl, l);
   // CHECK: atomicrmw add
   __nvvm_atom_add_gen_ll(&sll, ll);
 
   // CHECK: atomicrmw sub
__device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip,
   __nvvm_atom_add_gen_f(fp, f);
 
   // CHECK: call i32 @llvm.nvvm.atomic.load.inc.32.p0i32
   __nvvm_atom_inc_gen_ui(uip, ui);
 
   // CHECK: call i32 @llvm.nvvm.atomic.load.dec.32.p0i32
   __nvvm_atom_dec_gen_ui(uip, ui);
 
+
+  //////////////////////////////////////////////////////////////////
+  // 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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __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}}
+  __nvvm_atom_sys_cas_gen_ll(&sll, ll, 0);
+#endif
+
   // CHECK: ret
 }
 
 // CHECK-LABEL: nvvm_ldg
 __device__ void nvvm_ldg(const void *p) {
   // CHECK: call i8 @llvm.nvvm.ldg.global.i.i8.p0i8(i8* {{%[0-9]+}}, i32 1)
   // CHECK: call i8 @llvm.nvvm.ldg.global.i.i8.p0i8(i8* {{%[0-9]+}}, i32 1)
   __nvvm_ldg_c((const char *)p);