1a7d80c50Smax# RUN: %PYTHON %s | FileCheck %s 2a7d80c50Smax# This is just a smoke test that the dialect is functional. 3a7d80c50Smax 4a7d80c50Smaxfrom mlir.ir import * 5a7d80c50Smaxfrom mlir.dialects import nvvm 692233062Smaxfrom mlir.dialects import llvm 792233062Smaxfrom mlir.dialects import func 8a7d80c50Smax 9a7d80c50Smax 10a7d80c50Smaxdef constructAndPrintInModule(f): 11a7d80c50Smax print("\nTEST:", f.__name__) 12a7d80c50Smax with Context(), Location.unknown(): 13a7d80c50Smax module = Module.create() 14a7d80c50Smax with InsertionPoint(module.body): 15a7d80c50Smax f() 16a7d80c50Smax print(module) 17a7d80c50Smax return f 18a7d80c50Smax 19a7d80c50Smax 20a7d80c50Smax# CHECK-LABEL: testSmoke 21a7d80c50Smax@constructAndPrintInModule 22a7d80c50Smaxdef testSmoke(): 2392233062Smax i64 = IntegerType.get_signless(64) 2492233062Smax mat64f32_t = Type.parse( 2592233062Smax "!llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>" 2692233062Smax ) 2792233062Smax shape_attr = Attribute.parse("#nvvm.shape<m = 64, n = 32, k = 16>") 2892233062Smax # CHECK-LABEL: func @wgmma_f32_f16_f16(%arg0: i64, %arg1: i64) 2992233062Smax @func.FuncOp.from_py_func(i64, i64) 3092233062Smax def wgmma_f32_f16_f16(desc_a, desc_b): 31a7d80c50Smax # CHECK: nvvm.cp.async.wait.group 5 32a7d80c50Smax nvvm.CpAsyncWaitGroupOp(5) 3392233062Smax # CHECK: %0 = llvm.mlir.undef : [[MAT_T:.*]] 3492233062Smax result = llvm.UndefOp(mat64f32_t) 35*12c241b3SGuray Ozen # 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]] 3692233062Smax result1 = nvvm.WgmmaMmaAsyncOp( 3792233062Smax results_=mat64f32_t, 3892233062Smax inouts=result, 3992233062Smax descriptorA=desc_a, 4092233062Smax descriptorB=desc_b, 4192233062Smax shape=shape_attr, 4292233062Smax typeA=nvvm.WGMMATypes.f16, 4392233062Smax typeB=nvvm.WGMMATypes.f16, 44*12c241b3SGuray Ozen typeD=nvvm.WGMMATypes.f32, 4592233062Smax scaleD=nvvm.WGMMAScaleOut.zero, 4692233062Smax scaleA=nvvm.WGMMAScaleIn.neg, 4792233062Smax scaleB=nvvm.WGMMAScaleIn.neg, 4892233062Smax layoutA=nvvm.MMALayout.col, 4992233062Smax layoutB=nvvm.MMALayout.col, 5092233062Smax ) 51