[clang] [flang] [llvm] [AMDGPU] Add "lds-buffer-load-insts" attribute for all targets < gfx11 (PR #133055)

[Juan Manuel Martinez Caamaño via llvm-commits]

jmmartinez wrote:

> At the hardware level, GFX11 removed the ability for buffer, **scratch** and **global** instructions to return directly to LDS. So can we use one attribute that covers all three of those?

I've been "grep"ping in upstream (but not yet downstream):
* Scratch load lds are there after gfx942, but we have no llvm-ir builtin nor any other way than inline assembly to generate these.
* Global load lds is available on gfx9 and gfx10. However, the `__builtin_amdgcn_global_load_lds` builtin is restricted to platforms with the attribute `gfx940-insts`; while it could be available on gfx90a for example.

A single attribute doesn't match all the platforms where these instructions are.

What do you think about having a common attribute for buffer and global load lds (`mem-to-lds-loads`?):
* The advantage is that it would make __builtin_amdgcn_global_load_lds available on all gfx9 and gfx10 targets.
* The down-side is that it would restrict __builtin_amdgcn_raw_ptr_buffer_load_lds to gfx9 and gfx10; and miss gfx6,7,8 where it could be supported.
* In the future, `__builtin_amdgcn_scratch_load_lds` could use the `gfx940-insts` attribute

What do you think about this?


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