<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/61343>61343</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[MLIR]Crashed on 'convert-vector-to-gpu' with '<unknown>:0: error: MMAMatrixType elements must be F16 or F32'
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
BealZephyr
</td>
</tr>
</table>
<pre>
Test on [llvm16](https://github.com/llvm/llvm-project/tree/release/16.x)
steps to reproduce:
```
mlir-opt -convert-vector-to-gpu test.mlir
```
test case:
```
#map0 = affine_map<(d0, d1) -> (d1, d0)>
#map1 = affine_map<(d0, d1, d2) -> (d0, d2)>
#map2 = affine_map<(d0, d1, d2) -> (d1, d2)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
func.func @m16n8k32_int8_row_col_row(%arg0: memref<128x128xi8, 3>, %arg1: memref<128x128xi8, 3>, %arg2: memref<128x128xi32>) {
%cst_0 = arith.constant dense<0> : vector<32x8xi8>
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%cst = arith.constant 0 : i8
%cst0 = arith.constant 0 : i32
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi8, 3>, vector<16x32xi8>
%B = vector.transfer_read %arg1[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi8, 3>, vector<8x32xi8>
%C = vector.transfer_read %arg2[%c0, %c0], %cst0 {in_bounds = [true, true]} : memref<128x128xi32>, vector<16x8xi32>
%D = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B, %C : vector<16x32xi8>, vector<8x32xi8> into vector<16x8xi32>
vector.transfer_write %D, %arg2[%c0, %c0] {in_bounds = [true, true]} : vector<16x8xi32>, memref<128x128xi32>
return
}
```
crash trace:
```
<unknown>:0: error: MMAMatrixType elements must be F16 or F32
mlir-opt: /home/ty/llvm16/mlir/include/mlir/IR/StorageUniquerSupport.h:153: static ConcreteT mlir::detail::StorageUserBase<mlir::gpu::MMAMatrixType, mlir::Type, mlir::gpu::MMAMatrixStorageType, mlir::detail::TypeUniquer>::get(mlir::MLIRContext *, Args...) [ConcreteT = mlir::gpu::MMAMatrixType, BaseT = mlir::Type, StorageT = mlir::gpu::MMAMatrixStorageType, UniquerT = mlir::detail::TypeUniquer, Traits = <>, Args = <llvm::ArrayRef<long>, mlir::Type, llvm::StringRef>]: Assertion `succeeded(ConcreteT::verify(getDefaultDiagnosticEmitFn(ctx), args...))' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: /home/ty/llvm16/llvm/mlir/build/bin/mlir-opt -convert-vector-to-gpu test.mlir
#0 0x0000000000481d7d llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /home/ty/llvm16/llvm/lib/Support/Unix/Signals.inc:567:11
#1 0x000000000048220b PrintStackTraceSignalHandler(void*) /home/ty/llvm16/llvm/lib/Support/Unix/Signals.inc:641:1
#2 0x00000000004805a6 llvm::sys::RunSignalHandlers() /home/ty/llvm16/llvm/lib/Support/Signals.cpp:104:5
#3 0x0000000000482935 SignalHandler(int) /home/ty/llvm16/llvm/lib/Support/Unix/Signals.inc:412:1
#4 0x00007fbcf86b2980 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x12980)
#5 0x00007fbcf75a2e87 raise /build/glibc-CVJwZb/glibc-2.27/signal/../sysdeps/unix/sysv/linux/raise.c:51:0
#6 0x00007fbcf75a47f1 abort /build/glibc-CVJwZb/glibc-2.27/stdlib/abort.c:81:0
#7 0x00007fbcf75943fa __assert_fail_base /build/glibc-CVJwZb/glibc-2.27/assert/assert.c:89:0
#8 0x00007fbcf7594472 (/lib/x86_64-linux-gnu/libc.so.6+0x30472)
#9 0x0000000000c96bb5 mlir::gpu::MMAMatrixType mlir::detail::StorageUserBase<mlir::gpu::MMAMatrixType, mlir::Type, mlir::gpu::MMAMatrixStorageType, mlir::detail::TypeUniquer>::get<llvm::ArrayRef<long>, mlir::Type, llvm::StringRef>(mlir::MLIRContext*, llvm::ArrayRef<long>, mlir::Type, llvm::StringRef) /home/ty/llvm16/mlir/include/mlir/IR/StorageUniquerSupport.h:152:5
#10 0x0000000000c3b3c9 mlir::gpu::MMAMatrixType::get(llvm::ArrayRef<long>, mlir::Type, llvm::StringRef) /home/ty/llvm16/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp:55:10
#11 0x0000000003219152 convertTransferReadOp(mlir::vector::TransferReadOp, llvm::DenseMap<mlir::Value, mlir::Value, llvm::DenseMapInfo<mlir::Value, void>, llvm::detail::DenseMapPair<mlir::Value, mlir::Value>>&) /home/ty/llvm16/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp:484:7
#12 0x0000000003218bcc mlir::convertVectorToMMAOps(mlir::Operation*) /home/ty/llvm16/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp:989:5
#13 0x0000000003224446 (anonymous namespace)::ConvertVectorToGPUPass::runOnOperation() /home/ty/llvm16/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp:1071:3
#14 0x0000000003a697a7 mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /home/ty/llvm16/mlir/lib/Pass/Pass.cpp:471:21
#15 0x0000000003a69dbd mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /home/ty/llvm16/mlir/lib/Pass/Pass.cpp:534:16
#16 0x0000000003a6b6fc mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /home/ty/llvm16/mlir/lib/Pass/Pass.cpp:839:10
#17 0x0000000003a6b621 mlir::PassManager::run(mlir::Operation*) /home/ty/llvm16/mlir/lib/Pass/Pass.cpp:819:60
#18 0x0000000003a61f82 performActions(llvm::raw_ostream&, bool, bool, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, bool, bool) /home/ty/llvm16/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:91:17
#19 0x0000000003a61bdf processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, bool, bool, bool, bool, bool, bool, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, llvm::ThreadPool*) /home/ty/llvm16/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:139:12
#20 0x0000000003a61998 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool)::$_0::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) const /home/ty/llvm16/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:181:12
#21 0x0000000003a6188d mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool)::$_0>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) /home/ty/llvm16/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#22 0x0000000003b93ae9 llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) const /home/ty/llvm16/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#23 0x0000000003b930c5 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) /home/ty/llvm16/mlir/lib/Support/ToolUtilities.cpp:28:12
#24 0x0000000003a608c3 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool) /home/ty/llvm16/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:186:10
#25 0x0000000003a60a5f mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::PassPipelineCLParser const&, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool, bool) /home/ty/llvm16/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:209:10
#26 0x0000000003a616a1 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&, bool) /home/ty/llvm16/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:306:14
#27 0x00000000004116d9 main /home/ty/llvm16/mlir/tools/mlir-opt/mlir-opt.cpp:271:7
#28 0x00007fbcf7585c87 __libc_start_main /build/glibc-CVJwZb/glibc-2.27/csu/../csu/libc-start.c:344:0
#29 0x00000000004113ba _start (/home/ty/llvm16/llvm/mlir/build/bin/mlir-opt+0x4113ba)
Aborted (core dumped)
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzsWltz4ygW_jXkhYpLAl0f8uDY8WxvdSqpJD0P86JCCNtsy6AF1LH__RZIsi5xrpOZ6dmaqZ5IRofDd-4HJKI13wjGLkB4CcLlGanNVqqLS0bK31i1PaizXBaHiwemDZQCgvCyLH_s_AiES4CSrTGVBngO0Aqg1YabbZ3PqNwBtLJk7eW8UvI_jBqAVkYxBtBKsZIRbe_8aLYHKAXeEnjz5q82rNLQSKhYpWRRU2ZXaAgir_3nfu5Krs5lZeA5leIHU-b8B6NGqnMjzzdVDQ3TZmaJTs62TyG1MJ5hDxDekcqDAC8hWa-5YNmOVAAvAEoKD6AFLHyAUngO8BW0Y74b86w8-GrIw3-FxwIWaMTJO45NOKH3cvKf44Q_hskfcXJ_17WgM_sHgsDb-ZFIvmOUcWGSTMnHjMrSXgFKAAqJ2ngAz-GO7RRbA7zwUbK3__PEsseWNVrAhtJ_MyU6TYmRo0ohiC8brNBOoNpkrVkVN9sZlUIbIgwsmLDusPCczHgOG38CeIHR3q3cSd7wOcnEczO5KNh-SOufovVP02rzAuNkTPkSBoyGdnIz5o68kWtmFBF6zVSmGClgZ57w0snWapd6LtYXR2DxJRdZLmtRaMcLhJdG1cySuGu4BPESvsF0R-X60R6jJ9q9fA2p_-cjTU4BXbwGFL0M1Pt9SFsnH-vzON7DXA5hUimMIrRVUsH2XGxsFjgu3yauBqXLPP0tbuFzwxQxUmXmULHBTFQRRcqSlQC1s06MKFbU1HAp7FDD7zsXRcMF4RanHQJ4QYrCCmN1gMJ5q7rL9roYh-rQm07bDnJh5IvqmhrzUXHDnBKHOeeEVd9jypMA0OJZEzfYFDO1Em1cx8uTlYsqorfQGvj54oYXtfgu5KOwrPHcZWWmlAU0h9fX82tiFN8_HCoGWcl2TBgNd7U2MGdw5UdQKrjq8ktXiO1UgFZbubO13RzaBsCPAFq5OoxWXNCyLlg_8OUOoNW9kYps2DfB_1szdV9XlVRmtgV47ofYctWGGE7hQgqqmGEP0M3Gc4DnBTOEl819x0czdekq-6Kn21R1czMSzin8SPN05Omsdo2npEMg9mkrTaNfy4oZgJKe_vrrl7uFFIbtDQTIefVcbfRsNnMlK7zspbWO9BZJrNBT6u5Zh_tVZhMBWzGm856TFi3ggyLctN6PF61XW9G6Idccuplzpcjhznl7KcWmi4An4PsZ90ZxsXFTrmww4Tmca82UTSUQRJ6uKWWsYAVAyVGBzdQfTPH1AaBkw8ySrUldmiUnGyG14fRqx81KAJRQ4zpStIDkaA33L4ZrwktWzBqnv_16Nb-_grrOd9xAAvN6Y1tWqYxtXt_dGnOta6YBWkEiCtiGCTRbBptozgn97iK6Xf7eEPodFvWuOoa4NwNeeqvkRpGdBV-7qH0pKFsobSjmNS8Le-WiHXxHd22Ttge9vXf8L0j8Ii4GptMH3dzcKi6ME-DB5SiU9ESKPGZSG8XIDqDIlRlhXEC8LEPJc5tImtQB0Oqb4Hs7wDeClHrGBQV4HkaxzSl-D9mfQEbIy-EEX8PjX0QUpXXw5IfkhQvYzwEVBbbRHWBCE0xeSKJTaryrxQiadj32u0F1aGhlXcn3AquoHg2eaijFIZyq5BNtFPhorI6gBRCvc7pOohyliQezTDFtpGKZstkzOXLfJ1EWBeclF_X-fCPq5kFltrYhm2k58wC69Pa-5XLcedplwuEycUgQS2KoCNe28B9jY1PynJ4vfv3342_58SeaoRiglXZyALSyOWOlD7pglY3ouhFTH_QPh0bU9qdjPXNe6bsCfEQSTZAE8dqHJJdO0rchMUWjDTfLLZJMFonHi6QBXhOYZcTl0sxmuiwnb5e9mXe8aZZMx0sm0yWDGL3BdtRaLXJWw14Qo5HV0pF30jTK8_DVUvl37R4-sXI-04i0fcjnLPNCQvhwN4j65GSz97jiUJxjmr7eKQ3asT9L0sbBl5yUTbn_5fZbJ-4vt9_a8TYFh6HLw72UoyKFkZ_6IYJtVX5o9yh3jBQ31ciu3QbDoZ-QDeVYMqHZtTsC6if_Ssp64qLHoadTv4i1PD3dFctGl_20obN3LG6JnfsWBDYkrlxz8Ea1L5yqtNtsrn51WnmQjQUGv1rtB4ktgHGvfTTRfpJTOsDU2qFjdH09v7FZf2CHm8ruk93qLzYNH8eculQ7iAs8xoyCIIhsqiVCisNO1hoKsmO6ct1X2sBcjAX55fbbLdFtq6FqcSMGcrzYZ3xcDt-LbZnCvSDBSBASpTGJn8maN9WDvKks5nlBqqPnq1qMrOGEatLccyYaPpkLUh4019dEkE2zv8mlLO21Fu7wvHitQx3po1neXTqHczIjvxc6nApd5MW7hb7lFSu5YBNXtKRHWaJP18Jwkl3qi9BGuY2IVE-4jgmalZsHLfZbopgwNrlAd7b5ngB6qugQ28j2o17R0UTRebSmE3ydxEe12gbnpfh-UW0fxp7gdFIU4ifYkf8K9t-flU4A8y2wwTsTP5kA89cJghVTa6l2c3fkp1_Z8XWu1V21Kdqtz5YoVmSVUaNe6F7WirLrje2WOkeZuPZLPc66Fg5W1hy79ZO-yg2npLxjui7dLuO0el0lamvcGPobtfogZakHW26AVtclVzeVuSZcdDne7RIHdSmd6jkv1rBSkjKtL-v12m3MXlLzUa2167SeqPWa7aQ6dKwG9EVzeJIVrGSGPT-nLdRPDfra9Y-xTU_R9lt3bMO1UYdWH_2qD263eOvQvD043mRGvwlkdLQj8qZ2TNNk6Lk9i5_Gnn-Nfd7rRW0MNiwBCjKvuZVV896iaWMASv4KxU3NlzaJ6zM9zW35h57mTz0tSYZ9xdhkH7HxT6HI45aZkrLMCf2ercUI9j8R9ZkRdWV1aDfMf5G-ngbSK6eQ_aFDOzBfPgC0un_4umpVT8qrvVFEuzOHIJxE0Xg_mKeYsPT_IFr-llnxYyaNkolJ8dSkHh2eIuqq5GYuittJc_Vzp4efwn6vUnywZe5fItha-M3wkhvOuk0Jmhp4eozgJRT_02N9XkX41LYlGu900fQ4xCPh-ic33tgQ3XHG4ustUZqpZ_aon2mRP8IyyJucQaDp-YkfEf9Zy7TnQ3RLmpOg-bOn6W9VymcKhz3ndkEvXDx-8en7UZHCHeHitVXN0_W6206T7siv38ujyauxJKRJDLOs5DnNtCHKZN3Cb3kZR3XdvYVsbt1Dx8e9mMNB0L-Zs8unU1FxTmCzcPuC7sOfD7j3dg3H44u7eS6VYYVlTaVi7isGVvTfAY8_VTorLnCR4pScsQs_ipMQBzhCZ9sL7IckpWGOEclZwhBKUr9wDZkfrAmiZ_wCeQh72Pd9H6MAzWiSFB7xKaH-2o9ZAAKP7QgvZ1aSmVSbM_cZxkXk4wCflSRnpb5ovmkT7BG6h833amfqwn3AkdcbDQKv5NronovhpnTfVF9__XIHwuVCEb1lhfuEGsUnv6cAKIaP3GwtwWd9nIXis1qVFx__FsXp4X8BAAD__0UB7NY">