[all-commits] [llvm/llvm-project] ea17c7: clang/AMDGPU: Emit atomicrmw for __builtin_amdgcn_...

Matt Arsenault via All-commits all-commits at lists.llvm.org
Thu Aug 1 11:11:26 PDT 2024


  Branch: refs/heads/users/arsenm/amdgpu-clang-emit-atomicrmw-builtin-global-atomic-fadd
  Home:   https://github.com/llvm/llvm-project
  Commit: ea17c792053e32e39a7261e3bdf1673d98e4d94a
      https://github.com/llvm/llvm-project/commit/ea17c792053e32e39a7261e3bdf1673d98e4d94a
  Author: Matt Arsenault <Matthew.Arsenault at amd.com>
  Date:   2024-08-01 (Thu, 01 Aug 2024)

  Changed paths:
    M clang/lib/CodeGen/CGBuiltin.cpp
    M clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl

  Log Message:
  -----------
  clang/AMDGPU: Emit atomicrmw for __builtin_amdgcn_global_atomic_fadd_{f32|f64}

Need to emit syncscope and new metadata to get the native instruction,
most of the time.


  Commit: 81f345d6741a8326367b6b7df9816125f594b713
      https://github.com/llvm/llvm-project/commit/81f345d6741a8326367b6b7df9816125f594b713
  Author: Matt Arsenault <Matthew.Arsenault at amd.com>
  Date:   2024-08-01 (Thu, 01 Aug 2024)

  Changed paths:
    M clang/lib/CodeGen/CGBuiltin.cpp
    M clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx12.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx8.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl
    M clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx940.cl

  Log Message:
  -----------
  Use monotonic ordering


Compare: https://github.com/llvm/llvm-project/compare/9352ace1c597...81f345d6741a

To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications


More information about the All-commits mailing list