================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
