[libclc] Move mem_fence and barrier to clc library #151446
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
__clc_mem_fence and __clc_barrier function have two parameters Scope and MemorySemantics, which are defined in SPIR-V spec. The design allows the clc functions to implement SPIR-V ControlBarrier and MemoryBarrier functions in the future. The default memory ordering in clc is set to SequentiallyConsistent, which is also the default and strongest ordering in OpenCL and C++. The default memory scope in clc is set to memory_scope_device for amdgcn and ptx-nvidiacl since __opencl_c_atomic_scope_all_devices feature macro is not defined for these targets. llvm-diff shows no change to amdgcn--amdhsa.bc and nvptx64--nvidiacl.bc.
arsenm
reviewed
Aug 1, 2025
| #define __CLC_MEM_FENCE_CLC_MEM_SCOPE_SEMANTICS_H__ | ||
|
|
||
| // Scope values are defined in SPIR-V spec. | ||
| typedef enum Scope { |
There was a problem hiding this comment.
Should stick to OpenCL spec and terminology; I thought clang already had predefined enums for both of these
There was a problem hiding this comment.
Should stick to OpenCL spec and terminology; I thought clang already had predefined enums for both of these
done, thanks @arsenm
Changed to use clang macro __MEMORY_SCOPE_DEVICE/__MEMORY_SCOPE_WRKGRP and __ATOMIC_SEQ_CST
…s with clang macros
arsenm
approved these changes
Aug 5, 2025
|
|
||
| #include <clc/mem_fence/clc_mem_fence.h> | ||
|
|
||
| void __clc_amdgcn_s_waitcnt(unsigned flags); |
There was a problem hiding this comment.
We should get rid of this; it can be properly implemented with __builtin_amdgcn_fence
There was a problem hiding this comment.
We should get rid of this; it can be properly implemented with
__builtin_amdgcn_fence
I'll follow up
There was a problem hiding this comment.
We should get rid of this; it can be properly implemented with
__builtin_amdgcn_fenceI'll follow up
re-implement with __builtin_amdgcn_fence: #152275
wenju-he
added a commit
to wenju-he/llvm-project
that referenced
this pull request
Aug 6, 2025
…ltin_amdgcn_fence for AMDGPU It is necessary to add MemorySemantic argument 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 llvm#151446 (comment) llvm-diff shows no change to nvptx64--nvidiacl.bc.
wenju-he
added a commit
that referenced
this pull request
Sep 1, 2025
…ltin_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 #151446 (comment) llvm-diff shows no change to nvptx64--nvidiacl.bc.
llvm-sync bot
pushed a commit
to arm/arm-toolchain
that referenced
this pull request
Sep 1, 2025
…d 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 llvm/llvm-project#151446 (comment) llvm-diff shows no change to nvptx64--nvidiacl.bc.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
__clc_mem_fence and __clc_work_group_barrier function have two
parameters memory_scope and memory_order. The design allows the clc
functions to implement SPIR-V ControlBarrier and MemoryBarrier
functions in the future.
The default memory ordering in clc is set to __ATOMIC_SEQ_CST, which is
also the default and strongest ordering in OpenCL and C++.
OpenCL cl_mem_fence_flags parameter is converted to combination of
__MEMORY_SCOPE_DEVICE and __MEMORY_SCOPE_WRKGRP, which is passed to clc.
llvm-diff shows no change to nvptx64--nvidiacl.bc.
llvm-diff show a small change to amdgcn--amdhsa.bc and the number of
LLVM IR instruction is reduced by 1: https://alive2.llvm.org/ce/z/_Uhqvt