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

Matt Arsenault via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Tue Aug 8 06:05:50 PDT 2023


arsenm added a comment.

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


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D157388



More information about the llvm-commits mailing list