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

[Guray Ozen]

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

> How do you think this should look?

We've two options:

**Option 1:** Add nvgpu.grid_constant as an NVGPU arg attribute, and in the nvgpu-to-nvvm lowering translate it to nvvm.grid_constant (or the equivalent NVVM/LLVM arg attr). But you should


**Option 2:** Model it in the type system as !nvgpu.grid_constant<T> plus an nvgpu.unwrap_grid_constant op to get T. Then nvgpu-to-nvvm lowers the wrapper/unwrap and emits the desired NVVM/LLVM arg attributes.

```
gpu.func @k(%arg_ptr : !nvgpu.grid_constant<!llvm.ptr>) {
   
    %p = nvgpu.unwrap_grid_constant %arg_ptr : !nvgpu.grid_constant<!llvm.ptr> -> !llvm.ptr
   
  }

```

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