1# RUN: %PYTHON %s | FileCheck %s 2# This is just a smoke test that the dialect is functional. 3 4from mlir.ir import * 5from mlir.dialects import nvvm 6from mlir.dialects import llvm 7from mlir.dialects import func 8 9 10def constructAndPrintInModule(f): 11 print("\nTEST:", f.__name__) 12 with Context(), Location.unknown(): 13 module = Module.create() 14 with InsertionPoint(module.body): 15 f() 16 print(module) 17 return f 18 19 20# CHECK-LABEL: testSmoke 21@constructAndPrintInModule 22def testSmoke(): 23 i64 = IntegerType.get_signless(64) 24 mat64f32_t = Type.parse( 25 "!llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>" 26 ) 27 shape_attr = Attribute.parse("#nvvm.shape<m = 64, n = 32, k = 16>") 28 # CHECK-LABEL: func @wgmma_f32_f16_f16(%arg0: i64, %arg1: i64) 29 @func.FuncOp.from_py_func(i64, i64) 30 def wgmma_f32_f16_f16(desc_a, desc_b): 31 # CHECK: nvvm.cp.async.wait.group 5 32 nvvm.CpAsyncWaitGroupOp(5) 33 # CHECK: %0 = llvm.mlir.undef : [[MAT_T:.*]] 34 result = llvm.UndefOp(mat64f32_t) 35 # CHECK: %1 = nvvm.wgmma.mma_async %arg0, %arg1, %0, <m = 64, n = 32, k = 16>, D[<f32>, <zero>], A[<f16>, <neg>, <col>], B[<f16>, <neg>, <col>] : [[MAT_T]] -> [[MAT_T]] 36 result1 = nvvm.WgmmaMmaAsyncOp( 37 results_=mat64f32_t, 38 inouts=result, 39 descriptorA=desc_a, 40 descriptorB=desc_b, 41 shape=shape_attr, 42 typeA=nvvm.WGMMATypes.f16, 43 typeB=nvvm.WGMMATypes.f16, 44 typeD=nvvm.WGMMATypes.f32, 45 scaleD=nvvm.WGMMAScaleOut.zero, 46 scaleA=nvvm.WGMMAScaleIn.neg, 47 scaleB=nvvm.WGMMAScaleIn.neg, 48 layoutA=nvvm.MMALayout.col, 49 layoutB=nvvm.MMALayout.col, 50 ) 51