Index: include/clang/Basic/BuiltinsNVPTX.def =================================================================== --- include/clang/Basic/BuiltinsNVPTX.def +++ include/clang/Basic/BuiltinsNVPTX.def @@ -14,6 +14,10 @@ // 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") @@ -452,18 +456,28 @@ 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") @@ -478,97 +492,155 @@ 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") @@ -611,3 +683,4 @@ BUILTIN(__nvvm_ldg_d2, "E2dE2dC*", "") #undef BUILTIN +#undef TARGET_BUILTIN Index: lib/Basic/Targets.cpp =================================================================== --- lib/Basic/Targets.cpp +++ lib/Basic/Targets.cpp @@ -1830,8 +1830,19 @@ return llvm::makeArrayRef(BuiltinInfo, clang::NVPTX::LastTSBuiltin - Builtin::FirstTSBuiltin); } + bool + initFeatureMap(llvm::StringMap &Features, DiagnosticsEngine &Diags, + StringRef CPU, + const std::vector &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(Feature) + .Cases("ptx", "nvptx", true) + .Case("satom", GPU >= CudaArch::SM_60) // Atomics w/ scope. + .Default(false); } ArrayRef getGCCRegNames() const override; @@ -1886,6 +1897,8 @@ { #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" }; Index: lib/CodeGen/CGBuiltin.cpp =================================================================== --- lib/CodeGen/CGBuiltin.cpp +++ lib/CodeGen/CGBuiltin.cpp @@ -8105,7 +8105,13 @@ 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: @@ -8224,6 +8230,109 @@ 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; } Index: test/CodeGen/builtins-nvptx.c =================================================================== --- test/CodeGen/builtins-nvptx.c +++ test/CodeGen/builtins-nvptx.c @@ -1,8 +1,12 @@ // 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)) @@ -191,8 +195,9 @@ // 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 @@ -280,6 +285,255 @@ // 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 }