diff options
| author | MLIR Team <no-reply@google.com> | 2019-11-04 12:36:04 -0800 |
|---|---|---|
| committer | A. Unique TensorFlower <gardener@tensorflow.org> | 2019-11-04 12:36:37 -0800 |
| commit | 1f43d0d000312c1f004cd76302a8be1636fba413 (patch) | |
| tree | f9702e16793b0a48e2161666ba8b63b501e2b45b /mlir/test/Target | |
| parent | e4a912eb5a688dcb39036f8a809fc37977beaf1d (diff) | |
| download | bcm5719-llvm-1f43d0d000312c1f004cd76302a8be1636fba413.tar.gz bcm5719-llvm-1f43d0d000312c1f004cd76302a8be1636fba413.zip | |
[NVVM] Add mma.sync operation.
PiperOrigin-RevId: 278440547
Diffstat (limited to 'mlir/test/Target')
| -rw-r--r-- | mlir/test/Target/nvvmir.mlir | 9 |
1 files changed, 9 insertions, 0 deletions
diff --git a/mlir/test/Target/nvvmir.mlir b/mlir/test/Target/nvvmir.mlir index fea83e01fb4..2e63ecd68bc 100644 --- a/mlir/test/Target/nvvmir.mlir +++ b/mlir/test/Target/nvvmir.mlir @@ -64,6 +64,15 @@ llvm.func @nvvm_vote(%0 : !llvm.i32, %1 : !llvm.i1) -> !llvm.i32 { llvm.return %3 : !llvm.i32 } +llvm.func @nvvm_mma(%a0 : !llvm<"<2 x half>">, %a1 : !llvm<"<2 x half>">, + %b0 : !llvm<"<2 x half>">, %b1 : !llvm<"<2 x half>">, + %c0 : !llvm.float, %c1 : !llvm.float, %c2 : !llvm.float, %c3 : !llvm.float, + %c4 : !llvm.float, %c5 : !llvm.float, %c6 : !llvm.float, %c7 : !llvm.float) { + // CHECK: call { float, float, float, float, float, float, float, float } @llvm.nvvm.mma.m8n8k4.row.row.f32.f32 + %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="row", blayout="row"} : (!llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float) -> !llvm<"{ float, float, float, float, float, float, float, float }"> + llvm.return %0 : !llvm<"{ float, float, float, float, float, float, float, float }"> +} + // This function has the "kernel" attribute attached and should appear in the // NVVM annotations after conversion. llvm.func @kernel_func() attributes {gpu.kernel} { |
