diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def index eaf6e93796aa7..32f2807fc3164 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.def +++ b/clang/include/clang/Basic/BuiltinsNVPTX.def @@ -687,6 +687,12 @@ TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_xchg_gen_ll, "LLiLLiD*LLi", "n") 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_xchg_gen_f, "ffD*f", "n") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_f, "ffD*f", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_f, "ffD*f", "n", SM_60) +BUILTIN(__nvvm_atom_xchg_gen_d, "ddD*d", "n") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_d, "ddD*d", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_d, "ddD*d", "n", SM_60) BUILTIN(__nvvm_atom_max_gen_i, "iiD*i", "n") TARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i", "n", SM_60) @@ -785,6 +791,12 @@ TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_60) BUILTIN(__nvvm_atom_cas_gen_ll, "LLiLLiD*LLiLLi", "n") 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) +BUILTIN(__nvvm_atom_cas_gen_f, "ffD*ff", "n") +TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_f, "ffD*ff", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_f, "ffD*ff", "n", SM_60) +BUILTIN(__nvvm_atom_cas_gen_d, "ddD*dd", "n") +TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_d, "ddD*dd", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_d, "ddD*dd", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_acquire_add_gen_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_add_gen_i, "iiD*i", "n", SM_70) @@ -811,6 +823,12 @@ TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_gen_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_xchg_gen_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_gen_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_gen_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_max_gen_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_max_gen_i, "iiD*i", "n", SM_70) @@ -909,6 +927,12 @@ TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cas_gen_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_gen_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_gen_d, "ddD*dd", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_add_gen_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_add_gen_i, "iiD*i", "n", SM_70) @@ -935,6 +959,12 @@ TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_gen_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_xchg_gen_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_gen_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_gen_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_max_gen_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_max_gen_i, "iiD*i", "n", SM_70) @@ -1033,6 +1063,12 @@ TARGET_BUILTIN(__nvvm_atom_release_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cas_gen_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_cas_gen_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_cas_gen_d, "ddD*dd", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_add_gen_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_add_gen_i, "iiD*i", "n", SM_70) @@ -1059,6 +1095,12 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_gen_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_gen_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_gen_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_gen_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_gen_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_gen_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_max_gen_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_max_gen_i, "iiD*i", "n", SM_70) @@ -1157,6 +1199,12 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_gen_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_gen_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_gen_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_gen_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_gen_d, "ddD*dd", "n", SM_70) BUILTIN(__nvvm_atom_add_global_i, "iiD*i", "n") TARGET_BUILTIN(__nvvm_atom_cta_add_global_i, "iiD*i", "n", SM_60) @@ -1187,6 +1235,12 @@ TARGET_BUILTIN(__nvvm_atom_sys_xchg_global_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_xchg_global_ll, "LLiLLiD*LLi", "n") TARGET_BUILTIN(__nvvm_atom_cta_xchg_global_ll, "LLiLLiD*LLi", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_sys_xchg_global_ll, "LLiLLiD*LLi", "n", SM_60) +BUILTIN(__nvvm_atom_xchg_global_f, "ffD*f", "n") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_global_f, "ffD*f", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_global_f, "ffD*f", "n", SM_60) +BUILTIN(__nvvm_atom_xchg_global_d, "ddD*d", "n") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_global_d, "ddD*d", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_global_d, "ddD*d", "n", SM_60) BUILTIN(__nvvm_atom_max_global_i, "iiD*i", "n") TARGET_BUILTIN(__nvvm_atom_cta_max_global_i, "iiD*i", "n", SM_60) @@ -1285,6 +1339,12 @@ TARGET_BUILTIN(__nvvm_atom_sys_cas_global_l, "LiLiD*LiLi", "n", SM_60) BUILTIN(__nvvm_atom_cas_global_ll, "LLiLLiD*LLiLLi", "n") TARGET_BUILTIN(__nvvm_atom_cta_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_sys_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_60) +BUILTIN(__nvvm_atom_cas_global_f, "ffD*ff", "n") +TARGET_BUILTIN(__nvvm_atom_cta_cas_global_f, "ffD*ff", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_global_f, "ffD*ff", "n", SM_60) +BUILTIN(__nvvm_atom_cas_global_d, "ddD*dd", "n") +TARGET_BUILTIN(__nvvm_atom_cta_cas_global_d, "ddD*dd", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_global_d, "ddD*dd", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_acquire_add_global_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_add_global_i, "iiD*i", "n", SM_70) @@ -1311,6 +1371,12 @@ TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_global_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_xchg_global_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_global_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_global_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_max_global_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_max_global_i, "iiD*i", "n", SM_70) @@ -1409,6 +1475,12 @@ TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_global_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cas_global_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_global_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_global_d, "ddD*dd", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_add_global_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_add_global_i, "iiD*i", "n", SM_70) @@ -1435,6 +1507,12 @@ TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_global_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_xchg_global_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_global_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_global_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_max_global_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_max_global_i, "iiD*i", "n", SM_70) @@ -1533,6 +1611,12 @@ TARGET_BUILTIN(__nvvm_atom_release_sys_cas_global_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_sys_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cas_global_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_cas_global_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_cas_global_d, "ddD*dd", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_add_global_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_add_global_i, "iiD*i", "n", SM_70) @@ -1559,6 +1643,12 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_global_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_global_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_global_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_global_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_global_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_global_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_max_global_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_max_global_i, "iiD*i", "n", SM_70) @@ -1657,6 +1747,12 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_global_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_global_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_global_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_global_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_global_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_global_d, "ddD*dd", "n", SM_70) BUILTIN(__nvvm_atom_add_shared_i, "iiD*i", "n") TARGET_BUILTIN(__nvvm_atom_cta_add_shared_i, "iiD*i", "n", SM_60) @@ -1687,6 +1783,12 @@ TARGET_BUILTIN(__nvvm_atom_sys_xchg_shared_l, "LiLiD*Li", "n", SM_60) BUILTIN(__nvvm_atom_xchg_shared_ll, "LLiLLiD*LLi", "n") TARGET_BUILTIN(__nvvm_atom_cta_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_sys_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_60) +BUILTIN(__nvvm_atom_xchg_shared_f, "ffD*f", "n") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_shared_f, "ffD*f", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_shared_f, "ffD*f", "n", SM_60) +BUILTIN(__nvvm_atom_xchg_shared_d, "ddD*d", "n") +TARGET_BUILTIN(__nvvm_atom_cta_xchg_shared_d, "ddD*d", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_xchg_shared_d, "ddD*d", "n", SM_60) BUILTIN(__nvvm_atom_max_shared_i, "iiD*i", "n") TARGET_BUILTIN(__nvvm_atom_cta_max_shared_i, "iiD*i", "n", SM_60) @@ -1785,6 +1887,12 @@ TARGET_BUILTIN(__nvvm_atom_sys_cas_shared_l, "LiLiD*LiLi", "n", SM_60) BUILTIN(__nvvm_atom_cas_shared_ll, "LLiLLiD*LLiLLi", "n") TARGET_BUILTIN(__nvvm_atom_cta_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_sys_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_60) +BUILTIN(__nvvm_atom_cas_shared_f, "ffD*ff", "n") +TARGET_BUILTIN(__nvvm_atom_cta_cas_shared_f, "ffD*ff", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_shared_f, "ffD*ff", "n", SM_60) +BUILTIN(__nvvm_atom_cas_shared_d, "ddD*dd", "n") +TARGET_BUILTIN(__nvvm_atom_cta_cas_shared_d, "ddD*dd", "n", SM_60) +TARGET_BUILTIN(__nvvm_atom_sys_cas_shared_d, "ddD*dd", "n", SM_60) TARGET_BUILTIN(__nvvm_atom_acquire_add_shared_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_add_shared_i, "iiD*i", "n", SM_70) @@ -1811,6 +1919,12 @@ TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_shared_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_xchg_shared_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_xchg_shared_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_xchg_shared_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_max_shared_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_max_shared_i, "iiD*i", "n", SM_70) @@ -1909,6 +2023,12 @@ TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_shared_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cas_shared_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_cta_cas_shared_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acquire_sys_cas_shared_d, "ddD*dd", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_add_shared_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_add_shared_i, "iiD*i", "n", SM_70) @@ -1935,6 +2055,12 @@ TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_shared_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_xchg_shared_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_xchg_shared_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_xchg_shared_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_max_shared_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_max_shared_i, "iiD*i", "n", SM_70) @@ -2033,6 +2159,12 @@ TARGET_BUILTIN(__nvvm_atom_release_sys_cas_shared_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_cta_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_release_sys_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cas_shared_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_cta_cas_shared_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_release_sys_cas_shared_d, "ddD*dd", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_add_shared_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_add_shared_i, "iiD*i", "n", SM_70) @@ -2059,6 +2191,12 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_shared_l, "LiLiD*Li", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_shared_ll, "LLiLLiD*LLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_shared_f, "ffD*f", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_xchg_shared_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_xchg_shared_d, "ddD*d", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_xchg_shared_d, "ddD*d", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_max_shared_i, "iiD*i", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_max_shared_i, "iiD*i", "n", SM_70) @@ -2157,6 +2295,12 @@ TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_shared_l, "LiLiD*LiLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_shared_ll, "LLiLLiD*LLiLLi", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_shared_f, "ffD*ff", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cas_shared_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_cta_cas_shared_d, "ddD*dd", "n", SM_70) +TARGET_BUILTIN(__nvvm_atom_acq_rel_sys_cas_shared_d, "ddD*dd", "n", SM_70) #pragma push_macro("LD_VOLATILE_BUILTIN_TYPES") #define LD_VOLATILE_BUILTIN_TYPES(ADDR_SPACE) \ diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 34b2aab1849a2..a4c910d2d1e09 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -17672,6 +17672,14 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { AtomicOrdering::SequentiallyConsistent); } + case NVPTX::BI__nvvm_atom_xchg_gen_f: + case NVPTX::BI__nvvm_atom_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f); + + case NVPTX::BI__nvvm_atom_cas_gen_f: + case NVPTX::BI__nvvm_atom_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f); + case NVPTX::BI__nvvm_atom_inc_gen_ui: { Value *Ptr = EmitScalarExpr(E->getArg(0)); Value *Val = EmitScalarExpr(E->getArg(1)); @@ -17741,10 +17749,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { 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_cta_xchg_gen_f: + case NVPTX::BI__nvvm_atom_cta_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_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_sys_xchg_gen_f: + case NVPTX::BI__nvvm_atom_sys_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_sys); case NVPTX::BI__nvvm_atom_cta_max_gen_i: case NVPTX::BI__nvvm_atom_cta_max_gen_l: case NVPTX::BI__nvvm_atom_cta_max_gen_ll: @@ -17825,6 +17839,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_sys_cas_gen_l: case NVPTX::BI__nvvm_atom_sys_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_sys); + case NVPTX::BI__nvvm_atom_cta_cas_gen_f: + case NVPTX::BI__nvvm_atom_cta_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_cta); + case NVPTX::BI__nvvm_atom_sys_cas_gen_f: + case NVPTX::BI__nvvm_atom_sys_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_sys); case NVPTX::BI__nvvm_atom_acquire_add_gen_i: case NVPTX::BI__nvvm_atom_acquire_add_gen_l: case NVPTX::BI__nvvm_atom_acquire_add_gen_ll: @@ -17836,6 +17856,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_xchg_gen_l: case NVPTX::BI__nvvm_atom_acquire_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_acquire); + case NVPTX::BI__nvvm_atom_acquire_xchg_gen_f: + case NVPTX::BI__nvvm_atom_acquire_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_acquire); case NVPTX::BI__nvvm_atom_acquire_max_gen_i: case NVPTX::BI__nvvm_atom_acquire_max_gen_l: case NVPTX::BI__nvvm_atom_acquire_max_gen_ll: @@ -17876,6 +17899,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_cas_gen_l: case NVPTX::BI__nvvm_atom_acquire_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_acquire); + case NVPTX::BI__nvvm_atom_acquire_cas_gen_f: + case NVPTX::BI__nvvm_atom_acquire_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_acquire); case NVPTX::BI__nvvm_atom_acquire_cta_add_gen_i: case NVPTX::BI__nvvm_atom_acquire_cta_add_gen_l: case NVPTX::BI__nvvm_atom_acquire_cta_add_gen_ll: @@ -17894,10 +17920,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_cta_xchg_gen_l: case NVPTX::BI__nvvm_atom_acquire_cta_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_acquire_cta); + case NVPTX::BI__nvvm_atom_acquire_cta_xchg_gen_f: + case NVPTX::BI__nvvm_atom_acquire_cta_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_acquire_cta); case NVPTX::BI__nvvm_atom_acquire_sys_xchg_gen_i: case NVPTX::BI__nvvm_atom_acquire_sys_xchg_gen_l: case NVPTX::BI__nvvm_atom_acquire_sys_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_acquire_sys); + case NVPTX::BI__nvvm_atom_acquire_sys_xchg_gen_f: + case NVPTX::BI__nvvm_atom_acquire_sys_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_acquire_sys); case NVPTX::BI__nvvm_atom_acquire_cta_max_gen_i: case NVPTX::BI__nvvm_atom_acquire_cta_max_gen_l: case NVPTX::BI__nvvm_atom_acquire_cta_max_gen_ll: @@ -17978,6 +18010,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_sys_cas_gen_l: case NVPTX::BI__nvvm_atom_acquire_sys_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_acquire_sys); + case NVPTX::BI__nvvm_atom_acquire_cta_cas_gen_f: + case NVPTX::BI__nvvm_atom_acquire_cta_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_acquire_cta); + case NVPTX::BI__nvvm_atom_acquire_sys_cas_gen_f: + case NVPTX::BI__nvvm_atom_acquire_sys_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_acquire_sys); case NVPTX::BI__nvvm_atom_release_add_gen_i: case NVPTX::BI__nvvm_atom_release_add_gen_l: case NVPTX::BI__nvvm_atom_release_add_gen_ll: @@ -17989,6 +18027,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_xchg_gen_l: case NVPTX::BI__nvvm_atom_release_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_release); + case NVPTX::BI__nvvm_atom_release_xchg_gen_f: + case NVPTX::BI__nvvm_atom_release_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_release); case NVPTX::BI__nvvm_atom_release_max_gen_i: case NVPTX::BI__nvvm_atom_release_max_gen_l: case NVPTX::BI__nvvm_atom_release_max_gen_ll: @@ -18029,6 +18070,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_cas_gen_l: case NVPTX::BI__nvvm_atom_release_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_release); + case NVPTX::BI__nvvm_atom_release_cas_gen_f: + case NVPTX::BI__nvvm_atom_release_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_release); case NVPTX::BI__nvvm_atom_release_cta_add_gen_i: case NVPTX::BI__nvvm_atom_release_cta_add_gen_l: case NVPTX::BI__nvvm_atom_release_cta_add_gen_ll: @@ -18047,10 +18091,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_cta_xchg_gen_l: case NVPTX::BI__nvvm_atom_release_cta_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_release_cta); + case NVPTX::BI__nvvm_atom_release_cta_xchg_gen_f: + case NVPTX::BI__nvvm_atom_release_cta_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_release_cta); case NVPTX::BI__nvvm_atom_release_sys_xchg_gen_i: case NVPTX::BI__nvvm_atom_release_sys_xchg_gen_l: case NVPTX::BI__nvvm_atom_release_sys_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_release_sys); + case NVPTX::BI__nvvm_atom_release_sys_xchg_gen_f: + case NVPTX::BI__nvvm_atom_release_sys_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_release_sys); case NVPTX::BI__nvvm_atom_release_cta_max_gen_i: case NVPTX::BI__nvvm_atom_release_cta_max_gen_l: case NVPTX::BI__nvvm_atom_release_cta_max_gen_ll: @@ -18131,6 +18181,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_sys_cas_gen_l: case NVPTX::BI__nvvm_atom_release_sys_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_release_sys); + case NVPTX::BI__nvvm_atom_release_cta_cas_gen_f: + case NVPTX::BI__nvvm_atom_release_cta_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_release_cta); + case NVPTX::BI__nvvm_atom_release_sys_cas_gen_f: + case NVPTX::BI__nvvm_atom_release_sys_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_release_sys); case NVPTX::BI__nvvm_atom_acq_rel_add_gen_i: case NVPTX::BI__nvvm_atom_acq_rel_add_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_add_gen_ll: @@ -18142,6 +18198,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_xchg_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_acq_rel); + case NVPTX::BI__nvvm_atom_acq_rel_xchg_gen_f: + case NVPTX::BI__nvvm_atom_acq_rel_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_acq_rel); case NVPTX::BI__nvvm_atom_acq_rel_max_gen_i: case NVPTX::BI__nvvm_atom_acq_rel_max_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_max_gen_ll: @@ -18182,6 +18241,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_cas_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_acq_rel); + case NVPTX::BI__nvvm_atom_acq_rel_cas_gen_f: + case NVPTX::BI__nvvm_atom_acq_rel_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_acq_rel); case NVPTX::BI__nvvm_atom_acq_rel_cta_add_gen_i: case NVPTX::BI__nvvm_atom_acq_rel_cta_add_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_add_gen_ll: @@ -18200,10 +18262,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_acq_rel_cta); + case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_gen_f: + case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_acq_rel_cta); case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_gen_i: case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_gen_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_i_acq_rel_sys); + case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_gen_f: + case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_gen_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_gen_f_acq_rel_sys); case NVPTX::BI__nvvm_atom_acq_rel_cta_max_gen_i: case NVPTX::BI__nvvm_atom_acq_rel_cta_max_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_max_gen_ll: @@ -18284,6 +18352,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_gen_l: case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_gen_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_i_acq_rel_sys); + case NVPTX::BI__nvvm_atom_acq_rel_cta_cas_gen_f: + case NVPTX::BI__nvvm_atom_acq_rel_cta_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_acq_rel_cta); + case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_gen_f: + case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_gen_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_gen_f_acq_rel_sys); case NVPTX::BI__nvvm_atom_add_global_i: case NVPTX::BI__nvvm_atom_add_global_l: case NVPTX::BI__nvvm_atom_add_global_ll: @@ -18295,6 +18369,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_xchg_global_l: case NVPTX::BI__nvvm_atom_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i); + case NVPTX::BI__nvvm_atom_xchg_global_f: + case NVPTX::BI__nvvm_atom_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f); case NVPTX::BI__nvvm_atom_max_global_i: case NVPTX::BI__nvvm_atom_max_global_l: case NVPTX::BI__nvvm_atom_max_global_ll: @@ -18335,6 +18412,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_cas_global_l: case NVPTX::BI__nvvm_atom_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i); + case NVPTX::BI__nvvm_atom_cas_global_f: + case NVPTX::BI__nvvm_atom_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f); case NVPTX::BI__nvvm_atom_cta_add_global_i: case NVPTX::BI__nvvm_atom_cta_add_global_l: case NVPTX::BI__nvvm_atom_cta_add_global_ll: @@ -18353,10 +18433,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_cta_xchg_global_l: case NVPTX::BI__nvvm_atom_cta_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_cta); + case NVPTX::BI__nvvm_atom_cta_xchg_global_f: + case NVPTX::BI__nvvm_atom_cta_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_cta); case NVPTX::BI__nvvm_atom_sys_xchg_global_i: case NVPTX::BI__nvvm_atom_sys_xchg_global_l: case NVPTX::BI__nvvm_atom_sys_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_sys); + case NVPTX::BI__nvvm_atom_sys_xchg_global_f: + case NVPTX::BI__nvvm_atom_sys_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_sys); case NVPTX::BI__nvvm_atom_cta_max_global_i: case NVPTX::BI__nvvm_atom_cta_max_global_l: case NVPTX::BI__nvvm_atom_cta_max_global_ll: @@ -18437,6 +18523,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_sys_cas_global_l: case NVPTX::BI__nvvm_atom_sys_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_sys); + case NVPTX::BI__nvvm_atom_cta_cas_global_f: + case NVPTX::BI__nvvm_atom_cta_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_cta); + case NVPTX::BI__nvvm_atom_sys_cas_global_f: + case NVPTX::BI__nvvm_atom_sys_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_sys); case NVPTX::BI__nvvm_atom_acquire_add_global_i: case NVPTX::BI__nvvm_atom_acquire_add_global_l: case NVPTX::BI__nvvm_atom_acquire_add_global_ll: @@ -18448,6 +18540,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_xchg_global_l: case NVPTX::BI__nvvm_atom_acquire_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_acquire); + case NVPTX::BI__nvvm_atom_acquire_xchg_global_f: + case NVPTX::BI__nvvm_atom_acquire_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_acquire); case NVPTX::BI__nvvm_atom_acquire_max_global_i: case NVPTX::BI__nvvm_atom_acquire_max_global_l: case NVPTX::BI__nvvm_atom_acquire_max_global_ll: @@ -18488,6 +18583,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_cas_global_l: case NVPTX::BI__nvvm_atom_acquire_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_acquire); + case NVPTX::BI__nvvm_atom_acquire_cas_global_f: + case NVPTX::BI__nvvm_atom_acquire_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_acquire); case NVPTX::BI__nvvm_atom_acquire_cta_add_global_i: case NVPTX::BI__nvvm_atom_acquire_cta_add_global_l: case NVPTX::BI__nvvm_atom_acquire_cta_add_global_ll: @@ -18506,10 +18604,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_cta_xchg_global_l: case NVPTX::BI__nvvm_atom_acquire_cta_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_acquire_cta); + case NVPTX::BI__nvvm_atom_acquire_cta_xchg_global_f: + case NVPTX::BI__nvvm_atom_acquire_cta_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_acquire_cta); case NVPTX::BI__nvvm_atom_acquire_sys_xchg_global_i: case NVPTX::BI__nvvm_atom_acquire_sys_xchg_global_l: case NVPTX::BI__nvvm_atom_acquire_sys_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_acquire_sys); + case NVPTX::BI__nvvm_atom_acquire_sys_xchg_global_f: + case NVPTX::BI__nvvm_atom_acquire_sys_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_acquire_sys); case NVPTX::BI__nvvm_atom_acquire_cta_max_global_i: case NVPTX::BI__nvvm_atom_acquire_cta_max_global_l: case NVPTX::BI__nvvm_atom_acquire_cta_max_global_ll: @@ -18590,6 +18694,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_sys_cas_global_l: case NVPTX::BI__nvvm_atom_acquire_sys_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_acquire_sys); + case NVPTX::BI__nvvm_atom_acquire_cta_cas_global_f: + case NVPTX::BI__nvvm_atom_acquire_cta_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_acquire_cta); + case NVPTX::BI__nvvm_atom_acquire_sys_cas_global_f: + case NVPTX::BI__nvvm_atom_acquire_sys_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_acquire_sys); case NVPTX::BI__nvvm_atom_release_add_global_i: case NVPTX::BI__nvvm_atom_release_add_global_l: case NVPTX::BI__nvvm_atom_release_add_global_ll: @@ -18601,6 +18711,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_xchg_global_l: case NVPTX::BI__nvvm_atom_release_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_release); + case NVPTX::BI__nvvm_atom_release_xchg_global_f: + case NVPTX::BI__nvvm_atom_release_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_release); case NVPTX::BI__nvvm_atom_release_max_global_i: case NVPTX::BI__nvvm_atom_release_max_global_l: case NVPTX::BI__nvvm_atom_release_max_global_ll: @@ -18641,6 +18754,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_cas_global_l: case NVPTX::BI__nvvm_atom_release_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_release); + case NVPTX::BI__nvvm_atom_release_cas_global_f: + case NVPTX::BI__nvvm_atom_release_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_release); case NVPTX::BI__nvvm_atom_release_cta_add_global_i: case NVPTX::BI__nvvm_atom_release_cta_add_global_l: case NVPTX::BI__nvvm_atom_release_cta_add_global_ll: @@ -18659,10 +18775,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_cta_xchg_global_l: case NVPTX::BI__nvvm_atom_release_cta_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_release_cta); + case NVPTX::BI__nvvm_atom_release_cta_xchg_global_f: + case NVPTX::BI__nvvm_atom_release_cta_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_release_cta); case NVPTX::BI__nvvm_atom_release_sys_xchg_global_i: case NVPTX::BI__nvvm_atom_release_sys_xchg_global_l: case NVPTX::BI__nvvm_atom_release_sys_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_release_sys); + case NVPTX::BI__nvvm_atom_release_sys_xchg_global_f: + case NVPTX::BI__nvvm_atom_release_sys_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_release_sys); case NVPTX::BI__nvvm_atom_release_cta_max_global_i: case NVPTX::BI__nvvm_atom_release_cta_max_global_l: case NVPTX::BI__nvvm_atom_release_cta_max_global_ll: @@ -18743,6 +18865,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_sys_cas_global_l: case NVPTX::BI__nvvm_atom_release_sys_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_release_sys); + case NVPTX::BI__nvvm_atom_release_cta_cas_global_f: + case NVPTX::BI__nvvm_atom_release_cta_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_release_cta); + case NVPTX::BI__nvvm_atom_release_sys_cas_global_f: + case NVPTX::BI__nvvm_atom_release_sys_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_release_sys); case NVPTX::BI__nvvm_atom_acq_rel_add_global_i: case NVPTX::BI__nvvm_atom_acq_rel_add_global_l: case NVPTX::BI__nvvm_atom_acq_rel_add_global_ll: @@ -18754,6 +18882,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_xchg_global_l: case NVPTX::BI__nvvm_atom_acq_rel_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_acq_rel); + case NVPTX::BI__nvvm_atom_acq_rel_xchg_global_f: + case NVPTX::BI__nvvm_atom_acq_rel_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_acq_rel); case NVPTX::BI__nvvm_atom_acq_rel_max_global_i: case NVPTX::BI__nvvm_atom_acq_rel_max_global_l: case NVPTX::BI__nvvm_atom_acq_rel_max_global_ll: @@ -18794,6 +18925,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_cas_global_l: case NVPTX::BI__nvvm_atom_acq_rel_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_acq_rel); + case NVPTX::BI__nvvm_atom_acq_rel_cas_global_f: + case NVPTX::BI__nvvm_atom_acq_rel_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_acq_rel); case NVPTX::BI__nvvm_atom_acq_rel_cta_add_global_i: case NVPTX::BI__nvvm_atom_acq_rel_cta_add_global_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_add_global_ll: @@ -18812,10 +18946,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_global_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_acq_rel_cta); + case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_global_f: + case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_acq_rel_cta); case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_global_i: case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_global_l: case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_global_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_i_acq_rel_sys); + case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_global_f: + case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_global_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_global_f_acq_rel_sys); case NVPTX::BI__nvvm_atom_acq_rel_cta_max_global_i: case NVPTX::BI__nvvm_atom_acq_rel_cta_max_global_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_max_global_ll: @@ -18896,6 +19036,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_global_l: case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_global_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_i_acq_rel_sys); + case NVPTX::BI__nvvm_atom_acq_rel_cta_cas_global_f: + case NVPTX::BI__nvvm_atom_acq_rel_cta_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_acq_rel_cta); + case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_global_f: + case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_global_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_global_f_acq_rel_sys); case NVPTX::BI__nvvm_atom_add_shared_i: case NVPTX::BI__nvvm_atom_add_shared_l: case NVPTX::BI__nvvm_atom_add_shared_ll: @@ -18907,6 +19053,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_xchg_shared_l: case NVPTX::BI__nvvm_atom_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i); + case NVPTX::BI__nvvm_atom_xchg_shared_f: + case NVPTX::BI__nvvm_atom_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f); case NVPTX::BI__nvvm_atom_max_shared_i: case NVPTX::BI__nvvm_atom_max_shared_l: case NVPTX::BI__nvvm_atom_max_shared_ll: @@ -18947,6 +19096,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_cas_shared_l: case NVPTX::BI__nvvm_atom_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i); + case NVPTX::BI__nvvm_atom_cas_shared_f: + case NVPTX::BI__nvvm_atom_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f); case NVPTX::BI__nvvm_atom_cta_add_shared_i: case NVPTX::BI__nvvm_atom_cta_add_shared_l: case NVPTX::BI__nvvm_atom_cta_add_shared_ll: @@ -18965,10 +19117,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_cta_xchg_shared_l: case NVPTX::BI__nvvm_atom_cta_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_cta); + case NVPTX::BI__nvvm_atom_cta_xchg_shared_f: + case NVPTX::BI__nvvm_atom_cta_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_cta); case NVPTX::BI__nvvm_atom_sys_xchg_shared_i: case NVPTX::BI__nvvm_atom_sys_xchg_shared_l: case NVPTX::BI__nvvm_atom_sys_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_sys); + case NVPTX::BI__nvvm_atom_sys_xchg_shared_f: + case NVPTX::BI__nvvm_atom_sys_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_sys); case NVPTX::BI__nvvm_atom_cta_max_shared_i: case NVPTX::BI__nvvm_atom_cta_max_shared_l: case NVPTX::BI__nvvm_atom_cta_max_shared_ll: @@ -19049,6 +19207,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_sys_cas_shared_l: case NVPTX::BI__nvvm_atom_sys_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_sys); + case NVPTX::BI__nvvm_atom_cta_cas_shared_f: + case NVPTX::BI__nvvm_atom_cta_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_cta); + case NVPTX::BI__nvvm_atom_sys_cas_shared_f: + case NVPTX::BI__nvvm_atom_sys_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_sys); case NVPTX::BI__nvvm_atom_acquire_add_shared_i: case NVPTX::BI__nvvm_atom_acquire_add_shared_l: case NVPTX::BI__nvvm_atom_acquire_add_shared_ll: @@ -19060,6 +19224,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_xchg_shared_l: case NVPTX::BI__nvvm_atom_acquire_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_acquire); + case NVPTX::BI__nvvm_atom_acquire_xchg_shared_f: + case NVPTX::BI__nvvm_atom_acquire_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_acquire); case NVPTX::BI__nvvm_atom_acquire_max_shared_i: case NVPTX::BI__nvvm_atom_acquire_max_shared_l: case NVPTX::BI__nvvm_atom_acquire_max_shared_ll: @@ -19100,6 +19267,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_cas_shared_l: case NVPTX::BI__nvvm_atom_acquire_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_acquire); + case NVPTX::BI__nvvm_atom_acquire_cas_shared_f: + case NVPTX::BI__nvvm_atom_acquire_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_acquire); case NVPTX::BI__nvvm_atom_acquire_cta_add_shared_i: case NVPTX::BI__nvvm_atom_acquire_cta_add_shared_l: case NVPTX::BI__nvvm_atom_acquire_cta_add_shared_ll: @@ -19118,10 +19288,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_cta_xchg_shared_l: case NVPTX::BI__nvvm_atom_acquire_cta_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_acquire_cta); + case NVPTX::BI__nvvm_atom_acquire_cta_xchg_shared_f: + case NVPTX::BI__nvvm_atom_acquire_cta_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_acquire_cta); case NVPTX::BI__nvvm_atom_acquire_sys_xchg_shared_i: case NVPTX::BI__nvvm_atom_acquire_sys_xchg_shared_l: case NVPTX::BI__nvvm_atom_acquire_sys_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_acquire_sys); + case NVPTX::BI__nvvm_atom_acquire_sys_xchg_shared_f: + case NVPTX::BI__nvvm_atom_acquire_sys_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_acquire_sys); case NVPTX::BI__nvvm_atom_acquire_cta_max_shared_i: case NVPTX::BI__nvvm_atom_acquire_cta_max_shared_l: case NVPTX::BI__nvvm_atom_acquire_cta_max_shared_ll: @@ -19202,6 +19378,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acquire_sys_cas_shared_l: case NVPTX::BI__nvvm_atom_acquire_sys_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_acquire_sys); + case NVPTX::BI__nvvm_atom_acquire_cta_cas_shared_f: + case NVPTX::BI__nvvm_atom_acquire_cta_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_acquire_cta); + case NVPTX::BI__nvvm_atom_acquire_sys_cas_shared_f: + case NVPTX::BI__nvvm_atom_acquire_sys_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_acquire_sys); case NVPTX::BI__nvvm_atom_release_add_shared_i: case NVPTX::BI__nvvm_atom_release_add_shared_l: case NVPTX::BI__nvvm_atom_release_add_shared_ll: @@ -19213,6 +19395,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_xchg_shared_l: case NVPTX::BI__nvvm_atom_release_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_release); + case NVPTX::BI__nvvm_atom_release_xchg_shared_f: + case NVPTX::BI__nvvm_atom_release_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_release); case NVPTX::BI__nvvm_atom_release_max_shared_i: case NVPTX::BI__nvvm_atom_release_max_shared_l: case NVPTX::BI__nvvm_atom_release_max_shared_ll: @@ -19253,6 +19438,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_cas_shared_l: case NVPTX::BI__nvvm_atom_release_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_release); + case NVPTX::BI__nvvm_atom_release_cas_shared_f: + case NVPTX::BI__nvvm_atom_release_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_release); case NVPTX::BI__nvvm_atom_release_cta_add_shared_i: case NVPTX::BI__nvvm_atom_release_cta_add_shared_l: case NVPTX::BI__nvvm_atom_release_cta_add_shared_ll: @@ -19271,10 +19459,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_cta_xchg_shared_l: case NVPTX::BI__nvvm_atom_release_cta_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_release_cta); + case NVPTX::BI__nvvm_atom_release_cta_xchg_shared_f: + case NVPTX::BI__nvvm_atom_release_cta_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_release_cta); case NVPTX::BI__nvvm_atom_release_sys_xchg_shared_i: case NVPTX::BI__nvvm_atom_release_sys_xchg_shared_l: case NVPTX::BI__nvvm_atom_release_sys_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_release_sys); + case NVPTX::BI__nvvm_atom_release_sys_xchg_shared_f: + case NVPTX::BI__nvvm_atom_release_sys_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_release_sys); case NVPTX::BI__nvvm_atom_release_cta_max_shared_i: case NVPTX::BI__nvvm_atom_release_cta_max_shared_l: case NVPTX::BI__nvvm_atom_release_cta_max_shared_ll: @@ -19355,6 +19549,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_release_sys_cas_shared_l: case NVPTX::BI__nvvm_atom_release_sys_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_release_sys); + case NVPTX::BI__nvvm_atom_release_cta_cas_shared_f: + case NVPTX::BI__nvvm_atom_release_cta_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_release_cta); + case NVPTX::BI__nvvm_atom_release_sys_cas_shared_f: + case NVPTX::BI__nvvm_atom_release_sys_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_release_sys); case NVPTX::BI__nvvm_atom_acq_rel_add_shared_i: case NVPTX::BI__nvvm_atom_acq_rel_add_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_add_shared_ll: @@ -19366,6 +19566,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_xchg_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_acq_rel); + case NVPTX::BI__nvvm_atom_acq_rel_xchg_shared_f: + case NVPTX::BI__nvvm_atom_acq_rel_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_acq_rel); case NVPTX::BI__nvvm_atom_acq_rel_max_shared_i: case NVPTX::BI__nvvm_atom_acq_rel_max_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_max_shared_ll: @@ -19406,6 +19609,9 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_cas_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_acq_rel); + case NVPTX::BI__nvvm_atom_acq_rel_cas_shared_f: + case NVPTX::BI__nvvm_atom_acq_rel_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_acq_rel); case NVPTX::BI__nvvm_atom_acq_rel_cta_add_shared_i: case NVPTX::BI__nvvm_atom_acq_rel_cta_add_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_add_shared_ll: @@ -19424,10 +19630,16 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_acq_rel_cta); + case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_shared_f: + case NVPTX::BI__nvvm_atom_acq_rel_cta_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_acq_rel_cta); case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_shared_i: case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_shared_ll: return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_i_acq_rel_sys); + case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_shared_f: + case NVPTX::BI__nvvm_atom_acq_rel_sys_xchg_shared_d: + return MakeScopedAtomic(Intrinsic::nvvm_atomic_exch_shared_f_acq_rel_sys); case NVPTX::BI__nvvm_atom_acq_rel_cta_max_shared_i: case NVPTX::BI__nvvm_atom_acq_rel_cta_max_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_cta_max_shared_ll: @@ -19508,6 +19720,12 @@ CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) { case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_shared_l: case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_shared_ll: return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_i_acq_rel_sys); + case NVPTX::BI__nvvm_atom_acq_rel_cta_cas_shared_f: + case NVPTX::BI__nvvm_atom_acq_rel_cta_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_acq_rel_cta); + case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_shared_f: + case NVPTX::BI__nvvm_atom_acq_rel_sys_cas_shared_d: + return MakeScopedCasAtomic(Intrinsic::nvvm_atomic_cas_shared_f_acq_rel_sys); case NVPTX::BI__nvvm_match_all_sync_i32p: case NVPTX::BI__nvvm_match_all_sync_i64p: { Value *Mask = EmitScalarExpr(E->getArg(0)); diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index 368974095fad7..39c3e134f649b 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -254,6 +254,10 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, __nvvm_atom_xchg_gen_l(&dl, l); // CHECK: atomicrmw xchg i64* {{.*}} seq_cst, align 8 __nvvm_atom_xchg_gen_ll(&sll, ll); + // CHECK: call float @llvm.nvvm.atomic.exch.gen.f.f32.p0f32 + __nvvm_atom_xchg_gen_f(fp, f); + // CHECK: call double @llvm.nvvm.atomic.exch.gen.f.f64.p0f64 + __nvvm_atom_xchg_gen_d(dfp, df); // CHECK: atomicrmw max i32* {{.*}} seq_cst, align 4 __nvvm_atom_max_gen_i(ip, i); @@ -290,6 +294,10 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, // CHECK: cmpxchg i64* {{.*}} seq_cst seq_cst, align 8 // CHECK-NEXT: extractvalue { i64, i1 } {{%[0-9]+}}, 0 __nvvm_atom_cas_gen_ll(&sll, 0, ll); + // CHECK: call float @llvm.nvvm.atomic.cas.gen.f.f32.p0f32 + __nvvm_atom_cas_gen_f(fp, 0, f); + // CHECK: call double @llvm.nvvm.atomic.cas.gen.f.f64.p0f64 + __nvvm_atom_cas_gen_d(dfp, 0, df); // CHECK: atomicrmw fadd float* {{.*}} seq_cst, align 4 __nvvm_atom_add_gen_f(fp, f); @@ -425,6 +433,12 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64 // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature sm_60}} __nvvm_atom_cta_xchg_gen_ll(&sll, ll); + // CHECK: call float @llvm.nvvm.atomic.exch.gen.f.cta.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_f' needs target feature sm_60}} + __nvvm_atom_cta_xchg_gen_f(fp, f); + // CHECK: call double @llvm.nvvm.atomic.exch.gen.f.cta.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_d' needs target feature sm_60}} + __nvvm_atom_cta_xchg_gen_d(dfp, df); // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32 // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature sm_60}} @@ -436,6 +450,12 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64 // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature sm_60}} __nvvm_atom_sys_xchg_gen_ll(&sll, ll); + // CHECK: call float @llvm.nvvm.atomic.exch.gen.f.sys.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_f' needs target feature sm_60}} + __nvvm_atom_sys_xchg_gen_f(fp, f); + // CHECK: call double @llvm.nvvm.atomic.exch.gen.f.sys.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_d' needs target feature sm_60}} + __nvvm_atom_sys_xchg_gen_d(dfp, df); // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature sm_60}} @@ -1827,6 +1847,222 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_shared_l' needs target feature sm_70}} __nvvm_atom_acq_rel_cta_xchg_shared_l((__attribute__((address_space(3))) long *)&dl, l); + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.acquire.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_acquire_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_acquire_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.acquire.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_acquire_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_acquire_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.acquire.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_acquire_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_acquire_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.release.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_release_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_release_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.release.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_release_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_release_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.release.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_release_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_release_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.acq.rel.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.acq.rel.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.acq.rel.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.acquire.sys.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_acquire_sys_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_acquire_sys_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.acquire.sys.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_acquire_sys_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_acquire_sys_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.acquire.sys.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_acquire_sys_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_acquire_sys_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.release.sys.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_release_sys_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_release_sys_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.release.sys.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_release_sys_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_release_sys_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.release.sys.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_release_sys_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_release_sys_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.acq.rel.sys.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_sys_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_sys_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.acq.rel.sys.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_sys_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_sys_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.acq.rel.sys.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_sys_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_sys_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.acquire.cta.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_acquire_cta_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_acquire_cta_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.acquire.cta.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_acquire_cta_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_acquire_cta_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.acquire.cta.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_acquire_cta_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_acquire_cta_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.release.cta.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_release_cta_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_release_cta_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.release.cta.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_release_cta_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_release_cta_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.release.cta.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_release_cta_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_release_cta_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.gen.f.acq.rel.cta.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_gen_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_cta_xchg_gen_f(fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.global.f.acq.rel.cta.f32.p1f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_global_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_cta_xchg_global_f((__attribute__((address_space(1))) float *)fp, f); + + // CHECK_SM70_LP64: call float @llvm.nvvm.atomic.exch.shared.f.acq.rel.cta.f32.p3f32 + // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_shared_f' needs target feature sm_70}} + __nvvm_atom_acq_rel_cta_xchg_shared_f((__attribute__((address_space(3))) float *)fp, f); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.acquire.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_acquire_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_acquire_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.acquire.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_acquire_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_acquire_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.acquire.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_acquire_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_acquire_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.release.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_release_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_release_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.release.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_release_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_release_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.release.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_release_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_release_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.acq.rel.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.acq.rel.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.acq.rel.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.acquire.sys.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_acquire_sys_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_acquire_sys_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.acquire.sys.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_acquire_sys_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_acquire_sys_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.acquire.sys.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_acquire_sys_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_acquire_sys_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.release.sys.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_release_sys_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_release_sys_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.release.sys.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_release_sys_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_release_sys_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.release.sys.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_release_sys_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_release_sys_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.acq.rel.sys.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_sys_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_sys_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.acq.rel.sys.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_sys_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_sys_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.acq.rel.sys.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_sys_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_sys_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.acquire.cta.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_acquire_cta_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_acquire_cta_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.acquire.cta.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_acquire_cta_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_acquire_cta_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.acquire.cta.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_acquire_cta_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_acquire_cta_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.release.cta.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_release_cta_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_release_cta_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.release.cta.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_release_cta_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_release_cta_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.release.cta.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_release_cta_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_release_cta_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.gen.f.acq.rel.cta.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_gen_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_cta_xchg_gen_d(dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.global.f.acq.rel.cta.f64.p1f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_global_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_cta_xchg_global_d((__attribute__((address_space(1))) double *)dfp, df); + + // CHECK_SM70_LP64: call double @llvm.nvvm.atomic.exch.shared.f.acq.rel.cta.f64.p3f64 + // expected-error@+1 {{'__nvvm_atom_acq_rel_cta_xchg_shared_d' needs target feature sm_70}} + __nvvm_atom_acq_rel_cta_xchg_shared_d((__attribute__((address_space(3))) double *)dfp, df); + // CHECK_SM70_LP64: call i32 @llvm.nvvm.atomic.max.gen.i.acquire.i32.p0i32 // expected-error@+1 {{'__nvvm_atom_acquire_max_gen_i' needs target feature sm_70}} __nvvm_atom_acquire_max_gen_i(ip, i); diff --git a/libclc/ptx-nvidiacl/libspirv/atomic/atomic_cmpxchg.cl b/libclc/ptx-nvidiacl/libspirv/atomic/atomic_cmpxchg.cl index 5d9466bd1e5a4..88d72a856ed94 100644 --- a/libclc/ptx-nvidiacl/libspirv/atomic/atomic_cmpxchg.cl +++ b/libclc/ptx-nvidiacl/libspirv/atomic/atomic_cmpxchg.cl @@ -87,6 +87,8 @@ __CLC_NVVM_ATOMIC_CAS(int, i, int, i, cas, CompareExchange) __CLC_NVVM_ATOMIC_CAS(long, l, long, l, cas, CompareExchange) __CLC_NVVM_ATOMIC_CAS(unsigned int, j, int, i, cas, CompareExchange) __CLC_NVVM_ATOMIC_CAS(unsigned long, m, long, l, cas, CompareExchange) +__CLC_NVVM_ATOMIC_CAS(float, f, float, f, cas, CompareExchange) +__CLC_NVVM_ATOMIC_CAS(double, d, double, d, cas, CompareExchange) #undef __CLC_NVVM_ATOMIC_CAS_IMPL_ORDER #undef __CLC_NVVM_ATOMIC_CAS diff --git a/libclc/ptx-nvidiacl/libspirv/atomic/atomic_xchg.cl b/libclc/ptx-nvidiacl/libspirv/atomic/atomic_xchg.cl index e423b73c380c0..361d7dbb4a948 100644 --- a/libclc/ptx-nvidiacl/libspirv/atomic/atomic_xchg.cl +++ b/libclc/ptx-nvidiacl/libspirv/atomic/atomic_xchg.cl @@ -14,6 +14,8 @@ __CLC_NVVM_ATOMIC(int, i, int, i, xchg, _Z22__spirv_AtomicExchange) __CLC_NVVM_ATOMIC(long, l, long, l, xchg, _Z22__spirv_AtomicExchange) __CLC_NVVM_ATOMIC(unsigned int, j, int, i, xchg, _Z22__spirv_AtomicExchange) __CLC_NVVM_ATOMIC(unsigned long, m, long, l, xchg, _Z22__spirv_AtomicExchange) +__CLC_NVVM_ATOMIC(float, f, float, f, xchg, _Z22__spirv_AtomicExchange) +__CLC_NVVM_ATOMIC(double, d, double, d, xchg, _Z22__spirv_AtomicExchange) #undef __CLC_NVVM_ATOMIC_TYPES #undef __CLC_NVVM_ATOMIC diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index df5ea22e9b339..3389ea27de122 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -1375,12 +1375,16 @@ let TargetPrefix = "nvvm" in { defm _f: PTXAtomicWithScope2_sem; defm _i: PTXAtomicWithScope2_sem; } + multiclass PTXAtomicWithScope3_fi { + defm _f: PTXAtomicWithScope3_sem; + defm _i: PTXAtomicWithScope3_sem; + } defm int_nvvm_ld_gen : PTXLdWithScope_fi; defm int_nvvm_st_gen : PTXStWithScope_fi; defm int_nvvm_atomic_add_gen : PTXAtomicWithScope2_fi; defm int_nvvm_atomic_inc_gen_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_dec_gen_i : PTXAtomicWithScope2_sem; - defm int_nvvm_atomic_exch_gen_i: PTXAtomicWithScope2_sem; + defm int_nvvm_atomic_exch_gen: PTXAtomicWithScope2_fi; defm int_nvvm_atomic_xor_gen_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_max_gen_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_max_gen_ui: PTXAtomicWithScope2_sem; @@ -1388,14 +1392,14 @@ let TargetPrefix = "nvvm" in { defm int_nvvm_atomic_min_gen_ui: PTXAtomicWithScope2_sem; defm int_nvvm_atomic_or_gen_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_and_gen_i : PTXAtomicWithScope2_sem; - defm int_nvvm_atomic_cas_gen_i : PTXAtomicWithScope3_sem; + defm int_nvvm_atomic_cas_gen : PTXAtomicWithScope3_fi; defm int_nvvm_ld_shared : PTXLdWithScope_fi; defm int_nvvm_st_shared : PTXStWithScope_fi; defm int_nvvm_atomic_add_shared : PTXAtomicWithScope2_fi; defm int_nvvm_atomic_inc_shared_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_dec_shared_i : PTXAtomicWithScope2_sem; - defm int_nvvm_atomic_exch_shared_i: PTXAtomicWithScope2_sem; + defm int_nvvm_atomic_exch_shared: PTXAtomicWithScope2_fi; defm int_nvvm_atomic_xor_shared_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_max_shared_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_max_shared_ui: PTXAtomicWithScope2_sem; @@ -1403,14 +1407,14 @@ let TargetPrefix = "nvvm" in { defm int_nvvm_atomic_min_shared_ui: PTXAtomicWithScope2_sem; defm int_nvvm_atomic_or_shared_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_and_shared_i : PTXAtomicWithScope2_sem; - defm int_nvvm_atomic_cas_shared_i : PTXAtomicWithScope3_sem; + defm int_nvvm_atomic_cas_shared : PTXAtomicWithScope3_fi; defm int_nvvm_ld_global : PTXLdWithScope_fi; defm int_nvvm_st_global : PTXStWithScope_fi; defm int_nvvm_atomic_add_global : PTXAtomicWithScope2_fi; defm int_nvvm_atomic_inc_global_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_dec_global_i : PTXAtomicWithScope2_sem; - defm int_nvvm_atomic_exch_global_i: PTXAtomicWithScope2_sem; + defm int_nvvm_atomic_exch_global: PTXAtomicWithScope2_fi; defm int_nvvm_atomic_xor_global_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_max_global_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_max_global_ui: PTXAtomicWithScope2_sem; @@ -1418,7 +1422,7 @@ let TargetPrefix = "nvvm" in { defm int_nvvm_atomic_min_global_ui: PTXAtomicWithScope2_sem; defm int_nvvm_atomic_or_global_i : PTXAtomicWithScope2_sem; defm int_nvvm_atomic_and_global_i : PTXAtomicWithScope2_sem; - defm int_nvvm_atomic_cas_global_i : PTXAtomicWithScope3_sem; + defm int_nvvm_atomic_cas_global : PTXAtomicWithScope3_fi; // Bar.Sync diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index eaee35746eac1..0eb994ea8e35d 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -2264,6 +2264,8 @@ multiclass ATOM2_bitwise_impl { multiclass ATOM2_exch_impl { defm _b32 : ATOM2S_impl; defm _b64 : ATOM2S_impl; + defm _f32 : ATOM2S_impl; + defm _f64 : ATOM2S_impl; } // atom.{min,max} @@ -2286,6 +2288,8 @@ multiclass ATOM2_incdec_impl { multiclass ATOM3_cas_impl { defm _b32 : ATOM3S_impl; defm _b64 : ATOM3S_impl; + defm _f32 : ATOM3S_impl; + defm _f64 : ATOM3S_impl; } defm INT_PTX_LD : ATOM_ld_impl; diff --git a/llvm/test/CodeGen/NVPTX/atomics-with-semantics.ll b/llvm/test/CodeGen/NVPTX/atomics-with-semantics.ll index 2f565b71db79d..2c2fd2ff22de1 100644 --- a/llvm/test/CodeGen/NVPTX/atomics-with-semantics.ll +++ b/llvm/test/CodeGen/NVPTX/atomics-with-semantics.ll @@ -506,6 +506,168 @@ entry: ; CHECK: atom.acq_rel.cta.shared.exch.b64 %tmp161 = tail call i64 @llvm.nvvm.atomic.exch.shared.i.acq.rel.cta.i64.p3i64(i64 addrspace(3)* %llp3, i64 %ll); + ; CHECK: atom.acquire.exch.b32 + %tmp708 = tail call float @llvm.nvvm.atomic.exch.gen.f.acquire.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.acquire.global.exch.b32 + %tmp709 = tail call float @llvm.nvvm.atomic.exch.global.f.acquire.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.acquire.shared.exch.b32 + %tmp710 = tail call float @llvm.nvvm.atomic.exch.shared.f.acquire.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.release.exch.b32 + %tmp711 = tail call float @llvm.nvvm.atomic.exch.gen.f.release.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.release.global.exch.b32 + %tmp712 = tail call float @llvm.nvvm.atomic.exch.global.f.release.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.release.shared.exch.b32 + %tmp713 = tail call float @llvm.nvvm.atomic.exch.shared.f.release.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.acq_rel.exch.b32 + %tmp714 = tail call float @llvm.nvvm.atomic.exch.gen.f.acq.rel.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.acq_rel.global.exch.b32 + %tmp715 = tail call float @llvm.nvvm.atomic.exch.global.f.acq.rel.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.acq_rel.shared.exch.b32 + %tmp716 = tail call float @llvm.nvvm.atomic.exch.shared.f.acq.rel.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.acquire.sys.exch.b32 + %tmp717 = tail call float @llvm.nvvm.atomic.exch.gen.f.acquire.sys.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.acquire.sys.global.exch.b32 + %tmp718 = tail call float @llvm.nvvm.atomic.exch.global.f.acquire.sys.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.acquire.sys.shared.exch.b32 + %tmp719 = tail call float @llvm.nvvm.atomic.exch.shared.f.acquire.sys.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.release.sys.exch.b32 + %tmp720 = tail call float @llvm.nvvm.atomic.exch.gen.f.release.sys.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.release.sys.global.exch.b32 + %tmp721 = tail call float @llvm.nvvm.atomic.exch.global.f.release.sys.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.release.sys.shared.exch.b32 + %tmp722 = tail call float @llvm.nvvm.atomic.exch.shared.f.release.sys.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.acq_rel.sys.exch.b32 + %tmp723 = tail call float @llvm.nvvm.atomic.exch.gen.f.acq.rel.sys.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.acq_rel.sys.global.exch.b32 + %tmp724 = tail call float @llvm.nvvm.atomic.exch.global.f.acq.rel.sys.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.acq_rel.sys.shared.exch.b32 + %tmp725 = tail call float @llvm.nvvm.atomic.exch.shared.f.acq.rel.sys.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.acquire.cta.exch.b32 + %tmp726 = tail call float @llvm.nvvm.atomic.exch.gen.f.acquire.cta.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.acquire.cta.global.exch.b32 + %tmp727 = tail call float @llvm.nvvm.atomic.exch.global.f.acquire.cta.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.acquire.cta.shared.exch.b32 + %tmp728 = tail call float @llvm.nvvm.atomic.exch.shared.f.acquire.cta.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.release.cta.exch.b32 + %tmp729 = tail call float @llvm.nvvm.atomic.exch.gen.f.release.cta.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.release.cta.global.exch.b32 + %tmp730 = tail call float @llvm.nvvm.atomic.exch.global.f.release.cta.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.release.cta.shared.exch.b32 + %tmp731 = tail call float @llvm.nvvm.atomic.exch.shared.f.release.cta.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.acq_rel.cta.exch.b32 + %tmp732 = tail call float @llvm.nvvm.atomic.exch.gen.f.acq.rel.cta.f32.p0f32(float* %fp, float %f); + + ; CHECK: atom.acq_rel.cta.global.exch.b32 + %tmp733 = tail call float @llvm.nvvm.atomic.exch.global.f.acq.rel.cta.f32.p1f32(float addrspace(1)* %fp1, float %f); + + ; CHECK: atom.acq_rel.cta.shared.exch.b32 + %tmp734 = tail call float @llvm.nvvm.atomic.exch.shared.f.acq.rel.cta.f32.p3f32(float addrspace(3)* %fp3, float %f); + + ; CHECK: atom.acquire.exch.b64 + %tmp808 = tail call double @llvm.nvvm.atomic.exch.gen.f.acquire.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.acquire.global.exch.b64 + %tmp809 = tail call double @llvm.nvvm.atomic.exch.global.f.acquire.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.acquire.shared.exch.b64 + %tmp810 = tail call double @llvm.nvvm.atomic.exch.shared.f.acquire.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.release.exch.b64 + %tmp811 = tail call double @llvm.nvvm.atomic.exch.gen.f.release.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.release.global.exch.b64 + %tmp812 = tail call double @llvm.nvvm.atomic.exch.global.f.release.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.release.shared.exch.b64 + %tmp813 = tail call double @llvm.nvvm.atomic.exch.shared.f.release.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.acq_rel.exch.b64 + %tmp814 = tail call double @llvm.nvvm.atomic.exch.gen.f.acq.rel.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.acq_rel.global.exch.b64 + %tmp815 = tail call double @llvm.nvvm.atomic.exch.global.f.acq.rel.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.acq_rel.shared.exch.b64 + %tmp816 = tail call double @llvm.nvvm.atomic.exch.shared.f.acq.rel.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.acquire.sys.exch.b64 + %tmp817 = tail call double @llvm.nvvm.atomic.exch.gen.f.acquire.sys.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.acquire.sys.global.exch.b64 + %tmp818 = tail call double @llvm.nvvm.atomic.exch.global.f.acquire.sys.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.acquire.sys.shared.exch.b64 + %tmp819 = tail call double @llvm.nvvm.atomic.exch.shared.f.acquire.sys.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.release.sys.exch.b64 + %tmp820 = tail call double @llvm.nvvm.atomic.exch.gen.f.release.sys.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.release.sys.global.exch.b64 + %tmp821 = tail call double @llvm.nvvm.atomic.exch.global.f.release.sys.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.release.sys.shared.exch.b64 + %tmp822 = tail call double @llvm.nvvm.atomic.exch.shared.f.release.sys.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.acq_rel.sys.exch.b64 + %tmp823 = tail call double @llvm.nvvm.atomic.exch.gen.f.acq.rel.sys.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.acq_rel.sys.global.exch.b64 + %tmp824 = tail call double @llvm.nvvm.atomic.exch.global.f.acq.rel.sys.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.acq_rel.sys.shared.exch.b64 + %tmp825 = tail call double @llvm.nvvm.atomic.exch.shared.f.acq.rel.sys.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.acquire.cta.exch.b64 + %tmp826 = tail call double @llvm.nvvm.atomic.exch.gen.f.acquire.cta.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.acquire.cta.global.exch.b64 + %tmp827 = tail call double @llvm.nvvm.atomic.exch.global.f.acquire.cta.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.acquire.cta.shared.exch.b64 + %tmp828 = tail call double @llvm.nvvm.atomic.exch.shared.f.acquire.cta.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.release.cta.exch.b64 + %tmp829 = tail call double @llvm.nvvm.atomic.exch.gen.f.release.cta.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.release.cta.global.exch.b64 + %tmp830 = tail call double @llvm.nvvm.atomic.exch.global.f.release.cta.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.release.cta.shared.exch.b64 + %tmp831 = tail call double @llvm.nvvm.atomic.exch.shared.f.release.cta.f64.p3f64(double addrspace(3)* %dfp3, double %df); + + ; CHECK: atom.acq_rel.cta.exch.b64 + %tmp832 = tail call double @llvm.nvvm.atomic.exch.gen.f.acq.rel.cta.f64.p0f64(double* %dfp, double %df); + + ; CHECK: atom.acq_rel.cta.global.exch.b64 + %tmp833 = tail call double @llvm.nvvm.atomic.exch.global.f.acq.rel.cta.f64.p1f64(double addrspace(1)* %dfp1, double %df); + + ; CHECK: atom.acq_rel.cta.shared.exch.b64 + %tmp834 = tail call double @llvm.nvvm.atomic.exch.shared.f.acq.rel.cta.f64.p3f64(double addrspace(3)* %dfp3, double %df); + ; CHECK: atom.acquire.max.s32 %tmp162 = tail call i32 @llvm.nvvm.atomic.max.gen.i.acquire.i32.p0i32(i32* %ip, i32 %i); @@ -2126,6 +2288,168 @@ entry: ; CHECK: atom.acq_rel.cta.shared.cas.b64 %tmp701 = tail call i64 @llvm.nvvm.atomic.cas.shared.i.acq.rel.cta.i64.p3i64(i64 addrspace(3)* %llp3, i64 %ll, i64 %ll); + ; CHECK: atom.acquire.cas.b32 + %tmp748 = tail call float @llvm.nvvm.atomic.cas.gen.f.acquire.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.acquire.global.cas.b32 + %tmp749 = tail call float @llvm.nvvm.atomic.cas.global.f.acquire.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.acquire.shared.cas.b32 + %tmp750 = tail call float @llvm.nvvm.atomic.cas.shared.f.acquire.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.release.cas.b32 + %tmp751 = tail call float @llvm.nvvm.atomic.cas.gen.f.release.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.release.global.cas.b32 + %tmp752 = tail call float @llvm.nvvm.atomic.cas.global.f.release.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.release.shared.cas.b32 + %tmp753 = tail call float @llvm.nvvm.atomic.cas.shared.f.release.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.acq_rel.cas.b32 + %tmp754 = tail call float @llvm.nvvm.atomic.cas.gen.f.acq.rel.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.acq_rel.global.cas.b32 + %tmp755 = tail call float @llvm.nvvm.atomic.cas.global.f.acq.rel.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.acq_rel.shared.cas.b32 + %tmp756 = tail call float @llvm.nvvm.atomic.cas.shared.f.acq.rel.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.acquire.sys.cas.b32 + %tmp757 = tail call float @llvm.nvvm.atomic.cas.gen.f.acquire.sys.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.acquire.sys.global.cas.b32 + %tmp758 = tail call float @llvm.nvvm.atomic.cas.global.f.acquire.sys.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.acquire.sys.shared.cas.b32 + %tmp759 = tail call float @llvm.nvvm.atomic.cas.shared.f.acquire.sys.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.release.sys.cas.b32 + %tmp760 = tail call float @llvm.nvvm.atomic.cas.gen.f.release.sys.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.release.sys.global.cas.b32 + %tmp761 = tail call float @llvm.nvvm.atomic.cas.global.f.release.sys.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.release.sys.shared.cas.b32 + %tmp762 = tail call float @llvm.nvvm.atomic.cas.shared.f.release.sys.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.acq_rel.sys.cas.b32 + %tmp763 = tail call float @llvm.nvvm.atomic.cas.gen.f.acq.rel.sys.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.acq_rel.sys.global.cas.b32 + %tmp764 = tail call float @llvm.nvvm.atomic.cas.global.f.acq.rel.sys.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.acq_rel.sys.shared.cas.b32 + %tmp765 = tail call float @llvm.nvvm.atomic.cas.shared.f.acq.rel.sys.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.acquire.cta.cas.b32 + %tmp766 = tail call float @llvm.nvvm.atomic.cas.gen.f.acquire.cta.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.acquire.cta.global.cas.b32 + %tmp767 = tail call float @llvm.nvvm.atomic.cas.global.f.acquire.cta.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.acquire.cta.shared.cas.b32 + %tmp768 = tail call float @llvm.nvvm.atomic.cas.shared.f.acquire.cta.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.release.cta.cas.b32 + %tmp769 = tail call float @llvm.nvvm.atomic.cas.gen.f.release.cta.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.release.cta.global.cas.b32 + %tmp770 = tail call float @llvm.nvvm.atomic.cas.global.f.release.cta.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.release.cta.shared.cas.b32 + %tmp771 = tail call float @llvm.nvvm.atomic.cas.shared.f.release.cta.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.acq_rel.cta.cas.b32 + %tmp772 = tail call float @llvm.nvvm.atomic.cas.gen.f.acq.rel.cta.f32.p0f32(float* %fp, float %f, float %f); + + ; CHECK: atom.acq_rel.cta.global.cas.b32 + %tmp773 = tail call float @llvm.nvvm.atomic.cas.global.f.acq.rel.cta.f32.p1f32(float addrspace(1)* %fp1, float %f, float %f); + + ; CHECK: atom.acq_rel.cta.shared.cas.b32 + %tmp774 = tail call float @llvm.nvvm.atomic.cas.shared.f.acq.rel.cta.f32.p3f32(float addrspace(3)* %fp3, float %f, float %f); + + ; CHECK: atom.acquire.cas.b64 + %tmp848 = tail call double @llvm.nvvm.atomic.cas.gen.f.acquire.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.acquire.global.cas.b64 + %tmp849 = tail call double @llvm.nvvm.atomic.cas.global.f.acquire.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.acquire.shared.cas.b64 + %tmp850 = tail call double @llvm.nvvm.atomic.cas.shared.f.acquire.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.release.cas.b64 + %tmp851 = tail call double @llvm.nvvm.atomic.cas.gen.f.release.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.release.global.cas.b64 + %tmp852 = tail call double @llvm.nvvm.atomic.cas.global.f.release.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.release.shared.cas.b64 + %tmp853 = tail call double @llvm.nvvm.atomic.cas.shared.f.release.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.acq_rel.cas.b64 + %tmp854 = tail call double @llvm.nvvm.atomic.cas.gen.f.acq.rel.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.acq_rel.global.cas.b64 + %tmp855 = tail call double @llvm.nvvm.atomic.cas.global.f.acq.rel.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.acq_rel.shared.cas.b64 + %tmp856 = tail call double @llvm.nvvm.atomic.cas.shared.f.acq.rel.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.acquire.sys.cas.b64 + %tmp857 = tail call double @llvm.nvvm.atomic.cas.gen.f.acquire.sys.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.acquire.sys.global.cas.b64 + %tmp858 = tail call double @llvm.nvvm.atomic.cas.global.f.acquire.sys.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.acquire.sys.shared.cas.b64 + %tmp859 = tail call double @llvm.nvvm.atomic.cas.shared.f.acquire.sys.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.release.sys.cas.b64 + %tmp860 = tail call double @llvm.nvvm.atomic.cas.gen.f.release.sys.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.release.sys.global.cas.b64 + %tmp861 = tail call double @llvm.nvvm.atomic.cas.global.f.release.sys.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.release.sys.shared.cas.b64 + %tmp862 = tail call double @llvm.nvvm.atomic.cas.shared.f.release.sys.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.acq_rel.sys.cas.b64 + %tmp863 = tail call double @llvm.nvvm.atomic.cas.gen.f.acq.rel.sys.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.acq_rel.sys.global.cas.b64 + %tmp864 = tail call double @llvm.nvvm.atomic.cas.global.f.acq.rel.sys.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.acq_rel.sys.shared.cas.b64 + %tmp865 = tail call double @llvm.nvvm.atomic.cas.shared.f.acq.rel.sys.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.acquire.cta.cas.b64 + %tmp866 = tail call double @llvm.nvvm.atomic.cas.gen.f.acquire.cta.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.acquire.cta.global.cas.b64 + %tmp867 = tail call double @llvm.nvvm.atomic.cas.global.f.acquire.cta.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.acquire.cta.shared.cas.b64 + %tmp868 = tail call double @llvm.nvvm.atomic.cas.shared.f.acquire.cta.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.release.cta.cas.b64 + %tmp869 = tail call double @llvm.nvvm.atomic.cas.gen.f.release.cta.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.release.cta.global.cas.b64 + %tmp870 = tail call double @llvm.nvvm.atomic.cas.global.f.release.cta.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.release.cta.shared.cas.b64 + %tmp871 = tail call double @llvm.nvvm.atomic.cas.shared.f.release.cta.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + + ; CHECK: atom.acq_rel.cta.cas.b64 + %tmp872 = tail call double @llvm.nvvm.atomic.cas.gen.f.acq.rel.cta.f64.p0f64(double* %dfp, double %df, double %df); + + ; CHECK: atom.acq_rel.cta.global.cas.b64 + %tmp873 = tail call double @llvm.nvvm.atomic.cas.global.f.acq.rel.cta.f64.p1f64(double addrspace(1)* %dfp1, double %df, double %df); + + ; CHECK: atom.acq_rel.cta.shared.cas.b64 + %tmp874 = tail call double @llvm.nvvm.atomic.cas.shared.f.acq.rel.cta.f64.p3f64(double addrspace(3)* %dfp3, double %df, double %df); + ; CHECK: ld.relaxed.gpu.s32 %tmpldst0 = tail call i32 @llvm.nvvm.ld.gen.i.i32.p0(i32* %ip); @@ -2979,6 +3303,60 @@ declare i64 @llvm.nvvm.atomic.exch.shared.i.release.cta.i64.p3i64(i64 addrspace( declare i64 @llvm.nvvm.atomic.exch.gen.i.acq.rel.cta.i64.p0i64(i64* nocapture, i64) #1 declare i64 @llvm.nvvm.atomic.exch.global.i.acq.rel.cta.i64.p1i64(i64 addrspace(1)* nocapture, i64) #1 declare i64 @llvm.nvvm.atomic.exch.shared.i.acq.rel.cta.i64.p3i64(i64 addrspace(3)* nocapture, i64) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.acquire.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.acquire.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.acquire.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.release.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.release.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.release.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.acq.rel.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.acq.rel.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.acq.rel.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.acquire.sys.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.acquire.sys.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.acquire.sys.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.release.sys.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.release.sys.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.release.sys.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.acq.rel.sys.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.acq.rel.sys.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.acq.rel.sys.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.acquire.cta.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.acquire.cta.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.acquire.cta.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.release.cta.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.release.cta.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.release.cta.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.gen.f.acq.rel.cta.f32.p0f32(float* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.global.f.acq.rel.cta.f32.p1f32(float addrspace(1)* nocapture, float) #1 +declare float @llvm.nvvm.atomic.exch.shared.f.acq.rel.cta.f32.p3f32(float addrspace(3)* nocapture, float) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.acquire.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.acquire.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.acquire.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.release.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.release.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.release.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.acq.rel.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.acq.rel.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.acq.rel.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.acquire.sys.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.acquire.sys.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.acquire.sys.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.release.sys.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.release.sys.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.release.sys.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.acq.rel.sys.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.acq.rel.sys.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.acq.rel.sys.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.acquire.cta.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.acquire.cta.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.acquire.cta.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.release.cta.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.release.cta.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.release.cta.f64.p3f64(double addrspace(3)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.gen.f.acq.rel.cta.f64.p0f64(double* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.global.f.acq.rel.cta.f64.p1f64(double addrspace(1)* nocapture, double) #1 +declare double @llvm.nvvm.atomic.exch.shared.f.acq.rel.cta.f64.p3f64(double addrspace(3)* nocapture, double) #1 declare i32 @llvm.nvvm.atomic.max.gen.i.acquire.i32.p0i32(i32* nocapture, i32) #1 declare i32 @llvm.nvvm.atomic.max.global.i.acquire.i32.p1i32(i32 addrspace(1)* nocapture, i32) #1 declare i32 @llvm.nvvm.atomic.max.shared.i.acquire.i32.p3i32(i32 addrspace(3)* nocapture, i32) #1 @@ -3519,5 +3897,59 @@ declare i64 @llvm.nvvm.atomic.cas.shared.i.release.cta.i64.p3i64(i64 addrspace(3 declare i64 @llvm.nvvm.atomic.cas.gen.i.acq.rel.cta.i64.p0i64(i64* nocapture, i64, i64) #1 declare i64 @llvm.nvvm.atomic.cas.global.i.acq.rel.cta.i64.p1i64(i64 addrspace(1)* nocapture, i64, i64) #1 declare i64 @llvm.nvvm.atomic.cas.shared.i.acq.rel.cta.i64.p3i64(i64 addrspace(3)* nocapture, i64, i64) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.acquire.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.acquire.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.acquire.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.release.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.release.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.release.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.acq.rel.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.acq.rel.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.acq.rel.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.acquire.sys.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.acquire.sys.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.acquire.sys.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.release.sys.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.release.sys.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.release.sys.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.acq.rel.sys.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.acq.rel.sys.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.acq.rel.sys.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.acquire.cta.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.acquire.cta.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.acquire.cta.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.release.cta.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.release.cta.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.release.cta.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.gen.f.acq.rel.cta.f32.p0f32(float* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.global.f.acq.rel.cta.f32.p1f32(float addrspace(1)* nocapture, float, float) #1 +declare float @llvm.nvvm.atomic.cas.shared.f.acq.rel.cta.f32.p3f32(float addrspace(3)* nocapture, float, float) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.acquire.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.acquire.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.acquire.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.release.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.release.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.release.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.acq.rel.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.acq.rel.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.acq.rel.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.acquire.sys.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.acquire.sys.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.acquire.sys.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.release.sys.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.release.sys.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.release.sys.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.acq.rel.sys.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.acq.rel.sys.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.acq.rel.sys.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.acquire.cta.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.acquire.cta.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.acquire.cta.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.release.cta.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.release.cta.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.release.cta.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.gen.f.acq.rel.cta.f64.p0f64(double* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.global.f.acq.rel.cta.f64.p1f64(double addrspace(1)* nocapture, double, double) #1 +declare double @llvm.nvvm.atomic.cas.shared.f.acq.rel.cta.f64.p3f64(double addrspace(3)* nocapture, double, double) #1 attributes #1 = { argmemonly nounwind }