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

[Shilei Tian via cfe-commits]

================
@@ -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"> {
----------------
shiltian wrote:

I'd prefer to use a dedicated feature instead of `gfxXYZ-insts`. That would be more future-proof in case a future target doesn't fully support `gfx9-insts` (for example, not a full set to qualify for that feature) but still supports this specific functionality.

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