[PATCH] D74012: [mlir][spirv] Use spv.entry_point_abi in GPU to SPIR-V conversions

[Lei Zhang via Phabricator via llvm-commits]

antiagainst marked 5 inline comments as done.
antiagainst added a comment.

@herhut: +1 to what Mahesh said. Additionally, I'd like to tighten SPIR-V side to use attributes in general for passing in pattern configurations.

There is mismatch between GPU dialect and the SPIR-V side. For this one (`spv.entry_point_abi = {local_size = ...}`) probably we can push one layer upwards and propose a `gpu.workgroup_size` at the GPU dialect level; then the SPIR-V lowering can convert `gpu.workgroup_size` to `spv.entry_point_abi = {local_size = ...}`.  But as explained in https://llvm.discourse.group/t/using-attributes-to-specify-workgroup-configuration-when-lowering-to-gpu/496/20, for `gpu.launch` with non-constant workgroup sizes, we need to specify the `SpecId`s for them, which does not really make sense to appear at GPU dialect level; so likely we need the SPIR-V lowering path user (say, IREE) to attach something like `spv.entry_point_abi = {local_size_spec_id = ...}` before going to GPU dialect level and pass it all the way down the stack. (You can think this as part of the SPIR-V target in lieu of a proper target mechanism in SPIR-V. We have SPIR-V conversion target on SPIR-V side but that's covering different things than the ABI here. Complexities. ;-P) Again, we need to attach a SPIR-V specific attribute to the `gpu.func` eventually. I'd like to have consistency between the normal constant case and spec constant case.



================
Comment at: mlir/include/mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h:26
                                 SPIRVTypeConverter &typeConverter,
-                                OwningRewritePatternList &patterns,
-                                ArrayRef<int64_t> workGroupSize);
+                                OwningRewritePatternList &patterns);
 } // namespace mlir
----------------
mravishankar wrote:
> Can we leave the workgroup size as optional. If provided will be used for overriding the default way of using the attribute.
We probably don't want to do that. One of the purpose is to tighten the contract on SPIR-V lowerings and make them consistent. Having two ways is causing more confusion IMO.


================
Comment at: mlir/lib/Dialect/SPIRV/TargetAndABI.cpp:49
+spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) {
+  while (op && !op->hasTrait<OpTrait::FunctionLike>())
+    op = op->getParentOp();
----------------
mravishankar wrote:
> Cant you just do op->getParentOfType<FuncOp>(op) / op->getParentOfType<gpu::FuncOp>(op).
I think this is simpler and more composable given this is a utility for writing SPIR-V lowerings. There may exist some other funcs that one would like to lower towards SPIR-V so I think we can be a bit flexible here.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D74012/new/

https://reviews.llvm.org/D74012