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