[PATCH] D157388: [AMDGPU] Support FMin/FMax in AMDGPUAtomicOptimizer.

[Pravin Jagtap via Phabricator via llvm-commits]

pravinjagtap added a comment.

In D157388#4569250 <https://reviews.llvm.org/D157388#4569250>, @arsenm wrote:

> In D157388#4569245 <https://reviews.llvm.org/D157388#4569245>, @pravinjagtap wrote:
>
>> For FMin and FMax cases, `clang` itself is emitting CAS loop for both
>>
>>   __device__ inline float atomicMax(float* addr, float val) and
>>   __device__ inline float unsafeAtomicMax(float* addr, float val) 
>>
>> I am not sure how to potentially avoid this CAS loop before we reach atomic optimization pass.
>> CC: @b-sumner @arsenm
>
> Clang should not be expanding any atomics itself

All this is being implemented outside of llvm. The logic for creating CAS loop is inserted in `hipamd/include/hip/amd_detail/amd_hip_atomic.h`. How should we go about this? do I need to update hipamd repo ? CC: @b-sumner @arsenm


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D157388