Skip to content

[mlir][test] Fix filecheck annotation typos #92897

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
May 24, 2024
Merged

[mlir][test] Fix filecheck annotation typos #92897

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
f8280ce
moved filecheck fixes for mlir
May 21, 2024
2d5d9d5
few additional fixes
May 21, 2024
4291c34
Merge branch 'main' into filecheck-mlir
May 23, 2024
1d2b98b
review and test fix
May 23, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions mlir/test/Analysis/DataFlow/test-next-access.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,7 @@ func.func @branch(%arg0: memref<f32>, %arg1: f32, %arg2: i1) -> f32 {
return %phi : f32
}

// CHECK-LABEL @dead_branch
// CHECK-LABEL: @dead_branch
func.func @dead_branch(%arg0: memref<f32>, %arg1: f32) -> f32 {
// CHECK: name = "store"
// CHECK-SAME: next_access = ["unknown", ["load 2"]]
Expand Down Expand Up @@ -191,7 +191,7 @@ func.func @loop_cf(%arg0: memref<?xf32>, %arg1: f32, %arg2: index, %arg3: index,
return %8 : f32
}

// CHECK-LABEL @conditional_cf
// CHECK-LABEL: @conditional_cf
func.func @conditional_cf(%arg0: i1, %arg1: memref<f32>) {
// CHECK: name = "pre"
// CHECK-SAME: next_access = {{\[}}["then", "post"]]
Expand Down
8 changes: 4 additions & 4 deletions mlir/test/Conversion/BufferizationToMemRef/bufferization-to-memref.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,7 @@ func.func @conversion_dealloc_simple(%arg0: memref<2xf32>, %arg1: i1) {
return
}

// CHECk: scf.if [[ARG1]] {
// CHECk-NEXT: memref.dealloc [[ARG0]] : memref<2xf32>
// CHECk-NEXT: }
// CHECk-NEXT: return
// CHECK: scf.if [[ARG1]] {
// CHECK-NEXT: memref.dealloc [[ARG0]] : memref<2xf32>
// CHECK-NEXT: }
// CHECK-NEXT: return
4 changes: 2 additions & 2 deletions mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -778,11 +778,11 @@ func.func @create_tensor_map(%devicePtr2d : memref<64x128xf32>, %devicePtr1d : m
%crd0 = arith.constant 64 : index
%crd1 = arith.constant 128 : index
%devicePtr2d_unranked = memref.cast %devicePtr2d : memref<64x128xf32> to memref<*xf32>
// CHECK : llvm.call @mgpuTensorMapEncodeTiledMemref
// CHECK: llvm.call @mgpuTensorMapEncodeTiledMemref
%tensorMap2d = nvgpu.tma.create.descriptor %devicePtr2d_unranked box[%crd0, %crd1] : memref<*xf32> -> !tensorMap2d

%devicePtr1d_unranked = memref.cast %devicePtr1d : memref<128xf32> to memref<*xf32>
// CHECK : llvm.call @mgpuTensorMapEncodeTiledMemref
// CHECK: llvm.call @mgpuTensorMapEncodeTiledMemref
%tensorMap1d = nvgpu.tma.create.descriptor %devicePtr1d_unranked box[%crd1] : memref<*xf32> -> !tensorMap1d
func.return
}
Expand Down
12 changes: 6 additions & 6 deletions mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@ llvm.func @init_mbarrier(%barrier_gen : !llvm.ptr, %barrier : !llvm.ptr<3>, %cou
llvm.func @init_mbarrier_arrive_expect_tx(%barrier : !llvm.ptr<3>, %txcount : i32, %pred : i1) {
//CHECK: llvm.inline_asm has_side_effects asm_dialect = att "mbarrier.arrive.expect_tx.shared.b64 _, [$0], $1;", "r,r"
nvvm.mbarrier.arrive.expect_tx.shared %barrier, %txcount : !llvm.ptr<3>, i32
//CHECK : llvm.inline_asm has_side_effects asm_dialect = att "@$2 mbarrier.arrive.expect_tx.shared.b64 _, [$0], $1;", "r,r,b "
//CHECK: llvm.inline_asm has_side_effects asm_dialect = att "@$2 mbarrier.arrive.expect_tx.shared.b64 _, [$0], $1;", "r,r,b "
nvvm.mbarrier.arrive.expect_tx.shared %barrier, %txcount, predicate = %pred : !llvm.ptr<3>, i32, i1
llvm.return
}
Expand Down Expand Up @@ -129,7 +129,7 @@ func.func @tma_load_5d_all(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %bar
func.func @tma_load_1d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier: !llvm.ptr<3>, %crd0: i32, %p : i1) {
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.bulk.tensor.1d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2} ], [$3];", "r,l,r,r"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0] : !llvm.ptr<3>, !llvm.ptr
// CHECK : llvm.inline_asm has_side_effects asm_dialect = att "@$4 cp.async.bulk.tensor.1d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2} ], [$3];", "l,r,r,r,b"
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "@$4 cp.async.bulk.tensor.1d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2} ], [$3];", "l,r,r,r,b"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0] predicate=%p : !llvm.ptr<3>, !llvm.ptr
return
}
Expand All @@ -138,7 +138,7 @@ func.func @tma_load_1d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier
func.func @tma_load_2d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier: !llvm.ptr<3>, %crd0: i32, %crd1: i32, %p : i1) {
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3} ], [$4];", "r,l,r,r,r"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1] : !llvm.ptr<3>, !llvm.ptr
// CHECK : llvm.inline_asm has_side_effects asm_dialect = att "@$5 cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3} ], [$4];", "l,r,r,r,r,b"
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "@$5 cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3} ], [$4];", "l,r,r,r,r,b"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1] predicate=%p : !llvm.ptr<3>, !llvm.ptr
return
}
Expand All @@ -147,7 +147,7 @@ func.func @tma_load_2d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier
func.func @tma_load_3d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier: !llvm.ptr<3>, %crd0: i32, %crd1: i32, %crd2: i32, %p : i1) {
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.bulk.tensor.3d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4} ], [$5];", "r,l,r,r,r,r"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1,%crd2] : !llvm.ptr<3>, !llvm.ptr
// CHECK : llvm.inline_asm has_side_effects asm_dialect = att "@$6 cp.async.bulk.tensor.3d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4}], [$5];", "l,r,r,r,r,r,b"
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "@$6 cp.async.bulk.tensor.3d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4}], [$5];", "l,r,r,r,r,r,b"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1,%crd2] predicate=%p : !llvm.ptr<3>, !llvm.ptr
return
}
Expand All @@ -156,7 +156,7 @@ func.func @tma_load_3d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier
func.func @tma_load_4d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier: !llvm.ptr<3>, %crd0: i32, %crd1: i32, %crd2: i32, %crd3: i32, %p : i1) {
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.bulk.tensor.4d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4,$5} ], [$6];", "r,l,r,r,r,r,r"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1,%crd2,%crd3] : !llvm.ptr<3>, !llvm.ptr
// CHECK : llvm.inline_asm has_side_effects asm_dialect = att "@$7 cp.async.bulk.tensor.4d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4,$5}], [$6];", "l,r,r,r,r,r,r,b"
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "@$7 cp.async.bulk.tensor.4d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4,$5}], [$6];", "l,r,r,r,r,r,r,b"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1,%crd2,%crd3] predicate=%p : !llvm.ptr<3>, !llvm.ptr
return
}
Expand All @@ -165,7 +165,7 @@ func.func @tma_load_4d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier
func.func @tma_load_5d(%tmaDescriptor: !llvm.ptr, %dest : !llvm.ptr<3>, %barrier: !llvm.ptr<3>, %crd0: i32, %crd1: i32, %crd2: i32, %crd3: i32, %crd4: i32, %p : i1) {
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.bulk.tensor.5d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4,$5,$6} ], [$7];", "r,l,r,r,r,r,r,r"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1,%crd2,%crd3,%crd4] : !llvm.ptr<3>, !llvm.ptr
// CHECK : llvm.inline_asm has_side_effects asm_dialect = att "@$8 cp.async.bulk.tensor.5d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4,$5,$6}], [$7];", "l,r,r,r,r,r,r,r,b"
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "@$8 cp.async.bulk.tensor.5d.shared::cluster.global.mbarrier::complete_tx::bytes [$0], [$1, {$2,$3,$4,$5,$6}], [$7];", "l,r,r,r,r,r,r,r,b"
nvvm.cp.async.bulk.tensor.shared.cluster.global %dest, %tmaDescriptor, %barrier, box[%crd0,%crd1,%crd2,%crd3,%crd4] predicate=%p : !llvm.ptr<3>, !llvm.ptr
return
}
Expand Down
6 changes: 3 additions & 3 deletions mlir/test/Conversion/PDLToPDLInterp/pdl-to-pdl-interp-matcher.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -588,7 +588,7 @@ module @variadic_results_all {
// CHECK-DAG: %[[OPS:.*]] = pdl_interp.get_users of %[[VAL0]] : !pdl.value
// CHECK-DAG: pdl_interp.foreach %[[OP:.*]] : !pdl.operation in %[[OPS]]
// CHECK-DAG: %[[OPERANDS:.*]] = pdl_interp.get_operands of %[[OP]]
// CHECK-DAG pdl_interp.are_equal %[[VALS]], %[[OPERANDS]] -> ^{{.*}}, ^[[CONTINUE:.*]]
// CHECK-DAG: pdl_interp.are_equal %[[VALS]], %[[OPERANDS]] -> ^{{.*}}, ^[[CONTINUE:.*]]
// CHECK-DAG: pdl_interp.is_not_null %[[OP]]
// CHECK-DAG: pdl_interp.check_result_count of %[[OP]] is 0
pdl.pattern @variadic_results_all : benefit(1) {
Expand Down Expand Up @@ -701,7 +701,7 @@ module @common_connector {
// CHECK-DAG: pdl_interp.are_equal %[[ROOTA_OP]], %[[VAL0]] : !pdl.value
// CHECK-DAG: %[[ROOTB_OP:.*]] = pdl_interp.get_operand 0 of %[[ROOTB]]
// CHECK-DAG: pdl_interp.are_equal %[[ROOTB_OP]], %[[VAL0]] : !pdl.value
// CHECK-DAG } -> ^[[CONTA:.*]]
// CHECK-DAG: } -> ^[[CONTA:.*]]
pdl.pattern @common_connector : benefit(1) {
%type = type
%op = operation -> (%type, %type : !pdl.type, !pdl.type)
Expand Down Expand Up @@ -742,7 +742,7 @@ module @common_connector_range {
// CHECK-DAG: pdl_interp.are_equal %[[ROOTA_OPS]], %[[VALS0]] : !pdl.range<value>
// CHECK-DAG: %[[ROOTB_OPS:.*]] = pdl_interp.get_operands of %[[ROOTB]]
// CHECK-DAG: pdl_interp.are_equal %[[ROOTB_OPS]], %[[VALS0]] : !pdl.range<value>
// CHECK-DAG } -> ^[[CONTA:.*]]
// CHECK-DAG: } -> ^[[CONTA:.*]]
pdl.pattern @common_connector_range : benefit(1) {
%types = types
%op = operation -> (%types, %types : !pdl.range<type>, !pdl.range<type>)
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/spirv-storage-class-mapping.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,5 +91,5 @@ spirv.func @pointerCodeSectionINTEL(!spirv.ptr<i1, CodeSectionINTEL>) "None"
spirv.func @pointerDeviceOnlyINTEL(!spirv.ptr<i1, DeviceOnlyINTEL>) "None"

// CHECK-OPENCL: llvm.func @pointerHostOnlyINTEL(!llvm.ptr<6>)
// CHECK-UNKOWN: llvm.func @pointerHostOnlyINTEL(!llvm.ptr)
// CHECK-UNKNOWN: llvm.func @pointerHostOnlyINTEL(!llvm.ptr)
spirv.func @pointerHostOnlyINTEL(!spirv.ptr<i1, HostOnlyINTEL>) "None"
4 changes: 2 additions & 2 deletions mlir/test/Dialect/Arith/unsigned-when-equivalent.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// RUN: mlir-opt -arith-unsigned-when-equivalent %s | FileCheck %s

// CHECK-LABEL func @not_with_maybe_overflow
// CHECK-LABEL: func @not_with_maybe_overflow
// CHECK: arith.divsi
// CHECK: arith.ceildivsi
// CHECK: arith.floordivsi
Expand Down Expand Up @@ -32,7 +32,7 @@ func.func @not_with_maybe_overflow(%arg0 : i32) {
func.return
}

// CHECK-LABEL func @yes_with_no_overflow
// CHECK-LABEL: func @yes_with_no_overflow
// CHECK: arith.divui
// CHECK: arith.ceildivui
// CHECK: arith.divui
Expand Down
32 changes: 16 additions & 16 deletions mlir/test/Dialect/ArmSME/tile-allocation-liveness.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -366,15 +366,15 @@ func.func @avoidable_spill(%a: vector<[4]xf32>, %b: vector<[4]xf32>, %c: vector<

// CHECK-LIVE-RANGE-LABEL: @cond_branch_with_backedge
// CHECK-LIVE-RANGE: ^bb1:
// CHECK-LIVE-RANGE--NEXT: ||| | arith.cmpi
// CHECK-LIVE-RANGE--NEXT: EEE E cf.cond_br
// CHECK-LIVE-RANGE-NEXT: ||| | arith.cmpi
// CHECK-LIVE-RANGE-NEXT: EEE E cf.cond_br
//
// CHECK-LIVE-RANGE--NEXT: ^[[BB3_COPIES:[[:alnum:]]+]]:
// CHECK-LIVE-RANGE--NEXT: ||| ES arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: E|| |S arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: E| ||S arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: E |||S arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: EEEE cf.br
// CHECK-LIVE-RANGE-NEXT: ^[[BB3_COPIES:[[:alnum:]]+]]:
// CHECK-LIVE-RANGE-NEXT: ||| ES arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: E|| |S arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: E| ||S arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: E |||S arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: EEEE cf.br
//
// It is important to note that the first three live ranges in ^bb1 do not end
// at the `cf.cond_br` they are live-out via the backedge bb1 -> bb2 -> bb1.
Expand All @@ -389,15 +389,15 @@ func.func @avoidable_spill(%a: vector<[4]xf32>, %b: vector<[4]xf32>, %c: vector<
//
// CHECK-LIVE-RANGE: ========== Coalesced Live Ranges:
// CHECK-LIVE-RANGE: ^bb1:
// CHECK-LIVE-RANGE--NEXT: |||| arith.cmpi
// CHECK-LIVE-RANGE--NEXT: EEEE cf.cond_br
// CHECK-LIVE-RANGE-NEXT: |||| arith.cmpi
// CHECK-LIVE-RANGE-NEXT: EEEE cf.cond_br
//
// CHECK-LIVE-RANGE--NEXT: ^[[BB3_COPIES]]:
// CHECK-LIVE-RANGE--NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE--NEXT: EEEE cf.br
// CHECK-LIVE-RANGE-NEXT: ^[[BB3_COPIES]]:
// CHECK-LIVE-RANGE-NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: |||| arm_sme.copy_tile
// CHECK-LIVE-RANGE-NEXT: EEEE cf.br

// CHECK-LABEL: @cond_branch_with_backedge
// CHECK-NOT: tile_id = 16
Expand Down
8 changes: 4 additions & 4 deletions mlir/test/Dialect/Bufferization/Transforms/lower-deallocations-func.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,10 +9,10 @@ func.func @conversion_dealloc_simple(%arg0: memref<2xf32>, %arg1: i1) {
return
}

// CHECk: scf.if [[ARG1]] {
// CHECk-NEXT: memref.dealloc [[ARG0]] : memref<2xf32>
// CHECk-NEXT: }
// CHECk-NEXT: return
// CHECK: scf.if [[ARG1]] {
// CHECK-NEXT: memref.dealloc [[ARG0]] : memref<2xf32>
// CHECK-NEXT: }
// CHECK-NEXT: return

// -----

Expand Down
8 changes: 4 additions & 4 deletions mlir/test/Dialect/Bufferization/Transforms/lower-deallocations.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,10 +29,10 @@ func.func @conversion_dealloc_simple(%arg0: memref<2xf32>, %arg1: i1) {
return
}

// CHECk: scf.if [[ARG1]] {
// CHECk-NEXT: memref.dealloc [[ARG0]] : memref<2xf32>
// CHECk-NEXT: }
// CHECk-NEXT: return
// CHECK: scf.if [[ARG1]] {
// CHECK-NEXT: memref.dealloc [[ARG0]] : memref<2xf32>
// CHECK-NEXT: }
// CHECK-NEXT: return

// -----

Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Dialect/GPU/barrier-elimination.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,7 @@ func.func @write_in_a_loop(%arg0: memref<?xf32>, %arg1: f32) attributes {__paral
return
}

// CHECK-LABEL @read_read_write_loop
// CHECK-LABEL: @read_read_write_loop
func.func @read_read_write_loop(%arg0: memref<?xf32>, %arg1: f32) attributes {__parallel_region_boundary_for_test} {
%c0 = arith.constant 0 : index
%c42 = arith.constant 42 : index
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Dialect/GPU/ops.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -227,7 +227,7 @@ module attributes {gpu.container_module} {
gpu.return
}

// CHECK-LABEL gpu.func @printf_test
// CHECK-LABEL: gpu.func @printf_test
// CHECK: (%[[ARG0:.*]]: i32)
// CHECK: gpu.printf "Value: %d" %[[ARG0]] : i32
gpu.func @printf_test(%arg0 : i32) {
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Dialect/GPU/outlining.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -123,7 +123,7 @@ llvm.func @launch_from_llvm_func() {
llvm.return
}

// CHECK-DL-LABLE: gpu.module @launch_from_llvm_func_kernel attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<index, 32 : i32>>}
// CHECK-DL-LABEL: gpu.module @launch_from_llvm_func_kernel attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<index, 32 : i32>>}

// -----

Expand Down
14 changes: 7 additions & 7 deletions mlir/test/Dialect/LLVMIR/nvvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -464,24 +464,24 @@ llvm.func private @mbarrier_test_wait_shared(%barrier: !llvm.ptr<3>, %token : i6
llvm.return
}

// CHECK-LABEL : @wgmma_fence_aligned
// CHECK-LABEL: @wgmma_fence_aligned
func.func @wgmma_fence_aligned() {
// CHECK : nvvm.wgmma.fence.aligned
// CHECK: nvvm.wgmma.fence.aligned
nvvm.wgmma.fence.aligned
return
}

// CHECK-LABEL : @wgmma_commit_group_sync_aligned
// CHECK-LABEL: @wgmma_commit_group_sync_aligned
func.func @wgmma_commit_group_sync_aligned() {
// CHECK : nvvm.wgmma.commit.group.sync.aligned
// CHECK: nvvm.wgmma.commit.group.sync.aligned
nvvm.wgmma.commit.group.sync.aligned
return
}


// CHECK-LABEL : @wgmma_commit_group_sync_aligned
// CHECK-LABEL: @wgmma_commit_group_sync_aligned
func.func @wgmma_wait_group_sync_aligned() {
// CHECK : nvvm.wgmma.wait.group.sync.aligned
// CHECK: nvvm.wgmma.wait.group.sync.aligned
nvvm.wgmma.wait.group.sync.aligned 0
return
}
Expand All @@ -495,7 +495,7 @@ gpu.module @module_1 [#nvvm.target<chip = "sm_90", features = "+ptx70", link = [
gpu.module @module_2 [#nvvm.target<chip = "sm_90">, #nvvm.target<chip = "sm_80">, #nvvm.target<chip = "sm_70">] {
}

// CHECK-LABEL : nvvm.grid_constant
// CHECK-LABEL: nvvm.grid_constant
llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant}) attributes {nvvm.kernel} {
llvm.return
}
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Dialect/Linalg/data-layout-propagation.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -795,7 +795,7 @@ func.func @reduction_pack_transpose_inner_dims(%arg0: tensor<128x256x32xi32>,
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9]+]]
// CHECK: %[[ARG1_EMPTY:.+]] = tensor.empty() : tensor<4x16x16x32xi32>
// CHECK: %[[PACK_ARG1:.+]] = tensor.pack %[[ARG1]]
// CHECK-SME: inner_dims_pos = [1, 0] inner_tiles = [16, 32]
// CHECK-SAME: inner_dims_pos = [1, 0] inner_tiles = [16, 32]
// CHECK-SAME: into %[[ARG1_EMPTY]]
// CHECK: %[[ARG0_EMPTY:.+]] = tensor.empty() : tensor<4x16x32x16x32xi32>
// CHECK: %[[PACK_ARG0:.+]] = tensor.pack %[[ARG0]]
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Dialect/Math/expand-math.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -221,7 +221,7 @@ func.func @roundf_func(%a: f32) -> f32 {
// CHECK-LABEL: func @powf_func
// CHECK-SAME: ([[ARG0:%.+]]: f64, [[ARG1:%.+]]: f64)
func.func @powf_func(%a: f64, %b: f64) ->f64 {
// CHECK-DAG = [[CST0:%.+]] = arith.constant 0.000000e+00
// CHECK-DAG: [[CST0:%.+]] = arith.constant 0.000000e+00
// CHECK-DAG: [[TWO:%.+]] = arith.constant 2.000000e+00
// CHECK-DAG: [[NEGONE:%.+]] = arith.constant -1.000000e+00
// CHECK-DAG: [[SQR:%.+]] = arith.mulf [[ARG0]], [[ARG0]]
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Dialect/SCF/transform-ops.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
// CHECK: scf.for
// CHECK: arith.addi
//
// CHECK-LABEL @loop_outline_op
// CHECK-LABEL: @loop_outline_op
func.func @loop_outline_op(%arg0: index, %arg1: index, %arg2: index) {
// CHECK: scf.for
// CHECK-NOT: scf.for
Expand Down
12 changes: 6 additions & 6 deletions mlir/test/Dialect/SPIRV/IR/logical-ops.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -180,47 +180,47 @@ func.func @logicalUnary(%arg0 : i32)
func.func @select_op_bool(%arg0: i1) -> () {
%0 = spirv.Constant true
%1 = spirv.Constant false
// CHECK : spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i1
// CHECK: spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i1
%2 = spirv.Select %arg0, %0, %1 : i1, i1
return
}

func.func @select_op_int(%arg0: i1) -> () {
%0 = spirv.Constant 2 : i32
%1 = spirv.Constant 3 : i32
// CHECK : spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i32
// CHECK: spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i32
%2 = spirv.Select %arg0, %0, %1 : i1, i32
return
}

func.func @select_op_float(%arg0: i1) -> () {
%0 = spirv.Constant 2.0 : f32
%1 = spirv.Constant 3.0 : f32
// CHECK : spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, f32
// CHECK: spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, f32
%2 = spirv.Select %arg0, %0, %1 : i1, f32
return
}

func.func @select_op_ptr(%arg0: i1) -> () {
%0 = spirv.Variable : !spirv.ptr<f32, Function>
%1 = spirv.Variable : !spirv.ptr<f32, Function>
// CHECK : spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, !spirv.ptr<f32, Function>
// CHECK: spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, !spirv.ptr<f32, Function>
%2 = spirv.Select %arg0, %0, %1 : i1, !spirv.ptr<f32, Function>
return
}

func.func @select_op_vec(%arg0: i1) -> () {
%0 = spirv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
%1 = spirv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
// CHECK : spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, vector<3xf32>
// CHECK: spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, vector<3xf32>
%2 = spirv.Select %arg0, %0, %1 : i1, vector<3xf32>
return
}

func.func @select_op_vec_condn_vec(%arg0: vector<3xi1>) -> () {
%0 = spirv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
%1 = spirv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
// CHECK : spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi1>, vector<3xf32>
// CHECK: spirv.Select {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi1>, vector<3xf32>
%2 = spirv.Select %arg0, %0, %1 : vector<3xi1>, vector<3xf32>
return
}
Expand Down
Loading
Loading