
// Inline PTX: h u16 , r u32,  l u64, f f32, d f64
// Ops:

// binary operations

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_AND() \
template<class ctype> \
inline __device__ typename std::enable_if<sizeof(ctype)==4, ctype>::type device_atomic_fetch_and(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  uint32_t asm_value = reinterpret_cast<uint32_t&>(value); \
  uint32_t asm_result = 0u; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.and.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b32" " %0,[%1],%2;\n\t" \
          "@!p atom.and"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b32" " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : "=r"(asm_result) : "l"(dest),"r"(asm_value) : "memory"); \
  return reinterpret_cast<ctype&>(asm_result); \
} \
template<class ctype> \
inline __device__ typename std::enable_if<sizeof(ctype)==8, ctype>::type device_atomic_fetch_and(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  uint64_t asm_value = reinterpret_cast<uint64_t&>(value); \
  uint64_t asm_result = 0u; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.and.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b64" " %0,[%1],%2;\n\t" \
          "@!p atom.and"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b64" " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : "=l"(asm_result) : "l"(dest),"l"(asm_value) : "memory"); \
  return reinterpret_cast<ctype&>(asm_result); \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_OR() \
template<class ctype> \
inline __device__ typename std::enable_if<sizeof(ctype)==4, ctype>::type device_atomic_fetch_or(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  uint32_t asm_value = reinterpret_cast<uint32_t&>(value); \
  uint32_t asm_result = 0u; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.or.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b32" " %0,[%1],%2;\n\t" \
          "@!p atom.or"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b32" " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : "=r"(asm_result) : "l"(dest),"r"(asm_value) : "memory"); \
  return reinterpret_cast<ctype&>(asm_result); \
} \
template<class ctype> \
inline __device__ typename std::enable_if<sizeof(ctype)==8, ctype>::type device_atomic_fetch_or(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  uint64_t asm_value = reinterpret_cast<uint64_t&>(value); \
  uint64_t asm_result = 0u; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.or.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b64" " %0,[%1],%2;\n\t" \
          "@!p atom.or"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b64" " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : "=l"(asm_result) : "l"(dest),"l"(asm_value) : "memory"); \
  return reinterpret_cast<ctype&>(asm_result); \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_XOR() \
template<class ctype> \
inline __device__ typename std::enable_if<sizeof(ctype)==4, ctype>::type device_atomic_fetch_xor(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  uint32_t asm_value = reinterpret_cast<uint32_t&>(value); \
  uint32_t asm_result = 0u; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.xor.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b32" " %0,[%1],%2;\n\t" \
          "@!p atom.xor"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b32" " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : "=r"(asm_result) : "l"(dest),"r"(asm_value) : "memory"); \
  return reinterpret_cast<ctype&>(asm_result); \
} \
template<class ctype> \
inline __device__ typename std::enable_if<sizeof(ctype)==8, ctype>::type device_atomic_fetch_xor(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  uint64_t asm_value = reinterpret_cast<uint64_t&>(value); \
  uint64_t asm_result = 0u; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.xor.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b64" " %0,[%1],%2;\n\t" \
          "@!p atom.xor"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM ".b64" " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : "=l"(asm_result) : "l"(dest),"l"(asm_value) : "memory"); \
  return reinterpret_cast<ctype&>(asm_result); \
}

// Fetch atomics
#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(ctype,asm_ctype,reg_ctype,reg_ret_ctype) \
inline __device__ ctype device_atomic_fetch_add(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result=0; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.add.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.add"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(value) : "memory"); \
  return result; \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(ctype,asm_ctype,reg_ctype,reg_ret_ctype) \
inline __device__ ctype device_atomic_fetch_sub(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result=0; \
  ctype neg_value = -value; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.add.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.add"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(neg_value) : "memory"); \
  return result; \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MIN(ctype,asm_ctype,reg_ctype,reg_ret_ctype) \
inline __device__ ctype device_atomic_fetch_min(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result=0; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.min.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.min"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(value) : "memory"); \
  return result; \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MAX(ctype,asm_ctype,reg_ctype,reg_ret_ctype) \
inline __device__ ctype device_atomic_fetch_max(ctype* dest, ctype value, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result=0; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.max.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.max"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(value) : "memory"); \
  return result; \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_INC(ctype,asm_ctype,reg_ctype,reg_ret_ctype) \
inline __device__ ctype device_atomic_fetch_inc(ctype* dest, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result = 0; \
  ctype limit = desul::Impl::numeric_limits_max<ctype>::value; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.inc.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.inc"       __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(limit) : "memory"); \
  return result; \
} \
inline __device__ ctype device_atomic_fetch_inc_mod(ctype* dest, ctype limit, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result = 0; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.inc.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.inc"       __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(limit) : "memory"); \
  return result; \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_DEC(ctype,asm_ctype,reg_ctype,reg_ret_ctype) \
inline __device__ ctype device_atomic_fetch_dec(ctype* dest, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result = 0; \
  ctype limit = desul::Impl::numeric_limits_max<ctype>::value; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.dec.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.dec"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(limit) : "memory"); \
  return result; \
} \
inline __device__ ctype device_atomic_fetch_dec_mod(ctype* dest, ctype limit, __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER, __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE) { \
  ctype result = 0; \
  asm volatile( \
          "{\n\t" \
          ".reg .pred p;\n\t" \
          "isspacep.global p, %1;\n\t" \
          "@p  atom.dec.global" __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "@!p atom.dec"        __DESUL_IMPL_CUDA_ASM_MEMORY_ORDER_ASM __DESUL_IMPL_CUDA_ASM_MEMORY_SCOPE_ASM asm_ctype " %0,[%1],%2;\n\t" \
          "}\n\t" \
    : reg_ret_ctype(result) : "l"(dest),reg_ctype(limit) : "memory"); \
  return result; \
}

#define __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_BIN_OP() \
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_AND() \
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_OR() \
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_XOR()


// Instantiate Functions

// General comments:
//  - float/double only support add
//  - inc/dec only supported with uint32_t
//  - int64_t does not support add

// floating point types
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(float,".f32","f","=f")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(float,".f32","f","=f")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(double,".f64","d","=d")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(double,".f64","d","=d")

// uint32_t
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(uint32_t,".u32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(uint32_t,".u32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MIN(uint32_t,".u32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MAX(uint32_t,".u32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_INC(uint32_t,".u32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_DEC(uint32_t,".u32","r","=r")

// uint64_t
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(uint64_t,".u64","l","=l")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(uint64_t,".u64","l","=l")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MIN(uint64_t,".u64","l","=l")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MAX(uint64_t,".u64","l","=l")
//__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_INC(uint64_t,".u64","l","=l")
//__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_DEC(uint64_t,".u64","l","=l")

// int32_t
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(int32_t,".s32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(int32_t,".s32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MIN(int32_t,".s32","r","=r")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MAX(int32_t,".s32","r","=r")
//__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_INC(int32_t,".s32","r","=r")
//__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_DEC(int32_t,".s32","r","=r")

// int64_t note: add/sub is using unsigned register
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD(int64_t,".u64","l","=l")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_SUB(int64_t,".u64","l","=l")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MIN(int64_t,".s64","l","=l")
__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MAX(int64_t,".s64","l","=l")
//__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_INC(int64_t,".s64","l","=l")
//__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_DEC(int64_t,".s64","l","=l")

__DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_BIN_OP()

#undef __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_ADD
#undef __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MIN
#undef __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_MAX
#undef __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_INC
#undef __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_DEC
#undef __DESUL_IMPL_CUDA_ASM_ATOMIC_FETCH_AND

