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

At the high level, the question is "if I am a dialect attribute on a `memref` argument/result, and that `memref` is being lowered to a bare pointer, can I be copied to that pointer argument?"

Which isn't actually all that GPU-specific and should be logic that happens in the CPU pipeline as well.

And it is sort of a dialect interface - when you have that string of the form `dialect.foo`, there are already optional hooks you can implement on `dialect` to let it do some verifications. This would be a weirdly specific dialect interface - it'd fundamentally only get implemented by `llvm` and `nvvm`, - but it's an option.

If we're OK with minor abuses, this _could_ be a hook on `LLVMTranslationDialectInterface`

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