39 lines · plain
1// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s2 3// CHECK-LABEL: define i64 @tcgen05_mma_smem_desc_test(i32 %0, i32 %1, i32 %2, i8 %3, i1 %4, i8 %5) {4llvm.func @tcgen05_mma_smem_desc_test(%startAddr: i32, %leadingDimOffset: i32, %strideDimOffset: i32,5 %baseOffset: i8, %leadingDimMode: i1, %swizzleMode: i8) -> i64 {6 // CHECK-NEXT: %7 = and i32 %0, 163837 // CHECK-NEXT: %8 = zext i32 %7 to i648 // CHECK-NEXT: %9 = shl i64 %8, 09 // CHECK-NEXT: %10 = or i64 0, %910 // CHECK-NEXT: %11 = and i32 %1, 1638311 // CHECK-NEXT: %12 = zext i32 %11 to i6412 // CHECK-NEXT: %13 = shl i64 %12, 1613 // CHECK-NEXT: %14 = or i64 %10, %1314 // CHECK-NEXT: %15 = and i32 %2, 1638315 // CHECK-NEXT: %16 = zext i32 %15 to i6416 // CHECK-NEXT: %17 = shl i64 %16, 3217 // CHECK-NEXT: %18 = or i64 %14, %1718 // CHECK-NEXT: %19 = or i64 %18, 7036874417766419 // CHECK-NEXT: %20 = zext i8 %3 to i3220 // CHECK-NEXT: %21 = and i32 %20, 721 // CHECK-NEXT: %22 = zext i32 %21 to i6422 // CHECK-NEXT: %23 = shl i64 %22, 4923 // CHECK-NEXT: %24 = or i64 %19, %2324 // CHECK-NEXT: %25 = zext i1 %4 to i3225 // CHECK-NEXT: %26 = and i32 %25, 126 // CHECK-NEXT: %27 = zext i32 %26 to i6427 // CHECK-NEXT: %28 = shl i64 %27, 5228 // CHECK-NEXT: %29 = or i64 %24, %2829 // CHECK-NEXT: %30 = zext i8 %5 to i3230 // CHECK-NEXT: %31 = and i32 %30, 731 // CHECK-NEXT: %32 = zext i32 %31 to i6432 // CHECK-NEXT: %33 = shl i64 %32, 6133 // CHECK-NEXT: %34 = or i64 %29, %3334 // CHECK-NEXT: ret i64 %3435 // CHECK-NEXT: }36 %desc = nvvm.tcgen05.mma_smem_desc (%startAddr, %leadingDimOffset, %strideDimOffset, %baseOffset, %leadingDimMode, %swizzleMode) : (i32, i32, i32, i8, i1, i8) -> i6437 llvm.return %desc : i6438}39