[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:

Ok, yes, it _could_, but default-preserve is a more reasonable behavior on trivial argument copies than default-drop. If I've stuck a tag - *especially* one starting with something like `llvm.`, `nvvm.`, `rocdl.`  and so on on a function argument to a `gpu.func`, and then said `gpu.func` becomes an `llvm.func`, the semantics of said tag almost never change (unless the argument type does).

Yes, in principle there should be some sort of interface for querying this sort of thing.

Pragmatically, I claim we should just land copy over argument attributes for unchanged types because it's most likely what people want. (If there _is_ an attribute that needs to be modified during `gpu.func` => `llvm.func`, you can run the conversion pattern and then clean it up afterwards - see existing handling for maximum workitem ID ranges over in ROCDL.)

Or, alternatively ... not copying would require some example of something that shouldn't be copied.

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