[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 "../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"

   OwningRewritePatternList patterns;
   LLVMConversionTarget target(getContext());
 
+  populateVectorToLLVMConversionPatterns(converter, patterns);
   populateStdToLLVMConversionPatterns(converter, patterns);
   populateAsyncStructuralTypeConversionsAndLegality(&getContext(), converter,
                                                     patterns, target);

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

 }
 
 // CHECK: [1, 1, 1, 1, 1]
+// CHECK: ( 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>

   %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<2xf32>, vector<5xf32>
+  vector.print %val1: vector<2xf32>
   return
 }
 

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

 #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"

   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);