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

[Sameer Sahasrabuddhe via llvm-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"> {
----------------
ssahasra wrote:

Thanks for catching that! The intrinsics already use `FlatGlobalInsts`, so updated the builtins to check the same. This is a reasonable proxy feature when we want to say "support cache policy bits from gfx9 onwards for global and flat". Cache policy bits existed prior to gfx9 for flat insts, but there's no real reason to extend support earlier than gfx9, which introduced those bits for global insts.

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