[Mlir-commits] [mlir] [MLIR][GPU] Support grid constant, byval, byref on gpu.func (PR #172037)

[Krzysztof Drewniak]

================
@@ -369,7 +370,10 @@ GPUFuncOpLowering::matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, OpAdaptor adaptor,
     }
 
     if (lowersToPointer) {
+      copyPointerAttribute(mlir::NVVM::NVVMDialect::getGridConstantAttrName());
----------------
krzysz00 wrote:

@grypp 
> We've two options:

We have other options. The one I'm arguing for is to allow `nvvm.grid_constant` argument attributes to be copied (either as a default of how argument attributes on trivially-translated types are handled or by some interface) from `gpu.func` to `llvm.func`

... and, now that I think about it, I think there's a general point here, which is that function argument and result attributes should be preserved when the type and usage of a function argument/result is unchanged. (Heck, `builtin.module` should grow a `config` / `annotations` / `bikeshed` attribute dictionary that has those same "preserve me" semantics)

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