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