[Mlir-commits] [mlir] cccedd4 - [MLIR][NVGPU] Enable strict property assembly format (#217272)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Aug 19 04:02:52 PDT 2026
Author: Mehdi Amini
Date: 2026-08-19T13:02:46+02:00
New Revision: cccedd4c719f7782ba8a838d44bdc088bb35d196
URL: https://github.com/llvm/llvm-project/commit/cccedd4c719f7782ba8a838d44bdc088bb35d196
DIFF: https://github.com/llvm/llvm-project/commit/cccedd4c719f7782ba8a838d44bdc088bb35d196.diff
LOG: [MLIR][NVGPU] Enable strict property assembly format (#217272)
Enable strict property assembly format mode for the NVGPU dialect. Spell
ldmatrix, MMA, async copy, async wait, and warpgroup MMA properties
directly in the affected declarative assembly formats.
Refresh NVGPU dialect, conversion, integration, and Python tests so
those properties use direct syntax while ordinary attributes remain in
attr-dict.
Assisted-by: Codex
Added:
Modified:
mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
mlir/test/Conversion/NVGPUToNVVM/nvgpu-extf.mlir
mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm-mma-fp8.mlir
mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
mlir/test/Conversion/NVGPUToNVVM/nvgpu-truncf.mlir
mlir/test/Conversion/VectorToGPU/fold-arith-vector-to-mma-ops-mma-sync.mlir
mlir/test/Conversion/VectorToGPU/vector-to-mma-ops-mma-sync.mlir
mlir/test/Dialect/MemRef/extract-address-computations.mlir
mlir/test/Dialect/NVGPU/fold-memref-alias-ops.mlir
mlir/test/Dialect/NVGPU/invalid.mlir
mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32.mlir
mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32x3.mlir
mlir/test/Dialect/NVGPU/nvgpu-extf-invalid.mlir
mlir/test/Dialect/NVGPU/nvgpu-truncf-invalid.mlir
mlir/test/Dialect/NVGPU/optimize-shared-memory.mlir
mlir/test/Dialect/NVGPU/roundtrip.mlir
mlir/test/Dialect/NVGPU/transform-create-async-groups.mlir
mlir/test/Dialect/NVGPU/transform-matmul-to-nvvm.mlir
mlir/test/Dialect/NVGPU/transform-pipeline-shared.mlir
mlir/test/Examples/NVGPU/Ch3.py
mlir/test/Examples/NVGPU/Ch4.py
mlir/test/Examples/NVGPU/Ch5.py
mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sparse-mma-2-4-f16.mlir
mlir/test/Integration/GPU/CUDA/TensorCore/sm80/transform-mma-sync-matmul-f32.mlir
mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir
mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir
mlir/test/python/dialects/nvgpu.py
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
index 1c0d7bd1113ea..d43a8c38fd1e0 100644
--- a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
+++ b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
@@ -30,6 +30,7 @@ def NVGPU_Dialect : Dialect {
let useDefaultTypePrinterParser = 1;
let useDefaultAttributePrinterParser = 1;
+ let useStrictPropertiesInAssemblyFormat = 1;
let extraClassDeclaration = [{
/// Return true if the given MemRefType has an integer address
diff --git a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
index 69b6b57de54b8..87e525de1c184 100644
--- a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
+++ b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPUOps.td
@@ -48,7 +48,7 @@ def NVGPU_LdMatrixOp : NVGPU_Op<"ldmatrix", [
Example:
```mlir
- %0 = nvgpu.ldmatrix %sm[%c0, %c0] {numTiles = 4 : i32, transpose = false} :
+ %0 = nvgpu.ldmatrix %sm[%c0, %c0] numTiles = 4 transpose = false :
memref<?x?xf16, 3> -> vector<4x2xf16>
```
}];
@@ -58,7 +58,9 @@ def NVGPU_LdMatrixOp : NVGPU_Op<"ldmatrix", [
I32Attr:$numTiles);
let results = (outs AnyVectorOfNonZeroRank:$res);
let assemblyFormat = [{
- $srcMemref`[` $indices `]` attr-dict `:` type($srcMemref) `->` type($res)
+ $srcMemref`[` $indices `]` `numTiles` `=` $numTiles `transpose` `=`
+ $transpose attr-dict `:`
+ type($srcMemref) `->` type($res)
}];
let hasVerifier = 1;
@@ -98,7 +100,7 @@ def NVGPU_MmaSyncOp : NVGPU_MmaSyncOp<"mma.sync"> {
Example:
```mlir
- %res = nvgpu.mma.sync (%matrixA, %matrixB, %matrixC) {mmaShape = [16, 8, 16]} :
+ %res = nvgpu.mma.sync (%matrixA, %matrixB, %matrixC) mmaShape = [16, 8, 16] :
(vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
```
}];
@@ -123,7 +125,8 @@ def NVGPU_MmaSyncOp : NVGPU_MmaSyncOp<"mma.sync"> {
];
let assemblyFormat = [{
- `(` $matrixA`,` $matrixB`,` $matrixC `)` attr-dict
+ `(` $matrixA`,` $matrixB`,` $matrixC `)` `mmaShape` `=` $mmaShape
+ (`tf32Enabled` $tf32Enabled^)? attr-dict
`:` `(` type($matrixA) `,` type($matrixB) `,` type($matrixC) `)` `->` type($res)
}];
@@ -153,7 +156,7 @@ def NVGPU_MmaSparseSyncOp : NVGPU_MmaSyncOp<"mma.sp.sync"> {
Example (targetingthe f16 16x8x32 `mma.sp` PTX instruction):
```mlir
- nvgpu.mma.sp.sync (%a, %b, %c) metadata (%meta) {mmaShape = [16, 8, 32]} :
+ nvgpu.mma.sp.sync (%a, %b, %c) metadata (%meta) mmaShape = [16, 8, 32] :
(vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
```
}];
@@ -178,7 +181,9 @@ def NVGPU_MmaSparseSyncOp : NVGPU_MmaSyncOp<"mma.sp.sync"> {
];
let assemblyFormat = [{
- `(` $matrixA`,` $matrixB`,` $matrixC `)` `metadata` `(` $sparseMetadata `)` attr-dict
+ `(` $matrixA`,` $matrixB`,` $matrixC `)` `metadata` `(` $sparseMetadata `)`
+ `mmaShape` `=` $mmaShape (`sparsitySelector` `=` $sparsitySelector^)?
+ (`tf32Enabled` $tf32Enabled^)? attr-dict
`:` `(` type($matrixA) `,` type($matrixB) `,` type($matrixC) `)` `->` type($res)
}];
@@ -246,8 +251,9 @@ def NVGPU_DeviceAsyncCopyOp : NVGPU_Op<"device_async_copy", [
Optional<Index>:$srcElements,
OptionalAttr<UnitAttr>:$bypassL1);
let assemblyFormat = [{
- $src `[` $srcIndices `]` `,` $dst `[` $dstIndices `]` `,` $dstElements (`,` $srcElements^)?
- attr-dict `:` type($src) `to` type($dst)
+ $src `[` $srcIndices `]` `,` $dst `[` $dstIndices `]` `,`
+ $dstElements (`,` $srcElements^)? (`bypassL1` $bypassL1^)?
+ attr-dict `:` type($src) `to` type($dst)
}];
let hasVerifier = 1;
}
@@ -300,7 +306,7 @@ def NVGPU_DeviceAsyncWaitOp : NVGPU_Op<"device_async_wait", []> {
let arguments = (ins NVGPU_DeviceAsyncToken:$asyncDependencies,
OptionalAttr<I32Attr>:$numGroups);
let assemblyFormat = [{
- $asyncDependencies attr-dict
+ $asyncDependencies (`numGroups` `=` $numGroups^)? attr-dict
}];
}
@@ -340,7 +346,8 @@ def NVGPU_MBarrierGetOp : NVGPU_Op<"mbarrier.get", []> {
let arguments = (ins NVGPU_MBarrierGroup:$barriers, Index:$mbarId);
let results = (outs AnyTypeOf<[I32, I64]>:$mbarrierPointer);
let assemblyFormat = [{
- $barriers `[` $mbarId `]` attr-dict `:` type($barriers) `->` type($mbarrierPointer)
+ $barriers `[` $mbarId `]` attr-dict `:`
+ type($barriers) `->` type($mbarrierPointer)
}];
}
@@ -452,7 +459,7 @@ def NVGPU_MBarrierTryWaitParityOp : NVGPU_Op<"mbarrier.try_wait.parity", []> {
```
}];
let arguments = (ins NVGPU_MBarrierGroup:$barriers, I1:$phaseParity, Index:$ticks, Index:$mbarId);
- let assemblyFormat = "$barriers `[` $mbarId `]` `,` $phaseParity `,` $ticks attr-dict `:` type($barriers)";
+ let assemblyFormat = "$barriers `[` $mbarId `]` `,` $phaseParity `,` $ticks attr-dict `:` type($barriers)";
}
def NVGPU_TmaFenceOp : NVGPU_Op<"tma.fence.descriptor", []> {
@@ -504,7 +511,7 @@ def NVGPU_TmaAsyncLoadOp : NVGPU_Op<"tma.async.load", [AttrSizedOperandSegments]
`to` $dst
(`multicast_mask` `=` $multicastMask^ )?
(`,` `predicate` `=` $predicate^)?
- attr-dict `:` type($tensorMapDescriptor) `,` type($barriers)
+ attr-dict `:` type($tensorMapDescriptor) `,` type($barriers)
`->` type($dst)
}];
let hasVerifier = 1;
@@ -552,7 +559,8 @@ def NVGPU_TmaCreateDescriptorOp : NVGPU_Op<"tma.create.descriptor", []> {
Variadic<Index>:$boxDimensions);
let results = (outs NVGPU_TensorMapDescriptor:$tensorMap);
let assemblyFormat = [{
- $tensor `box` `[` $boxDimensions `]` attr-dict `:` type($tensor) `->` type($tensorMap)
+ $tensor `box` `[` $boxDimensions `]` attr-dict `:`
+ type($tensor) `->` type($tensorMap)
}];
let hasVerifier = 1;
}
@@ -612,7 +620,9 @@ def NVGPU_WarpgroupMmaOp : NVGPU_Op<"warpgroup.mma"> {
NVGPU_WarpgroupAccumulator:$matrixC);
let results = (outs NVGPU_WarpgroupAccumulator:$matrixD);
let assemblyFormat = [{
- $descriptorA`,` $descriptorB`,` $matrixC attr-dict
+ $descriptorA`,` $descriptorB`,` $matrixC
+ (`waitGroup` `=` $waitGroup^)? (`transposeA` $transposeA^)?
+ (`transposeB` $transposeB^)? attr-dict
`:` type($descriptorA) `,` type($descriptorB) `,` type($matrixC) `->` type($matrixD)
}];
let hasVerifier = 1;
@@ -633,7 +643,8 @@ def NVGPU_WarpgroupMmaStoreOp : NVGPU_Op<"warpgroup.mma.store"> {
Arg<AnyMemRef, "", [MemWrite]>:$dstMemref);
let assemblyFormat = [{
- $matrixD `,` $dstMemref attr-dict `:` type($matrixD) `to` type($dstMemref)
+ $matrixD `,` $dstMemref attr-dict `:`
+ type($matrixD) `to` type($dstMemref)
}];
let hasVerifier = 1;
}
@@ -666,7 +677,7 @@ def NVGPU_RcpOp : NVGPU_Op<"rcp", [Pure,
DefaultValuedAttr<BoolAttr, "false">:$ftz);
let results = (outs VectorOfNonZeroRankOf<[F32]>:$out);
let assemblyFormat = [{
- $in attr-dict `:` type($out)
+ $in prop-dict attr-dict `:` type($out)
}];
let hasVerifier = 1;
}
@@ -713,7 +724,10 @@ def NVGPU_TruncfOp : NVGPU_Op<"truncf", [Pure]> {
Optional<I32>:$random_bits
);
let results = (outs AnyTypeOf<[AnyFloat, VectorOfNonZeroRankOf<[AnyFloat]>]>:$out);
- let assemblyFormat = "$in (`,` $random_bits^)? attr-dict `:` type($in) `to` type($out)";
+ let assemblyFormat = [{
+ $in (`,` $random_bits^)? prop-dict attr-dict
+ `:` type($in) `to` type($out)
+ }];
let hasVerifier = 1;
}
@@ -742,7 +756,9 @@ def NVGPU_ExtfOp : NVGPU_Op<"extf", [Pure]> {
DefaultValuedAttr<FPRoundingModeAttr, "NVVM::FPRoundingMode::RN">:$rnd,
DefaultValuedAttr<BoolAttr, "false">:$relu);
let results = (outs AnyTypeOf<[AnyFloat, VectorOfNonZeroRankOf<[AnyFloat]>]>:$out);
- let assemblyFormat = "$in attr-dict `:` type($in) `to` type($out)";
+ let assemblyFormat = [{
+ $in prop-dict attr-dict `:` type($in) `to` type($out)
+ }];
let hasVerifier = 1;
}
diff --git a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-extf.mlir b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-extf.mlir
index 8c91cfa963fa9..f0b535a994516 100644
--- a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-extf.mlir
+++ b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-extf.mlir
@@ -431,7 +431,7 @@ func.func @fpext_v3x1_f8_to_f16(%in : vector<3x1xf8E4M3FN>) -> vector<3x1xf16> {
func.func @fpext_f8_to_f16_relu(%in : vector<8xf8E4M3FN>) {
// CHECK: nvvm.convert.f8x2.to.f16x2
// CHECK-SAME: relu
- %out = nvgpu.extf %in {relu = true}
+ %out = nvgpu.extf %in <{relu = true}>
: vector<8xf8E4M3FN> to vector<8xf16>
return
}
diff --git a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm-mma-fp8.mlir b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm-mma-fp8.mlir
index 43980109ec7c2..8106d26401fc8 100644
--- a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm-mma-fp8.mlir
+++ b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm-mma-fp8.mlir
@@ -6,7 +6,7 @@ func.func @fp8_mma_e4m3_e4m3_m16n8k16(%arg0: vector<2x4xf8E4M3FN>, %arg1: vector
// CHECK-SAME: shape = <m = 16, n = 8, k = 16>
// CHECK-SAME: multiplicand_a_ptx_type = e4m3
// CHECK-SAME: multiplicand_b_ptx_type = e4m3
- %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<2x4xf8E4M3FN>, vector<1x4xf8E4M3FN>, vector<2x2xf32>) -> vector<2x2xf32>
+ %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<2x4xf8E4M3FN>, vector<1x4xf8E4M3FN>, vector<2x2xf32>) -> vector<2x2xf32>
return %0 : vector<2x2xf32>
}
@@ -18,7 +18,7 @@ func.func @fp8_mma_e4m3_e4m3_m16n8k32(%arg0: vector<4x4xf8E4M3FN>, %arg1: vector
// CHECK-SAME: shape = <m = 16, n = 8, k = 32>
// CHECK-SAME: multiplicand_a_ptx_type = e4m3
// CHECK-SAME: multiplicand_b_ptx_type = e4m3
- %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xf8E4M3FN>, vector<2x4xf8E4M3FN>, vector<2x2xf32>) -> vector<2x2xf32>
+ %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) mmaShape = [16, 8, 32] : (vector<4x4xf8E4M3FN>, vector<2x4xf8E4M3FN>, vector<2x2xf32>) -> vector<2x2xf32>
return %0 : vector<2x2xf32>
}
@@ -30,7 +30,7 @@ func.func @fp8_mma_e5m2_e5m2_m16n8k16(%arg0: vector<2x4xf8E5M2>, %arg1: vector<1
// CHECK-SAME: shape = <m = 16, n = 8, k = 16>
// CHECK-SAME: multiplicand_a_ptx_type = e5m2
// CHECK-SAME: multiplicand_b_ptx_type = e5m2
- %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<2x4xf8E5M2>, vector<1x4xf8E5M2>, vector<2x2xf32>) -> vector<2x2xf32>
+ %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<2x4xf8E5M2>, vector<1x4xf8E5M2>, vector<2x2xf32>) -> vector<2x2xf32>
return %0 : vector<2x2xf32>
}
@@ -42,6 +42,6 @@ func.func @fp8_mma_e5m2_e5m2_m16n8k32(%arg0: vector<4x4xf8E5M2>, %arg1: vector<2
// CHECK-SAME: shape = <m = 16, n = 8, k = 32>
// CHECK-SAME: multiplicand_a_ptx_type = e5m2
// CHECK-SAME: multiplicand_b_ptx_type = e5m2
- %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xf8E5M2>, vector<2x4xf8E5M2>, vector<2x2xf32>) -> vector<2x2xf32>
+ %0 = nvgpu.mma.sync(%arg0, %arg1, %arg2) mmaShape = [16, 8, 32] : (vector<4x4xf8E5M2>, vector<2x4xf8E5M2>, vector<2x2xf32>) -> vector<2x2xf32>
return %0 : vector<2x2xf32>
}
diff --git a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
index 380a93adfe739..7d32d8db81313 100644
--- a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
+++ b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
@@ -14,7 +14,7 @@ func.func @m16n8k16_fp16(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2:
// CHECK-NOT: llvm.extractvalue
// CHECK: [[d:%.+]] = nvvm.mma.sync
// CHECK-SAME: shape = <m = 16, n = 8, k = 16>
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
// CHECK-DAG: llvm.extractvalue [[d]][0] : !llvm.struct<(vector<2xf16>, vector<2xf16>)>
// CHECK-DAG: llvm.extractvalue [[d]][1] : !llvm.struct<(vector<2xf16>, vector<2xf16>)>
// CHECK: llvm.mlir.poison : !llvm.array<2 x vector<2xf16>>
@@ -31,7 +31,7 @@ func.func @m16n8k16_fp16_fp32(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %a
// CHECK: [[d:%.+]] = nvvm.mma.sync
// CHECK-SAME: shape = <m = 16, n = 8, k = 16>
// CHECK-SAME: (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32)>
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
// CHECK: [[undef:%.+]] = llvm.mlir.poison : vector<2xf32>
// CHECK-DAG: llvm.extractvalue [[d]][0] : !llvm.struct<(f32, f32, f32, f32)>
// CHECK-DAG: llvm.extractvalue [[d]][1] : !llvm.struct<(f32, f32, f32, f32)>
@@ -68,7 +68,7 @@ func.func @m16n8k16_bf16_fp32(%arg0: vector<4x2xbf16>, %arg1: vector<2x2xbf16>,
// CHECK-SAME: multiplicand_a_ptx_type = bf16
// CHECK-SAME: multiplicand_b_ptx_type = bf16
// CHECK-SAME: (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xbf16>, vector<2x2xbf16>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x2xbf16>, vector<2x2xbf16>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
@@ -82,7 +82,7 @@ func.func @m16n8k8_fp16(%arg0: vector<2x2xf16>, %arg1: vector<1x2xf16>, %arg2: v
// CHECK-NOT: llvm.extractvalue
// CHECK: [[d:%.+]] = nvvm.mma.sync
// CHECK-SAME: shape = <m = 16, n = 8, k = 8>
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<2x2xf16>, vector<1x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 8] : (vector<2x2xf16>, vector<1x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
// CHECK-DAG: llvm.extractvalue [[d]][0] : !llvm.struct<(vector<2xf16>, vector<2xf16>)>
// CHECK-DAG: llvm.extractvalue [[d]][1] : !llvm.struct<(vector<2xf16>, vector<2xf16>)>
// CHECK: llvm.mlir.poison : !llvm.array<2 x vector<2xf16>>
@@ -113,7 +113,7 @@ func.func @m16n8k32_int8(%arg0: vector<4x4xi8>, %arg1: vector<2x4xi8>, %arg2: ve
// CHECK-SAME: int_overflow = satfinite
// CHECK-SAME: multiplicand_a_ptx_type = s8
// CHECK-SAME: multiplicand_b_ptx_type = s8
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 32] : (vector<4x4xi8>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
@@ -132,7 +132,7 @@ func.func @m16n8k32_i4(%arg0: vector<2x8xi4>, %arg1: vector<1x8xi4>, %arg2: vect
// CHECK-SAME: int_overflow = satfinite
// CHECK-SAME: multiplicand_a_ptx_type = s4
// CHECK-SAME: multiplicand_b_ptx_type = s4
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<2x8xi4>, vector<1x8xi4>, vector<2x2xi32>) -> vector<2x2xi32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 32] : (vector<2x8xi4>, vector<1x8xi4>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
@@ -157,7 +157,7 @@ func.func @m16n8k64_i4(%arg0: vector<4x8xi4>, %arg1: vector<2x8xi4>, %arg2: vect
// CHECK-SAME: int_overflow = satfinite
// CHECK-SAME: multiplicand_a_ptx_type = s4
// CHECK-SAME: multiplicand_b_ptx_type = s4
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 64]} : (vector<4x8xi4>, vector<2x8xi4>, vector<2x2xi32>) -> vector<2x2xi32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 64] : (vector<4x8xi4>, vector<2x8xi4>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
@@ -168,7 +168,7 @@ func.func @m8n8k4_f64(%arg0: vector<1x1xf64>, %arg1: vector<1x1xf64>, %arg2: vec
// CHECK: llvm.extractvalue
// CHECK: [[d:%.+]] = nvvm.mma.sync A[{{%.+}}] B[{{%.+}}] C[{{%.+}}, {{%.+}}]
// CHECK-SAME: shape = <m = 8, n = 8, k = 4>
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [8, 8, 4]} : (vector<1x1xf64>, vector<1x1xf64>, vector<1x2xf64>) -> vector<1x2xf64>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [8, 8, 4] : (vector<1x1xf64>, vector<1x1xf64>, vector<1x2xf64>) -> vector<1x2xf64>
// CHECK: llvm.mlir.poison : vector<2xf64>
// CHECK-DAG: llvm.extractvalue [[d]][0] : !llvm.struct<(f64, f64)>
// CHECK-DAG: llvm.extractvalue [[d]][1] : !llvm.struct<(f64, f64)>
@@ -183,7 +183,7 @@ func.func @m8n8k4_f64(%arg0: vector<1x1xf64>, %arg1: vector<1x1xf64>, %arg2: vec
func.func @ldmatrix_x4(%arg0: memref<128x128xf16, 3>) -> vector<4x2xf16> {
%c0 = arith.constant 0 : index
// CHECK: nvvm.ldmatrix {{%.+}}, num = 4, layout = <row>, shape = <m = 8, n = 8>, element_type = <b16> : {{.*}} -> !llvm.struct<(i32, i32, i32, i32)>
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf16, 3> -> vector<4x2xf16>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf16, 3> -> vector<4x2xf16>
// CHECK: llvm.extractvalue
// CHECK: llvm.bitcast
// CHECK: llvm.insertvalue
@@ -203,7 +203,7 @@ func.func @ldmatrix_x4(%arg0: memref<128x128xf16, 3>) -> vector<4x2xf16> {
func.func @ldmatrix_x1(%arg0: memref<128x128xf16, 3>) -> vector<1x2xf16> {
%c0 = arith.constant 0 : index
// CHECK: nvvm.ldmatrix {{%.+}}, num = 1, layout = <row>, shape = <m = 8, n = 8>, element_type = <b16> : {{.*}} -> i32
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 1 : i32} : memref<128x128xf16, 3> -> vector<1x2xf16>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 1 transpose = false : memref<128x128xf16, 3> -> vector<1x2xf16>
// CHECK: llvm.bitcast
// CHECK: llvm.insertvalue
return %a : vector<1x2xf16>
@@ -224,7 +224,7 @@ func.func @m16n8k4_tf32(%arg0: vector<2x1xf32>, %arg1: vector<1x1xf32>, %arg2: v
// CHECK-SAME: multiplicand_a_ptx_type = tf32
// CHECK-SAME: multiplicand_b_ptx_type = tf32
// CHECK-SAME: -> !llvm.struct<(f32, f32, f32, f32)>
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 4], tf32Enabled} : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 4] tf32Enabled : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
// CHECK: [[undef:%.+]] = llvm.mlir.poison : vector<2xf32>
// CHECK-DAG: llvm.extractvalue [[d]][0] : !llvm.struct<(f32, f32, f32, f32)>
// CHECK-DAG: llvm.extractvalue [[d]][1] : !llvm.struct<(f32, f32, f32, f32)>
@@ -266,10 +266,10 @@ func.func @async_cp(
// CHECK: nvvm.cp.async.commit.group
%1 = nvgpu.device_async_create_group %0
// CHECK: nvvm.cp.async.wait.group 1
- nvgpu.device_async_wait %1 { numGroups = 1 : i32 }
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg
- %2 = nvgpu.device_async_copy %src[%i, %i], %dst[%i, %i, %i], 4 {bypassL1}: memref<128x128xf32> to memref<3x16x128xf32, 3>
+ %2 = nvgpu.device_async_copy %src[%i, %i], %dst[%i, %i, %i], 4 bypassL1: memref<128x128xf32> to memref<3x16x128xf32, 3>
return
}
@@ -320,11 +320,11 @@ func.func @async_cp_zfill_f32_align4(
// CHECK-DAG: %[[c4:.*]] = llvm.mul %[[c2]], %[[c3]] : i32
// CHECK-DAG: %[[c5:.*]] = llvm.lshr %[[c4]], %[[c1]] : i32
// CHECK-DAG: nvvm.cp.async.shared.global %[[ADDRESSDST]], %[[CAST2]], 16, cache = cg, %[[c5]]
- %0 = nvgpu.device_async_copy %src[%i, %i], %dst[%i, %i, %i], 4, %srcElements {bypassL1}: memref<128x128xf32> to memref<3x16x128xf32, 3>
+ %0 = nvgpu.device_async_copy %src[%i, %i], %dst[%i, %i, %i], 4, %srcElements bypassL1: memref<128x128xf32> to memref<3x16x128xf32, 3>
// CHECK: nvvm.cp.async.commit.group
%1 = nvgpu.device_async_create_group %0
// CHECK: nvvm.cp.async.wait.group 1
- nvgpu.device_async_wait %1 { numGroups = 1 : i32 }
+ nvgpu.device_async_wait %1 numGroups = 1
return
}
@@ -359,7 +359,7 @@ func.func @async_cp_zfill_f32_align1(
// CHECK: nvvm.cp.async.commit.group
%1 = nvgpu.device_async_create_group %0
// CHECK: nvvm.cp.async.wait.group 1
- nvgpu.device_async_wait %1 { numGroups = 1 : i32 }
+ nvgpu.device_async_wait %1 numGroups = 1
return
}
@@ -393,7 +393,7 @@ func.func @mma_sp_sync_f16_16832(%arg0: vector<4x2xf16>,
// CHECK-SAME: %[[sparseMetadata]] :
// CHECK-SAME: -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 32]} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 32] :
(vector<4x2xf16>, vector<4x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
// CHECK-DAG: llvm.extractvalue %[[d]][0] : !llvm.struct<(vector<2xf16>, vector<2xf16>)>
@@ -429,7 +429,7 @@ func.func @mma_sp_sync_f16_16816(%arg0: vector<2x2xf16>,
// CHECK-SAME: %[[sparseMetadata]] :
// CHECK-SAME: -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 16]} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 16] :
(vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -449,8 +449,7 @@ func.func @mma_sp_sync_f16_16816_01(%arg0: vector<2x2xf16>,
// CHECK-SAME: %[[sparseMetadata]] :
// CHECK-SAME: -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3)
- {mmaShape = [16, 8, 16], sparsitySelector = 1 : i32} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 16] sparsitySelector = 1 :
(vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -488,7 +487,7 @@ func.func @mma_sp_sync_i8_16864(%arg0: vector<4x4xi8>,
// CHECK-SAME: %[[sparseMetadata]] :
// CHECK-SAME: -> !llvm.struct<(i32, i32, i32, i32)
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 64]} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 64] :
(vector<4x4xi8>, vector<4x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
@@ -977,7 +976,7 @@ func.func @warpgroup_mma_128_128_64(
// CHECK: %[[S41:.+]] = llvm.insertvalue %[[S38]], %[[S40]][1] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)>
// CHECK: nvvm.wgmma.commit.group.sync.aligned
// CHECK: nvvm.wgmma.wait.group.sync.aligned 1
- %wgmmaResult = nvgpu.warpgroup.mma %descA, %descB, %acc {transposeB}:
+ %wgmmaResult = nvgpu.warpgroup.mma %descA, %descB, %acc transposeB:
!nvgpu.warpgroup.descriptor<tensor = memref<128x64xf16, 3>>,
!nvgpu.warpgroup.descriptor<tensor = memref<64x128xf16, 3>>,
!nvgpu.warpgroup.accumulator<fragmented = vector<128x128xf32>>
@@ -1327,7 +1326,7 @@ func.func @warpgroup_matrix_multiply_m128n128k64(
!nvgpu.warpgroup.accumulator<fragmented = vector<128x128xf32>>
// GEMM
- %matrixD = nvgpu.warpgroup.mma %descA, %descB, %matrixC {transposeB}:
+ %matrixD = nvgpu.warpgroup.mma %descA, %descB, %matrixC transposeB:
!nvgpu.warpgroup.descriptor<tensor = memref<128x64xf16, 3>>,
!nvgpu.warpgroup.descriptor<tensor = memref<64x128xf16, 3>>,
!nvgpu.warpgroup.accumulator<fragmented = vector<128x128xf32>>
@@ -1407,6 +1406,6 @@ func.func @rcp_approx_ftz_f32(%in: vector<32x16xf32>) {
// CHECK: %[[ELEM_RCP0:.*]] = nvvm.rcp.approx.ftz.f %[[ELEM_0]] : f32
// CHECK: llvm.insertelement %[[ELEM_RCP0]], %[[OUT1DVEC]][%[[IDX_0]] : i64] : vector<16xf32>
// CHECK-COUNT-511: nvvm.rcp.approx.ftz.f
- %out = nvgpu.rcp %in {approx = true, ftz = true} : vector<32x16xf32>
+ %out = nvgpu.rcp %in <{approx = true, ftz = true}> : vector<32x16xf32>
return
}
diff --git a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-truncf.mlir b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-truncf.mlir
index a648ea6868601..56de1b933f660 100644
--- a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-truncf.mlir
+++ b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-truncf.mlir
@@ -134,7 +134,7 @@ func.func @cvt_float_f32_to_e8m0_rz(%in : vector<8xf32>) {
// CHECK: nvvm.convert.f32x2.to.f8x2
// CHECK-SAME: rnd = <rz>
// CHECK-SAME: : i16(f8E8M0FNU)
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rz>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rz>}>
: vector<8xf32> to vector<8xf8E8M0FNU>
return
}
@@ -145,7 +145,7 @@ func.func @cvt_float_f32_to_e8m0_rp(%in : vector<8xf32>) {
// CHECK: nvvm.convert.f32x2.to.f8x2
// CHECK-SAME: rnd = <rp>
// CHECK-SAME: : i16(f8E8M0FNU)
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rp>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rp>}>
: vector<8xf32> to vector<8xf8E8M0FNU>
return
}
@@ -287,7 +287,7 @@ func.func @fptrunc_f64_to_f8(%arg0: vector<4xf64>) -> vector<4xf8E4M3FN> {
func.func @fptrunc_f32_to_f8_satfinite(%in : vector<8xf32>) {
// CHECK: nvvm.convert.f32x2.to.f8x2
// CHECK-SAME: sat = <satfinite>
- %out = nvgpu.truncf %in {sat = #nvvm.sat_mode<satfinite>}
+ %out = nvgpu.truncf %in <{sat = #nvvm.sat_mode<satfinite>}>
: vector<8xf32> to vector<8xf8E4M3FN>
return
}
@@ -296,7 +296,7 @@ func.func @fptrunc_f32_to_f8_satfinite(%in : vector<8xf32>) {
func.func @fptrunc_f32_to_f16_relu(%in : vector<4xf32>) {
// CHECK: nvvm.convert.f32x2.to.f16x2
// CHECK-SAME: relu
- %out = nvgpu.truncf %in {relu = true}
+ %out = nvgpu.truncf %in <{relu = true}>
: vector<4xf32> to vector<4xf16>
return
}
@@ -323,7 +323,7 @@ func.func @fptrunc_f32_to_f16_default_sat(%in : vector<4xf32>) {
func.func @fptrunc_f32_to_f16_explicit_none(%in : vector<4xf32>) {
// CHECK: nvvm.convert.f32x2.to.f16x2
// CHECK-NOT: satfinite
- %out = nvgpu.truncf %in {sat = #nvvm.sat_mode<none>}
+ %out = nvgpu.truncf %in <{sat = #nvvm.sat_mode<none>}>
: vector<4xf32> to vector<4xf16>
return
}
@@ -334,7 +334,7 @@ func.func @fptrunc_f32_to_f16_explicit_none(%in : vector<4xf32>) {
func.func @fptrunc_f32_to_f16_rs(%in : vector<4xf32>, %rbits : i32) {
// CHECK: nvvm.convert.f32x2.to.f16x2
// CHECK-SAME: rnd = <rs>
- %out = nvgpu.truncf %in, %rbits {rnd = #nvvm.fp_rnd_mode<rs>}
+ %out = nvgpu.truncf %in, %rbits <{rnd = #nvvm.fp_rnd_mode<rs>}>
: vector<4xf32> to vector<4xf16>
return
}
@@ -343,7 +343,7 @@ func.func @fptrunc_f32_to_f16_rs(%in : vector<4xf32>, %rbits : i32) {
func.func @fptrunc_f32_to_bf16_rs(%in : vector<4xf32>, %rbits : i32) {
// CHECK: nvvm.convert.f32x2.to.bf16x2
// CHECK-SAME: rnd = <rs>
- %out = nvgpu.truncf %in, %rbits {rnd = #nvvm.fp_rnd_mode<rs>}
+ %out = nvgpu.truncf %in, %rbits <{rnd = #nvvm.fp_rnd_mode<rs>}>
: vector<4xf32> to vector<4xbf16>
return
}
@@ -352,7 +352,7 @@ func.func @fptrunc_f32_to_bf16_rs(%in : vector<4xf32>, %rbits : i32) {
func.func @fptrunc_f32_to_f16_rz(%in : vector<4xf32>) {
// CHECK: nvvm.convert.f32x2.to.f16x2
// CHECK-SAME: rnd = <rz>
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rz>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rz>}>
: vector<4xf32> to vector<4xf16>
return
}
@@ -361,7 +361,7 @@ func.func @fptrunc_f32_to_f16_rz(%in : vector<4xf32>) {
func.func @fptrunc_f32_to_bf16_rz(%in : vector<4xf32>) {
// CHECK: nvvm.convert.f32x2.to.bf16x2
// CHECK-SAME: rnd = <rz>
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rz>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rz>}>
: vector<4xf32> to vector<4xbf16>
return
}
diff --git a/mlir/test/Conversion/VectorToGPU/fold-arith-vector-to-mma-ops-mma-sync.mlir b/mlir/test/Conversion/VectorToGPU/fold-arith-vector-to-mma-ops-mma-sync.mlir
index 0afaa19d59d15..a23bc5d867b7b 100644
--- a/mlir/test/Conversion/VectorToGPU/fold-arith-vector-to-mma-ops-mma-sync.mlir
+++ b/mlir/test/Conversion/VectorToGPU/fold-arith-vector-to-mma-ops-mma-sync.mlir
@@ -16,12 +16,12 @@ func.func @m16n8k16_mmasync16816_f16_f16_f32_row_row_row(%arg0: memref<42x32xf16
%cst_f16 = arith.constant 0.000000e+00 : f16
%cst_f32 = arith.constant 0.000000e+00 : f32
- // CHECK-DAG: nvgpu.ldmatrix %arg0[%{{.*}}, %{{.*}}] {numTiles = 4 : i32, transpose = false}
+ // CHECK-DAG: nvgpu.ldmatrix %arg0[%{{.*}}, %{{.*}}] numTiles = 4 transpose = false
%A = vector.transfer_read %arg0[%c0, %c0], %cst_f16 {in_bounds = [true, true]} : memref<42x32xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
%A_f32 = arith.extf %A : vector<16x16xf16> to vector<16x16xf32>
- // CHECK-DAG: nvgpu.ldmatrix %arg1[%{{.*}}, %{{.*}}] {numTiles = 4 : i32, transpose = true}
+ // CHECK-DAG: nvgpu.ldmatrix %arg1[%{{.*}}, %{{.*}}] numTiles = 4 transpose = true
%B = vector.transfer_read %arg1[%c0, %c0], %cst_f16 {permutation_map = #map0, in_bounds = [true, true]} : memref<32x64xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
%C = vector.transfer_read %arg2[%c0, %c0], %cst_f32 {in_bounds = [true, true]} : memref<42x64xf32, #gpu.address_space<workgroup>>, vector<16x16xf32>
@@ -29,7 +29,7 @@ func.func @m16n8k16_mmasync16816_f16_f16_f32_row_row_row(%arg0: memref<42x32xf16
%B0_f32 = arith.extf %B0 : vector<8x16xf16> to vector<8x16xf32>
%C0 = vector.extract_strided_slice %C {offsets = [0, 0], sizes = [16, 8], strides = [1, 1]} : vector<16x16xf32> to vector<16x8xf32>
- // CHECK-DAG: nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
+ // CHECK-DAG: nvgpu.mma.sync({{.*}}) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
%D0 = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A_f32, %B0_f32, %C0 : vector<16x16xf32>, vector<8x16xf32> into vector<16x8xf32>
vector.transfer_write %D0, %arg2[%c0, %c0] {in_bounds = [true, true]} : vector<16x8xf32>, memref<42x64xf32, #gpu.address_space<workgroup>>
@@ -38,7 +38,7 @@ func.func @m16n8k16_mmasync16816_f16_f16_f32_row_row_row(%arg0: memref<42x32xf16
%B1_f32 = arith.extf %B1 : vector<8x16xf16> to vector<8x16xf32>
%C1 = vector.extract_strided_slice %C {offsets = [0, 8], sizes = [16, 8], strides = [1, 1]} : vector<16x16xf32> to vector<16x8xf32>
- // CHECK-DAG: nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
+ // CHECK-DAG: nvgpu.mma.sync({{.*}}) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>
%D1 = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A_f32, %B1_f32, %C1 : vector<16x16xf32>, vector<8x16xf32> into vector<16x8xf32>
vector.transfer_write %D1, %arg2[%c0, %c0] {in_bounds = [true, true]} : vector<16x8xf32>, memref<42x64xf32, #gpu.address_space<workgroup>>
diff --git a/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops-mma-sync.mlir b/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops-mma-sync.mlir
index 912f7fba59e60..a5dcfd813b3a5 100644
--- a/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops-mma-sync.mlir
+++ b/mlir/test/Conversion/VectorToGPU/vector-to-mma-ops-mma-sync.mlir
@@ -44,7 +44,7 @@ func.func @m16n8k32_int8_row_row_row(%arg0: memref<128x128xi8, #gpu.address_spac
// CHECK: [[m_coord:%.+]] = affine.apply [[$strided_map]]()[{{%.+}}]
// CHECK: [[k_coord:%.+]] = affine.apply [[$contiguous_map]]()[{{%.+}}]
- // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false} : memref<128x128xi8, #gpu.address_space<workgroup>> -> vector<4x4xi8>
+ // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4 transpose = false : memref<128x128xi8, #gpu.address_space<workgroup>> -> vector<4x4xi8>
// Verify that the operandB load is lowered to scalar load to be able
// to transpose at 8-bit granularity. ldmatrix can only transpose at
@@ -88,7 +88,7 @@ func.func @m16n8k32_int8_row_row_row(%arg0: memref<128x128xi8, #gpu.address_spac
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi8, #gpu.address_space<workgroup>>, vector<16x32xi8>
%B = vector.transfer_read %arg1[%c39, %c40], %cst {in_bounds = [true, true], permutation_map = #map0} : memref<128x128xi8, #gpu.address_space<workgroup>>, vector<8x32xi8>
%C = vector.transfer_read %arg2[%c49, %c40], %cst0 {in_bounds = [true, true]} : memref<128x128xi32>, vector<16x8xi32>
- // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
+ // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) mmaShape = [16, 8, 32] : (vector<4x4xi8>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
%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>
// CHECK: [[row:%.+]] = affine.apply [[$rowC0_map]]()[{{%.+}}]
@@ -153,7 +153,7 @@ func.func @m8n8k4_f64_row_row_row(%arg0: memref<128x128xf64>, %arg1: memref<128x
%A = vector.transfer_read %arg0[%c1, %c1], %cst {in_bounds = [true, true]} : memref<128x128xf64>, vector<8x4xf64>
%B = vector.transfer_read %arg1[%c39, %c40], %cst {in_bounds = [true, true], permutation_map = #map0} : memref<128x128xf64>, vector<8x4xf64>
%C = vector.transfer_read %arg2[%c49, %c40], %cst0 {in_bounds = [true, true]} : memref<128x128xf64>, vector<8x8xf64>
- // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [8, 8, 4]} : (vector<1x1xf64>, vector<1x1xf64>, vector<1x2xf64>) -> vector<1x2xf64>
+ // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) mmaShape = [8, 8, 4] : (vector<1x1xf64>, vector<1x1xf64>, vector<1x2xf64>) -> vector<1x2xf64>
%D = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B, %C : vector<8x4xf64>, vector<8x4xf64> into vector<8x8xf64>
// CHECK-DAG: [[row:%.+]] = affine.apply [[$rowC0_map]]
@@ -185,10 +185,10 @@ func.func @m16n8k16_fp16_row_row_row(%arg0: memref<20x20xf16, #gpu.address_space
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4 transpose = false
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$strided_map]]
- // CHECK: nvgpu.ldmatrix %arg1[[[k_coord]], [[n_coord]]] {numTiles = 2 : i32, transpose = true}
+ // CHECK: nvgpu.ldmatrix %arg1[[[k_coord]], [[n_coord]]] numTiles = 2 transpose = true
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<20x20xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
%B = vector.transfer_read %arg1[%c0, %c0], %cst {permutation_map = #map0, in_bounds = [true, true]} : memref<20x20xf16, #gpu.address_space<workgroup>>, vector<8x16xf16>
%C = vector.transfer_read %arg2[%c0, %c0], %cst {in_bounds = [true, true]} : memref<20x20xf16, #gpu.address_space<workgroup>>, vector<16x8xf16>
@@ -219,22 +219,22 @@ func.func @m16n16k16_mmasync16816_fp16_f16_row_row_row(%arg0: memref<42x32xf16,
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK: [[fragmentA:%.+]] = nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK: [[fragmentA:%.+]] = nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4 transpose = false
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<42x32xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$strided_map]]
- // CHECK-DAG: [[fragmentB:%.+]] = nvgpu.ldmatrix %arg1[[[k_coord]], [[n_coord]]] {numTiles = 4 : i32, transpose = true}
+ // CHECK-DAG: [[fragmentB:%.+]] = nvgpu.ldmatrix %arg1[[[k_coord]], [[n_coord]]] numTiles = 4 transpose = true
%B = vector.transfer_read %arg1[%c0, %c0], %cst {permutation_map = #map0, in_bounds = [true, true]} : memref<32x64xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK-DAG: [[fragmentC:%.*]] = nvgpu.ldmatrix %arg2[[[m_coord]], [[n_coord]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK-DAG: [[fragmentC:%.*]] = nvgpu.ldmatrix %arg2[[[m_coord]], [[n_coord]]] numTiles = 4 transpose = false
%C = vector.transfer_read %arg2[%c0, %c0], %cst {in_bounds = [true, true]} : memref<42x64xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[fragmentB0:%.+]] = vector.extract_strided_slice [[fragmentB]] {offsets = [0, 0], sizes = [2, 2], strides = [1, 1]} : vector<4x2xf16> to vector<2x2xf16>
// CHECK-DAG: [[fragmentC0:%.+]] = vector.extract_strided_slice [[fragmentC]] {offsets = [0, 0], sizes = [2, 2], strides = [1, 1]} : vector<4x2xf16> to vector<2x2xf16>
- // CHECK: nvgpu.mma.sync([[fragmentA]], [[fragmentB0]], [[fragmentC0]]) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ // CHECK: nvgpu.mma.sync([[fragmentA]], [[fragmentB0]], [[fragmentC0]]) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
%B0 = vector.extract_strided_slice %B {offsets = [0, 0], sizes = [8, 16], strides = [1, 1]} : vector<16x16xf16> to vector<8x16xf16>
%C0 = vector.extract_strided_slice %C {offsets = [0, 0], sizes = [16, 8], strides = [1, 1]} : vector<16x16xf16> to vector<16x8xf16>
%D0 = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B0, %C0 : vector<16x16xf16>, vector<8x16xf16> into vector<16x8xf16>
@@ -242,7 +242,7 @@ func.func @m16n16k16_mmasync16816_fp16_f16_row_row_row(%arg0: memref<42x32xf16,
// CHECK-DAG: [[fragmentB1:%.+]] = vector.extract_strided_slice [[fragmentB]] {offsets = [2, 0], sizes = [2, 2], strides = [1, 1]} : vector<4x2xf16> to vector<2x2xf16>
// CHECK-DAG: [[fragmentC1:%.+]] = vector.extract_strided_slice [[fragmentC]] {offsets = [2, 0], sizes = [2, 2], strides = [1, 1]} : vector<4x2xf16> to vector<2x2xf16>
- // CHECK: nvgpu.mma.sync([[fragmentA]], [[fragmentB1]], [[fragmentC1]]) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ // CHECK: nvgpu.mma.sync([[fragmentA]], [[fragmentB1]], [[fragmentC1]]) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
%B1 = vector.extract_strided_slice %B {offsets = [8, 0], sizes = [8, 16], strides = [1, 1]} : vector<16x16xf16> to vector<8x16xf16>
%C1 = vector.extract_strided_slice %C {offsets = [0, 8], sizes = [16, 8], strides = [1, 1]} : vector<16x16xf16> to vector<16x8xf16>
%D1 = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B1, %C1 : vector<16x16xf16>, vector<8x16xf16> into vector<16x8xf16>
@@ -275,22 +275,22 @@ func.func @multi_dim_m16n8k16_fp16_row_row_row(%arg0: memref<4x32x1x32xf16, #gpu
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK: [[fragmentA:%.+]] = nvgpu.ldmatrix %arg0[[[c0]], [[m_coord]], [[c0]], [[k_coord]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK: [[fragmentA:%.+]] = nvgpu.ldmatrix %arg0[[[c0]], [[m_coord]], [[c0]], [[k_coord]]] numTiles = 4 transpose = false
%A = vector.transfer_read %arg0[%c0, %c0, %c0, %c0], %cst {in_bounds = [true, true], permutation_map = #map_a} : memref<4x32x1x32xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$strided_map]]
- // CHECK-DAG: [[fragmentB:%.+]] = nvgpu.ldmatrix %arg1[[[c0]], [[c0]], [[k_coord]], [[n_coord]]] {numTiles = 4 : i32, transpose = true}
+ // CHECK-DAG: [[fragmentB:%.+]] = nvgpu.ldmatrix %arg1[[[c0]], [[c0]], [[k_coord]], [[n_coord]]] numTiles = 4 transpose = true
%B = vector.transfer_read %arg1[%c0, %c0, %c0, %c0], %cst {in_bounds = [true, true], permutation_map = #map_b} : memref<4x1x32x32xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK-DAG: [[fragmentC:%.*]] = nvgpu.ldmatrix %arg2[[[c0]], [[m_coord]], [[n_coord]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK-DAG: [[fragmentC:%.*]] = nvgpu.ldmatrix %arg2[[[c0]], [[m_coord]], [[n_coord]]] numTiles = 4 transpose = false
%C = vector.transfer_read %arg2[%c0, %c0, %c0], %cst {in_bounds = [true, true]} : memref<1x32x40xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[fragmentB0:%.+]] = vector.extract_strided_slice [[fragmentB]] {offsets = [0, 0], sizes = [2, 2], strides = [1, 1]} : vector<4x2xf16> to vector<2x2xf16>
// CHECK-DAG: [[fragmentC0:%.+]] = vector.extract_strided_slice [[fragmentC]] {offsets = [0, 0], sizes = [2, 2], strides = [1, 1]} : vector<4x2xf16> to vector<2x2xf16>
- // CHECK: nvgpu.mma.sync([[fragmentA]], [[fragmentB0]], [[fragmentC0]]) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ // CHECK: nvgpu.mma.sync([[fragmentA]], [[fragmentB0]], [[fragmentC0]]) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
%B0 = vector.extract_strided_slice %B {offsets = [0, 0], sizes = [8, 16], strides = [1, 1]} : vector<16x16xf16> to vector<8x16xf16>
%C0 = vector.extract_strided_slice %C {offsets = [0, 0], sizes = [16, 8], strides = [1, 1]} : vector<16x16xf16> to vector<16x8xf16>
%D0 = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B0, %C0 : vector<16x16xf16>, vector<8x16xf16> into vector<16x8xf16>
@@ -318,17 +318,17 @@ func.func @batch_m16n8k16_fp16_row_row_row(%arg0: memref<2x20x20xf16, #gpu.addre
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK: nvgpu.ldmatrix %arg0[[[C0]], [[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false} : memref<2x20x20xf16, #gpu.address_space<workgroup>> -> vector<4x2xf16>
+ // CHECK: nvgpu.ldmatrix %arg0[[[C0]], [[m_coord]], [[k_coord]]] numTiles = 4 transpose = false : memref<2x20x20xf16, #gpu.address_space<workgroup>> -> vector<4x2xf16>
%A = vector.transfer_read %arg0[%c0, %c0, %c0], %cst {in_bounds = [true, true]} : memref<2x20x20xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$strided_map]]
- // CHECK: nvgpu.ldmatrix %arg1[[[C0]], [[k_coord]], [[n_coord]]] {numTiles = 2 : i32, transpose = true} : memref<2x20x20xf16, #gpu.address_space<workgroup>> -> vector<2x2xf16>
+ // CHECK: nvgpu.ldmatrix %arg1[[[C0]], [[k_coord]], [[n_coord]]] numTiles = 2 transpose = true : memref<2x20x20xf16, #gpu.address_space<workgroup>> -> vector<2x2xf16>
%B = vector.transfer_read %arg1[%c0, %c0, %c0], %cst {permutation_map = #map0, in_bounds = [true, true]} : memref<2x20x20xf16, #gpu.address_space<workgroup>>, vector<8x16xf16>
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK: nvgpu.ldmatrix %arg2[[[C0]], [[m_coord]], [[n_coord]]] {numTiles = 2 : i32, transpose = false} : memref<2x20x20xf16, #gpu.address_space<workgroup>> -> vector<2x2xf16>
+ // CHECK: nvgpu.ldmatrix %arg2[[[C0]], [[m_coord]], [[n_coord]]] numTiles = 2 transpose = false : memref<2x20x20xf16, #gpu.address_space<workgroup>> -> vector<2x2xf16>
%C = vector.transfer_read %arg2[%c0, %c0, %c0], %cst {in_bounds = [true, true]} : memref<2x20x20xf16, #gpu.address_space<workgroup>>, vector<16x8xf16>
%D = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B, %C : vector<16x16xf16>, vector<8x16xf16> into vector<16x8xf16>
vector.transfer_write %D, %arg2[%c0, %c0, %c0] {in_bounds = [true, true]} : vector<16x8xf16>, memref<2x20x20xf16, #gpu.address_space<workgroup>>
@@ -360,17 +360,17 @@ func.func @m16n8k16_fp16_row_col_row(%arg0: memref<20x20xf16, #gpu.address_space
%cst = arith.constant 0.000000e+00 : f16
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_ldmatrix_x4_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x4_map]]
- // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32
+ // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4
// CHECK-SAME: transpose = false
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$strided_ldmatrix_x2_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x2_map]]
- // CHECK: nvgpu.ldmatrix %arg1[[[n_coord]], [[k_coord]]] {numTiles = 2 : i32
+ // CHECK: nvgpu.ldmatrix %arg1[[[n_coord]], [[k_coord]]] numTiles = 2
// CHECK-SAME: transpose = false
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_ldmatrix_x4_map]]
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x4_map]]
- // CHECK: nvgpu.ldmatrix %arg2[[[m_coord]], [[n_coord]]] {numTiles = 2 : i32
+ // CHECK: nvgpu.ldmatrix %arg2[[[m_coord]], [[n_coord]]] numTiles = 2
// CHECK-SAME: transpose = false
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<20x20xf16, #gpu.address_space<workgroup>>, vector<16x16xf16>
%B = vector.transfer_read %arg1[%c0, %c0], %cst {in_bounds = [true, true]} : memref<20x20xf16, #gpu.address_space<workgroup>>, vector<8x16xf16>
@@ -413,7 +413,7 @@ func.func @m16n8k4_tf32_f32_row_row_row(%arg0: memref<20x20xf32, #gpu.address_sp
// CHECK-DAG: [[row:%.+]] = affine.apply [[$rowA_map]]
// CHECK-DAG: [[col:%.+]] = affine.apply [[$colA_map]]
- // CHECK: [[a_frag:%.+]] = nvgpu.ldmatrix %arg0[[[row]], [[col]]] {numTiles = 2 : i32, transpose = false}
+ // CHECK: [[a_frag:%.+]] = nvgpu.ldmatrix %arg0[[[row]], [[col]]] numTiles = 2 transpose = false
// b and c are not loaded by ldmatrix in this test.
// CHECK-NOT: nvgpu.ldmatrix
@@ -474,7 +474,7 @@ func.func @m16n8k8_tf32_f32_row_row_row(%arg0: memref<20x20xf32, #gpu.address_sp
// CHECK-DAG: [[row:%.+]] = affine.apply [[$rowA_map]]
// CHECK-DAG: [[col:%.+]] = affine.apply [[$colA_map]]
- // CHECK: [[a_frag:%.+]] = nvgpu.ldmatrix %arg0[[[row]], [[col]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK: [[a_frag:%.+]] = nvgpu.ldmatrix %arg0[[[row]], [[col]]] numTiles = 4 transpose = false
// b and c are not loaded by ldmatrix in this test.
// CHECK-NOT: nvgpu.ldmatrix
@@ -554,7 +554,7 @@ func.func @m16n8k8_tf32_f32_col_col_row(%arg0: memref<20x20xf32, #gpu.address_sp
// CHECK-DAG: [[row:%.+]] = affine.apply [[$rowB0_map]]
// CHECK-DAG: [[col:%.+]] = affine.apply [[$colB0_map]]
- // CHECK: [[b_frag:%.+]] = nvgpu.ldmatrix %arg1[[[row]], [[col]]] {numTiles = 2 : i32, transpose = false}
+ // CHECK: [[b_frag:%.+]] = nvgpu.ldmatrix %arg1[[[row]], [[col]]] numTiles = 2 transpose = false
// CHECK: [[d_frag:%.+]] = nvgpu.mma.sync([[a_frag]], [[b_frag]], [[c_frag]])
// CHECK-SAME: mmaShape = [16, 8, 8]
@@ -608,12 +608,12 @@ func.func @m16n8k64_int4_row_col_row(%arg0: memref<128x128xi4, #gpu.address_spac
// CHECK: [[lane:%.+]] = gpu.lane_id
// CHECK: [[m_coord:%.+]] = affine.apply [[$strided_ldmatrix_x4_map]]()[[[lane]]]
// CHECK: [[k_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x4_map]]()[[[lane]]]
- // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false} : memref<128x128xi4, #gpu.address_space<workgroup>> -> vector<4x8xi4>
+ // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4 transpose = false : memref<128x128xi4, #gpu.address_space<workgroup>> -> vector<4x8xi4>
// CHECK: [[lane:%.+]] = gpu.lane_id
// CHECK: [[n_coord:%.+]] = affine.apply [[$strided_ldmatrix_x2_map]]()[[[lane]]]
// CHECK: [[k_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x2_map]]()[[[lane]]]
- // CHECK: nvgpu.ldmatrix %arg1[[[n_coord]], [[k_coord]]] {numTiles = 2 : i32, transpose = false} : memref<128x128xi4, #gpu.address_space<workgroup>> -> vector<2x8xi4>
+ // CHECK: nvgpu.ldmatrix %arg1[[[n_coord]], [[k_coord]]] numTiles = 2 transpose = false : memref<128x128xi4, #gpu.address_space<workgroup>> -> vector<2x8xi4>
// CHECK: [[lane:%.+]] = gpu.lane_id
// CHECK: [[row:%.+]] = affine.apply [[$rowC0_map]]()[{{%.+}}]
@@ -628,7 +628,7 @@ func.func @m16n8k64_int4_row_col_row(%arg0: memref<128x128xi4, #gpu.address_spac
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi4, #gpu.address_space<workgroup>>, vector<16x64xi4>
%B = vector.transfer_read %arg1[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi4, #gpu.address_space<workgroup>>, vector<8x64xi4>
%C = vector.transfer_read %arg2[%c0, %c0], %cst0 {in_bounds = [true, true]} : memref<128x128xi32>, vector<16x8xi32>
- // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 64]} : (vector<4x8xi4>, vector<2x8xi4>, vector<2x2xi32>) -> vector<2x2xi32>
+ // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) mmaShape = [16, 8, 64] : (vector<4x8xi4>, vector<2x8xi4>, vector<2x2xi32>) -> vector<2x2xi32>
%D = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"], kind = #vector.kind<add>} %A, %B, %C : vector<16x64xi4>, vector<8x64xi4> into vector<16x8xi32>
// CHECK: [[lane:%.+]] = gpu.lane_id
@@ -679,12 +679,12 @@ func.func @m16n8k32_int8_row_col_row(%arg0: memref<128x128xi8, #gpu.address_spac
// CHECK: [[lane:%.+]] = gpu.lane_id
// CHECK: [[m_coord:%.+]] = affine.apply [[$strided_ldmatrix_x4_map]]()[[[lane]]]
// CHECK: [[k_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x4_map]]()[[[lane]]]
- // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false} : memref<128x128xi8, #gpu.address_space<workgroup>> -> vector<4x4xi8>
+ // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4 transpose = false : memref<128x128xi8, #gpu.address_space<workgroup>> -> vector<4x4xi8>
// CHECK: [[lane:%.+]] = gpu.lane_id
// CHECK: [[n_coord:%.+]] = affine.apply [[$strided_ldmatrix_x2_map]]()[[[lane]]]
// CHECK: [[k_coord:%.+]] = affine.apply [[$contiguous_ldmatrix_x2_map]]()[[[lane]]]
- // CHECK: nvgpu.ldmatrix %arg1[[[n_coord]], [[k_coord]]] {numTiles = 2 : i32, transpose = false} : memref<128x128xi8, #gpu.address_space<workgroup>> -> vector<2x4xi8>
+ // CHECK: nvgpu.ldmatrix %arg1[[[n_coord]], [[k_coord]]] numTiles = 2 transpose = false : memref<128x128xi8, #gpu.address_space<workgroup>> -> vector<2x4xi8>
// CHECK: [[lane:%.+]] = gpu.lane_id
// CHECK: [[m_coord:%.+]] = affine.apply [[$rowC0_map]]()[[[lane]]]
@@ -698,7 +698,7 @@ func.func @m16n8k32_int8_row_col_row(%arg0: memref<128x128xi8, #gpu.address_spac
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi8, #gpu.address_space<workgroup>>, vector<16x32xi8>
%B = vector.transfer_read %arg1[%c0, %c0], %cst {in_bounds = [true, true]} : memref<128x128xi8, #gpu.address_space<workgroup>>, vector<8x32xi8>
%C = vector.transfer_read %arg2[%c0, %c0], %cst0 {in_bounds = [true, true]} : memref<128x128xi32>, vector<16x8xi32>
- // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
+ // CHECK: [[d:%.+]] = nvgpu.mma.sync({{.*}}) mmaShape = [16, 8, 32] : (vector<4x4xi8>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
%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>
// CHECK: [[lane:%.+]] = gpu.lane_id
@@ -738,10 +738,10 @@ func.func @strided_memref_read_write(%arg0: !smem_type,
// CHECK-DAG: [[m_coord:%.+]] = affine.apply [[$strided_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$contiguous_map]]
- // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] {numTiles = 4 : i32, transpose = false}
+ // CHECK: nvgpu.ldmatrix %arg0[[[m_coord]], [[k_coord]]] numTiles = 4 transpose = false
// CHECK-DAG: [[n_coord:%.+]] = affine.apply [[$contiguous_map]]
// CHECK-DAG: [[k_coord:%.+]] = affine.apply [[$strided_map]]
- // CHECK: nvgpu.ldmatrix %arg1[[[k_coord]], [[n_coord]]] {numTiles = 2 : i32, transpose = true}
+ // CHECK: nvgpu.ldmatrix %arg1[[[k_coord]], [[n_coord]]] numTiles = 2 transpose = true
%A = vector.transfer_read %arg0[%c0, %c0], %cst {in_bounds = [true, true]} : !smem_type, vector<16x16xf16>
%B = vector.transfer_read %arg1[%c0, %c0], %cst {permutation_map = #map0, in_bounds = [true, true]} : !smem_type, vector<8x16xf16>
%C = vector.transfer_read %arg2[%c0, %c0], %cst {in_bounds = [true, true]} : !smem_type, vector<16x8xf16>
diff --git a/mlir/test/Dialect/MemRef/extract-address-computations.mlir b/mlir/test/Dialect/MemRef/extract-address-computations.mlir
index adf95bb1d7d54..1e0af79c6222f 100644
--- a/mlir/test/Dialect/MemRef/extract-address-computations.mlir
+++ b/mlir/test/Dialect/MemRef/extract-address-computations.mlir
@@ -314,14 +314,13 @@ module attributes {transform.with_named_sequence} {
// CHECK-SAME: %[[DYN_OFFSET2:[^:]*]]: index)
// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
// CHECK-DAG: %[[SUBVIEW:.*]] = memref.subview %[[BASE]][%[[DYN_OFFSET0]], %[[DYN_OFFSET1]], %[[DYN_OFFSET2]]] [1, 1, 8] [1, 1, 1] : memref<4x32x32xf16, 3> to memref<1x1x8xf16, strided<[1024, 32, 1], offset: ?>, 3>
-// CHECK: %[[LOADED_VAL:.*]] = nvgpu.ldmatrix %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] {numTiles = 4 : i32, transpose = false} : memref<1x1x8xf16, strided<[1024, 32, 1], offset: ?>, 3> -> vector<4x2xf16>
+// CHECK: %[[LOADED_VAL:.*]] = nvgpu.ldmatrix %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] numTiles = 4 transpose = false : memref<1x1x8xf16, strided<[1024, 32, 1], offset: ?>, 3> -> vector<4x2xf16>
// CHECK: return %[[LOADED_VAL]] : vector<4x2xf16>
func.func @test_ldmatrix(%base : memref<4x32x32xf16, 3>,
%offset0 : index, %offset1: index, %offset2: index)
-> vector<4x2xf16> {
%loaded_val = nvgpu.ldmatrix
- %base[%offset0, %offset1, %offset2]
- {numTiles = 4 : i32, transpose = false}
+ %base[%offset0, %offset1, %offset2] numTiles = 4 transpose = false
: memref<4x32x32xf16, 3> -> vector<4x2xf16>
return %loaded_val : vector<4x2xf16>
}
@@ -351,14 +350,13 @@ module attributes {transform.with_named_sequence} {
// CHECK-DAG: {{.*}}, {{.*}}, %[[SIZES:.*]]:3, {{.*}} = memref.extract_strided_metadata %[[BASE]]
// CHECK-DAG: %[[DYN_SIZE:.*]] = affine.min #[[$MIN_8_MAP]]()[%[[SIZES]]#2, %[[DYN_OFFSET2]]]
// CHECK-DAG: %[[SUBVIEW:.*]] = memref.subview %[[BASE]][%[[DYN_OFFSET0]], %[[DYN_OFFSET1]], %[[DYN_OFFSET2]]] [1, 1, %[[DYN_SIZE]]] [1, 1, 1] : memref<?x?x?xf16, 3> to memref<1x1x?xf16, strided<[?, ?, 1], offset: ?>, 3>
-// CHECK: %[[LOADED_VAL:.*]] = nvgpu.ldmatrix %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] {numTiles = 4 : i32, transpose = false} : memref<1x1x?xf16, strided<[?, ?, 1], offset: ?>, 3> -> vector<4x2xf16>
+// CHECK: %[[LOADED_VAL:.*]] = nvgpu.ldmatrix %[[SUBVIEW]][%[[C0]], %[[C0]], %[[C0]]] numTiles = 4 transpose = false : memref<1x1x?xf16, strided<[?, ?, 1], offset: ?>, 3> -> vector<4x2xf16>
// CHECK: return %[[LOADED_VAL]] : vector<4x2xf16>
func.func @test_dynamic_ldmatrix(%base : memref<?x?x?xf16, 3>,
%offset0 : index, %offset1: index, %offset2: index)
-> vector<4x2xf16> {
%loaded_val = nvgpu.ldmatrix
- %base[%offset0, %offset1, %offset2]
- {numTiles = 4 : i32, transpose = false}
+ %base[%offset0, %offset1, %offset2] numTiles = 4 transpose = false
: memref<?x?x?xf16, 3> -> vector<4x2xf16>
return %loaded_val : vector<4x2xf16>
}
diff --git a/mlir/test/Dialect/NVGPU/fold-memref-alias-ops.mlir b/mlir/test/Dialect/NVGPU/fold-memref-alias-ops.mlir
index 44dcffcc1f00d..4b0da5580d17f 100644
--- a/mlir/test/Dialect/NVGPU/fold-memref-alias-ops.mlir
+++ b/mlir/test/Dialect/NVGPU/fold-memref-alias-ops.mlir
@@ -4,7 +4,7 @@ func.func @fold_nvgpu_device_async_copy_zero_sub_idx(%gmem_memref_3d : memref<2x
%c0 = arith.constant 0 : index
%smem_memref_4d = memref.alloc() : memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
%gmem_memref_subview_2d = memref.subview %gmem_memref_3d[%idx_1, %idx_2, %idx_3] [1, 1, 8] [1, 1, 1] : memref<2x128x768xf16> to memref<1x8xf16, strided<[98304, 1], offset: ?>>
- %async_token = nvgpu.device_async_copy %gmem_memref_subview_2d[%c0, %c0], %smem_memref_4d[%c0, %c0, %c0, %c0], 8 {bypassL1} : memref<1x8xf16, strided<[98304, 1], offset: ?>> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
+ %async_token = nvgpu.device_async_copy %gmem_memref_subview_2d[%c0, %c0], %smem_memref_4d[%c0, %c0, %c0, %c0], 8 bypassL1 : memref<1x8xf16, strided<[98304, 1], offset: ?>> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
return
}
@@ -12,7 +12,7 @@ func.func @fold_nvgpu_device_async_copy_zero_sub_idx(%gmem_memref_3d : memref<2x
// CHECK-SAME: (%[[GMEM_MEMREF_3d:.+]]: memref<2x128x768xf16>, %[[IDX_1:.+]]: index, %[[IDX_2:.+]]: index, %[[IDX_3:.+]]: index)
// CHECK-DAG: %[[c0:.+]] = arith.constant 0 : index
// CHECK-DAG: %[[SMEM_MEMREF_4d:.+]] = memref.alloc() : memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
-// CHECK: nvgpu.device_async_copy %[[GMEM_MEMREF_3d]][%[[IDX_1]], %[[IDX_2]], %[[IDX_3]]], %[[SMEM_MEMREF_4d]][%[[c0]], %[[c0]], %[[c0]], %[[c0]]], 8 {bypassL1} : memref<2x128x768xf16> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
+// CHECK: nvgpu.device_async_copy %[[GMEM_MEMREF_3d]][%[[IDX_1]], %[[IDX_2]], %[[IDX_3]]], %[[SMEM_MEMREF_4d]][%[[c0]], %[[c0]], %[[c0]], %[[c0]]], 8 bypassL1 : memref<2x128x768xf16> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
// -----
@@ -21,7 +21,7 @@ func.func @fold_src_nvgpu_device_async_copy(%gmem_memref_3d : memref<2x128x768xf
%c0 = arith.constant 0 : index
%smem_memref_4d = memref.alloc() : memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
%gmem_memref_subview_2d = memref.subview %gmem_memref_3d[%src_idx_0, %src_idx_1, %src_idx_2] [1, 1, 8] [1, 1, 1] : memref<2x128x768xf16> to memref<1x8xf16, strided<[98304, 1], offset: ?>>
- %async_token = nvgpu.device_async_copy %gmem_memref_subview_2d[%src_sub_idx_0, %src_sub_idx_1], %smem_memref_4d[%c0, %c0, %c0, %c0], 8 {bypassL1} : memref<1x8xf16, strided<[98304, 1], offset: ?>> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
+ %async_token = nvgpu.device_async_copy %gmem_memref_subview_2d[%src_sub_idx_0, %src_sub_idx_1], %smem_memref_4d[%c0, %c0, %c0, %c0], 8 bypassL1 : memref<1x8xf16, strided<[98304, 1], offset: ?>> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
return
}
@@ -31,7 +31,7 @@ func.func @fold_src_nvgpu_device_async_copy(%gmem_memref_3d : memref<2x128x768xf
// CHECK-DAG: %[[c0:.+]] = arith.constant 0 : index
// CHECK-DAG: %[[RESOLVED_SRC_IDX_0:.+]] = affine.apply #[[MAP]]()[%[[SRC_IDX_0]], %[[SRC_SUB_IDX_0]]]
// CHECK-DAG: %[[RESOLVED_SRC_IDX_1:.+]] = affine.apply #[[MAP]]()[%[[SRC_IDX_2]], %[[SRC_SUB_IDX_1]]]
-// CHECK-DAG: nvgpu.device_async_copy %[[GMEM_MEMREF_3d]][%[[RESOLVED_SRC_IDX_0]], %[[SRC_IDX_1]], %[[RESOLVED_SRC_IDX_1]]], %[[SMEM_MEMREF_4d]][%[[c0]], %[[c0]], %[[c0]], %[[c0]]], 8 {bypassL1} : memref<2x128x768xf16> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
+// CHECK-DAG: nvgpu.device_async_copy %[[GMEM_MEMREF_3d]][%[[RESOLVED_SRC_IDX_0]], %[[SRC_IDX_1]], %[[RESOLVED_SRC_IDX_1]]], %[[SMEM_MEMREF_4d]][%[[c0]], %[[c0]], %[[c0]], %[[c0]]], 8 bypassL1 : memref<2x128x768xf16> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
// -----
@@ -41,7 +41,7 @@ func.func @fold_src_fold_dest_nvgpu_device_async_copy(%gmem_memref_3d : memref<2
%smem_memref_4d = memref.alloc() : memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
%gmem_memref_subview_2d = memref.subview %gmem_memref_3d[%src_idx_0, %src_idx_1, %src_idx_2] [1, 1, 8] [1, 1, 1] : memref<2x128x768xf16> to memref<1x8xf16, strided<[98304, 1], offset: ?>>
%smem_memref_2d = memref.subview %smem_memref_4d[%dest_idx_0, %dest_idx_1, %dest_idx_2, %dest_idx_3] [1, 1, 1, 8] [1, 1, 1, 1] : memref<5x1x64x64xf16, #gpu.address_space<workgroup>> to memref<1x8xf16, strided<[4096, 1], offset: ?>, #gpu.address_space<workgroup>>
- %async_token = nvgpu.device_async_copy %gmem_memref_subview_2d[%src_sub_idx_0, %src_sub_idx_1], %smem_memref_2d[%dest_sub_idx_0, %dest_sub_idx_1], 8 {bypassL1} : memref<1x8xf16, strided<[98304, 1], offset: ?>> to memref<1x8xf16, strided<[4096, 1], offset: ?>, #gpu.address_space<workgroup>>
+ %async_token = nvgpu.device_async_copy %gmem_memref_subview_2d[%src_sub_idx_0, %src_sub_idx_1], %smem_memref_2d[%dest_sub_idx_0, %dest_sub_idx_1], 8 bypassL1 : memref<1x8xf16, strided<[98304, 1], offset: ?>> to memref<1x8xf16, strided<[4096, 1], offset: ?>, #gpu.address_space<workgroup>>
return
}
@@ -52,7 +52,7 @@ func.func @fold_src_fold_dest_nvgpu_device_async_copy(%gmem_memref_3d : memref<2
// CHECK-DAG: %[[RESOLVED_SRC_IDX_1:.+]] = affine.apply #[[MAP]]()[%[[SRC_IDX_2]], %[[SRC_SUB_IDX_1]]]
// CHECK-DAG: %[[RESOLVED_DST_IDX_1:.+]] = affine.apply #[[MAP]]()[%[[DEST_IDX_1]], %[[DEST_SUB_IDX_0]]]
// CHECK-DAG: %[[RESOLVED_DST_IDX_3:.+]] = affine.apply #[[MAP]]()[%[[DEST_IDX_3]], %[[DEST_SUB_IDX_1]]]
-// CHECK-DAG: nvgpu.device_async_copy %[[GMEM_MEMREF_3d]][%[[RESOLVED_SRC_IDX_0]], %[[SRC_IDX_1]], %[[RESOLVED_SRC_IDX_1]]], %[[SMEM_MEMREF_4d]][%[[DEST_IDX_0]], %[[RESOLVED_DST_IDX_1]], %[[DEST_IDX_2]], %[[RESOLVED_DST_IDX_3]]], 8 {bypassL1} : memref<2x128x768xf16> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
+// CHECK-DAG: nvgpu.device_async_copy %[[GMEM_MEMREF_3d]][%[[RESOLVED_SRC_IDX_0]], %[[SRC_IDX_1]], %[[RESOLVED_SRC_IDX_1]]], %[[SMEM_MEMREF_4d]][%[[DEST_IDX_0]], %[[RESOLVED_DST_IDX_1]], %[[DEST_IDX_2]], %[[RESOLVED_DST_IDX_3]]], 8 bypassL1 : memref<2x128x768xf16> to memref<5x1x64x64xf16, #gpu.address_space<workgroup>>
// -----
@@ -65,7 +65,7 @@ func.func @test_ldmatrix(%arg0: memref<4x32x32xf16, 3>, %arg1: index, %arg2: ind
%1 = affine.apply #map1()[%arg2]
%2 = affine.apply #map1()[%arg3]
%subview = memref.subview %arg0[%arg1, %arg2, %arg3] [%0, %1, %2] [1, 1, 1] : memref<4x32x32xf16, 3> to memref<?x?x?xf16, strided<[1024, 32, 1], offset: ?>, 3>
- %3 = nvgpu.ldmatrix %subview[%c0, %c0, %c0] {numTiles = 4 : i32, transpose = false} : memref<?x?x?xf16, strided<[1024, 32, 1], offset: ?>, 3> -> vector<4x2xf16>
+ %3 = nvgpu.ldmatrix %subview[%c0, %c0, %c0] numTiles = 4 transpose = false : memref<?x?x?xf16, strided<[1024, 32, 1], offset: ?>, 3> -> vector<4x2xf16>
return %3 : vector<4x2xf16>
}
@@ -74,13 +74,13 @@ func.func @test_ldmatrix(%arg0: memref<4x32x32xf16, 3>, %arg1: index, %arg2: ind
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index
-// CHECK: nvgpu.ldmatrix %[[ARG0]][%[[ARG1]], %[[ARG2]], %[[ARG3]]] {numTiles = 4 : i32, transpose = false} : memref<4x32x32xf16, 3> -> vector<4x2xf16>
+// CHECK: nvgpu.ldmatrix %[[ARG0]][%[[ARG1]], %[[ARG2]], %[[ARG3]]] numTiles = 4 transpose = false : memref<4x32x32xf16, 3> -> vector<4x2xf16>
// -----
func.func @ldmatrix_expand(%arg0: memref<4096xf16, 3>, %arg1: index, %arg2: index, %arg3: index) -> vector<4x2xf16> {
%exp = memref.expand_shape %arg0 [[0, 1, 2]] output_shape [4, 32, 32] : memref<4096xf16, 3> into memref<4x32x32xf16, 3>
- %3 = nvgpu.ldmatrix %exp[%arg1, %arg2, %arg3] {numTiles = 4 : i32, transpose = false} : memref<4x32x32xf16, 3> -> vector<4x2xf16>
+ %3 = nvgpu.ldmatrix %exp[%arg1, %arg2, %arg3] numTiles = 4 transpose = false : memref<4x32x32xf16, 3> -> vector<4x2xf16>
return %3 : vector<4x2xf16>
}
@@ -90,4 +90,4 @@ func.func @ldmatrix_expand(%arg0: memref<4096xf16, 3>, %arg1: index, %arg2: inde
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG3:[a-zA-Z0-9_]+]]: index
// CHECK: %[[LIN:[a-zA-Z0-9_]+]] = affine.linearize_index disjoint [%[[ARG1]], %[[ARG2]], %[[ARG3]]] by (4, 32, 32)
-// CHECK: nvgpu.ldmatrix %[[ARG0]][%[[LIN]]] {numTiles = 4 : i32, transpose = false} : memref<4096xf16, 3> -> vector<4x2xf16>
+// CHECK: nvgpu.ldmatrix %[[ARG0]][%[[LIN]]] numTiles = 4 transpose = false : memref<4096xf16, 3> -> vector<4x2xf16>
diff --git a/mlir/test/Dialect/NVGPU/invalid.mlir b/mlir/test/Dialect/NVGPU/invalid.mlir
index 0341a08a2c778..46e17eec3f932 100644
--- a/mlir/test/Dialect/NVGPU/invalid.mlir
+++ b/mlir/test/Dialect/NVGPU/invalid.mlir
@@ -3,7 +3,7 @@
func.func @ldmatrix_address_space_f16_x4(%arg0: memref<128x128xf16, 2>) -> vector<4x1xf16> {
%c0 = arith.constant 0 : index
// expected-error @below {{expected nvgpu.ldmatrix srcMemref must have a memory space attribute of IntegerAttr(3) or gpu::AddressSpaceAttr(Workgroup)}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf16, 2> -> vector<4x1xf16>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf16, 2> -> vector<4x1xf16>
return %a : vector<4x1xf16>
}
// -----
@@ -11,7 +11,7 @@ func.func @ldmatrix_address_space_f16_x4(%arg0: memref<128x128xf16, 2>) -> vect
func.func @ldmatrix_num_elements_f16_x4(%arg0: memref<128x128xf16, 3>) -> vector<4x1xf16> {
%c0 = arith.constant 0 : index
// expected-error @+1 {{expected vector register shape[1] = 2}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf16, 3> -> vector<4x1xf16>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf16, 3> -> vector<4x1xf16>
return %a : vector<4x1xf16>
}
// -----
@@ -19,7 +19,7 @@ func.func @ldmatrix_num_elements_f16_x4(%arg0: memref<128x128xf16, 3>) -> vecto
func.func @ldmatrix_num_tiles_f16_x4(%arg0: memref<128x128xf16, 3>) -> vector<2x2xf16> {
%c0 = arith.constant 0 : index
// expected-error @+1 {{expected vector register shape[0] and numTiles to match}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf16, 3> -> vector<2x2xf16>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf16, 3> -> vector<2x2xf16>
return %a : vector<2x2xf16>
}
// -----
@@ -27,7 +27,7 @@ func.func @ldmatrix_num_tiles_f16_x4(%arg0: memref<128x128xf16, 3>) -> vector<2
func.func @ldmatrix_num_tiles_f32_x4(%arg0: memref<128x128xf32, 3>) -> vector<4x2xf32> {
%c0 = arith.constant 0 : index
// expected-error @+1 {{expected vector register shape[1] = 1}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf32, 3> -> vector<4x2xf32>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf32, 3> -> vector<4x2xf32>
return %a : vector<4x2xf32>
}
// -----
@@ -35,7 +35,7 @@ func.func @ldmatrix_num_tiles_f32_x4(%arg0: memref<128x128xf32, 3>) -> vector<4
func.func @ldmatrix_trans_f32_x4(%arg0: memref<128x128xf32, 3>) -> vector<4x1xf32> {
%c0 = arith.constant 0 : index
// expected-error @+1 {{nvgpu.ldmatrix transpose works only at 16b granularity}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = true, numTiles = 4 : i32} : memref<128x128xf32, 3> -> vector<4x1xf32>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = true : memref<128x128xf32, 3> -> vector<4x1xf32>
return %a : vector<4x1xf32>
}
// -----
@@ -43,7 +43,7 @@ func.func @ldmatrix_trans_f32_x4(%arg0: memref<128x128xf32, 3>) -> vector<4x1xf
func.func @ldmatrix_trans_f32_x4(%arg0: memref<128x128xf32, 3>) -> vector<4xf32> {
%c0 = arith.constant 0 : index
// expected-error @+1 {{results must be 2 dimensional vector}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf32, 3> -> vector<4xf32>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf32, 3> -> vector<4xf32>
return %a : vector<4xf32>
}
// -----
@@ -51,84 +51,84 @@ func.func @ldmatrix_trans_f32_x4(%arg0: memref<128x128xf32, 3>) -> vector<4xf32
func.func @ldmatrix_type_x4(%arg0: memref<128x128xf32, 3>) -> vector<4x2xf16> {
%c0 = arith.constant 0 : index
// expected-error @+1 {{'nvgpu.ldmatrix' op failed to verify that srcMemref and res have same element type}}
- %a = nvgpu.ldmatrix %arg0[%c0, %c0] {transpose = false, numTiles = 4 : i32} : memref<128x128xf32, 3> -> vector<4x2xf16>
+ %a = nvgpu.ldmatrix %arg0[%c0, %c0] numTiles = 4 transpose = false : memref<128x128xf32, 3> -> vector<4x2xf16>
return %a : vector<4x2xf16>
}
// -----
func.func @m16n8k16_fp16_vector_shape_a(%arg0: vector<4x4xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{expected 256 warp-wide matrix A elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x4xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x4xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
// -----
func.func @m16n8k16_fp16_vector_shape_b(%arg0: vector<4x2xf16>, %arg1: vector<2x4xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{expected 128 warp-wide matrix B elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x4xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x4xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
// -----
func.func @m16n8k16_fp16_vector_shape_c(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x4xf16>) -> vector<2x4xf16> {
// expected-error @+1 {{expected 128 warp-wide matrix C elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x4xf16>) -> vector<2x4xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x4xf16>) -> vector<2x4xf16>
return %d : vector<2x4xf16>
}
// -----
func.func @m16n8k16_fp16_vector_shape_a_extended(%arg0: vector<2x4xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{expected matrix A to be shaped (4 x 2)}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<2x4xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<2x4xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
// -----
func.func @m16n8k16_fp16_tf32Enabled(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{expected tf32 tensor cores only for F32 operands}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16], tf32Enabled} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] tf32Enabled : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
// -----
func.func @m16n8k8_fp32_vector_shape_a(%arg0: vector<4x2xf32>, %arg1: vector<2x1xf32>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
// expected-error @+1 {{expected 128 warp-wide matrix A elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<4x2xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 8] : (vector<4x2xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
// -----
func.func @m16n8k8_fp32_vector_shape_a_extended(%arg0: vector<1x4xf32>, %arg1: vector<2x1xf32>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
// expected-error @+1 {{expected matrix A to be shaped (4 x 1)}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<1x4xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 8] : (vector<1x4xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
// -----
func.func @m8n8k4_fp64_vector_shape_a(%arg0: vector<1x2xf64>, %arg1: vector<1x1xf64>, %arg2: vector<1x2xf64>) -> vector<1x2xf64> {
// expected-error @+1 {{expected 32 warp-wide matrix A elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [8, 8, 4]} : (vector<1x2xf64>, vector<1x1xf64>, vector<1x2xf64>) -> vector<1x2xf64>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [8, 8, 4] : (vector<1x2xf64>, vector<1x1xf64>, vector<1x2xf64>) -> vector<1x2xf64>
return %d : vector<1x2xf64>
}
// -----
func.func @m8n8k4_fp64_vector_shape_c_extended(%arg0: vector<1x1xf64>, %arg1: vector<1x1xf64>, %arg2: vector<2x1xf64>) -> vector<2x1xf64> {
// expected-error @+1 {{expected matrix C to be shaped (1 x 2)}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [8, 8, 4]} : (vector<1x1xf64>, vector<1x1xf64>, vector<2x1xf64>) -> vector<2x1xf64>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [8, 8, 4] : (vector<1x1xf64>, vector<1x1xf64>, vector<2x1xf64>) -> vector<2x1xf64>
return %d : vector<2x1xf64>
}
// -----
func.func @m16n8k32_int8_vector_shape_b(%arg0: vector<4x4xi8>, %arg1: vector<4x4xi8>, %arg2: vector<2x2xi32>) -> vector<2x2xi32> {
// expected-error @+1 {{expected 256 warp-wide matrix B elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xi8>, vector<4x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 32] : (vector<4x4xi8>, vector<4x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
// -----
func.func @m16n8k32_int32_datatype(%arg0: vector<4x4xi32>, %arg1: vector<2x4xi8>, %arg2: vector<2x2xi32>) -> vector<2x2xi32> {
// expected-error @+1 {{op failed to verify that matrixA and matrixB have same element type}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 32]} : (vector<4x4xi32>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 32] : (vector<4x4xi32>, vector<2x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
// -----
@@ -189,7 +189,7 @@ func.func @mma_sp_sync_f16_16816(%arg0: vector<2x2xf16>,
%arg2: vector<2x2xf16>,
%arg3: vector<2xi16>) -> vector<2x2xf16> {
// expected-error @+1 {{'nvgpu.mma.sp.sync' op sparsity selector should be 0 or 1}}
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 16], sparsitySelector = 42 : i32} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 16] sparsitySelector = 42 :
(vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -199,7 +199,7 @@ func.func @mma_sp_sync_f16_16816(%arg0: vector<2x2xf16>,
func.func @async_cp_zfill_f32_align1(
%src: memref<128x128xf32>, %dst: memref<3x16x128xf32, 3>, %i : index, %srcElements : index) {
// expected-error @+1 {{'nvgpu.device_async_copy' op bypassL1 does not satify alignment for 'memref<3x16x128xf32, 3>' with destination element 1. Unset bypassL1, or set destination element to 4}}
- %0 = nvgpu.device_async_copy %src[%i, %i], %dst[%i, %i, %i], 1, %srcElements {bypassL1} : memref<128x128xf32> to memref<3x16x128xf32, 3>
+ %0 = nvgpu.device_async_copy %src[%i, %i], %dst[%i, %i, %i], 1, %srcElements bypassL1 : memref<128x128xf32> to memref<3x16x128xf32, 3>
return
}
@@ -340,26 +340,26 @@ func.func @tma_generate_descriptor_incorrect_last_dim(%desc: !desc, %buffer2: m
func.func @rcp_unsupported_rounding_0(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.rcp' op has a limitation. #nvvm.fp_rnd_mode<rn> is not supported yet.}}
- %out = nvgpu.rcp %in {rounding = #nvvm.fp_rnd_mode<rn>, approx = true, ftz = true} : vector<16xf32>
+ %out = nvgpu.rcp %in <{rounding = #nvvm.fp_rnd_mode<rn>, approx = true, ftz = true}> : vector<16xf32>
}
// -----
func.func @rcp_unsupported_rounding_1(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.rcp' op has a limitation. non-approx or non-ftz is not supported yet.}}
- %out = nvgpu.rcp %in {ftz = true} : vector<16xf32>
+ %out = nvgpu.rcp %in <{ftz = true}> : vector<16xf32>
}
// -----
func.func @rcp_unsupported_ftz(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.rcp' op has a limitation. non-approx or non-ftz is not supported yet.}}
- %out = nvgpu.rcp %in {approx = true} : vector<16xf32>
+ %out = nvgpu.rcp %in <{approx = true}> : vector<16xf32>
}
// -----
func.func @check_matrixA_dim(%arg0: vector<16xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{matrixA must be 2 dimensional vector}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<16xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<16xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -367,7 +367,7 @@ func.func @check_matrixA_dim(%arg0: vector<16xf16>, %arg1: vector<2x2xf16>, %arg
func.func @check_matrixB_dim(%arg0: vector<4x4xf16>, %arg1: vector<4xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{matrixB must be 2 dimensional vector}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x4xf16>, vector<4xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x4xf16>, vector<4xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -375,7 +375,7 @@ func.func @check_matrixB_dim(%arg0: vector<4x4xf16>, %arg1: vector<4xf16>, %arg2
func.func @check_matrixC_dim(%arg0: vector<4x4xf16>, %arg1: vector<2x2xf16>, %arg2: vector<4xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{matrixC must be 2 dimensional vector}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x4xf16>, vector<2x2xf16>, vector<4xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x4xf16>, vector<2x2xf16>, vector<4xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -394,7 +394,7 @@ func.func @tma_last_dim_bytes(%desc: !desc, %buffer: memref<32x8xi8,3>, %mbarrie
func.func @mma_sync_invalid_shape_2_elements(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{mmaShape must have exactly 3 elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -402,7 +402,7 @@ func.func @mma_sync_invalid_shape_2_elements(%arg0: vector<4x2xf16>, %arg1: vect
func.func @mma_sync_invalid_shape_4_elements(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
// expected-error @+1 {{mmaShape must have exactly 3 elements}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16, 4]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16, 4] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -410,7 +410,7 @@ func.func @mma_sync_invalid_shape_4_elements(%arg0: vector<4x2xf16>, %arg1: vect
func.func @mma_sparse_sync_invalid_shape_2_elements(%arg0: vector<2x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>, %arg3: vector<2xi16>) -> vector<2x2xf16> {
// expected-error @+1 {{mmaShape must have exactly 3 elements}}
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8], sparsitySelector = 0 : i32} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8] sparsitySelector = 0 :
(vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -419,7 +419,7 @@ func.func @mma_sparse_sync_invalid_shape_2_elements(%arg0: vector<2x2xf16>, %arg
func.func @mma_sparse_sync_invalid_shape_4_elements(%arg0: vector<2x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>, %arg3: vector<2xi16>) -> vector<2x2xf16> {
// expected-error @+1 {{mmaShape must have exactly 3 elements}}
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 16, 4], sparsitySelector = 0 : i32} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 16, 4] sparsitySelector = 0 :
(vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
diff --git a/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32.mlir b/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32.mlir
index 6d0cf348273c3..586e1afbcee31 100644
--- a/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32.mlir
+++ b/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32.mlir
@@ -4,7 +4,7 @@
func.func @m16n8k4_tf32(%arg0: vector<2x1xf32>, %arg1: vector<1x1xf32>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
// CHECK: nvgpu.mma.sync
// CHECK-SAME: tf32Enabled
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 4]} : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 4] : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
@@ -14,7 +14,7 @@ func.func @m16n8k4_tf32(%arg0: vector<2x1xf32>, %arg1: vector<1x1xf32>, %arg2: v
func.func @m16n8k8_tf32(%arg0: vector<4x1xf32>, %arg1: vector<2x1xf32>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
// CHECK: nvgpu.mma.sync
// CHECK-SAME: tf32Enabled
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<4x1xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 8] : (vector<4x1xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
// -----
@@ -24,6 +24,6 @@ func.func @m16n8k8_tf32(%arg0: vector<4x1xf32>, %arg1: vector<2x1xf32>, %arg2: v
// CHECK-NOT: tf32Enabled
// CHECK: return
func.func @mma_sync_f16(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
diff --git a/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32x3.mlir b/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32x3.mlir
index 7db60254d2ec8..07afc8e094331 100644
--- a/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32x3.mlir
+++ b/mlir/test/Dialect/NVGPU/mma-sync-f32-to-tf32x3.mlir
@@ -3,7 +3,7 @@
// CHECK-LABEL: m16n8k4_tf32
func.func @m16n8k4_tf32(%arg0: vector<2x1xf32>, %arg1: vector<1x1xf32>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
// expected-error @+1 {{TF32x3 is not supported at the moment for nvgpu.mma.sync on f32 datatype}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 4]} : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 4] : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
@@ -12,7 +12,7 @@ func.func @m16n8k4_tf32(%arg0: vector<2x1xf32>, %arg1: vector<1x1xf32>, %arg2: v
// CHECK-LABEL: m16n8k8_tf32
func.func @m16n8k8_tf32(%arg0: vector<4x1xf32>, %arg1: vector<2x1xf32>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
// expected-error @+1 {{TF32x3 is not supported at the moment for nvgpu.mma.sync on f32 datatype}}
- %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 8]} : (vector<4x1xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+ %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) mmaShape = [16, 8, 8] : (vector<4x1xf32>, vector<2x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
return %d : vector<2x2xf32>
}
// -----
diff --git a/mlir/test/Dialect/NVGPU/nvgpu-extf-invalid.mlir b/mlir/test/Dialect/NVGPU/nvgpu-extf-invalid.mlir
index ef659efb96f8e..d81c438aa29e9 100644
--- a/mlir/test/Dialect/NVGPU/nvgpu-extf-invalid.mlir
+++ b/mlir/test/Dialect/NVGPU/nvgpu-extf-invalid.mlir
@@ -28,7 +28,7 @@ func.func @fpext_e8m0_to_f16(%in : vector<16xf8E8M0FNU>) {
func.func @fpext_bad_rounding(%in : vector<16xf8E5M2>) {
// expected-error @+1 {{'nvgpu.extf' op expects RN rounding mode, but got #nvvm.fp_rnd_mode<rz>}}
- %out = nvgpu.extf %in {rnd = #nvvm.fp_rnd_mode<rz>}
+ %out = nvgpu.extf %in <{rnd = #nvvm.fp_rnd_mode<rz>}>
: vector<16xf8E5M2> to vector<16xf16>
return
}
@@ -37,7 +37,7 @@ func.func @fpext_bad_rounding(%in : vector<16xf8E5M2>) {
func.func @fpext_relu_bf16(%in : vector<8xf8E5M2>) {
// expected-error @+1 {{'nvgpu.extf' op relu is not supported for bf16 destination}}
- %out = nvgpu.extf %in {relu = true} : vector<8xf8E5M2> to vector<8xbf16>
+ %out = nvgpu.extf %in <{relu = true}> : vector<8xf8E5M2> to vector<8xbf16>
return
}
diff --git a/mlir/test/Dialect/NVGPU/nvgpu-truncf-invalid.mlir b/mlir/test/Dialect/NVGPU/nvgpu-truncf-invalid.mlir
index 1b2fb3db5ceca..c99f5279c0315 100644
--- a/mlir/test/Dialect/NVGPU/nvgpu-truncf-invalid.mlir
+++ b/mlir/test/Dialect/NVGPU/nvgpu-truncf-invalid.mlir
@@ -20,7 +20,7 @@ func.func @fptrunc_src_bitwidth(%in : vector<16xf8E5M2>) {
func.func @fptrunc_e8m0_bad_rounding(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.truncf' op expects RZ or RP rounding mode when result type is e8m0, but got #nvvm.fp_rnd_mode<rn>}}
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rn>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rn>}>
: vector<16xf32> to vector<16xf8E8M0FNU>
return
}
@@ -29,7 +29,7 @@ func.func @fptrunc_e8m0_bad_rounding(%in : vector<16xf32>) {
func.func @fptrunc_unsupported_sat_mode(%in : vector<8xf32>) {
// expected-error @+1 {{'nvgpu.truncf' op attribute 'sat' failed to satisfy constraint: Describes the saturation mode whose value is one of {none, satfinite}}}
- %out = nvgpu.truncf %in {sat = #nvvm.sat_mode<sat>}
+ %out = nvgpu.truncf %in <{sat = #nvvm.sat_mode<sat>}>
: vector<8xf32> to vector<8xf8E4M3FN>
return
}
@@ -38,7 +38,7 @@ func.func @fptrunc_unsupported_sat_mode(%in : vector<8xf32>) {
func.func @fptrunc_f32_to_f8_rz(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.truncf' op expects RN rounding mode, but got #nvvm.fp_rnd_mode<rz>}}
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rz>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rz>}>
: vector<16xf32> to vector<16xf8E4M3FN>
return
}
@@ -47,7 +47,7 @@ func.func @fptrunc_f32_to_f8_rz(%in : vector<16xf32>) {
func.func @fptrunc_f64_to_f16_rz(%in : vector<4xf64>) {
// expected-error @+1 {{'nvgpu.truncf' op expects RN rounding mode for f64 input, but got #nvvm.fp_rnd_mode<rz>}}
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rz>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rz>}>
: vector<4xf64> to vector<4xf16>
return
}
@@ -56,7 +56,7 @@ func.func @fptrunc_f64_to_f16_rz(%in : vector<4xf64>) {
func.func @fptrunc_rs_unsupported_types(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.truncf' op RS (stochastic) rounding is only supported for f32->f16/bf16, got 'f32' -> 'f8E4M3FN'}}
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rs>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rs>}>
: vector<16xf32> to vector<16xf8E4M3FN>
return
}
@@ -65,7 +65,7 @@ func.func @fptrunc_rs_unsupported_types(%in : vector<16xf32>) {
func.func @fptrunc_rs_no_random_bits(%in : vector<4xf32>) {
// expected-error @+1 {{'nvgpu.truncf' op random_bits operand is required with RS rounding}}
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rs>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rs>}>
: vector<4xf32> to vector<4xf16>
return
}
@@ -74,7 +74,7 @@ func.func @fptrunc_rs_no_random_bits(%in : vector<4xf32>) {
func.func @fptrunc_bad_rounding(%in : vector<16xf32>) {
// expected-error @+1 {{'nvgpu.truncf' op expects RN rounding mode, but got #nvvm.fp_rnd_mode<rp>}}
- %out = nvgpu.truncf %in {rnd = #nvvm.fp_rnd_mode<rp>}
+ %out = nvgpu.truncf %in <{rnd = #nvvm.fp_rnd_mode<rp>}>
: vector<16xf32> to vector<16xf8E4M3FN>
return
}
diff --git a/mlir/test/Dialect/NVGPU/optimize-shared-memory.mlir b/mlir/test/Dialect/NVGPU/optimize-shared-memory.mlir
index 8c4381c2bcbf6..2682ad52b8268 100644
--- a/mlir/test/Dialect/NVGPU/optimize-shared-memory.mlir
+++ b/mlir/test/Dialect/NVGPU/optimize-shared-memory.mlir
@@ -20,7 +20,7 @@ func.func @optimize_128x32xf16_32x128xf16(%arg0: memref<128x128xf16>,
%0 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shm[%stRow, %stCol], 8
: memref<128x128xf16> to memref<128x32xf16, 3>
%1 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: [[c6:%.+]] = arith.constant 6 : index
// CHECK: [[srcBits:%.+]] = arith.andi [[fragRow]], [[c6]]
@@ -28,7 +28,7 @@ func.func @optimize_128x32xf16_32x128xf16(%arg0: memref<128x128xf16>,
// CHECK: [[xorBits:%.+]] = arith.shli [[srcBits]], [[c2]]
// CHECK: [[fragColPerm:%.+]] = arith.xori [[fragCol]], [[xorBits]]
// CHECK: nvgpu.ldmatrix [[shm]][[[fragRow]], [[fragColPerm]]]
- %mat = nvgpu.ldmatrix %shm[%fragRow, %fragCol] {numTiles = 4 : i32, transpose = false}
+ %mat = nvgpu.ldmatrix %shm[%fragRow, %fragCol] numTiles = 4 transpose = false
: memref<128x32xf16, 3> -> vector<4x2xf16>
// CHECK: [[c15:%.+]] = arith.constant 15 : index
@@ -40,7 +40,7 @@ func.func @optimize_128x32xf16_32x128xf16(%arg0: memref<128x128xf16>,
%2 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shmB[%stRow, %stCol], 8
: memref<128x128xf16> to memref<32x128xf16, 3>
%3 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: [[c15:%.+]] = arith.constant 15 : index
// CHECK: [[srcBits:%.+]] = arith.andi [[fragRow]], [[c15]]
@@ -48,7 +48,7 @@ func.func @optimize_128x32xf16_32x128xf16(%arg0: memref<128x128xf16>,
// CHECK: [[xorBits:%.+]] = arith.shli [[srcBits]], [[c3]]
// CHECK: [[fragColPerm:%.+]] = arith.xori [[fragCol]], [[xorBits]]
// CHECK: nvgpu.ldmatrix [[shmB]][[[fragRow]], [[fragColPerm]]]
- %matB = nvgpu.ldmatrix %shmB[%fragRow, %fragCol] {numTiles = 4 : i32, transpose = false}
+ %matB = nvgpu.ldmatrix %shmB[%fragRow, %fragCol] numTiles = 4 transpose = false
: memref<32x128xf16, 3> -> vector<4x2xf16>
return %mat, %matB: vector<4x2xf16>, vector<4x2xf16>
@@ -77,7 +77,7 @@ func.func @optimize_64x16xf32_16x64xf32(%arg0: memref<128x128xf32>,
%0 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shm[%stRow, %stCol], 4
: memref<128x128xf32> to memref<64x16xf32, 3>
%1 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: [[c6:%.+]] = arith.constant 6 : index
// CHECK: [[srcBits:%.+]] = arith.andi [[fragRow]], [[c6]]
@@ -85,7 +85,7 @@ func.func @optimize_64x16xf32_16x64xf32(%arg0: memref<128x128xf32>,
// CHECK: [[xorBits:%.+]] = arith.shli [[srcBits]], [[c1]]
// CHECK: [[fragColPerm:%.+]] = arith.xori [[fragCol]], [[xorBits]]
// CHECK: nvgpu.ldmatrix [[shm]][[[fragRow]], [[fragColPerm]]]
- %mat = nvgpu.ldmatrix %shm[%fragRow, %fragCol] {numTiles = 4 : i32, transpose = false}
+ %mat = nvgpu.ldmatrix %shm[%fragRow, %fragCol] numTiles = 4 transpose = false
: memref<64x16xf32, 3> -> vector<4x1xf32>
// CHECK: [[c6:%.+]] = arith.constant 6 : index
@@ -133,7 +133,7 @@ func.func @optimize_64x16xf32_16x64xf32(%arg0: memref<128x128xf32>,
%2 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shmB[%stRow, %stCol], 4
: memref<128x128xf32> to memref<16x64xf32, 3>
%3 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: [[c15:%.+]] = arith.constant 15 : index
// CHECK: [[srcBits:%.+]] = arith.andi [[fragRow]], [[c15]]
@@ -141,7 +141,7 @@ func.func @optimize_64x16xf32_16x64xf32(%arg0: memref<128x128xf32>,
// CHECK: [[xorBits:%.+]] = arith.shli [[srcBits]], [[c2]]
// CHECK: [[fragColPerm:%.+]] = arith.xori [[fragCol]], [[xorBits]]
// CHECK: nvgpu.ldmatrix [[shmB]][[[fragRow]], [[fragColPerm]]]
- %matB = nvgpu.ldmatrix %shmB[%fragRow, %fragCol] {numTiles = 4 : i32, transpose = false}
+ %matB = nvgpu.ldmatrix %shmB[%fragRow, %fragCol] numTiles = 4 transpose = false
: memref<16x64xf32, 3> -> vector<4x1xf32>
// CHECK: [[c15:%.+]] = arith.constant 15 : index
@@ -178,7 +178,7 @@ func.func @small_column_size_f64(%arg0: memref<32x32xf64>,
%0 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shm[%stRow, %stCol], 2
: memref<32x32xf64> to memref<32x4xf64, 3>
%1 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: [[c6:%.+]] = arith.constant 4 : index
// CHECK: [[srcBits:%.+]] = arith.andi [[fragRow]], [[c6]]
@@ -204,10 +204,10 @@ func.func @too_small_column_size_f16(%arg0: memref<128x128xf16>,
%0 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shm[%stRow, %stCol], 8
: memref<128x128xf16> to memref<128x8xf16, 3>
%1 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: nvgpu.ldmatrix [[shm]][[[fragRow]], [[fragCol]]]
- %mat = nvgpu.ldmatrix %shm[%fragRow, %fragCol] {numTiles = 1 : i32, transpose = false}
+ %mat = nvgpu.ldmatrix %shm[%fragRow, %fragCol] numTiles = 1 transpose = false
: memref<128x8xf16, 3> -> vector<1x2xf16>
return %mat: vector<1x2xf16>
@@ -230,10 +230,10 @@ func.func @abort_if_subview(%arg0: memref<128x128xf16>,
%0 = nvgpu.device_async_copy %arg0[%ldRow, %ldCol], %shm[%stRow, %stCol], 8
: memref<128x128xf16> to memref<128x32xf16, 3>
%1 = nvgpu.device_async_create_group %0
- nvgpu.device_async_wait %1 { numGroups = 1 : i32}
+ nvgpu.device_async_wait %1 numGroups = 1
// CHECK: nvgpu.ldmatrix [[shmView]][[[fragRow]], [[fragCol]]]
- %mat = nvgpu.ldmatrix %shmView[%fragRow, %fragCol] {numTiles = 1 : i32, transpose = false}
+ %mat = nvgpu.ldmatrix %shmView[%fragRow, %fragCol] numTiles = 1 transpose = false
: memref<64x32xf16, 3> -> vector<1x2xf16>
return %mat: vector<1x2xf16>
diff --git a/mlir/test/Dialect/NVGPU/roundtrip.mlir b/mlir/test/Dialect/NVGPU/roundtrip.mlir
index 77586fc9b7abf..93cfc5ba5b9a9 100644
--- a/mlir/test/Dialect/NVGPU/roundtrip.mlir
+++ b/mlir/test/Dialect/NVGPU/roundtrip.mlir
@@ -3,8 +3,8 @@
// CHECK-LABEL: func @ldmatrix(
func.func @ldmatrix(%arg0: memref<?x?xf16, 3>, %x: index, %y: index) {
// CHECK: nvgpu.ldmatrix %{{.*}}[%{{.*}}, %{{.*}}]
-// CHECK-SAME: {numTiles = 4 : i32, transpose = false} : memref<?x?xf16, 3> -> vector<4x2xf16>
- %l = nvgpu.ldmatrix %arg0[%x, %y] {numTiles = 4 : i32, transpose = false} :
+// CHECK-SAME: numTiles = 4 transpose = false : memref<?x?xf16, 3> -> vector<4x2xf16>
+ %l = nvgpu.ldmatrix %arg0[%x, %y] numTiles = 4 transpose = false :
memref<?x?xf16, 3> -> vector<4x2xf16>
return
}
@@ -13,8 +13,8 @@ func.func @ldmatrix(%arg0: memref<?x?xf16, 3>, %x: index, %y: index) {
func.func @mma_sync(%arg0: vector<4x2xf16>,
%arg1: vector<2x2xf16>,
%arg2: vector<2x2xf16>) -> vector<2x2xf16> {
-// CHECK: nvgpu.mma.sync(%{{.*}}, %{{.*}}, %{{.*}}) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
- %d = nvgpu.mma.sync(%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} :
+// CHECK: nvgpu.mma.sync(%{{.*}}, %{{.*}}, %{{.*}}) mmaShape = [16, 8, 16] : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ %d = nvgpu.mma.sync(%arg0, %arg1, %arg2) mmaShape = [16, 8, 16] :
(vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -24,10 +24,9 @@ func.func @mma_sp_sync_f16_16832(%arg0: vector<4x2xf16>,
%arg1: vector<4x2xf16>,
%arg2: vector<2x2xf16>,
%arg3: vector<2xi16>) -> vector<2x2xf16> {
- // CHECK: nvgpu.mma.sp.sync(%{{.*}}, %{{.*}}, %{{.*}}) metadata(%{{.+}}) {
- // CHECK-SAME: mmaShape = [16, 8, 32]
+ // CHECK: nvgpu.mma.sp.sync(%{{.*}}, %{{.*}}, %{{.*}}) metadata(%{{.+}}) mmaShape = [16, 8, 32]
// CHECK-SAME: (vector<4x2xf16>, vector<4x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 32]} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 32] :
(vector<4x2xf16>, vector<4x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -37,10 +36,9 @@ func.func @mma_sp_sync_f16_16816(%arg0: vector<2x2xf16>,
%arg1: vector<2x2xf16>,
%arg2: vector<2x2xf16>,
%arg3: vector<2xi16>) -> vector<2x2xf16> {
- // CHECK: nvgpu.mma.sp.sync(%{{.*}}, %{{.*}}, %{{.*}}) metadata(%{{.+}}) {
- // CHECK-SAME: mmaShape = [16, 8, 16]
+ // CHECK: nvgpu.mma.sp.sync(%{{.*}}, %{{.*}}, %{{.*}}) metadata(%{{.+}}) mmaShape = [16, 8, 16]
// CHECK-SAME: (vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 16]} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 16] :
(vector<2x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
return %d : vector<2x2xf16>
}
@@ -50,10 +48,9 @@ func.func @mma_sp_sync_i8_16864(%arg0: vector<4x4xi8>,
%arg1: vector<4x4xi8>,
%arg2: vector<2x2xi32>,
%arg3: vector<2xi16>) -> vector<2x2xi32> {
- // CHECK: nvgpu.mma.sp.sync(%{{.*}}, %{{.*}}, %{{.*}}) metadata(%{{.+}}) {
- // CHECK-SAME: mmaShape = [16, 8, 64]
+ // CHECK: nvgpu.mma.sp.sync(%{{.*}}, %{{.*}}, %{{.*}}) metadata(%{{.+}}) mmaShape = [16, 8, 64]
// CHECK-SAME: (vector<4x4xi8>, vector<4x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
- %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) {mmaShape = [16, 8, 64]} :
+ %d = nvgpu.mma.sp.sync(%arg0, %arg1, %arg2) metadata(%arg3) mmaShape = [16, 8, 64] :
(vector<4x4xi8>, vector<4x4xi8>, vector<2x2xi32>) -> vector<2x2xi32>
return %d : vector<2x2xi32>
}
@@ -65,8 +62,8 @@ func.func @async_cp(%dst : memref<2x7x5xf32, 3>, %src : memref<4x5xf32>){
%0 = nvgpu.device_async_copy %src[%c0, %c0], %dst[%c0, %c0, %c0], 4 : memref<4x5xf32> to memref<2x7x5xf32, 3>
// CHECK: %{{.*}} = nvgpu.device_async_create_group
%token = nvgpu.device_async_create_group %0
- // CHECK: nvgpu.device_async_wait %{{.*}} {numGroups = 1 : i32}
- nvgpu.device_async_wait %token {numGroups = 1 : i32}
+ // CHECK: nvgpu.device_async_wait %{{.*}} numGroups = 1
+ nvgpu.device_async_wait %token numGroups = 1
return
}
diff --git a/mlir/test/Dialect/NVGPU/transform-create-async-groups.mlir b/mlir/test/Dialect/NVGPU/transform-create-async-groups.mlir
index 8290001c45856..e7129c70442d5 100644
--- a/mlir/test/Dialect/NVGPU/transform-create-async-groups.mlir
+++ b/mlir/test/Dialect/NVGPU/transform-create-async-groups.mlir
@@ -9,7 +9,7 @@ builtin.module {
%c4 = arith.constant 4 : index
%cst_0 = arith.constant 0.000000e+00 : f32
// Make sure we emit the bypassL1.
- // CHECK: %[[CP0:.*]] = nvgpu.device_async_copy {{.*}}, {{.*}}, 4 {bypassL1} :
+ // CHECK: %[[CP0:.*]] = nvgpu.device_async_copy {{.*}}, {{.*}}, 4 bypassL1 :
%1 = vector.transfer_read %a[%c0, %c0], %cst_0 {in_bounds = [true]} : memref<1024x1024xf32>, vector<4xf32>
vector.transfer_write %1, %0[%c0, %c0, %c0] {in_bounds = [true]} : vector<4xf32>, memref<4x32x16xf32, #gpu.address_space<workgroup>>
// CHECK-NOT: nvgpu.device_async_create_group
@@ -145,7 +145,7 @@ builtin.module {
%0 = memref.alloc() : memref<4x32x16xf32, #gpu.address_space<workgroup>>
%c0 = arith.constant 0 : index
%cst_0 = arith.constant 0.000000e+00 : f32
- // CHECK: nvgpu.device_async_copy {{.*}}, {{.*}}, 4, %[[sz]] {bypassL1} :
+ // CHECK: nvgpu.device_async_copy {{.*}}, {{.*}}, 4, %[[sz]] bypassL1 :
%mask = vector.create_mask %sz : vector<4xi1>
%1 = vector.transfer_read %a[%c0, %c0], %cst_0, %mask {in_bounds = [true]} : memref<1024x1024xf32>, vector<4xf32>
vector.transfer_write %1, %0[%c0, %c0, %c0] {in_bounds = [true]} : vector<4xf32>, memref<4x32x16xf32, #gpu.address_space<workgroup>>
@@ -178,15 +178,15 @@ builtin.module {
// CHECK: %[[cmpi0:.*]] = arith.cmpi slt, %[[c0]], %[[sz0]]
// CHECK: %[[s0:.*]] = arith.select %[[cmpi0]], %[[sz1]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c0]]], {{.*}}, 4, %[[s0]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c0]]], {{.*}}, 4, %[[s0]] bypassL1
// CHECK: %[[cmpi1:.*]] = arith.cmpi slt, %[[c1]], %[[sz0]]
// CHECK: %[[s1:.*]] = arith.select %[[cmpi1]], %[[sz1]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c0]]], {{.*}}, 4, %[[s1]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c0]]], {{.*}}, 4, %[[s1]] bypassL1
// CHECK: %[[cmpi2:.*]] = arith.cmpi slt, %[[c2]], %[[sz0]]
// CHECK: %[[s2:.*]] = arith.select %[[cmpi2]], %[[sz1]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c2]], %[[c0]]], {{.*}}, 4, %[[s2]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c2]], %[[c0]]], {{.*}}, 4, %[[s2]] bypassL1
%mask = vector.create_mask %sz0, %sz1 : vector<3x4xi1>
%1 = vector.transfer_read %a[%c0, %c0], %cst_0, %mask {in_bounds = [true, true]} : memref<1024x1024xf32>, vector<3x4xf32>
vector.transfer_write %1, %0[%c0, %c0, %c0] {in_bounds = [true, true]} : vector<3x4xf32>, memref<4x32x16xf32, #gpu.address_space<workgroup>>
@@ -226,30 +226,30 @@ builtin.module {
// CHECK: %[[cmpi1:.*]] = arith.cmpi slt, %[[c0]], %[[sz1]]
// CHECK: %[[cond0:.*]] = arith.andi %[[cmpi1]], %[[cmpi0]]
// CHECK: %[[s0:.*]] = arith.select %[[cond0]], %[[sz2]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c0]], %[[c0]]], {{.*}}, 4, %[[s0]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c0]], %[[c0]]], {{.*}}, 4, %[[s0]] bypassL1
// CHECK: %[[cmpi2:.*]] = arith.cmpi slt, %[[c1]], %[[sz1]]
// CHECK: %[[cond1:.*]] = arith.andi %[[cmpi2]], %[[cmpi0]]
// CHECK: %[[s1:.*]] = arith.select %[[cond1]], %[[sz2]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c1]], %[[c0]]], {{.*}}, 4, %[[s1]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c1]], %[[c0]]], {{.*}}, 4, %[[s1]] bypassL1
// CHECK: %[[cmpi3:.*]] = arith.cmpi slt, %[[c2]], %[[sz1]]
// CHECK: %[[cond2:.*]] = arith.andi %[[cmpi3]], %[[cmpi0]]
// CHECK: %[[s2:.*]] = arith.select %[[cond2]], %[[sz2]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c2]], %[[c0]]], {{.*}}, 4, %[[s2]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c0]], %[[c2]], %[[c0]]], {{.*}}, 4, %[[s2]] bypassL1
// CHECK: %[[cmpi4:.*]] = arith.cmpi slt, %[[c1]], %[[sz0]]
// CHECK: %[[cond3:.*]] = arith.andi %[[cmpi1]], %[[cmpi4]]
// CHECK: %[[s3:.*]] = arith.select %[[cond3]], %[[sz2]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c0]], %[[c0]]], {{.*}}, 4, %[[s3]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c0]], %[[c0]]], {{.*}}, 4, %[[s3]] bypassL1
// CHECK: %[[cond4:.*]] = arith.andi %[[cmpi2]], %[[cmpi4]]
// CHECK: %[[s4:.*]] = arith.select %[[cond4]], %[[sz2]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c1]], %[[c0]]], {{.*}}, 4, %[[s4]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c1]], %[[c0]]], {{.*}}, 4, %[[s4]] bypassL1
// CHECK: %[[cond5:.*]] = arith.andi %[[cmpi3]], %[[cmpi4]]
// CHECK: %[[s5:.*]] = arith.select %[[cond5]], %[[sz2]], %[[c0]]
- // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c2]], %[[c0]]], {{.*}}, 4, %[[s5]] {bypassL1}
+ // CHECK: nvgpu.device_async_copy %[[a]][%[[c1]], %[[c2]], %[[c0]]], {{.*}}, 4, %[[s5]] bypassL1
%mask = vector.create_mask %sz0, %sz1, %sz2 : vector<2x3x4xi1>
%1 = vector.transfer_read %a[%c0, %c0, %c0], %cst_0, %mask {in_bounds = [true, true, true]} : memref<1024x1024x1024xf32>, vector<2x3x4xf32>
vector.transfer_write %1, %0[%c0, %c0, %c0] {in_bounds = [true, true, true]} : vector<2x3x4xf32>, memref<4x32x16xf32, #gpu.address_space<workgroup>>
diff --git a/mlir/test/Dialect/NVGPU/transform-matmul-to-nvvm.mlir b/mlir/test/Dialect/NVGPU/transform-matmul-to-nvvm.mlir
index c6aa2039a9511..d9b6a06c02c41 100644
--- a/mlir/test/Dialect/NVGPU/transform-matmul-to-nvvm.mlir
+++ b/mlir/test/Dialect/NVGPU/transform-matmul-to-nvvm.mlir
@@ -48,7 +48,7 @@ func.func @matmul_16x8x4xf32_global(
// CHECK: %[[VAL_33:.*]] = vector.insert %[[VAL_26]], %[[VAL_32]] [1, 0] : f32 into vector<2x2xf32>
// CHECK: %[[RES:.*]] = vector.insert %[[VAL_29]], %[[VAL_33]] [1, 1] : f32 into vector<2x2xf32>
//
-// CHECK: %[[VAL_35:.*]] = nvgpu.mma.sync(%[[LHS]], %[[RHS]], %[[RES]]) {mmaShape = [16, 8, 4], tf32Enabled} : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
+// CHECK: %[[VAL_35:.*]] = nvgpu.mma.sync(%[[LHS]], %[[RHS]], %[[RES]]) mmaShape = [16, 8, 4] tf32Enabled : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
//
// CHECK: %[[VAL_36:.*]] = vector.extract %[[VAL_35]][0, 0] : f32 from vector<2x2xf32>
// CHECK: %[[VAL_37:.*]] = vector.extract %[[VAL_35]][0, 1] : f32 from vector<2x2xf32>
@@ -96,7 +96,7 @@ func.func @matmul_16x8x16xf16_global(
// CHECK-COUNT-4: memref.load {{.*}} : memref<16x8xf16>
// CHECK-COUNT-4: vector.insert {{.*}} : f16 into vector<2x2xf16>
//
- // CHECK: nvgpu.mma.sync(%{{.*}}) {mmaShape = [16, 8, 16]}
+ // CHECK: nvgpu.mma.sync(%{{.*}}) mmaShape = [16, 8, 16]
// CHECK-SAME: : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
//
// CHECK-COUNT-4: vector.extract %{{.*}} : f16 from vector<2x2xf16>
diff --git a/mlir/test/Dialect/NVGPU/transform-pipeline-shared.mlir b/mlir/test/Dialect/NVGPU/transform-pipeline-shared.mlir
index e959949babd9e..75bcd349d9b82 100644
--- a/mlir/test/Dialect/NVGPU/transform-pipeline-shared.mlir
+++ b/mlir/test/Dialect/NVGPU/transform-pipeline-shared.mlir
@@ -100,7 +100,7 @@ func.func @async_depth_2_predicated(%global: memref<?xf32>, %alloc_size: index)
scf.for %i = %c0 to %c98 step %c4 {
// Condition for the predication "select" below.
// CHECK: %[[CMP0:.+]] = arith.cmpi slt, %[[I]], %[[C90]]
- // CHECK: nvgpu.device_async_wait %[[ITER_ARG0]] {numGroups = 1
+ // CHECK: nvgpu.device_async_wait %[[ITER_ARG0]] numGroups = 1
// Original "select" with updated induction variable.
// CHECK: %[[I_PLUS_8:.+]] = arith.addi %[[I]], %[[C8]]
// CHECK: %[[CMP1:.+]] = arith.cmpi slt, %[[I_PLUS_8]], %[[C96]]
@@ -156,12 +156,12 @@ func.func @async_depth_2_peeled(%global: memref<?xf32>) {
// CHECK: nvgpu.device_async_copy
// CHECK: nvgpu.device_async_copy
// CHECK: scf.for
- // CHECK: nvgpu.device_async_wait %{{.*}} {numGroups = 1
+ // CHECK: nvgpu.device_async_wait %{{.*}} numGroups = 1
// CHECK: arith.select
// CHECK: nvgpu.device_async_copy
// CHECK: scf.yield
- // CHECK: nvgpu.device_async_wait %{{.*}} {numGroups = 1
- // CHECK: nvgpu.device_async_wait %{{.*}} {numGroups = 0
+ // CHECK: nvgpu.device_async_wait %{{.*}} numGroups = 1
+ // CHECK: nvgpu.device_async_wait %{{.*}} numGroups = 0
scf.for %i = %c0 to %c98 step %c4 {
%c96 = arith.constant 96 : index
%cond = arith.cmpi slt, %i, %c96 : index
diff --git a/mlir/test/Examples/NVGPU/Ch3.py b/mlir/test/Examples/NVGPU/Ch3.py
index 815e16dc68553..8289069812612 100644
--- a/mlir/test/Examples/NVGPU/Ch3.py
+++ b/mlir/test/Examples/NVGPU/Ch3.py
@@ -189,7 +189,7 @@ def gemm_tma_kernel():
# DUMPIR: %[[WG_ACC:.*]] = nvgpu.warpgroup.mma.init.accumulator -> <fragmented = vector<128x128xf32>>
# DUMPIR: %[[GEN0:.*]] = nvgpu.warpgroup.generate.descriptor %[[VIEW]], %[[TMA0]] : memref<128x64xf16, #gpu.address_space<workgroup>>, <tensor = memref<128x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none> -> <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>
# DUMPIR: %[[GEN1:.*]] = nvgpu.warpgroup.generate.descriptor %[[VIEW_15]], %[[TMA1]] : memref<64x128xf16, #gpu.address_space<workgroup>>, <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none> -> <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>
-# DUMPIR: %[[MMA:.*]] = nvgpu.warpgroup.mma %[[GEN0]], %[[GEN1]], %[[WG_ACC]] {transposeB} : <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>, <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
+# DUMPIR: %[[MMA:.*]] = nvgpu.warpgroup.mma %[[GEN0]], %[[GEN1]], %[[WG_ACC]] transposeB : <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>, <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
# DUMPIR: nvgpu.warpgroup.mma.store %[[MMA]], %{{.*}} : <fragmented = vector<128x128xf32>> to memref<128x128xf32>
# DUMPIR: gpu.terminator
# DUMPIR: }
diff --git a/mlir/test/Examples/NVGPU/Ch4.py b/mlir/test/Examples/NVGPU/Ch4.py
index c66259d141336..d96abb7a0d36e 100644
--- a/mlir/test/Examples/NVGPU/Ch4.py
+++ b/mlir/test/Examples/NVGPU/Ch4.py
@@ -397,7 +397,7 @@ def gemm_multistage_kernel():
# DUMPIR: %[[VIEW_B_L:.*]] = memref.view %[[SMEM_B_L]][%[[OFF_B_L]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<64x128xf16, #gpu.address_space<workgroup>>
# DUMPIR: %[[DESC_A_L:.*]] = nvgpu.warpgroup.generate.descriptor %[[VIEW_A_L]], %{{.*}} : memref<128x64xf16, #gpu.address_space<workgroup>>, <tensor = memref<128x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none> -> <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>
# DUMPIR: %[[DESC_B_L:.*]] = nvgpu.warpgroup.generate.descriptor %[[VIEW_B_L]], %{{.*}} : memref<64x128xf16, #gpu.address_space<workgroup>>, <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none> -> <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>
-# DUMPIR: %[[ACC_L:.*]] = nvgpu.warpgroup.mma %[[DESC_A_L]], %[[DESC_B_L]], %arg16 {transposeB} : <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>, <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
+# DUMPIR: %[[ACC_L:.*]] = nvgpu.warpgroup.mma %[[DESC_A_L]], %[[DESC_B_L]], %arg16 transposeB : <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>, <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
# DUMPIR: %[[C6_NEXT:.*]] = arith.constant 6 : index
# DUMPIR: %[[ITER_NEXT:.*]] = arith.addi %arg15, %[[C6_NEXT]] : index
# DUMPIR: %[[C16_CMP:.*]] = arith.constant 16 : index
diff --git a/mlir/test/Examples/NVGPU/Ch5.py b/mlir/test/Examples/NVGPU/Ch5.py
index 4f06f97142620..f6af8cae90b09 100644
--- a/mlir/test/Examples/NVGPU/Ch5.py
+++ b/mlir/test/Examples/NVGPU/Ch5.py
@@ -435,7 +435,7 @@ def gemm_warp_specialized_kernel():
# DUMPIR: %[[VIEW_B_CONS:.*]] = memref.view %[[SMEM_CONS_2]][%[[OFF_B_CONS]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<64x128xf16, #gpu.address_space<workgroup>>
# DUMPIR: %[[DESC_A:.*]] = nvgpu.warpgroup.generate.descriptor %[[VIEW_A_CONS]], %{{.*}} : memref<128x64xf16, #gpu.address_space<workgroup>>, <tensor = memref<128x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none> -> <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>
# DUMPIR: %[[DESC_B:.*]] = nvgpu.warpgroup.generate.descriptor %[[VIEW_B_CONS]], %{{.*}} : memref<64x128xf16, #gpu.address_space<workgroup>>, <tensor = memref<64x64xf16, 3>, swizzle = swizzle_128b, l2promo = none, oob = zero, interleave = none> -> <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>
-# DUMPIR: %[[ACC:.*]] = nvgpu.warpgroup.mma %[[DESC_A]], %[[DESC_B]], %arg16 {transposeB} : <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>, <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
+# DUMPIR: %[[ACC:.*]] = nvgpu.warpgroup.mma %[[DESC_A]], %[[DESC_B]], %arg16 transposeB : <tensor = memref<128x64xf16, #gpu.address_space<workgroup>>>, <tensor = memref<64x128xf16, #gpu.address_space<workgroup>>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
# DUMPIR: %[[C0_CMP:.*]] = arith.constant 0 : index
# DUMPIR: %[[IS_NOT_FIRST:.*]] = arith.cmpi ugt, %arg15, %[[C0_CMP]] : index
# DUMPIR: %[[ARRIVE_PRED:.*]] = arith.andi %[[IS_NOT_FIRST]], %[[IS_PRIMARY_2]] : i1
diff --git a/mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sparse-mma-2-4-f16.mlir b/mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sparse-mma-2-4-f16.mlir
index d8d7c1c39db91..ae06a666e0dcd 100644
--- a/mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sparse-mma-2-4-f16.mlir
+++ b/mlir/test/Integration/Dialect/SparseTensor/GPU/CUDA/sparse-mma-2-4-f16.mlir
@@ -193,8 +193,7 @@ module attributes {gpu.container_module} {
// Sparsity selector. For 16x8x32, the default "0" means threads T0/T1
// within each group of four threads contribute metadata.
%d = nvgpu.mma.sp.sync(%A_data, %B_data, %accum)
- metadata(%meta)
- {mmaShape = [16, 8, 32]} : (vector<4x2xf16>, vector<4x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
+ metadata(%meta) mmaShape = [16, 8, 32] : (vector<4x2xf16>, vector<4x2xf16>, vector<2x2xf16>) -> vector<2x2xf16>
//===----------------------------------------------------------------------===//
// Write back results to gpu global memory
diff --git a/mlir/test/Integration/GPU/CUDA/TensorCore/sm80/transform-mma-sync-matmul-f32.mlir b/mlir/test/Integration/GPU/CUDA/TensorCore/sm80/transform-mma-sync-matmul-f32.mlir
index 2eef2ff8f3564..10a84760e8dba 100644
--- a/mlir/test/Integration/GPU/CUDA/TensorCore/sm80/transform-mma-sync-matmul-f32.mlir
+++ b/mlir/test/Integration/GPU/CUDA/TensorCore/sm80/transform-mma-sync-matmul-f32.mlir
@@ -3,7 +3,7 @@
// RUN: | FileCheck %s --check-prefix=CHECK-MMA-SYNC
// CHECK-MMA-SYNC-LABEL: func @main() {
-// CHECK-MMA-SYNC: nvgpu.mma.sync(%{{.*}}) {mmaShape = [16, 8, 4], tf32Enabled}
+// CHECK-MMA-SYNC: nvgpu.mma.sync(%{{.*}}) mmaShape = [16, 8, 4] tf32Enabled
// CHECK-MMA-SYNC-SAME: : (vector<2x1xf32>, vector<1x1xf32>, vector<2x2xf32>) -> vector<2x2xf32>
// Tested to run locally in 1.7s.
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir b/mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir
index 22474cbcd39f3..069cdf0d6194e 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/gemm_f32_f16_f16_128x128x128.mlir
@@ -206,7 +206,7 @@ func.func @main() {
%dA = nvgpu.warpgroup.generate.descriptor %lhsSlice, %descA : memref<128x64xf16, strided<[64, 1], offset: ?>, #gpu.address_space<workgroup>>, !lhsTensorMap -> !nvgpu.warpgroup.descriptor<tensor=memref<128x64xf16, 3>>
%dB = nvgpu.warpgroup.generate.descriptor %rhsSlice, %descB : memref<64x128xf16, strided<[128, 1], offset: ?>, #gpu.address_space<workgroup>>, !rhsTensorMap -> !nvgpu.warpgroup.descriptor<tensor=memref<64x128xf16, 3>>
// Perform WGMMA 128x128x64
- %md = nvgpu.warpgroup.mma %dA, %dB, %mc {transposeB} : <tensor = memref<128x64xf16,3>>, <tensor = memref<64x128xf16,3>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
+ %md = nvgpu.warpgroup.mma %dA, %dB, %mc transposeB : <tensor = memref<128x64xf16,3>>, <tensor = memref<64x128xf16,3>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
scf.yield %md : !nvgpu.warpgroup.accumulator<fragmented = vector<128x128xf32>>
}
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir b/mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir
index 39bad38f36468..7fe041965fb37 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/gemm_pred_f32_f16_f16_128x128x128.mlir
@@ -214,7 +214,7 @@ func.func @main() {
%dA = nvgpu.warpgroup.generate.descriptor %lhsSlice, %descA : memref<128x64xf16, strided<[64, 1], offset: ?>, #gpu.address_space<workgroup>>, !lhsTensorMap -> !nvgpu.warpgroup.descriptor<tensor=memref<128x64xf16, 3>>
%dB = nvgpu.warpgroup.generate.descriptor %rhsSlice, %descB : memref<64x128xf16, strided<[128, 1], offset: ?>, #gpu.address_space<workgroup>>, !rhsTensorMap -> !nvgpu.warpgroup.descriptor<tensor=memref<64x128xf16, 3>>
// Perform WGMMA 128x128x64
- %md = nvgpu.warpgroup.mma %dA, %dB, %mc {transposeB} : <tensor = memref<128x64xf16,3>>, <tensor = memref<64x128xf16,3>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
+ %md = nvgpu.warpgroup.mma %dA, %dB, %mc transposeB : <tensor = memref<128x64xf16,3>>, <tensor = memref<64x128xf16,3>>, <fragmented = vector<128x128xf32>> -> <fragmented = vector<128x128xf32>>
scf.yield %md : !nvgpu.warpgroup.accumulator<fragmented = vector<128x128xf32>>
}
diff --git a/mlir/test/python/dialects/nvgpu.py b/mlir/test/python/dialects/nvgpu.py
index 6df32bdd3c273..fbf8f84b1bc9a 100644
--- a/mlir/test/python/dialects/nvgpu.py
+++ b/mlir/test/python/dialects/nvgpu.py
@@ -39,5 +39,5 @@ def testSmoke():
mem_t = MemRefType.get((10, 10), F32Type.get(), memory_space=Attribute.parse("3"))
vec_t = VectorType.get((4, 1), F32Type.get())
mem = memref.AllocOp(mem_t, [], [])
- # CHECK: %0 = nvgpu.ldmatrix %alloc[%c42, %c42] {numTiles = 4 : i32, transpose = false} : memref<10x10xf32, 3> -> vector<4x1xf32>
+ # CHECK: %0 = nvgpu.ldmatrix %alloc[%c42, %c42] numTiles = 4 transpose = false : memref<10x10xf32, 3> -> vector<4x1xf32>
nvgpu.LdMatrixOp(vec_t, mem, [cst, cst], False, 4)
More information about the Mlir-commits
mailing list