ssahasra wrote: > I feel like this is retreading the same ground from > https://discourse.llvm.org/t/rfc-add-clang-atomic-control-options-and-pragmas/80641 > > Given that, would something like this make sense? > > ```c++ > template <typename T> > __device__ __forceinline__ T loadFifo(T *p) { > [[clang::atomic(nontemporal)]] { > return __scoped_atomic_load_n(p, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM); > } > } > ```
Superficially, it may seem that the only requirement is to add `!nontemporal` to an atomic. But note that the expectation is even more ... this feature is supposed to allow the compiler to "expand" or widen the operation. (It is not clear if arbitrary loads can be fused, or only other calls to the same function.) This all seems backwards, where the design started by first determining the required ISA and then figuring out what hack in the compiler can produce that ISA. We need a forward design: What are the requirements of this protocal load: 1. Does it need to be unsplittable? (I don't want to call it atomic just yet). 2. Does it contra-indicate temporal locality, so that it should not be retained in the cache? 3. Does it need to be merged with other loads whenever possible? 4. Can it be elided with matching operations? 5. Does it have ordering semantics? 6. Is it required to complete at a specified scope? Does that affect correctness or performance or both? 7. Can it pair with a fence so that two threads can establish a synchronize-with relationship using this load? (That's the other aspect of something being atomic.) https://github.com/llvm/llvm-project/pull/229207 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
