[clang] [llvm] [AMDGPU[ Expand `av` built-ins and intrinsics to b8/b16/b32/b64 (PR #229429)

[Fabian Ritter via cfe-commits]

================
@@ -5496,28 +5496,37 @@ returns the bit at the position of the current lane. It is almost equivalent to
 `(mask & (1 << lane_id)) != 0`, except that its behavior is only defined if
 the given mask has the same value for all active lanes of the current wave.
 
-#### \_\_builtin_amdgcn_av\_{load,store}\_b128
+#### \_\_builtin_amdgcn_av\_{load,store}\_{b8,b16,b32,b64,b128}
 
 Signature:
 
 ```c
+typedef __attribute__((__vector_size__(2 * sizeof(unsigned int)))) unsigned int v2u;
 typedef __attribute__((__vector_size__(4 * sizeof(unsigned int)))) unsigned int v4u;
 
+unsigned char __builtin_amdgcn_av_load_b8(unsigned char *src, int scope);
+unsigned short __builtin_amdgcn_av_load_b16(unsigned short *src, int scope);
+unsigned int __builtin_amdgcn_av_load_b32(unsigned int *src, int scope);
+v2u __builtin_amdgcn_av_load_b64(v2u *src, int scope);
 v4u __builtin_amdgcn_av_load_b128(v4u *src, int scope);
 
+void __builtin_amdgcn_av_store_b8(unsigned char *dst, unsigned char data, int scope);
+void __builtin_amdgcn_av_store_b16(unsigned short *dst, unsigned short data, int scope);
+void __builtin_amdgcn_av_store_b32(unsigned int *dst, unsigned int data, int scope);
+void __builtin_amdgcn_av_store_b64(v2u *dst, v2u data, int scope);
 void __builtin_amdgcn_av_store_b128(v4u *dst, v4u data, int scope);
 ```
 
-Load or store a vector of 4 unsigned integers from or to memory with cache
-behavior specified by `scope`, which is one of the `__MEMORY_SCOPE_*` macros
-defined for {ref}`scoped atomic builtins <langext-c11-atomic>`.
+Load or store 8, 16, 32, 64 or 128 bits from or to memory with cache behavior
+specified by `scope`, which is one of the `__MEMORY_SCOPE_*` macros defined for
+{ref}`scoped atomic builtins <langext-c11-atomic>`.
----------------
ritter-x2a wrote:

I wonder how this interacts with C/C++'s strict aliasing rules. If you'd cast a pointer to a `uint64_t` (or a `double`, or a vector of 2 `float`s) object to `v2u` and pass it to `__builtin_amdgcn_av_load_b64`, wouldn't you access that object with the wrong data type and therefore cause UB? (To be fair, that would already be a concern with the `b128` variant, if you wanted to use it for anything other than unsigned ints.)

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