1// RUN: mlir-opt --convert-nvvm-to-llvm --split-input-file -verify-diagnostics %s 2 3!mat64f32 = !llvm.struct<(f32, f32, f32, f32, f32, f32, f32)> 4func.func @wgmma_f32_f16_f16(%descA : i64, %descB : i64) -> !mat64f32{ 5 %result = llvm.mlir.undef : !mat64f32 6 // expected-error @+1 {{'nvvm.wgmma.mma_async' op results 64, however output struct has 7 elements}} 7 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 8 #nvvm.shape<m = 64, n = 128, k = 16>, 9 D [<f32>, <zero>], 10 A [<f16>, #nvvm.wgmma_scale_in<neg>, <col>], 11 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 12 : !mat64f32 -> !mat64f32 13 return %res : !mat64f32 14} 15 16// ----- 17 18func.func @wgmma_f32_satfinite(%descA : i64, %descB : i64) { 19 %result = llvm.mlir.undef : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 20 // expected-error @+1 {{`satfinite` can be only used with s32 accumulator, however the current accumulator is f32}} 21 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 22 #nvvm.shape<m = 64, n = 16, k = 16>, 23 D [<f32>, <zero>, <satfinite>], 24 A [<f16>, #nvvm.wgmma_scale_in<neg>, <col>], 25 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 26 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 27 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 28 return 29} 30 31// ----- 32 33func.func @wgmma_f32_m32(%descA : i64, %descB : i64) { 34 %result = llvm.mlir.undef : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 35 // expected-error @+1 {{shape 'm' must be 64}} 36 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 37 #nvvm.shape<m = 32, n = 16, k = 16>, 38 D [<f32>, <zero>], 39 A [<f16>, #nvvm.wgmma_scale_in<neg>, <col>], 40 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 41 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 42 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 43 return 44} 45 46// ----- 47 48func.func @wgmma_f32_m32(%descA : i64, %descB : i64) { 49 %result = llvm.mlir.undef : !llvm.struct<(f32, f32, i32, f32, f32, f32, f32, f32)> 50 // expected-error @+1 {{op all elements in struct must be same type but there is 'i32'}} 51 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 52 #nvvm.shape<m = 64, n = 16, k = 16>, 53 D [<f32>, <zero>], 54 A [<f16>, #nvvm.wgmma_scale_in<neg>, <col>], 55 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 56 : !llvm.struct<(f32, f32, i32, f32, f32, f32, f32, f32)> 57 -> !llvm.struct<(f32, f32, i32, f32, f32, f32, f32, f32)> 58 return 59} 60 61// ----- 62 63func.func @wgmma_f32_m32(%descA : i64, %descB : i64) { 64 %result = llvm.mlir.undef : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 65 // expected-error @+1 {{op shape 'k' must be 16 for input type f16}} 66 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 67 #nvvm.shape<m = 64, n = 16, k = 3>, 68 D [<f32>, <zero>], 69 A [<f16>, #nvvm.wgmma_scale_in<neg>, <col>], 70 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 71 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 72 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 73 return 74} 75 76// ----- 77 78func.func @wgmma_transpose(%descA : i64, %descB : i64) { 79 %result = llvm.mlir.undef : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 80 // expected-error @+1 {{op given layouts layout_a = col and layout_b = col for input types tf32 and tf32 requires transpose. However, this is only supported for: f16 and bf16}} 81 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 82 #nvvm.shape<m = 64, n = 16, k = 8>, 83 D [<f32>, <zero>], 84 A [<tf32>, #nvvm.wgmma_scale_in<neg>, <col>], 85 B [<tf32>, #nvvm.wgmma_scale_in<neg>, <col>] 86 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 87 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 88 return 89} 90 91// ----- 92 93func.func @wgmma_transpose(%descA : i64, %descB : i64) { 94 %result = llvm.mlir.undef : !llvm.struct<(f16, f16, f16, f16)> 95 // expected-error @+1 {{'nvvm.wgmma.mma_async' op f16 += tf32 * tf32, it is not supported.}} 96 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 97 #nvvm.shape<m = 64, n = 16, k = 8>, 98 D [<f16>, <zero>], 99 A [<tf32>, #nvvm.wgmma_scale_in<neg>, <col>], 100 B [<tf32>, #nvvm.wgmma_scale_in<neg>, <col>] 101 :!llvm.struct<(f16, f16, f16, f16)> 102 -> !llvm.struct<(f16, f16, f16, f16)> 103 return 104} 105 106// ----- 107 108func.func @wgmma_f32_m32(%descA : i64, %descB : i64) { 109 %result = llvm.mlir.undef : !llvm.struct<(i32, i32, i32, i32)> 110 // expected-error @+1 {{input struct and result struct must be the same type}} 111 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 112 #nvvm.shape<m = 64, n = 8, k = 16>, 113 D [<f16>, <zero>], 114 A [<f16>, #nvvm.wgmma_scale_in<neg>, <col>], 115 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 116 : !llvm.struct<(i32, i32, i32, i32)> 117 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 118 return 119} 120 121// ----- 122 123func.func @wgmma_f32_m32(%descA : i64, %descB : i64) { 124 %result = llvm.mlir.undef : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 125 // expected-error @+1 {{op f32 += bf16 * f16, it is not supported}} 126 %res = nvvm.wgmma.mma_async %descA, %descB, %result, 127 #nvvm.shape<m = 64, n = 8, k = 16>, 128 D [<f32>, <zero>], 129 A [<bf16>, #nvvm.wgmma_scale_in<neg>, <col>], 130 B [<f16>, #nvvm.wgmma_scale_in<neg>, <col>] 131 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 132 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 133 return 134} 135// ----- 136 137func.func @set_max_register() { 138 // expected-error @+1 {{new register size must be in between 24 to 256}} 139 nvvm.setmaxregister decrease 8 140 func.return 141} 142 143// ----- 144 145func.func @set_max_register() { 146 // expected-error @+1 {{new register size must be multiple of 8}} 147 nvvm.setmaxregister decrease 51 148 func.return 149} 150 151// ----- 152 153func.func @fence_proxy() { 154 // expected-error @+1 {{op only async_shared fence can have space attribute}} 155 nvvm.fence.proxy { kind = #nvvm.proxy_kind<async>, space = #nvvm.shared_space<cluster>} 156 func.return 157} 158 159// ----- 160 161func.func @fence_proxy() { 162 // expected-error @+1 {{op async_shared fence requires space attribute}} 163 nvvm.fence.proxy { kind = #nvvm.proxy_kind<async.shared>} 164 func.return 165} 166