[clang] [AMDGPU][Clang] add __builtin_amdgcn_av_(load|store)_b128 (PR #199176)

Shilei Tian via cfe-commits cfe-commits at lists.llvm.org
Tue Jun 2 06:10:21 PDT 2026


================
@@ -280,6 +280,19 @@ def __builtin_amdgcn_raw_ptr_buffer_load_async_lds : AMDGPUBuiltin<"void(__amdgp
 def __builtin_amdgcn_struct_ptr_buffer_load_lds : AMDGPUBuiltin<"void(__amdgpu_buffer_rsrc_t, void address_space<3> *, _Constant unsigned int, int, int, int, _Constant int, _Constant int)", [], "vmem-to-lds-load-insts">;
 def __builtin_amdgcn_struct_ptr_buffer_load_async_lds : AMDGPUBuiltin<"void(__amdgpu_buffer_rsrc_t, void address_space<3> *, _Constant unsigned int, int, int, int, _Constant int, _Constant int)", [], "vmem-to-lds-load-insts">;
 
+//===----------------------------------------------------------------------===//
+// Global Available/Visible memory accesses.
+//===----------------------------------------------------------------------===//
+
+def __builtin_amdgcn_av_load_b128
+    : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(_ExtVector<4, unsigned int> *, int)", [], "gfx9-insts"> {
+  let Documentation = [DocAVLoadB128];
+}
+def __builtin_amdgcn_av_store_b128
+    : AMDGPUBuiltin<"void(_ExtVector<4, unsigned int> *, _ExtVector<4, unsigned int>, int)", [], "gfx9-insts"> {
+  let Documentation = [DocAVStoreB128];
----------------
shiltian wrote:

We can add `let Args = ...` here to make the document easier to follow, when we mentioned the argument names in the doc string.

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


More information about the cfe-commits mailing list