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

[Sameer Sahasrabuddhe via cfe-commits]

================
@@ -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