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