[llvm] [AMDGPU] Make globally-addressable-scratch opt-in (PR #189555)

[Pierre van Houtryve via llvm-commits]

================
@@ -1292,9 +1292,16 @@ defm XF32Insts : AMDGPUSubtargetFeature<"xf32-insts",
   "v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32"
 >;
 
-defm GloballyAddressableScratch : AMDGPUSubtargetFeature<"globally-addressable-scratch",
-  "FLAT instructions can access scratch memory for any thread in any wave",
-  /*GenPredicate=*/0
+def FeatureGloballyAddressableScratchSupport : SubtargetFeature<"globally-addressable-scratch-support",
----------------
Pierre-vh wrote:

What is the right way to add such an attribute that'd be propagated from kernels to all callees?
Do I just add something like `amdgpu-globally-addressable-scratch` in `AMDGPUAttributes.def` ?

I'd also need the feature to be removed if GAS is not supported I assume. Does that happen in the attributor?

I am a bit worried about making this an attribute because it affects the memory model. e.g. if a kernel A doesn't use GAS, and a kernel B uses it, then B cannot acquire values from A.

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