[clang] [clang][AMDGPU] Clean-up handling of named barrier type (PR #207687)

Sameer Sahasrabuddhe via cfe-commits cfe-commits at lists.llvm.org
Thu Jul 23 23:54:30 PDT 2026


================
@@ -30,3 +30,38 @@ Please note that the specific architecture and feature names will vary depending
 
 Clang exposes AMDGPU hardware intrinsics as target-specific builtins with the
 `__builtin_amdgcn_` prefix. These are documented in {doc}`AMDGPUBuiltinReference`.
+
+## Target-Specific Types
+
+### Named Workgroup Barrier Type
+
+The `__amdgpu_named_workgroup_barrier_t` type is used to represent the GFX12.5 named barriers.
+Example usage:
+
+```c
+__amdgpu_named_workgroup_barrier_t x;
+__amdgpu_named_workgroup_barrier_t arr[2]; // Arrays are also fine
+
+void foo(int a)
+{
+  __builtin_amdgcn_s_barrier_init(&x, a);
+}
+```
+
+A "named barrier wrapper" is a class that contains exactly one field, which is either
+a single value or an array of values of one of the following types:
+
+* `__amdgpu_named_workgroup_barrier_t`.
+* Another "named barrier wrapper".
+
+Named barrier wrappers let users add helper methods around named barrier objects.
----------------
ssahasra wrote:

This breaks the flow of the text. It should be moved earlier as an introduction to wrapper classes.

https://github.com/llvm/llvm-project/pull/207687


More information about the cfe-commits mailing list