[all-commits] [llvm/llvm-project] a247da: [libclc] update __clc_mem_fence: add MemorySemanti...

[Wenju He via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: a247da4f9363116c54b91a37755edd994c56dbf8
      https://github.com/llvm/llvm-project/commit/a247da4f9363116c54b91a37755edd994c56dbf8
  Author: Wenju He <wenju.he at intel.com>
  Date:   2025-09-01 (Mon, 01 Sep 2025)

  Changed paths:
    M libclc/clc/include/clc/mem_fence/clc_mem_fence.h
    A libclc/clc/include/clc/mem_fence/clc_mem_semantic.h
    M libclc/clc/include/clc/synchronization/clc_work_group_barrier.h
    M libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl
    M libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl
    M libclc/clc/lib/ptx-nvidiacl/mem_fence/clc_mem_fence.cl
    M libclc/clc/lib/ptx-nvidiacl/synchronization/clc_work_group_barrier.cl
    M libclc/opencl/include/clc/opencl/synchronization/utils.h
    M libclc/opencl/lib/amdgcn/mem_fence/fence.cl
    M libclc/opencl/lib/amdgcn/synchronization/barrier.cl
    M libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl
    M libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl

  Log Message:
  -----------
  [libclc] update __clc_mem_fence: add MemorySemantic arg and use __builtin_amdgcn_fence for AMDGPU (#152275)

It is necessary to add MemorySemantic argument for AMDGPU which means
the memory or address space to which the memory ordering is applied.

The MemorySemantic is also necessary for implementing the SPIR-V
MemoryBarrier instruction. Additionally, the implementation of
__clc_mem_fence on Intel GPUs requires the MemorySemantic argument.

Using __builtin_amdgcn_fence for AMDGPU is follow-up of
https://github.com/llvm/llvm-project/pull/151446#discussion_r2254006508

llvm-diff shows no change to nvptx64--nvidiacl.bc.



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