26 lines · plain
1// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s2 3// -----4 5llvm.func @tma_reduce_0d(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %ch : i64) {6 // expected-error @below {{expects coordinates between 1 to 5 dimension}}7 nvvm.cp.async.bulk.tensor.reduce %tma_desc, %src, box[] {redKind = #nvvm.tma_redux_kind<add>}: !llvm.ptr, !llvm.ptr<3>8 llvm.return9}10 11// -----12 13llvm.func @tma_reduce_2d_im2col(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %d0 : i32, %d1 : i32, %ch : i64) {14 // expected-error @below {{to use im2col mode, the tensor has to be at least 3-dimensional}}15 nvvm.cp.async.bulk.tensor.reduce %tma_desc, %src, box[%d0, %d1] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col>}: !llvm.ptr, !llvm.ptr<3>16 llvm.return17}18 19// -----20 21llvm.func @tma_store_reduce_scatter(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %d0 : i32, %ch : i64) {22 // expected-error @below {{Scatter mode unsupported for CpAsyncBulkTensorReduceOp}}23 nvvm.cp.async.bulk.tensor.reduce %tma_desc, %src, box[%d0] {redKind = #nvvm.tma_redux_kind<add>, mode = #nvvm.tma_store_mode<tile_scatter4>} : !llvm.ptr, !llvm.ptr<3>24 llvm.return25}26