================
@@ -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.
+
+In C++, a class that inherits from a named barrier wrapper is also considered a
----------------
ssahasra wrote:
I managed to wrangle the following from Claude Opus. To me it reads
subjectively better than the current wording. It covers both the cases --- a
containing class or a derived class.
A "named barrier wrapper" is a
[standard-layout](https://en.cppreference.com/w/cpp/language/classes#Standard-layout_class)
class (see std::is_standard_layout) that has exactly one direct subobject —
either a single non-static data member or a single base class — whose type is
one of the following:
__amdgpu_named_workgroup_barrier_t, or an array thereof.
Another "named barrier wrapper", or an array thereof.
Named barrier wrappers let users add helper methods around named barrier
objects.
Because a standard-layout class places this single subobject at offset zero and
is pointer-interconvertible with it, a named barrier wrapper has the same
layout as the named barrier it ultimately wraps. This covers two cases with one
rule:
A class whose sole non-static data member is a named barrier (or another
wrapper), e.g. struct W { __amdgpu_named_workgroup_barrier_t x; };.
A class that adds no data members of its own but derives from a named
barrier wrapper, e.g. struct D : W {};.
The standard-layout requirement forbids virtual functions and virtual base
classes, and — together with the "exactly one subobject" rule — forbids any
additional data members, whether declared or inherited.
https://github.com/llvm/llvm-project/pull/207687
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
