[PATCH] D80804: [AMDGPU] Introduce Clang builtins to be mapped to AMDGCN atomic inc/dec intrinsics

Fri May 29 22:50:25 PDT 2020

sameerds added a comment.

In D80804#2063522 <https://reviews.llvm.org/D80804#2063522>, @saiislam wrote:

> In D80804#2063451 <https://reviews.llvm.org/D80804#2063451>, @sameerds wrote:
>
> > Actually, the question really is about why inc/dec are needed as separate operations either as IR intrinsics or Clang builtins. Why not just expose a __builtin_amdgcn_atomicrmw that takes a scope, and map it to the LLVM atomicrmw? That would be way cleaner. The language can provide convenience functions for inc/dec that internally call the rmw builtin.
>
>
> At the moment, atomic inc/dec exist along with atomicrmw. This patch only aims to devise a way to make it (inc/dec) accessible from the language level.

Just to clarify that a bit, the point is that atomic inc/dec are not available as operations in atomicrmw, and cannot be treated as an optimization of (add 1) or (sub 1). So the intrinsics are required until atomicrmw is enhanced, which is out of scope for this patch.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D80804/new/

https://reviews.llvm.org/D80804