[mlir] Add vector op support to cuda-runner including vector.print

[herhut]

mlir/test/mlir-cuda-runner/gpu-to-cubin.mlir

19

Could you instead have an additional test case?

[ThomasRaoux]

mlir/test/mlir-cuda-runner/gpu-to-cubin.mlir

19

I added a vector.print without changing the rest of the test. Does that make sense to you or do you think it should be a different test? I didn't to copy paste this test but if you think it is clearer I can change it.

[ThomasRaoux]

rebase.

[kiranchandramohan]

Failing a buildbot https://lab.llvm.org/buildbot/#/builders/33/builds/3150.

ConvertLaunchFuncToRuntimeCalls.cpp:(.text._ZN12_GLOBAL__N_123GpuToLLVMConversionPass14runOnOperationEv+0x84): undefined reference to `mlir::populateVectorToLLVMConversionPatterns(mlir::LLVMTypeConverter&, mlir::OwningRewritePatternList&, bool, bool)'

[ThomasRaoux]

In D97346#2636065, @kiranchandramohan wrote:

Failing a buildbot https://lab.llvm.org/buildbot/#/builders/33/builds/3150.

ConvertLaunchFuncToRuntimeCalls.cpp:(.text._ZN12_GLOBAL__N_123GpuToLLVMConversionPass14runOnOperationEv+0x84): undefined reference to `mlir::populateVectorToLLVMConversionPatterns(mlir::LLVMTypeConverter&, mlir::OwningRewritePatternList&, bool, bool)'

Should be fixed by: https://github.com/llvm/llvm-project/commit/44f24f3996e8a32d0bb3d6d79a66643c36f088da

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToRuntimeCalls.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
 
 #include "../PassDetail.h"
 #include "mlir/Conversion/AsyncToLLVM/AsyncToLLVM.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
+#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 
 #include "llvm/ADT/STLExtras.h"

 };
 } // namespace
 
 void GpuToLLVMConversionPass::runOnOperation() {
   LLVMTypeConverter converter(&getContext());
   OwningRewritePatternList patterns;
   LLVMConversionTarget target(getContext());
 
+  populateVectorToLLVMConversionPatterns(converter, patterns);
   populateStdToLLVMConversionPatterns(converter, patterns);
   populateAsyncStructuralTypeConversionsAndLegality(&getContext(), converter,
                                                     patterns, target);
   populateGpuToLLVMConversionPatterns(converter, patterns, gpuBinaryAnnotation);
 
   if (failed(
           applyPartialConversion(getOperation(), target, std::move(patterns))))
     signalPassFailure();

mlir/test/mlir-cuda-runner/gpu-to-cubin.mlir

func @other_func(%arg0 : f32, %arg1 : memref<?xf32>) {
   gpu.launch blocks(%bx, %by, %bz) in (%grid_x = %cst, %grid_y = %cst, %grid_z = %cst)
              threads(%tx, %ty, %tz) in (%block_x = %cst2, %block_y = %cst, %block_z = %cst) {
     store %arg0, %arg1[%tx] : memref<?xf32>
     gpu.terminator
   }
   return
 }
 
-// CHECK: [1, 1, 1, 1, 1]
+// CHECK: ( 1, 1, 1, 1, 1 )
 func @main() {
+  %v0 = constant 0.0 : f32
+  %c0 = constant 0: index
   %arg0 = alloc() : memref<5xf32>
   %21 = constant 5 : i32
   %22 = memref_cast %arg0 : memref<5xf32> to memref<?xf32>
   %23 = memref_cast %22 : memref<?xf32> to memref<*xf32>
   gpu.host_register %23 : memref<*xf32>
-  call @print_memref_f32(%23) : (memref<*xf32>) -> ()
+  %val = vector.transfer_read %arg0[%c0], %v0: memref<5xf32>, vector<5xf32>
+  vector.print %val: vector<5xf32>
   %24 = constant 1.0 : f32
   call @other_func(%24, %22) : (f32, memref<?xf32>) -> ()
-  call @print_memref_f32(%23) : (memref<*xf32>) -> ()
+  %val1 = vector.transfer_read %arg0[%c0], %v0: memref<5xf32>, vector<5xf32>
+  vector.print %val1: vector<5xf32>
   return
 }
-
-func private @print_memref_f32(%ptr : memref<*xf32>)

mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp

 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
 #include "mlir/Dialect/Async/IR/Async.h"
 #include "mlir/Dialect/Async/Passes.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/GPU/Passes.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
+#include "mlir/Dialect/Vector/VectorOps.h"
 #include "mlir/ExecutionEngine/JitRunner.h"
 #include "mlir/ExecutionEngine/OptUtils.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
 #include "mlir/Target/LLVMIR.h"
 #include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
 #include "mlir/Target/LLVMIR/Export.h"
int main(int argc, char **argv) {
   mlir::initializeLLVMPasses();
 
   mlir::JitRunnerConfig jitRunnerConfig;
   jitRunnerConfig.mlirTransformer = runMLIRPasses;
 
   mlir::DialectRegistry registry;
   registry.insert<mlir::LLVM::LLVMDialect, mlir::NVVM::NVVMDialect,
                   mlir::async::AsyncDialect, mlir::gpu::GPUDialect,
-                  mlir::StandardOpsDialect>();
+                  mlir::vector::VectorDialect, mlir::StandardOpsDialect>();
   registry.addDialectInterface<NVVM::NVVMDialect,
                                NVVMDialectLLVMIRTranslationInterface>();
   mlir::registerLLVMDialectTranslation(registry);
 
   return mlir::JitRunnerMain(argc, argv, registry, jitRunnerConfig);
 }