166 lines · plain
1// RUN: mlir-opt --convert-nvvm-to-llvm --split-input-file -verify-diagnostics %s2 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 : !mat64f326 // 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 -> !mat64f3213 return %res : !mat64f3214}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 8140 func.return141}142 143// -----144 145func.func @set_max_register() {146 // expected-error @+1 {{new register size must be multiple of 8}}147 nvvm.setmaxregister decrease 51148 func.return149}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.return157}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.return165}166