[PATCH] D74270: [mlir][GPUToSPIRV] Modify the lowering of gpu.block_dim to be consistent with Vulkan SPEC

[Stephan Herhut via Phabricator via llvm-commits]

herhut added inline comments.


================
Comment at: mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp:261
   }
-  int32_t index = 0;
   if (dimAttr.getValue() == "x") {
+    return 0;
----------------
Maybe llvm::StringSwitch?


================
Comment at: mlir/test/Conversion/GPUToSPIRV/builtins.mlir:83
       attributes {gpu.kernel} {
-      // CHECK: [[ADDRESS:%.*]] = spv._address_of [[WORKGROUPSIZE]]
-      // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
-      // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
+      // The constant value is obtained fomr the command line option above.
+      // CHECK: spv.constant 32 : i32
----------------
nit: from 


================
Comment at: mlir/test/Conversion/GPUToSPIRV/builtins.mlir:96
+    %c0 = constant 1 : index
+    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_y", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
+    return
----------------
I think we should either require that the `gpu.launch` has the same constant argument than what the SPIR-V lowering uses or update the description of the operation to state that the provided dimensions are just guidelines. 

I do not understand where the need to pass the workgroup size comes from. Wouldn't it be sufficient to require that the arguments to the launch are constants? The pass that constructs the launch and earlier passes that do the tiling need to know this value anyway. So they could manifest them in IR as constant arguments to launch.


================
Comment at: mlir/test/Conversion/GPUToSPIRV/builtins.mlir:116
+  func @builtin() {
+    %c0 = constant 1 : index
+    "gpu.launch_func"(%c0, %c0, %c0, %c0, %c0, %c0) {kernel = "builtin_workgroup_size_z", kernel_module = @kernels} : (index, index, index, index, index, index) -> ()
----------------
Maybe use a constant other than 1 here to differentiate from the default.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D74270