relaxed load and store

[jinz2014]

Can you please provide a solution that is portable for relaxed load and store operations ? Currently, the assembly codes (e.g. load) are required. Thanks.

static __device__ __forceinline__ uint4 load_relaxed(uint4 const *ptr)
{
#ifdef PTX
  uint4 retval;
  NV_IF_TARGET(NV_PROVIDES_SM_70,
               (asm volatile("ld.relaxed.gpu.v4.u32 {%0, %1, %2, %3}, [%4];"
                             : "=r"(retval.x), "=r"(retval.y), "=r"(retval.z), "=r"(retval.w)
                             : _CUB_ASM_PTR_(ptr)
                             : "memory");),
               (asm volatile("ld.cg.v4.u32 {%0, %1, %2, %3}, [%4];"
                             : "=r"(retval.x), "=r"(retval.y), "=r"(retval.z), "=r"(retval.w)
                             : _CUB_ASM_PTR_(ptr)
                             : "memory");));
  return retval;
#else
 return *ptr;
#endif
}

[miscco]

Could you provide some more details. What does portable mean in this context? Different architectures, other vendors?

[jinz2014]

Yes. A CUDA program without assembly codes for vendors.

[jrhemstad]

Using cuda::atomic_ref::load/store(cuda::memory_order_relaxed) is probably your easiest solution.

https://godbolt.org/z/39cKqhnW3

[jinz2014]

Thank you for the link. However, I find that atomic_ref does not support types larger than 8 bytes.