diff options
Diffstat (limited to 'mlir/test')
125 files changed, 3647 insertions, 362 deletions
diff --git a/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir b/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir index dbff233..455f886 100644 --- a/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir +++ b/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir @@ -1,4 +1,5 @@ -// RUN: mlir-opt %s -convert-math-to-rocdl -allow-unregistered-dialect -split-input-file | FileCheck %s +// RUN: mlir-opt %s -allow-unregistered-dialect -split-input-file -pass-pipeline='builtin.module(convert-math-to-rocdl{chipset=gfx803})' | FileCheck %s --check-prefix=PRE9 +// RUN: mlir-opt %s -allow-unregistered-dialect -split-input-file -pass-pipeline='builtin.module(convert-math-to-rocdl{chipset=gfx942})' | FileCheck %s --check-prefix=POST9 module @test_module { // CHECK: llvm.func @__ocml_fmod_f16(f16, f16) -> f16 @@ -596,3 +597,76 @@ module @test_module { func.return %result : vector<2x2xf16> } } + +// ----- + +// f16 clamp → rocdl.fmed3 on gfx9+ +// CHECK-LABEL: func.func @clampf_f16 +func.func @clampf_f16(%x: f16, %lo: f16, %hi: f16) -> f16 { + %r = math.clampf %x to [%lo, %hi] : f16 + return %r : f16 + // POST9: rocdl.fmed3 {{.*}} : f16 + // PRE9-NOT: rocdl.fmed3 + // PRE9: math.clampf {{.*}} : f16 +} + +// f32 clamp → rocdl.fmed3 on gfx9+ +// CHECK-LABEL: func.func @clampf_f32 +func.func @clampf_f32(%x: f32, %lo: f32, %hi: f32) -> f32 { + %r = math.clampf %x to [%lo, %hi] : f32 + return %r : f32 + // POST9: rocdl.fmed3 {{.*}} : f32 + // PRE9-NOT: rocdl.fmed3 + // PRE9: math.clampf {{.*}} : f32 +} + +// ----- + +// Vector f16 clamp → rocdl.fmed3 on gfx9+ +// CHECK-LABEL: func.func @clampf_vector_f16 +func.func @clampf_vector_f16(%x: vector<2xf16>, %lo: vector<2xf16>, %hi: vector<2xf16>) -> vector<2xf16> { + %r = math.clampf %x to [%lo, %hi] : vector<2xf16> + return %r : vector<2xf16> + // POST9: rocdl.fmed3 {{.*}} : vector<2xf16> + // PRE9-NOT: rocdl.fmed3 + // PRE9: math.clampf {{.*}} : vector<2xf16> +} + +// ----- + +// Vector f32 clamp → rocdl.fmed3 on gfx9+ +// CHECK-LABEL: func.func @clampf_vector_f32 +func.func @clampf_vector_f32(%x: vector<2xf32>, %lo: vector<2xf32>, %hi: vector<2xf32>) -> vector<2xf32> { + %r = math.clampf %x to [%lo, %hi] : vector<2xf32> + return %r : vector<2xf32> + // POST9: rocdl.fmed3 {{.*}} : vector<2xf32> + // PRE9-NOT: rocdl.fmed3 + // PRE9: math.clampf {{.*}} : vector<2xf32> +} + +// ----- + +// Multi-dimensional vector f16 clamp → rocdl.fmed3 on gfx9+ (unrolled to 1D vectors) +// CHECK-LABEL: func.func @clampf_vector_2d_f16 +func.func @clampf_vector_2d_f16(%x: vector<2x2xf16>, %lo: vector<2x2xf16>, %hi: vector<2x2xf16>) -> vector<2x2xf16> { + %r = math.clampf %x to [%lo, %hi] : vector<2x2xf16> + return %r : vector<2x2xf16> + // POST9: builtin.unrealized_conversion_cast {{.*}} : vector<2x2xf16> to !llvm.array<2 x vector<2xf16>> + // POST9: llvm.extractvalue {{.*}} : !llvm.array<2 x vector<2xf16>> + // POST9: rocdl.fmed3 {{.*}} : vector<2xf16> + // POST9: llvm.insertvalue {{.*}} : !llvm.array<2 x vector<2xf16>> + // POST9: llvm.extractvalue {{.*}} : !llvm.array<2 x vector<2xf16>> + // POST9: rocdl.fmed3 {{.*}} : vector<2xf16> + // POST9: llvm.insertvalue {{.*}} : !llvm.array<2 x vector<2xf16>> + // PRE9-NOT: rocdl.fmed3 + // PRE9: math.clampf {{.*}} : vector<2x2xf16> +} + +// ----- +// CHECK-LABEL: func.func @clampf_bf16 +func.func @clampf_bf16(%x: bf16, %lo: bf16, %hi: bf16) -> bf16 { + %r = math.clampf %x to [%lo, %hi] : bf16 + return %r : bf16 + // CHECK: math.clampf {{.*}} : bf16 + // CHECK-NOT: rocdl.fmed3 +} diff --git a/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir b/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir index 2d33888..d669a3b 100644 --- a/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir +++ b/mlir/test/Conversion/VectorToLLVM/vector-to-llvm.mlir @@ -76,6 +76,18 @@ func.func @broadcast_vec1d_from_f32(%arg0: f32) -> vector<2xf32> { // ----- +func.func @broadcast_single_elem_vec1d_from_f32(%arg0: f32) -> vector<1xf32> { + %0 = vector.broadcast %arg0 : f32 to vector<1xf32> + return %0 : vector<1xf32> +} +// CHECK-LABEL: @broadcast_single_elem_vec1d_from_f32 +// CHECK-SAME: %[[A:.*]]: f32) +// CHECK: %[[T0:.*]] = llvm.insertelement %[[A]] +// CHECK-NOT: llvm.shufflevector +// CHECK: return %[[T0]] : vector<1xf32> + +// ----- + func.func @broadcast_vec1d_from_f32_scalable(%arg0: f32) -> vector<[2]xf32> { %0 = vector.broadcast %arg0 : f32 to vector<[2]xf32> return %0 : vector<[2]xf32> diff --git a/mlir/test/Conversion/XeGPUToXeVM/dpas.mlir b/mlir/test/Conversion/XeGPUToXeVM/dpas.mlir index e6f22f0..a9ab0be 100644 --- a/mlir/test/Conversion/XeGPUToXeVM/dpas.mlir +++ b/mlir/test/Conversion/XeGPUToXeVM/dpas.mlir @@ -1,17 +1,13 @@ // RUN: mlir-opt -convert-xegpu-to-xevm %s | FileCheck %s -#sg_map_a_f16 = #xegpu.layout<lane_layout = [1, 16], lane_data = [1, 1]> -#sg_map_b_f16 = #xegpu.layout<lane_layout = [1, 16], lane_data = [2, 1]> -#sg_map_c_f32 = #xegpu.layout<lane_layout = [1, 16], lane_data = [1, 1]> - -gpu.module @load_store_check { +gpu.module @test_kernel { // CHECK-LABEL: func.func @dpas( // CHECK-SAME: %[[ARG0:.*]]: vector<8xf16>, %[[ARG1:.*]]: vector<16xf16>, %[[ARG2:.*]]: vector<8xf32> func.func @dpas(%a_loaded: vector<8xf16>, %b_loaded: vector<16xf16>, %c_loaded: vector<8xf32>) -> vector<8xf32> { // Loads are checked in a separate test. // CHECK: %[[D:.*]] = xevm.mma %[[ARG0]], %[[ARG1]], %[[ARG2]] {shape = <m = 8, n = 16, k = 16>, types = <d = f32, a = f16, b = f16, c = f32>} // CHECK-SAME: : (vector<8xf16>, vector<16xf16>, vector<8xf32>) -> vector<8xf32> - %d = xegpu.dpas %a_loaded, %b_loaded, %c_loaded {a_layout = #sg_map_a_f16, b_layout = #sg_map_b_f16, c_layout = #sg_map_c_f32} + %d = xegpu.dpas %a_loaded, %b_loaded, %c_loaded : vector<8xf16>, vector<16xf16>, vector<8xf32> -> vector<8xf32> return %d : vector<8xf32> } diff --git a/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir b/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir new file mode 100644 index 0000000..d4cb493 --- /dev/null +++ b/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir @@ -0,0 +1,201 @@ +// RUN: mlir-opt -split-input-file -convert-xegpu-to-xevm -cse %s | FileCheck %s + +gpu.module @test_kernel [#xevm.target<chip = "pvc">] { + + // e.g. for mem_desc<32x32xf16, @strides=[1, 16]> + // its memory layout tuple is (blocked shape = [1,1,32,32],strides=[1024,1024,32,1]) + //CHECK-LABEL: load_store_matrix_1 + gpu.func @load_store_matrix_1(%arg0: memref<4096xi8, 3>) -> f32 { + %0 = xegpu.create_mem_desc %arg0 : memref<4096xi8, 3> -> !xegpu.mem_desc<32x32xf32> + + //CHECK: %[[TID:.*]] = gpu.thread_id x + //CHECK: %[[C1:.*]] = arith.constant 1 : index + //CHECK: %[[MUL1:.*]] = arith.muli %[[TID]], %[[C1]] : index + //CHECK: %[[C4:.*]] = arith.constant 4 : i32 + //CHECK: %[[MUL2:.*]] = arith.muli {{.*}}, %[[C4]] : i32 + //CHECK: llvm.load {{.*}} : !llvm.ptr<3> -> f32 + + %tid_x = gpu.thread_id x + %c0 = arith.constant 0 : index + %1 = xegpu.load_matrix %0[%c0, %tid_x]: !xegpu.mem_desc<32x32xf32>, index, index -> f32 + + //CHECK: llvm.store {{.*}}, {{.*}} : f32, !llvm.ptr<3> + + xegpu.store_matrix %1, %0[%c0, %tid_x]: f32, !xegpu.mem_desc<32x32xf32>, index, index + + gpu.return %1: f32 + } + +// e.g. for mem_desc<32x64xf16, @block=[16, 16], @strides=[1, 32]> + // its memory layout tuple is ([2,4,16,16],[256,512,1,16]) + //CHECK-LABEL: load_store_matrix_2 + gpu.func @load_store_matrix_2(%arg0: memref<4096xi8, 3>) -> f16 { + %0 = xegpu.create_mem_desc %arg0 : memref<4096xi8, 3> -> !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<stride = [1, 32], block = [16, 16]>> + //CHECK: %[[c0:.*]] = arith.constant 0 : index + //CHECK: %[[tid_x:.*]] = gpu.thread_id x + //CHECK: %[[c13:.*]] = arith.constant 13 : index + //CHECK: %[[c16:.*]] = arith.constant 16 : index + //CHECK: %[[offsetx_0:.*]] = arith.divsi %[[c13]], %[[c16]] : index + //CHECK: %[[offsetx_1:.*]] = arith.remsi %[[c13]], %[[c16]] : index + //CHECK: %[[offsety_0:.*]] = arith.divsi %[[tid_x]], %[[c16]] : index + //CHECK: %[[offsety_1:.*]] = arith.remsi %[[tid_x]], %[[c16]] : index + + //CHECK: %[[c256:.*]] = arith.constant 256 : index + //CHECK: %[[mul0:.*]] = arith.muli %[[offsetx_0]], %[[c256]] : index + //CHECK: %[[add0:.*]] = arith.addi %[[mul0]], %[[c0]] : index + //CHECK: %[[c512:.*]] = arith.constant 512 : index + //CHECK: %[[mul1:.*]] = arith.muli %[[offsety_0]], %[[c512]] : index + //CHECK: %[[add1:.*]] = arith.addi %[[mul1]], %[[add0]] : index + //CHECK: %[[c1:.*]] = arith.constant 1 : index + //CHECK: %[[mul2:.*]] = arith.muli %[[offsetx_1]], %[[c1]] : index + //CHECK: %[[add2:.*]] = arith.addi %[[mul2]], %[[add1]] : index + //CHECK: %[[mul3:.*]] = arith.muli %[[offsety_1]], %[[c16]] : index + //CHECK: %[[add3:.*]] = arith.addi %[[mul3]], %[[add2]] : index + + //CHECK: %[[loaded:.*]] = llvm.load {{.*}}: !llvm.ptr<3> -> f16 + + + %tid_x = gpu.thread_id x + %c13 = arith.constant 13 : index + %1 = xegpu.load_matrix %0[%c13, %tid_x]: !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<stride = [1, 32], block = [16, 16]>>, index, index -> f16 + + //CHECK: llvm.store %[[loaded]], {{.*}} : f16, !llvm.ptr<3> + + xegpu.store_matrix %1, %0[%c13, %tid_x]: f16, !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<stride = [1, 32], block = [16, 16]>>, index, index + gpu.return %1: f16 + } + + + // e.g. for mem_desc<32x64xf16, @block=[16, 16]> + // its memory layout tuple is ([2,4,16,16],[1024,256,16,1]) + //CHECK-LABEL: load_store_matrix_3 + gpu.func @load_store_matrix_3(%arg0: memref<4096xi8, 3>) -> f16 { + //CHECK: %[[c0:.*]] = arith.constant 0 : index + //CHECK: %[[view:.*]] = memref.view %arg0[%[[c0]]][] : memref<4096xi8, 3> to memref<2048xf16, 3> + %0 = xegpu.create_mem_desc %arg0 : memref<4096xi8, 3> -> !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<block = [16, 16]>> + + //CHECK: %[[tid_x:.*]] = gpu.thread_id x + //CHECK: %[[c19:.*]] = arith.constant 19 : index + %tid_x = gpu.thread_id x + %c19 = arith.constant 19: index + + //CHECK: %[[intptr:.*]] = memref.extract_aligned_pointer_as_index %[[view]] : memref<2048xf16, 3> -> index + //CHECK: %[[basePtrI64:.*]] = arith.index_castui %[[intptr]] : index to i32 + //CHECK: %[[c16:.*]] = arith.constant 16 : index + //CHECK: %[[offsetx_0:.*]] = arith.divsi %[[c19]], %[[c16]] : index + //CHECK: %[[offsetx_1:.*]] = arith.remsi %[[c19]], %[[c16]] : index + //CHECK: %[[offsety_0:.*]] = arith.divsi %[[tid_x]], %[[c16]] : index + //CHECK: %[[offsety_1:.*]] = arith.remsi %[[tid_x]], %[[c16]] : index + //CHECK: %[[c1024:.*]] = arith.constant 1024 : index + //CHECK: %[[mul0:.*]] = arith.muli %[[offsetx_0]], %[[c1024]] : index + //CHECK: %[[add0:.*]] = arith.addi %[[mul0]], %[[c0]] : index + //CHECK: %[[c256:.*]] = arith.constant 256 : index + //CHECK: %[[mul1:.*]] = arith.muli %[[offsety_0]], %[[c256]] : index + //CHECK: %[[add1:.*]] = arith.addi %[[mul1]], %[[add0]] : index + //CHECK: %[[mul2:.*]] = arith.muli %[[offsetx_1]], %[[c16]] : index + //CHECK: %[[add2:.*]] = arith.addi %[[mul2]], %[[add1]] : index + //CHECK: %[[c1:.*]] = arith.constant 1 : index + //CHECK: %[[mul3:.*]] = arith.muli %[[offsety_1]], %[[c1]] : index + //CHECK: %[[add3:.*]] = arith.addi %[[mul3]], %[[add2]] : index + + //CHECK: %[[loaded:.*]] = llvm.load {{.*}} : !llvm.ptr<3> -> f16 + %1 = xegpu.load_matrix %0[%c19, %tid_x]: !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<block = [16, 16]>>, index, index -> f16 + + //CHECK: llvm.store %[[loaded]], {{.*}} : f16, !llvm.ptr<3> + xegpu.store_matrix %1, %0[%c19, %tid_x]: f16, !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<block = [16, 16]>>, index, index + + //CHECK: gpu.return %[[loaded]] : f16 + gpu.return %1: f16 + } + + // e.g. for mem_desc<32x64xf16, @block=[16, 16], @strides=[1, 16]> + // its memory layout tuple is ([2,4,16,16],[256,512,1,16]) + //CHECK-LABEL: load_store_matrix_4 + gpu.func @load_store_matrix_4(%arg0: memref<4096xi8, 3>) -> vector<8xf16> { + %0 = xegpu.create_mem_desc %arg0 : memref<4096xi8, 3> -> !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<stride = [1, 32], block = [16, 16]>> + + //CHECK: %[[c0:.*]] = arith.constant 0 : index + //CHECK: %[[tid_x:.*]] = gpu.thread_id x + + //CHECK: %[[c16:.*]] = arith.constant 16 : index + //CHECK: %[[offsetx_0:.*]] = arith.divsi %[[c16]], %[[c16]] : index + //CHECK: %[[offsetx_1:.*]] = arith.remsi %[[c16]], %[[c16]] : index + //CHECK: %[[offsety_0:.*]] = arith.divsi %[[tid_x]], %[[c16]] : index + //CHECK: %[[offsety_1:.*]] = arith.remsi %[[tid_x]], %[[c16]] : index + + //CHECK: %[[c256:.*]] = arith.constant 256 : index + //CHECK: %[[mul0:.*]] = arith.muli %[[offsetx_0]], %[[c256]] : index + //CHECK: %[[add0:.*]] = arith.addi %[[mul0]], %[[c0]] : index + //CHECK: %[[c512:.*]] = arith.constant 512 : index + //CHECK: %[[mul1:.*]] = arith.muli %[[offsety_0]], %[[c512]] : index + //CHECK: %[[add1:.*]] = arith.addi %[[mul1]], %[[add0]] : index + //CHECK: %[[c1:.*]] = arith.constant 1 : index + //CHECK: %[[mul2:.*]] = arith.muli %[[offsetx_1]], %[[c1]] : index + //CHECK: %[[add2:.*]] = arith.addi %[[mul2]], %[[add1]] : index + //CHECK: %[[mul3:.*]] = arith.muli %[[offsety_1]], %[[c16]] : index + //CHECK: %[[add3:.*]] = arith.addi %[[mul3]], %[[add2]] : index + + //CHECK: %[[loaded:.*]] = llvm.load {{.*}}: !llvm.ptr<3> -> vector<8xf16> + + %tid_x = gpu.thread_id x + %c16 = arith.constant 16 : index + %1 = xegpu.load_matrix %0[%c16, %tid_x] : !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<stride = [1, 32], block = [16, 16]>>, index, index -> vector<8xf16> + + //CHECK: llvm.store %[[loaded]], {{.*}} : vector<8xf16>, !llvm.ptr<3> + xegpu.store_matrix %1, %0[%c16, %tid_x] : vector<8xf16>, !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<stride = [1, 32], block = [16, 16]>>, index, index + + gpu.return %1: vector<8xf16> + } + + + // e.g. for mem_desc<32x64xf16, @block=[16, 16]> + // its memory layout tuple is ([2,4,16,16],[1024,256,16,1]) + //CHECK-LABEL: load_store_matrix_5 + gpu.func @load_store_matrix_5(%arg0: memref<4096xi8, 3>) -> vector<8xf16> { + //CHECK: %[[c0:.*]] = arith.constant 0 : index + //CHECK: %[[view:.*]] = memref.view %arg0[%[[c0]]][] : memref<4096xi8, 3> to memref<2048xf16, 3> + + %0 = xegpu.create_mem_desc %arg0 : memref<4096xi8, 3> -> !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<block = [16, 16]>> + + //CHECK: %[[c16:.*]] = arith.constant 16 : index + //CHECK: %[[c48:.*]] = arith.constant 48 : index + + %c16 = arith.constant 16 : index + %c48 = arith.constant 48 : index + + //CHECK: %[[intptr:.*]] = memref.extract_aligned_pointer_as_index %[[view]] : memref<2048xf16, 3> -> index + //CHECK: %[[basePtrI64:.*]] = arith.index_castui %[[intptr]] : index to i32 + //CHECK: %[[offset0:.*]] = arith.divsi %[[c16]], %[[c16]] : index + //CHECK: %[[offset1:.*]] = arith.remsi %[[c16]], %[[c16]] : index + //CHECK: %[[offset2:.*]] = arith.divsi %[[c48]], %[[c16]] : index + //CHECK: %[[offset3:.*]] = arith.remsi %[[c48]], %[[c16]] : index + //CHECK: %[[c1024:.*]] = arith.constant 1024 : index + //CHECK: %[[mul0:.*]] = arith.muli %[[offset0]], %[[c1024]] : index + //CHECK: %[[add0:.*]] = arith.addi %[[mul0]], %[[c0]] : index + //CHECK: %[[c256:.*]] = arith.constant 256 : index + //CHECK: %[[mul1:.*]] = arith.muli %[[offset2]], %[[c256]] : index + //CHECK: %[[add1:.*]] = arith.addi %[[mul1]], %[[add0]] : index + //CHECK: %[[mul2:.*]] = arith.muli %[[offset1]], %[[c16]] : index + //CHECK: %[[add2:.*]] = arith.addi %[[mul2]], %[[add1]] : index + //CHECK: %[[c1:.*]] = arith.constant 1 : index + //CHECK: %[[mul3:.*]] = arith.muli %[[offset3]], %[[c1]] : index + //CHECK: %[[linearOffset:.*]] = arith.addi %[[mul3]], %[[add2]] : index + //CHECK: %[[linearOffsetI64:.*]] = arith.index_castui %[[linearOffset]] : index to i32 + //CHECK: %[[c2:.*]] = arith.constant 2 : i32 + //CHECK: %[[byteOffset:.*]] = arith.muli %[[linearOffsetI64]], %[[c2]] : i32 + //CHECK: %[[finalPtr:.*]] = arith.addi %[[basePtrI64]], %[[byteOffset]] : i32 + //CHECK: %[[ptr:.*]] = llvm.inttoptr %[[finalPtr]] : i32 to !llvm.ptr<3> + //CHECK: %[[loadedI16:.*]] = xevm.blockload %[[ptr]] : (!llvm.ptr<3>) -> vector<8xi16> + //CHECK: %[[loaded:.*]] = vector.bitcast %[[loadedI16]] : vector<8xi16> to vector<8xf16> + + %1 = xegpu.load_matrix %0[%c16, %c48] {subgroup_block_io}: !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<block = [16, 16]>>, index, index -> vector<8xf16> + + //CHECK: %[[storeDataI16:.*]] = vector.bitcast %[[loaded]] : vector<8xf16> to vector<8xi16> + //CHECK: xevm.blockstore %[[ptr]], %[[storeDataI16]] : (!llvm.ptr<3>, vector<8xi16>) + + xegpu.store_matrix %1, %0[%c16, %c48] {subgroup_block_io}: vector<8xf16>, !xegpu.mem_desc<32x64xf16, #xegpu.mem_layout<block = [16, 16]>>, index, index + + gpu.return %1: vector<8xf16> + } + +} diff --git a/mlir/test/Conversion/XeGPUToXeVM/loadstoreprefetch.mlir b/mlir/test/Conversion/XeGPUToXeVM/loadstoreprefetch.mlir index 0b150e9..9c552d8 100644 --- a/mlir/test/Conversion/XeGPUToXeVM/loadstoreprefetch.mlir +++ b/mlir/test/Conversion/XeGPUToXeVM/loadstoreprefetch.mlir @@ -14,19 +14,36 @@ gpu.func @load_gather_i64_src_value_offset(%src: i64, %offset: vector<1xindex>) // CHECK: %[[VAR4:.*]] = arith.addi %[[ARG0]], %[[VAR3]] : i64 // CHECK: %[[VAR5:.*]] = llvm.inttoptr %[[VAR4]] : i64 to !llvm.ptr<1> // CHECK: %[[VAR6:.*]] = scf.if %[[VAR2]] -> (f16) { - // CHECK: %[[VAR7:.*]] = llvm.load %[[VAR5]] {cache_control = #xevm.load_cache_control<L1c_L2uc_L3uc>} : !llvm.ptr<1> -> vector<1xf16> - // CHECK: %[[VAR8:.*]] = vector.extract %[[VAR7]][0] : f16 from vector<1xf16> - // CHECK: scf.yield %[[VAR8]] : f16 - // CHECK: } else { - // CHECK: %[[CST_0:.*]] = arith.constant dense<0.000000e+00> : vector<1xf16> - // CHECK: %[[VAR7:.*]] = vector.extract %[[CST_0]][0] : f16 from vector<1xf16> + // CHECK: %[[VAR7:.*]] = llvm.load %[[VAR5]] {cache_control = #xevm.load_cache_control<L1c_L2uc_L3uc>} : !llvm.ptr<1> -> f16 // CHECK: scf.yield %[[VAR7]] : f16 + // CHECK: } else { + // CHECK: %[[CST_0:.*]] = arith.constant 0.000000e+00 : f16 + // CHECK: scf.yield %[[CST_0]] : f16 // CHECK: } %3 = xegpu.load %src[%offset], %1 <{l1_hint = #xegpu.cache_hint<cached>, l2_hint = #xegpu.cache_hint<uncached>}> : i64, vector<1xindex>, vector<1xi1> -> vector<1xf16> gpu.return } } + +// ----- +gpu.module @test { +// CHECK-LABEL: @source_materialize_single_elem_vec +// CHECK-SAME: %[[ARG0:.*]]: i64, %[[ARG1:.*]]: vector<1xindex>, %[[ARG2:.*]]: memref<1xf16> +gpu.func @source_materialize_single_elem_vec(%src: i64, %offset: vector<1xindex>, %dst: memref<1xf16>) { + %1 = arith.constant dense<1>: vector<1xi1> + %3 = xegpu.load %src[%offset], %1 <{l1_hint = #xegpu.cache_hint<cached>, l2_hint = #xegpu.cache_hint<uncached>}> + : i64, vector<1xindex>, vector<1xi1> -> vector<1xf16> + // CHECK: %[[VAR_IF:.*]] = scf.if + // CHECK: %[[VAR_RET:.*]] = vector.broadcast %[[VAR_IF]] : f16 to vector<1xf16> + // CHECK: %[[C0:.*]] = arith.constant 0 : index + // CHECK: vector.store %[[VAR_RET]], %[[ARG2]][%[[C0]]] : memref<1xf16>, vector<1xf16> + %c0 = arith.constant 0 : index + vector.store %3, %dst[%c0] : memref<1xf16>, vector<1xf16> + gpu.return +} +} + // ----- gpu.module @test { diff --git a/mlir/test/Dialect/Affine/canonicalize.mlir b/mlir/test/Dialect/Affine/canonicalize.mlir index e56079c..1169cd1 100644 --- a/mlir/test/Dialect/Affine/canonicalize.mlir +++ b/mlir/test/Dialect/Affine/canonicalize.mlir @@ -2235,6 +2235,136 @@ func.func @affine_leading_zero_no_outer_bound(%arg0: index, %arg1: index) -> ind // ----- +// CHECK-LABEL: func @delin_apply_cancel_exact +// CHECK-SAME: (%[[ARG0:.+]]: index, %[[ARG1:.+]]: memref<?xindex>) +// CHECK-COUNT-6: memref.store %[[ARG0]], %[[ARG1]][%[[ARG0]]] +// CHECK-NOT: memref.store +// CHECK: return +func.func @delin_apply_cancel_exact(%arg0: index, %arg1: memref<?xindex>) { + %a:3 = affine.delinearize_index %arg0 into (4, 5) : index, index, index + %b:3 = affine.delinearize_index %arg0 into (3, 4, 5) : index, index, index + %c:2 = affine.delinearize_index %arg0 into (20) : index, index + + %t1 = affine.apply affine_map<()[s0, s1, s2] -> (s0 + s1 * 5 + s2 * 20)>()[%a#2, %a#1, %a#0] + memref.store %t1, %arg1[%t1] : memref<?xindex> + + %t2 = affine.apply affine_map<()[s0, s1, s2] -> (s0 + s2 * 20 + s1 * 5)>()[%a#2, %a#1, %a#0] + memref.store %t2, %arg1[%t2] : memref<?xindex> + + %t3 = affine.apply affine_map<()[s0, s1, s2] -> (s1 * 20 + s2 * 5 + s0)>()[%a#2, %a#0, %a#1] + memref.store %t3, %arg1[%t3] : memref<?xindex> + + %t4 = affine.apply affine_map<()[s0, s1, s2] -> (s0 + s1 * 5 + s2 * 20)>()[%b#2, %b#1, %b#0] + memref.store %t4, %arg1[%t4] : memref<?xindex> + + %t5 = affine.apply affine_map<()[s0, s1] -> (s0 + s1 * 20)>()[%c#1, %c#0] + memref.store %t5, %arg1[%t5] : memref<?xindex> + + %t6 = affine.apply affine_map<()[s0, s1] -> (s1 * 20 + s0)>()[%c#1, %c#0] + memref.store %t6, %arg1[%t5] : memref<?xindex> + + return +} + +// ----- + +// CHECK-LABEL: func @delin_apply_cancel_exact_dim +// CHECK: affine.for %[[arg1:.+]] = 0 to 256 +// CHECK: memref.store %[[arg1]] +// CHECK: return +func.func @delin_apply_cancel_exact_dim(%arg0: memref<?xindex>) { + affine.for %arg1 = 0 to 256 { + %a:3 = affine.delinearize_index %arg1 into (2, 2, 64) : index, index, index + %i = affine.apply affine_map<(d0, d1, d2) -> (d0 + d1 * 128 + d2 * 64)>(%a#2, %a#0, %a#1) + memref.store %i, %arg0[%i] : memref<?xindex> + } + return +} + +// ----- + +// CHECK-DAG: #[[$MAP:.+]] = affine_map<()[s0] -> (s0 + 512)> +// CHECK-LABEL: func @delin_apply_cancel_const_term +// CHECK-SAME: (%[[ARG0:.+]]: index, %[[ARG1:.+]]: memref<?xindex>) +// CHECK: affine.apply #[[$MAP]]()[%[[ARG0]]] +// CHECK: return +func.func @delin_apply_cancel_const_term(%arg0: index, %arg1: memref<?xindex>) { + %a:3 = affine.delinearize_index %arg0 into (2, 2, 64) : index, index, index + + %t1 = affine.apply affine_map<()[s0, s1, s2] -> (s0 + s1 * 128 + s2 * 64 + 512)>()[%a#2, %a#0, %a#1] + memref.store %t1, %arg1[%t1] : memref<?xindex> + + return +} + +// ----- + +// CHECK-DAG: #[[$MAP:.+]] = affine_map<()[s0, s1] -> (s0 + s1 + 512)> +// CHECK-LABEL: func @delin_apply_cancel_var_term +// CHECK-SAME: (%[[ARG0:.+]]: index, %[[ARG1:.+]]: memref<?xindex>, %[[ARG2:.+]]: index) +// CHECK: affine.apply #[[$MAP]]()[%[[ARG2]], %[[ARG0]]] +// CHECK: return +func.func @delin_apply_cancel_var_term(%arg0: index, %arg1: memref<?xindex>, %arg2: index) { + %a:3 = affine.delinearize_index %arg0 into (2, 2, 64) : index, index, index + + %t1 = affine.apply affine_map<()[s0, s1, s2, s3] -> (s0 + s1 * 128 + s2 * 64 + s3 + 512)>()[%a#2, %a#0, %a#1, %arg2] + memref.store %t1, %arg1[%t1] : memref<?xindex> + + return +} + +// ----- + +// CHECK-DAG: #[[$MAP:.+]] = affine_map<()[s0] -> (s0 * 2 + s0 ceildiv 4)> +// CHECK-LABEL: func @delin_apply_cancel_nested_exprs +// CHECK-SAME: (%[[ARG0:.+]]: index, %[[ARG1:.+]]: memref<?xindex>) +// CHECK: affine.apply #[[$MAP]]()[%[[ARG0]]] +// CHECK: return +func.func @delin_apply_cancel_nested_exprs(%arg0: index, %arg1: memref<?xindex>) { + %a:2 = affine.delinearize_index %arg0 into (20) : index, index + + %t1 = affine.apply affine_map<()[s0, s1] -> ((s0 + s1 * 20) ceildiv 4 + (s1 * 20 + s0) * 2)>()[%a#1, %a#0] + memref.store %t1, %arg1[%t1] : memref<?xindex> + + return +} + +// ----- + +// CHECK-DAG: #[[$MAP:.+]] = affine_map<()[s0, s1] -> (s0 + s1)> +// CHECK-LABEL: func @delin_apply_cancel_preserve_rotation +// CHECK-SAME: (%[[ARG0:.+]]: index, %[[ARG1:.+]]: memref<?xindex>) +// CHECK: %[[A:.+]]:2 = affine.delinearize_index %[[ARG0]] into (20) +// CHECK: affine.apply #[[$MAP]]()[%[[A]]#1, %[[ARG0]]] +// CHECK: return +func.func @delin_apply_cancel_preserve_rotation(%arg0: index, %arg1: memref<?xindex>) { + %a:2 = affine.delinearize_index %arg0 into (20) : index, index + + %t1 = affine.apply affine_map<()[s0, s1] -> (s0 + s1 * 20 + s0)>()[%a#1, %a#0] + memref.store %t1, %arg1[%t1] : memref<?xindex> + + return +} + +// ----- + +// CHECK-DAG: #[[$MAP:.+]] = affine_map<()[s0, s1] -> (s0 + s1 * 5)> +// CHECK-LABEL: func @delin_apply_dont_cancel_partial +// CHECK-SAME: (%[[ARG0:.+]]: index, %[[ARG1:.+]]: memref<?xindex>) +// CHECK: %[[A:.+]]:3 = affine.delinearize_index %[[ARG0]] into (3, 4, 5) +// CHECK: affine.apply #[[$MAP]]()[%[[A]]#2, %[[A]]#1] +// CHECK: return +func.func @delin_apply_dont_cancel_partial(%arg0: index, %arg1: memref<?xindex>) { + %a:3 = affine.delinearize_index %arg0 into (3, 4, 5) : index, index, index + + %t1 = affine.apply affine_map<()[s0, s1] -> (s0 + s1 * 5)>()[%a#2, %a#1] + memref.store %t1, %arg1[%t1] : memref<?xindex> + + return +} + +// ----- + // CHECK-LABEL: @cst_value_to_cst_attr_basis_delinearize_index // CHECK-SAME: (%[[ARG0:.*]]: index) // CHECK: %[[RET:.*]]:3 = affine.delinearize_index %[[ARG0]] into (3, 4, 2) : index, index diff --git a/mlir/test/Dialect/Bufferization/Transforms/one-shot-non-module-bufferize.mlir b/mlir/test/Dialect/Bufferization/Transforms/one-shot-non-module-bufferize.mlir index e2ab876..b52612d 100644 --- a/mlir/test/Dialect/Bufferization/Transforms/one-shot-non-module-bufferize.mlir +++ b/mlir/test/Dialect/Bufferization/Transforms/one-shot-non-module-bufferize.mlir @@ -24,10 +24,46 @@ // CHECK-NOT: copy // CHECK: %[[call:.*]]:2 = call @inner_func(%[[arg0]]) %0, %1 = call @inner_func(%t0) : (tensor<?xf32>) -> (tensor<?xf32>, f32) - // CHECK: return %[[call]]#1, %[[call]]#0 : f32, memref<?xf32,{{.*}}> + // CHECK: return %[[call]]#1, %[[call]]#0 : f32, memref<?xf32{{.*}}> return %1, %0 : f32, tensor<?xf32> } "test.finish" () : () -> () }) : () -> () +// ----- +#enc1 = #test.tensor_encoding<"hello"> +#enc2 = #test.tensor_encoding<"not hello"> + +"test.symbol_scope_isolated"() ({ + // CHECK: func @inner_func( + // CHECK-SAME: %[[arg0:.*]]: memref<?xf32, #test.memref_layout<"hello">>) + // CHECK-SAME: -> memref<?xf32, #test.memref_layout<"hello">> + func.func @inner_func(%t: tensor<?xf32, #enc1>) + -> tensor<?xf32, #enc1> { + // CHECK: return %[[arg0]] + return %t : tensor<?xf32, #enc1> + } + + // CHECK: func @outer_func( + // CHECK-SAME: %[[arg0:.*]]: memref<?xf32, #test.memref_layout<"hello">>) + // CHECK-SAME: -> (memref<?xf32, #test.memref_layout<"hello">>, + // CHECK-SAME: memref<?xf32, #test.memref_layout<"not hello">>) + func.func @outer_func(%t0: tensor<?xf32, #enc1>) + -> (tensor<?xf32, #enc1>, tensor<?xf32, #enc2>) { + // CHECK: %[[call:.*]] = call @inner_func(%[[arg0]]) + %0 = call @inner_func(%t0) + : (tensor<?xf32, #enc1>) -> (tensor<?xf32, #enc1>) + + // CHECK: %[[local:.*]] = "test.create_memref_op"() : () + // CHECK-SAME: -> memref<?xf32, #test.memref_layout<"not hello">> + %local = "test.create_tensor_op"() : () -> tensor<?xf32, #enc2> + // CHECK: %[[dummy:.*]] = "test.dummy_memref_op"(%[[local]]) + %1 = "test.dummy_tensor_op"(%local) : (tensor<?xf32, #enc2>) + -> tensor<?xf32, #enc2> + + // CHECK: return %[[call]], %[[dummy]] + return %0, %1 : tensor<?xf32, #enc1>, tensor<?xf32, #enc2> + } + "test.finish" () : () -> () +}) : () -> () diff --git a/mlir/test/Dialect/LLVMIR/canonicalize.mlir b/mlir/test/Dialect/LLVMIR/canonicalize.mlir index 8accf6e..755e3a3 100644 --- a/mlir/test/Dialect/LLVMIR/canonicalize.mlir +++ b/mlir/test/Dialect/LLVMIR/canonicalize.mlir @@ -235,6 +235,17 @@ llvm.func @fold_gep_canon(%x : !llvm.ptr) -> !llvm.ptr { // ----- +// CHECK-LABEL: fold_shufflevector +// CHECK-SAME: %[[ARG1:[[:alnum:]]+]]: vector<1xf32>, %[[ARG2:[[:alnum:]]+]]: vector<1xf32> +llvm.func @fold_shufflevector(%v1 : vector<1xf32>, %v2 : vector<1xf32>) -> vector<1xf32> { + // CHECK-NOT: llvm.shufflevector + %c = llvm.shufflevector %v1, %v2 [0] : vector<1xf32> + // CHECK: llvm.return %[[ARG1]] + llvm.return %c : vector<1xf32> +} + +// ----- + // Check that LLVM constants participate in cross-dialect constant folding. The // resulting constant is created in the arith dialect because the last folded // operation belongs to it. diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir index 358bd33..242c04f 100644 --- a/mlir/test/Dialect/LLVMIR/rocdl.mlir +++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir @@ -1035,6 +1035,20 @@ llvm.func @rocdl.s.wait.expcnt() { llvm.return } +llvm.func @rocdl.s.wait.asynccnt() { + // CHECK-LABEL: rocdl.s.wait.asynccnt + // CHECK: rocdl.s.wait.asynccnt 0 + rocdl.s.wait.asynccnt 0 + llvm.return +} + +llvm.func @rocdl.s.wait.tensorcnt() { + // CHECK-LABEL: rocdl.s.wait.tensorcnt + // CHECK: rocdl.s.wait.tensorcnt 0 + rocdl.s.wait.tensorcnt 0 + llvm.return +} + // ----- llvm.func @rocdl.readfirstlane(%src : f32) -> f32 { diff --git a/mlir/test/Dialect/Linalg/vectorization/linalg-ops-with-patterns.mlir b/mlir/test/Dialect/Linalg/vectorization/linalg-ops-with-patterns.mlir index 35f520a..93a0336 100644 --- a/mlir/test/Dialect/Linalg/vectorization/linalg-ops-with-patterns.mlir +++ b/mlir/test/Dialect/Linalg/vectorization/linalg-ops-with-patterns.mlir @@ -1,5 +1,9 @@ // RUN: mlir-opt %s -transform-interpreter -split-input-file | FileCheck %s +///---------------------------------------------------------------------------------------- +/// Tests for linalg.dot +///---------------------------------------------------------------------------------------- + // CHECK-LABEL: contraction_dot func.func @contraction_dot(%A: memref<1584xf32>, %B: memref<1584xf32>, %C: memref<f32>) { @@ -20,6 +24,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.matvec +///---------------------------------------------------------------------------------------- + // CHECK-LABEL: contraction_matvec func.func @contraction_matvec(%A: memref<1584x1584xf32>, %B: memref<1584xf32>, %C: memref<1584xf32>) { @@ -41,6 +49,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.matmul +///---------------------------------------------------------------------------------------- + // CHECK-LABEL: contraction_matmul func.func @contraction_matmul(%A: memref<1584x1584xf32>, %B: memref<1584x1584xf32>, %C: memref<1584x1584xf32>) { // CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<1584x1584x1584xf32> @@ -138,6 +150,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.batch_matmul +///---------------------------------------------------------------------------------------- + // CHECK-LABEL: contraction_batch_matmul func.func @contraction_batch_matmul(%A: memref<1584x1584x1584xf32>, %B: memref<1584x1584x1584xf32>, %C: memref<1584x1584x1584xf32>) { // CHECK: arith.mulf %{{.*}}, %{{.*}} : vector<1584x1584x1584x1584xf32> @@ -159,6 +175,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.cantract +///---------------------------------------------------------------------------------------- + // CHECK-LABEL: @matmul_as_contract // CHECK-SAME: %[[A:.*]]: tensor<24x12xf32> // CHECK-SAME: %[[B:.*]]: tensor<12x25xf32> @@ -220,6 +240,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.fill +///---------------------------------------------------------------------------------------- + // CHECK-LABEL: func @test_vectorize_fill func.func @test_vectorize_fill(%A : memref<8x16xf32>, %arg0 : f32) { // CHECK: %[[V:.*]] = vector.broadcast {{.*}} : f32 to vector<8x16xf32> @@ -259,70 +283,14 @@ module attributes {transform.with_named_sequence} { // ----- -// CHECK-LABEL: func @test_vectorize_copy -func.func @test_vectorize_copy(%A : memref<8x16xf32>, %B : memref<8x16xf32>) { - // CHECK: %[[V:.*]] = vector.transfer_read {{.*}} : memref<8x16xf32>, vector<8x16xf32> - // CHECK: vector.transfer_write %[[V]], {{.*}} : vector<8x16xf32>, memref<8x16xf32> - memref.copy %A, %B : memref<8x16xf32> to memref<8x16xf32> - return -} +///---------------------------------------------------------------------------------------- +/// Tests for linalg.pack +///---------------------------------------------------------------------------------------- -module attributes {transform.with_named_sequence} { - transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) { - %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!transform.any_op) -> !transform.any_op - %1 = transform.get_parent_op %0 {isolated_from_above} : (!transform.any_op) -> !transform.any_op - %2 = transform.structured.vectorize_children_and_apply_patterns %1 : (!transform.any_op) -> !transform.any_op - transform.yield - } -} +// Note, see a similar test in: +// * vectorization.mlir. -// ----- - -// CHECK-LABEL: func @test_vectorize_copy_0d -func.func @test_vectorize_copy_0d(%A : memref<f32>, %B : memref<f32>) { - // CHECK-SAME: (%[[A:.*]]: memref<f32>, %[[B:.*]]: memref<f32>) - // CHECK: %[[V:.*]] = vector.transfer_read %[[A]][]{{.*}} : memref<f32>, vector<f32> - // CHECK: %[[val:.*]] = vector.extract %[[V]][] : f32 from vector<f32> - // CHECK: %[[VV:.*]] = vector.broadcast %[[val]] : f32 to vector<f32> - // CHECK: vector.transfer_write %[[VV]], %[[B]][] : vector<f32>, memref<f32> - memref.copy %A, %B : memref<f32> to memref<f32> - return -} - -module attributes {transform.with_named_sequence} { - transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) { - %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!transform.any_op) -> !transform.any_op - %1 = transform.get_parent_op %0 {isolated_from_above} : (!transform.any_op) -> !transform.any_op - %2 = transform.structured.vectorize_children_and_apply_patterns %1 : (!transform.any_op) -> !transform.any_op - transform.yield - } -} - -// ----- - -// CHECK-LABEL: func @test_vectorize_copy_complex -// CHECK-NOT: vector< -func.func @test_vectorize_copy_complex(%A : memref<8x16xcomplex<f32>>, %B : memref<8x16xcomplex<f32>>) { - memref.copy %A, %B : memref<8x16xcomplex<f32>> to memref<8x16xcomplex<f32>> - return -} - -module attributes {transform.with_named_sequence} { - transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) { - %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!transform.any_op) -> !transform.any_op - %1 = transform.get_parent_op %0 {isolated_from_above} : (!transform.any_op) -> !transform.any_op - %2 = transform.structured.vectorize_children_and_apply_patterns %1 : (!transform.any_op) -> !transform.any_op - transform.yield - } -} - -// ----- - -// Input identical as the test in vectorization.mlir. Output is different - -// vector sizes are inferred (rather than user-specified) and hence _no_ -// masking was used. - -func.func @test_vectorize_pack(%arg0: tensor<32x8x16xf32>, %arg1: tensor<4x1x32x16x2xf32>) -> tensor<4x1x32x16x2xf32> { +func.func @pack_no_padding(%arg0: tensor<32x8x16xf32>, %arg1: tensor<4x1x32x16x2xf32>) -> tensor<4x1x32x16x2xf32> { %pack = linalg.pack %arg0 outer_dims_perm = [1, 2, 0] inner_dims_pos = [2, 1] inner_tiles = [16, 2] into %arg1 : tensor<32x8x16xf32> -> tensor<4x1x32x16x2xf32> return %pack : tensor<4x1x32x16x2xf32> } @@ -336,7 +304,7 @@ module attributes {transform.with_named_sequence} { } } -// CHECK-LABEL: func.func @test_vectorize_pack( +// CHECK-LABEL: func.func @pack_no_padding( // CHECK-SAME: %[[VAL_0:.*]]: tensor<32x8x16xf32>, // CHECK-SAME: %[[VAL_1:.*]]: tensor<4x1x32x16x2xf32>) -> tensor<4x1x32x16x2xf32> { // CHECK-DAG: %[[VAL_2:.*]] = ub.poison : f32 @@ -349,13 +317,16 @@ module attributes {transform.with_named_sequence} { // ----- -func.func @test_vectorize_padded_pack(%arg0: tensor<32x7x15xf32>, %arg1: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { +// Note, see a similar test in: +// * vectorization.mlir. + +func.func @pack_with_padding(%arg0: tensor<32x7x15xf32>, %arg1: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { %pad = arith.constant 0.000000e+00 : f32 %pack = linalg.pack %arg0 padding_value(%pad : f32) inner_dims_pos = [2, 1] inner_tiles = [16, 2] into %arg1 : tensor<32x7x15xf32> -> tensor<32x4x1x16x2xf32> return %pack : tensor<32x4x1x16x2xf32> } -// CHECK-LABEL: func.func @test_vectorize_padded_pack( +// CHECK-LABEL: func.func @pack_with_padding( // CHECK-SAME: %[[VAL_0:.*]]: tensor<32x7x15xf32>, // CHECK-SAME: %[[VAL_1:.*]]: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { // CHECK: %[[VAL_2:.*]] = arith.constant 0.000000e+00 : f32 @@ -377,6 +348,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.map +///---------------------------------------------------------------------------------------- + func.func @vectorize_map(%arg0: memref<64xf32>, %arg1: memref<64xf32>, %arg2: memref<64xf32>) { linalg.map ins(%arg0, %arg1 : memref<64xf32>, memref<64xf32>) @@ -403,6 +378,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.transpose +///---------------------------------------------------------------------------------------- + func.func @vectorize_transpose(%arg0: memref<16x32x64xf32>, %arg1: memref<32x64x16xf32>) { linalg.transpose ins(%arg0 : memref<16x32x64xf32>) @@ -424,6 +403,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.reduce +///---------------------------------------------------------------------------------------- + func.func @vectorize_reduce(%arg0: memref<16x32x64xf32>, %arg1: memref<16x64xf32>) { linalg.reduce ins(%arg0 : memref<16x32x64xf32>) @@ -449,6 +432,10 @@ module attributes {transform.with_named_sequence} { // ----- +///---------------------------------------------------------------------------------------- +/// Tests for linalg.generic +///---------------------------------------------------------------------------------------- + #matmul_trait = { indexing_maps = [ affine_map<(m, n, k) -> (m, k)>, @@ -1446,6 +1433,8 @@ module attributes {transform.with_named_sequence} { // ----- +// TODO: Two Linalg Ops in one tests - either split or document "why". + // CHECK-DAG: #[[$M6:.*]] = affine_map<(d0, d1) -> (d0, 0)> // CHECK-LABEL: func @fused_broadcast_red_2d @@ -1896,3 +1885,65 @@ module attributes {transform.with_named_sequence} { } } +// ----- + +///---------------------------------------------------------------------------------------- +/// Tests for memref.copy +///---------------------------------------------------------------------------------------- + +// CHECK-LABEL: func @test_vectorize_copy +func.func @test_vectorize_copy(%A : memref<8x16xf32>, %B : memref<8x16xf32>) { + // CHECK: %[[V:.*]] = vector.transfer_read {{.*}} : memref<8x16xf32>, vector<8x16xf32> + // CHECK: vector.transfer_write %[[V]], {{.*}} : vector<8x16xf32>, memref<8x16xf32> + memref.copy %A, %B : memref<8x16xf32> to memref<8x16xf32> + return +} + +module attributes {transform.with_named_sequence} { + transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) { + %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.get_parent_op %0 {isolated_from_above} : (!transform.any_op) -> !transform.any_op + %2 = transform.structured.vectorize_children_and_apply_patterns %1 : (!transform.any_op) -> !transform.any_op + transform.yield + } +} + +// ----- + +// CHECK-LABEL: func @test_vectorize_copy_0d +func.func @test_vectorize_copy_0d(%A : memref<f32>, %B : memref<f32>) { + // CHECK-SAME: (%[[A:.*]]: memref<f32>, %[[B:.*]]: memref<f32>) + // CHECK: %[[V:.*]] = vector.transfer_read %[[A]][]{{.*}} : memref<f32>, vector<f32> + // CHECK: %[[val:.*]] = vector.extract %[[V]][] : f32 from vector<f32> + // CHECK: %[[VV:.*]] = vector.broadcast %[[val]] : f32 to vector<f32> + // CHECK: vector.transfer_write %[[VV]], %[[B]][] : vector<f32>, memref<f32> + memref.copy %A, %B : memref<f32> to memref<f32> + return +} + +module attributes {transform.with_named_sequence} { + transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) { + %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.get_parent_op %0 {isolated_from_above} : (!transform.any_op) -> !transform.any_op + %2 = transform.structured.vectorize_children_and_apply_patterns %1 : (!transform.any_op) -> !transform.any_op + transform.yield + } +} + +// ----- + +// CHECK-LABEL: func @test_vectorize_copy_complex +// CHECK-NOT: vector< +func.func @test_vectorize_copy_complex(%A : memref<8x16xcomplex<f32>>, %B : memref<8x16xcomplex<f32>>) { + memref.copy %A, %B : memref<8x16xcomplex<f32>> to memref<8x16xcomplex<f32>> + return +} + +module attributes {transform.with_named_sequence} { + transform.named_sequence @__transform_main(%arg1: !transform.any_op {transform.readonly}) { + %0 = transform.structured.match ops{["memref.copy"]} in %arg1 : (!transform.any_op) -> !transform.any_op + %1 = transform.get_parent_op %0 {isolated_from_above} : (!transform.any_op) -> !transform.any_op + %2 = transform.structured.vectorize_children_and_apply_patterns %1 : (!transform.any_op) -> !transform.any_op + transform.yield + } +} diff --git a/mlir/test/Dialect/Linalg/vectorization/linalg-ops.mlir b/mlir/test/Dialect/Linalg/vectorization/linalg-ops.mlir index 11bea8d..1304a90 100644 --- a/mlir/test/Dialect/Linalg/vectorization/linalg-ops.mlir +++ b/mlir/test/Dialect/Linalg/vectorization/linalg-ops.mlir @@ -1307,14 +1307,17 @@ func.func @test_vectorize_unpack_no_vector_sizes_permute(%source: tensor<4x7x4xf /// Tests for linalg.pack ///---------------------------------------------------------------------------------------- -// Input identical as the test in vectorization-with-patterns.mlir. Output is -// different - vector sizes are inferred (rather than user-specified) and hence -// masking was used. +// This packing requires no padding, so no out-of-bounds read/write vector Ops. -// CHECK-LABEL: func @test_vectorize_pack +// Note, see a similar test in: +// * vectorization-with-patterns.mlir +// The output is identical (the input vector sizes == the inferred vector +// sizes, i.e. the tensor sizes). + +// CHECK-LABEL: func @pack_no_padding // CHECK-SAME: %[[SRC:.*]]: tensor<32x8x16xf32>, // CHECK-SAME: %[[DEST:.*]]: tensor<4x1x32x16x2xf32> -func.func @test_vectorize_pack(%src: tensor<32x8x16xf32>, %dest: tensor<4x1x32x16x2xf32>) -> tensor<4x1x32x16x2xf32> { +func.func @pack_no_padding(%src: tensor<32x8x16xf32>, %dest: tensor<4x1x32x16x2xf32>) -> tensor<4x1x32x16x2xf32> { %pack = linalg.pack %src outer_dims_perm = [1, 2, 0] inner_dims_pos = [2, 1] inner_tiles = [16, 2] into %dest : tensor<32x8x16xf32> -> tensor<4x1x32x16x2xf32> return %pack : tensor<4x1x32x16x2xf32> } @@ -1325,9 +1328,9 @@ func.func @test_vectorize_pack(%src: tensor<32x8x16xf32>, %dest: tensor<4x1x32x1 // CHECK: %[[SC:.*]] = vector.shape_cast %[[READ]] : vector<32x8x16xf32> to vector<32x4x2x1x16xf32> // CHECK: %[[TR:.*]] = vector.transpose %[[SC]], [1, 3, 0, 4, 2] : vector<32x4x2x1x16xf32> to vector<4x1x32x16x2xf32> // CHECK-DAG: %[[C0_1:.*]] = arith.constant 0 : index -// CHECK: %[[write:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] +// CHECK: %[[WRITE:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] // CHECK-SAME: {in_bounds = [true, true, true, true, true]} : vector<4x1x32x16x2xf32>, tensor<4x1x32x16x2xf32> -// CHECK: return %[[write]] : tensor<4x1x32x16x2xf32> +// CHECK: return %[[WRITE]] : tensor<4x1x32x16x2xf32> module attributes {transform.with_named_sequence} { transform.named_sequence @__transform_main(%src: !transform.any_op {transform.readonly}) { @@ -1339,14 +1342,18 @@ module attributes {transform.with_named_sequence} { // ----- -// Input identical as the test in vectorization-with-patterns.mlir. Output is -// different - vector sizes are inferred (rather than user-specified) and hence -// masking was used. +// This packing does require padding, so there are out-of-bounds read/write +// vector Ops. + +// Note, see a similar test in: +// * vectorization-with-patterns.mlir. +// The output is different (the input vector sizes != inferred vector sizes, +// i.e. the tensor sizes). -// CHECK-LABEL: func @test_vectorize_padded_pack +// CHECK-LABEL: func @pack_with_padding // CHECK-SAME: %[[SRC:.*]]: tensor<32x7x15xf32>, // CHECK-SAME: %[[DEST:.*]]: tensor<32x4x1x16x2xf32> -func.func @test_vectorize_padded_pack(%src: tensor<32x7x15xf32>, %dest: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { +func.func @pack_with_padding(%src: tensor<32x7x15xf32>, %dest: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { %pad = arith.constant 0.000000e+00 : f32 %pack = linalg.pack %src padding_value(%pad : f32) inner_dims_pos = [2, 1] inner_tiles = [16, 2] into %dest : tensor<32x7x15xf32> -> tensor<32x4x1x16x2xf32> return %pack : tensor<32x4x1x16x2xf32> @@ -1364,9 +1371,9 @@ func.func @test_vectorize_padded_pack(%src: tensor<32x7x15xf32>, %dest: tensor<3 // CHECK: %[[SC:.*]] = vector.shape_cast %[[READ]] : vector<32x8x16xf32> to vector<32x4x2x1x16xf32> // CHECK: %[[TR:.*]] = vector.transpose %[[SC]], [0, 1, 3, 4, 2] : vector<32x4x2x1x16xf32> to vector<32x4x1x16x2xf32> // CHECK-DAG: %[[C0_1:.*]] = arith.constant 0 : index -// CHECK: %[[write:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] +// CHECK: %[[WRITE:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] // CHECK-SAME: {in_bounds = [true, true, true, true, true]} : vector<32x4x1x16x2xf32>, tensor<32x4x1x16x2xf32> -// CHECK: return %[[write]] : tensor<32x4x1x16x2xf32> +// CHECK: return %[[WRITE]] : tensor<32x4x1x16x2xf32> module attributes {transform.with_named_sequence} { transform.named_sequence @__transform_main(%arg0: !transform.any_op {transform.readonly}) { @@ -1378,10 +1385,46 @@ module attributes {transform.with_named_sequence} { // ----- -// CHECK-LABEL: func @test_vectorize_dynamic_pack +// This packing does require padding, so there are out-of-bounds read/write +// vector Ops. + +// Note, see a similar test in: +// * vectorization-with-patterns.mlir. +// The output is identical (in both cases the vector sizes are inferred). + +// CHECK-LABEL: func @pack_with_padding_no_vector_sizes +// CHECK-SAME: %[[SRC:.*]]: tensor<32x7x15xf32>, +// CHECK-SAME: %[[DEST:.*]]: tensor<32x4x1x16x2xf32> +func.func @pack_with_padding_no_vector_sizes(%src: tensor<32x7x15xf32>, %dest: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { + %pad = arith.constant 0.000000e+00 : f32 + %pack = linalg.pack %src padding_value(%pad : f32) inner_dims_pos = [2, 1] inner_tiles = [16, 2] into %dest : tensor<32x7x15xf32> -> tensor<32x4x1x16x2xf32> + return %pack : tensor<32x4x1x16x2xf32> +} +// CHECK-DAG: %[[CST:.*]] = arith.constant 0.000000e+00 : f32 +// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index +// CHECK: %[[READ:.*]] = vector.transfer_read %{{.*}}[%[[C0]], %[[C0]], %[[C0]]], %[[CST]] +// CHECK-SAME: {in_bounds = [true, false, false]} : tensor<32x7x15xf32>, vector<32x8x16xf32> +// CHECK: %[[SC:.*]] = vector.shape_cast %[[READ]] : vector<32x8x16xf32> to vector<32x4x2x1x16xf32> +// CHECK: %[[TR:.*]] = vector.transpose %[[SC]], [0, 1, 3, 4, 2] : vector<32x4x2x1x16xf32> to vector<32x4x1x16x2xf32> +// CHECK-DAG: %[[C0_1:.*]] = arith.constant 0 : index +// CHECK: %[[WRITE:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] +// CHECK-SAME: {in_bounds = [true, true, true, true, true]} : vector<32x4x1x16x2xf32>, tensor<32x4x1x16x2xf32> +// CHECK: return %[[WRITE]] : tensor<32x4x1x16x2xf32> + +module attributes {transform.with_named_sequence} { + transform.named_sequence @__transform_main(%arg0: !transform.any_op {transform.readonly}) { + %0 = transform.structured.match ops{["linalg.pack"]} in %arg0 : (!transform.any_op) -> !transform.any_op + transform.structured.vectorize %0 : !transform.any_op + transform.yield + } +} + +// ----- + +// CHECK-LABEL: func @pack_with_dynamic_dims // CHECK-SAME: %[[SRC:.*]]: tensor<?x?xf32>, // CHECK-SAME: %[[DEST:.*]]: tensor<?x?x16x2xf32> -func.func @test_vectorize_dynamic_pack(%src: tensor<?x?xf32>, %dest: tensor<?x?x16x2xf32>) -> tensor<?x?x16x2xf32> { +func.func @pack_with_dynamic_dims(%src: tensor<?x?xf32>, %dest: tensor<?x?x16x2xf32>) -> tensor<?x?x16x2xf32> { %pack = linalg.pack %src inner_dims_pos = [1, 0] inner_tiles = [16, 2] into %dest : tensor<?x?xf32> -> tensor<?x?x16x2xf32> return %pack : tensor<?x?x16x2xf32> } @@ -1418,64 +1461,6 @@ module attributes {transform.with_named_sequence} { } } -// ----- - -// CHECK-LABEL: func @test_vectorize_pack_no_vector_sizes -// CHECK-SAME: %[[SRC:.*]]: tensor<64x4xf32>, -// CHECK-SAME: %[[DEST:.*]]: tensor<2x4x16x2xf32> -func.func @test_vectorize_pack_no_vector_sizes(%src: tensor<64x4xf32>, %dest: tensor<2x4x16x2xf32>) -> tensor<2x4x16x2xf32> { - %pack = linalg.pack %src outer_dims_perm = [1, 0] inner_dims_pos = [0, 1] inner_tiles = [16, 2] into %dest : tensor<64x4xf32> -> tensor<2x4x16x2xf32> - return %pack : tensor<2x4x16x2xf32> -} -// CHECK-DAG: %[[CST:.*]] = ub.poison : f32 -// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index -// CHECK: %[[READ:.*]] = vector.transfer_read %{{.*}}[%[[C0]], %[[C0]]], %[[CST]] -// CHECK-SAME: {in_bounds = [true, true]} : tensor<64x4xf32>, vector<64x4xf32> -// CHECK: %[[SC:.*]] = vector.shape_cast %[[READ]] : vector<64x4xf32> to vector<4x16x2x2xf32> -// CHECK: %[[TR:.*]] = vector.transpose %[[SC]], [2, 0, 1, 3] : vector<4x16x2x2xf32> to vector<2x4x16x2xf32> -// CHECK-DAG: %[[C0_1:.*]] = arith.constant 0 : index -// CHECK: %[[WRITE:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] -// CHECK-SAME: {in_bounds = [true, true, true, true]} : vector<2x4x16x2xf32>, tensor<2x4x16x2xf32> -// CHECK: return %[[WRITE]] : tensor<2x4x16x2xf32> - -module attributes {transform.with_named_sequence} { - transform.named_sequence @__transform_main(%arg0: !transform.any_op {transform.readonly}) { - %0 = transform.structured.match ops{["linalg.pack"]} in %arg0 : (!transform.any_op) -> !transform.any_op - transform.structured.vectorize %0 : !transform.any_op - transform.yield - } -} - -// ----- - -// CHECK-LABEL: test_vectorize_padded_pack_no_vector_sizes -// CHECK-SAME: %[[SRC:.*]]: tensor<32x7x15xf32>, -// CHECK-SAME: %[[DEST:.*]]: tensor<32x4x1x16x2xf32> -func.func @test_vectorize_padded_pack_no_vector_sizes(%src: tensor<32x7x15xf32>, %dest: tensor<32x4x1x16x2xf32>) -> tensor<32x4x1x16x2xf32> { - %pad = arith.constant 0.000000e+00 : f32 - %pack = linalg.pack %src padding_value(%pad : f32) inner_dims_pos = [2, 1] inner_tiles = [16, 2] into %dest : tensor<32x7x15xf32> -> tensor<32x4x1x16x2xf32> - return %pack : tensor<32x4x1x16x2xf32> -} -// CHECK-DAG: %[[CST:.*]] = arith.constant 0.000000e+00 : f32 -// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index -// CHECK: %[[READ:.*]] = vector.transfer_read %{{.*}}[%[[C0]], %[[C0]], %[[C0]]], %[[CST]] -// CHECK-SAME: {in_bounds = [true, false, false]} : tensor<32x7x15xf32>, vector<32x8x16xf32> -// CHECK: %[[SC:.*]] = vector.shape_cast %[[READ]] : vector<32x8x16xf32> to vector<32x4x2x1x16xf32> -// CHECK: %[[TR:.*]] = vector.transpose %[[SC]], [0, 1, 3, 4, 2] : vector<32x4x2x1x16xf32> to vector<32x4x1x16x2xf32> -// CHECK-DAG: %[[C0_1:.*]] = arith.constant 0 : index -// CHECK: %[[WRITE:.*]] = vector.transfer_write %[[TR]], %[[DEST]][%[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]], %[[C0_1]]] -// CHECK-SAME: {in_bounds = [true, true, true, true, true]} : vector<32x4x1x16x2xf32>, tensor<32x4x1x16x2xf32> -// CHECK: return %[[WRITE]] : tensor<32x4x1x16x2xf32> - -module attributes {transform.with_named_sequence} { - transform.named_sequence @__transform_main(%arg0: !transform.any_op {transform.readonly}) { - %0 = transform.structured.match ops{["linalg.pack"]} in %arg0 : (!transform.any_op) -> !transform.any_op - transform.structured.vectorize %0 : !transform.any_op - transform.yield - } -} - - ///---------------------------------------------------------------------------------------- /// Tests for other Ops ///---------------------------------------------------------------------------------------- diff --git a/mlir/test/Dialect/MemRef/canonicalize.mlir b/mlir/test/Dialect/MemRef/canonicalize.mlir index 16b7a5c..7160b52 100644 --- a/mlir/test/Dialect/MemRef/canonicalize.mlir +++ b/mlir/test/Dialect/MemRef/canonicalize.mlir @@ -911,6 +911,21 @@ func.func @reinterpret_noop(%arg : memref<2x3x4xf32>) -> memref<2x3x4xf32> { // ----- +// CHECK-LABEL: func @reinterpret_constant_fold +// CHECK-SAME: (%[[ARG:.*]]: memref<f32>) +// CHECK: %[[RES:.*]] = memref.reinterpret_cast %[[ARG]] to offset: [0], sizes: [100, 100], strides: [100, 1] +// CHECK: %[[CAST:.*]] = memref.cast %[[RES]] +// CHECK: return %[[CAST]] +func.func @reinterpret_constant_fold(%arg0: memref<f32>) -> memref<?x?xf32, strided<[?, ?], offset: ?>> { + %c0 = arith.constant 0 : index + %c1 = arith.constant 1 : index + %c100 = arith.constant 100 : index + %reinterpret_cast = memref.reinterpret_cast %arg0 to offset: [%c0], sizes: [%c100, %c100], strides: [%c100, %c1] : memref<f32> to memref<?x?xf32, strided<[?, ?], offset: ?>> + return %reinterpret_cast : memref<?x?xf32, strided<[?, ?], offset: ?>> +} + +// ----- + // CHECK-LABEL: func @reinterpret_of_reinterpret // CHECK-SAME: (%[[ARG:.*]]: memref<?xi8>, %[[SIZE1:.*]]: index, %[[SIZE2:.*]]: index) // CHECK: %[[RES:.*]] = memref.reinterpret_cast %[[ARG]] to offset: [0], sizes: [%[[SIZE2]]], strides: [1] @@ -996,10 +1011,9 @@ func.func @reinterpret_of_extract_strided_metadata_same_type(%arg0 : memref<?x?x // when the strides don't match. // CHECK-LABEL: func @reinterpret_of_extract_strided_metadata_w_different_stride // CHECK-SAME: (%[[ARG:.*]]: memref<8x2xf32>) -// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index -// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index -// CHECK: %[[RES:.*]] = memref.reinterpret_cast %[[ARG]] to offset: [%[[C0]]], sizes: [4, 2, 2], strides: [1, 1, %[[C1]]] -// CHECK: return %[[RES]] +// CHECK: %[[RES:.*]] = memref.reinterpret_cast %[[ARG]] to offset: [0], sizes: [4, 2, 2], strides: [1, 1, 1] +// CHECK: %[[CAST:.*]] = memref.cast %[[RES]] +// CHECK: return %[[CAST]] func.func @reinterpret_of_extract_strided_metadata_w_different_stride(%arg0 : memref<8x2xf32>) -> memref<?x?x?xf32, strided<[?, ?, ?], offset: ?>> { %base, %offset, %sizes:2, %strides:2 = memref.extract_strided_metadata %arg0 : memref<8x2xf32> -> memref<f32>, index, index, index, index, index %m2 = memref.reinterpret_cast %base to offset: [%offset], sizes: [4, 2, 2], strides: [1, 1, %strides#1] : memref<f32> to memref<?x?x?xf32, strided<[?, ?, ?], offset: ?>> @@ -1011,11 +1025,9 @@ func.func @reinterpret_of_extract_strided_metadata_w_different_stride(%arg0 : me // when the offset doesn't match. // CHECK-LABEL: func @reinterpret_of_extract_strided_metadata_w_different_offset // CHECK-SAME: (%[[ARG:.*]]: memref<8x2xf32>) -// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : index -// CHECK-DAG: %[[C2:.*]] = arith.constant 2 : index -// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index -// CHECK: %[[RES:.*]] = memref.reinterpret_cast %[[ARG]] to offset: [1], sizes: [%[[C8]], %[[C2]]], strides: [%[[C2]], %[[C1]]] -// CHECK: return %[[RES]] +// CHECK: %[[RES:.*]] = memref.reinterpret_cast %[[ARG]] to offset: [1], sizes: [8, 2], strides: [2, 1] +// CHECK: %[[CAST:.*]] = memref.cast %[[RES]] +// CHECK: return %[[CAST]] func.func @reinterpret_of_extract_strided_metadata_w_different_offset(%arg0 : memref<8x2xf32>) -> memref<?x?xf32, strided<[?, ?], offset: ?>> { %base, %offset, %sizes:2, %strides:2 = memref.extract_strided_metadata %arg0 : memref<8x2xf32> -> memref<f32>, index, index, index, index, index %m2 = memref.reinterpret_cast %base to offset: [1], sizes: [%sizes#0, %sizes#1], strides: [%strides#0, %strides#1] : memref<f32> to memref<?x?xf32, strided<[?, ?], offset: ?>> diff --git a/mlir/test/Dialect/OpenACC/pointer-like-interface-alloc.mlir b/mlir/test/Dialect/OpenACC/pointer-like-interface-alloc.mlir index 603ace8..3d4bec7 100644 --- a/mlir/test/Dialect/OpenACC/pointer-like-interface-alloc.mlir +++ b/mlir/test/Dialect/OpenACC/pointer-like-interface-alloc.mlir @@ -3,7 +3,7 @@ func.func @test_static_memref_alloc() { %0 = memref.alloca() {test.ptr} : memref<10x20xf32> // CHECK: Successfully generated alloc for operation: %[[ORIG:.*]] = memref.alloca() {test.ptr} : memref<10x20xf32> - // CHECK: Generated: %{{.*}} = memref.alloca() : memref<10x20xf32> + // CHECK: Generated: %{{.*}} = memref.alloca() {acc.var_name = #acc.var_name<"test_alloc">} : memref<10x20xf32> return } @@ -19,6 +19,6 @@ func.func @test_dynamic_memref_alloc() { // CHECK: Generated: %[[DIM0:.*]] = memref.dim %[[ORIG]], %[[C0]] : memref<?x?xf32> // CHECK: Generated: %[[C1:.*]] = arith.constant 1 : index // CHECK: Generated: %[[DIM1:.*]] = memref.dim %[[ORIG]], %[[C1]] : memref<?x?xf32> - // CHECK: Generated: %{{.*}} = memref.alloc(%[[DIM0]], %[[DIM1]]) : memref<?x?xf32> + // CHECK: Generated: %{{.*}} = memref.alloc(%[[DIM0]], %[[DIM1]]) {acc.var_name = #acc.var_name<"test_alloc">} : memref<?x?xf32> return } diff --git a/mlir/test/Dialect/OpenACC/recipe-populate-firstprivate.mlir b/mlir/test/Dialect/OpenACC/recipe-populate-firstprivate.mlir index 35355c6..8846c9e 100644 --- a/mlir/test/Dialect/OpenACC/recipe-populate-firstprivate.mlir +++ b/mlir/test/Dialect/OpenACC/recipe-populate-firstprivate.mlir @@ -2,7 +2,7 @@ // CHECK: acc.firstprivate.recipe @firstprivate_scalar : memref<f32> init { // CHECK: ^bb0(%{{.*}}: memref<f32>): -// CHECK: %[[ALLOC:.*]] = memref.alloca() : memref<f32> +// CHECK: %[[ALLOC:.*]] = memref.alloca() {acc.var_name = #acc.var_name<"scalar">} : memref<f32> // CHECK: acc.yield %[[ALLOC]] : memref<f32> // CHECK: } copy { // CHECK: ^bb0(%[[SRC:.*]]: memref<f32>, %[[DST:.*]]: memref<f32>): @@ -20,7 +20,7 @@ func.func @test_scalar() { // CHECK: acc.firstprivate.recipe @firstprivate_static_2d : memref<10x20xf32> init { // CHECK: ^bb0(%{{.*}}: memref<10x20xf32>): -// CHECK: %[[ALLOC:.*]] = memref.alloca() : memref<10x20xf32> +// CHECK: %[[ALLOC:.*]] = memref.alloca() {acc.var_name = #acc.var_name<"static_2d">} : memref<10x20xf32> // CHECK: acc.yield %[[ALLOC]] : memref<10x20xf32> // CHECK: } copy { // CHECK: ^bb0(%[[SRC:.*]]: memref<10x20xf32>, %[[DST:.*]]: memref<10x20xf32>): @@ -42,7 +42,7 @@ func.func @test_static_2d() { // CHECK: %[[DIM0:.*]] = memref.dim %[[ARG]], %[[C0]] : memref<?x?xf32> // CHECK: %[[C1:.*]] = arith.constant 1 : index // CHECK: %[[DIM1:.*]] = memref.dim %[[ARG]], %[[C1]] : memref<?x?xf32> -// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM0]], %[[DIM1]]) : memref<?x?xf32> +// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM0]], %[[DIM1]]) {acc.var_name = #acc.var_name<"dynamic_2d">} : memref<?x?xf32> // CHECK: acc.yield %[[ALLOC]] : memref<?x?xf32> // CHECK: } copy { // CHECK: ^bb0(%[[SRC:.*]]: memref<?x?xf32>, %[[DST:.*]]: memref<?x?xf32>): @@ -65,7 +65,7 @@ func.func @test_dynamic_2d(%arg0: index, %arg1: index) { // CHECK: ^bb0(%[[ARG:.*]]: memref<10x?xf32>): // CHECK: %[[C1:.*]] = arith.constant 1 : index // CHECK: %[[DIM1:.*]] = memref.dim %[[ARG]], %[[C1]] : memref<10x?xf32> -// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM1]]) : memref<10x?xf32> +// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM1]]) {acc.var_name = #acc.var_name<"mixed_dims">} : memref<10x?xf32> // CHECK: acc.yield %[[ALLOC]] : memref<10x?xf32> // CHECK: } copy { // CHECK: ^bb0(%[[SRC:.*]]: memref<10x?xf32>, %[[DST:.*]]: memref<10x?xf32>): @@ -86,7 +86,7 @@ func.func @test_mixed_dims(%arg0: index) { // CHECK: acc.firstprivate.recipe @firstprivate_scalar_int : memref<i32> init { // CHECK: ^bb0(%{{.*}}: memref<i32>): -// CHECK: %[[ALLOC:.*]] = memref.alloca() : memref<i32> +// CHECK: %[[ALLOC:.*]] = memref.alloca() {acc.var_name = #acc.var_name<"scalar_int">} : memref<i32> // CHECK: acc.yield %[[ALLOC]] : memref<i32> // CHECK: } copy { // CHECK: ^bb0(%[[SRC:.*]]: memref<i32>, %[[DST:.*]]: memref<i32>): diff --git a/mlir/test/Dialect/OpenACC/recipe-populate-private.mlir b/mlir/test/Dialect/OpenACC/recipe-populate-private.mlir index 8403ee8..3d5a918 100644 --- a/mlir/test/Dialect/OpenACC/recipe-populate-private.mlir +++ b/mlir/test/Dialect/OpenACC/recipe-populate-private.mlir @@ -2,7 +2,7 @@ // CHECK: acc.private.recipe @private_scalar : memref<f32> init { // CHECK: ^bb0(%{{.*}}: memref<f32>): -// CHECK: %[[ALLOC:.*]] = memref.alloca() : memref<f32> +// CHECK: %[[ALLOC:.*]] = memref.alloca() {acc.var_name = #acc.var_name<"scalar">} : memref<f32> // CHECK: acc.yield %[[ALLOC]] : memref<f32> // CHECK: } // CHECK-NOT: destroy @@ -16,7 +16,7 @@ func.func @test_scalar() { // CHECK: acc.private.recipe @private_static_2d : memref<10x20xf32> init { // CHECK: ^bb0(%{{.*}}: memref<10x20xf32>): -// CHECK: %[[ALLOC:.*]] = memref.alloca() : memref<10x20xf32> +// CHECK: %[[ALLOC:.*]] = memref.alloca() {acc.var_name = #acc.var_name<"static_2d">} : memref<10x20xf32> // CHECK: acc.yield %[[ALLOC]] : memref<10x20xf32> // CHECK: } // CHECK-NOT: destroy @@ -34,7 +34,7 @@ func.func @test_static_2d() { // CHECK: %[[DIM0:.*]] = memref.dim %[[ARG]], %[[C0]] : memref<?x?xf32> // CHECK: %[[C1:.*]] = arith.constant 1 : index // CHECK: %[[DIM1:.*]] = memref.dim %[[ARG]], %[[C1]] : memref<?x?xf32> -// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM0]], %[[DIM1]]) : memref<?x?xf32> +// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM0]], %[[DIM1]]) {acc.var_name = #acc.var_name<"dynamic_2d">} : memref<?x?xf32> // CHECK: acc.yield %[[ALLOC]] : memref<?x?xf32> // CHECK: } destroy { // CHECK: ^bb0(%{{.*}}: memref<?x?xf32>, %[[VAL:.*]]: memref<?x?xf32>): @@ -53,7 +53,7 @@ func.func @test_dynamic_2d(%arg0: index, %arg1: index) { // CHECK: ^bb0(%[[ARG:.*]]: memref<10x?xf32>): // CHECK: %[[C1:.*]] = arith.constant 1 : index // CHECK: %[[DIM1:.*]] = memref.dim %[[ARG]], %[[C1]] : memref<10x?xf32> -// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM1]]) : memref<10x?xf32> +// CHECK: %[[ALLOC:.*]] = memref.alloc(%[[DIM1]]) {acc.var_name = #acc.var_name<"mixed_dims">} : memref<10x?xf32> // CHECK: acc.yield %[[ALLOC]] : memref<10x?xf32> // CHECK: } destroy { // CHECK: ^bb0(%{{.*}}: memref<10x?xf32>, %[[VAL:.*]]: memref<10x?xf32>): @@ -70,7 +70,7 @@ func.func @test_mixed_dims(%arg0: index) { // CHECK: acc.private.recipe @private_scalar_int : memref<i32> init { // CHECK: ^bb0(%{{.*}}: memref<i32>): -// CHECK: %[[ALLOC:.*]] = memref.alloca() : memref<i32> +// CHECK: %[[ALLOC:.*]] = memref.alloca() {acc.var_name = #acc.var_name<"scalar_int">} : memref<i32> // CHECK: acc.yield %[[ALLOC]] : memref<i32> // CHECK: } // CHECK-NOT: destroy diff --git a/mlir/test/Dialect/Tensor/one-shot-bufferize.mlir b/mlir/test/Dialect/Tensor/one-shot-bufferize.mlir index b6c72be..f66cf7a 100644 --- a/mlir/test/Dialect/Tensor/one-shot-bufferize.mlir +++ b/mlir/test/Dialect/Tensor/one-shot-bufferize.mlir @@ -490,3 +490,32 @@ func.func @collapse_shape_regression( tensor.collapse_shape %0[[0, 1]] : tensor<5x6xf32> into tensor<30xf32> return } + +// ----- + +// CHECK-LABEL: func private @mult_return_callee( +// CHECK-SAME: %[[T:.*]]: memref<?xf32, strided<[?], offset: ?>>, %[[COND:.*]]: i1, +// CHECK-SAME: %[[A:.*]]: index, %[[B:.*]]: index) -> index { +// CHECK: cf.cond_br %[[COND]], ^bb1, ^bb2 +// CHECK: ^bb1: +// CHECK: return %[[A]] : index +// CHECK: ^bb2: +// CHECK: return %[[B]] : index +func.func private @mult_return_callee(%t: tensor<?xf32>, %cond:i1, %a: index, %b: index) -> (tensor<10xf32>, index) { + %casted = tensor.cast %t : tensor<?xf32> to tensor<10xf32> + cf.cond_br %cond,^a, ^b +^a: + return %casted, %a : tensor<10xf32>, index +^b: + return %casted, %b : tensor<10xf32>, index +} + +// CHECK-LABEL: func @mult_return( +// CHECK-SAME: %[[T:.*]]: memref<?xf32, strided<[?], offset: ?>>, %[[COND:.*]]: i1, +// CHECK-SAME: %[[A:.*]]: index, %[[B:.*]]: index) -> (memref<?xf32, strided<[?], offset: ?>>, index) { +func.func @mult_return(%t: tensor<?xf32>, %cond:i1, %a: index, %b: index) -> (tensor<10xf32>, index) { + // CHECK: %[[RET:.*]] = call @mult_return_callee(%[[T]], %[[COND]], %[[A]], %[[B]]) : (memref<?xf32, strided<[?], offset: ?>>, i1, index, index) -> index + // CHECK: return %[[T]], %[[RET]] : memref<?xf32, strided<[?], offset: ?>>, index + %t_res, %v = func.call @mult_return_callee(%t, %cond, %a, %b) : (tensor<?xf32>, i1, index, index) -> (tensor<10xf32>, index) + return %t_res, %v : tensor<10xf32>, index +} diff --git a/mlir/test/Dialect/Tosa/tosa-attach-target.mlir b/mlir/test/Dialect/Tosa/tosa-attach-target.mlir index d6c886c..a0c59c0 100644 --- a/mlir/test/Dialect/Tosa/tosa-attach-target.mlir +++ b/mlir/test/Dialect/Tosa/tosa-attach-target.mlir @@ -1,12 +1,14 @@ // RUN: mlir-opt %s -split-input-file -tosa-attach-target="profiles=pro_int,pro_fp extensions=int16,int4,bf16,fp8e4m3,fp8e5m2,fft,variable,controlflow,doubleround,inexactround,dynamic level=none" | FileCheck %s --check-prefix=CHECK-ALL // RUN: mlir-opt %s -split-input-file -tosa-attach-target="level=8k" | FileCheck %s --check-prefix=CHECK-LVL-8K // RUN: mlir-opt %s -split-input-file -tosa-attach-target | FileCheck %s --check-prefix=CHECK-DEFAULT +// RUN: mlir-opt %s -split-input-file -tosa-attach-target="specification_version=1.1.draft" | FileCheck %s --check-prefix=CHECK-VERSION-1P1 // ----- -// CHECK-ALL: module attributes {tosa.target_env = #tosa.target_env<level = none, profiles = [pro_int, pro_fp], extensions = [int16, int4, bf16, fp8e4m3, fp8e5m2, fft, variable, controlflow, doubleround, inexactround, dynamic]>} -// CHECK-LVL-8K: module attributes {tosa.target_env = #tosa.target_env<level = "8k", profiles = [], extensions = []>} -// CHECK-DEFAULT: module attributes {tosa.target_env = #tosa.target_env<level = "8k", profiles = [], extensions = []>} +// CHECK-ALL: module attributes {tosa.target_env = #tosa.target_env<specification_version = "1.0", level = none, profiles = [pro_int, pro_fp], extensions = [int16, int4, bf16, fp8e4m3, fp8e5m2, fft, variable, controlflow, doubleround, inexactround, dynamic]>} +// CHECK-LVL-8K: module attributes {tosa.target_env = #tosa.target_env<specification_version = "1.0", level = "8k", profiles = [], extensions = []>} +// CHECK-DEFAULT: module attributes {tosa.target_env = #tosa.target_env<specification_version = "1.0", level = "8k", profiles = [], extensions = []>} +// CHECK-VERSION-1P1: module attributes {tosa.target_env = #tosa.target_env<specification_version = "1.1.draft", level = "8k", profiles = [], extensions = []>} // CHECK-LABEL: test_simple func.func @test_simple(%arg0 : tensor<1x1x1x1xf32>, %arg1 : tensor<1x1x1x1xf32>) -> tensor<1x1x1x1xf32> { %1 = tosa.add %arg0, %arg1 : (tensor<1x1x1x1xf32>, tensor<1x1x1x1xf32>) -> tensor<1x1x1x1xf32> diff --git a/mlir/test/Dialect/Tosa/tosa-validation-version-1p0-invalid.mlir b/mlir/test/Dialect/Tosa/tosa-validation-version-1p0-invalid.mlir new file mode 100644 index 0000000..51089df --- /dev/null +++ b/mlir/test/Dialect/Tosa/tosa-validation-version-1p0-invalid.mlir @@ -0,0 +1,21 @@ +// RUN: mlir-opt %s -split-input-file -verify-diagnostics -tosa-attach-target="specification_version=1.0 profiles=pro_int,pro_fp extensions=int16,int4,bf16,fp8e4m3,fp8e5m2,fft,variable,controlflow,dynamic,doubleround,inexactround" -tosa-validate="strict-op-spec-alignment" + +// ----- + +func.func @test_matmul_fp8_mixed_precision_operands(%arg0: tensor<1x14x19xf8E4M3FN>, %arg1: tensor<1x19x28xf8E5M2>) -> tensor<1x14x28xf16> { + %azp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E4M3FN>}> : () -> tensor<1xf8E4M3FN> + %bzp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E5M2>}> : () -> tensor<1xf8E5M2> + // expected-error@+1 {{'tosa.matmul' op illegal: the target specification version (1.0) is not backwards compatible with the op compliance specification version (1.1)}} + %0 = tosa.matmul %arg0, %arg1, %azp0, %bzp0 : (tensor<1x14x19xf8E4M3FN>, tensor<1x19x28xf8E5M2>, tensor<1xf8E4M3FN>, tensor<1xf8E5M2>) -> tensor<1x14x28xf16> + return %0 : tensor<1x14x28xf16> +} + +// ----- + +func.func @test_matmul_fp8_input_fp32_acc_type(%arg0: tensor<1x14x19xf8E4M3FN>, %arg1: tensor<1x19x28xf8E4M3FN>) -> tensor<1x14x28xf32> { + %azp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E4M3FN>}> : () -> tensor<1xf8E4M3FN> + %bzp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E4M3FN>}> : () -> tensor<1xf8E4M3FN> + // expected-error@+1 {{'tosa.matmul' op illegal: the target specification version (1.0) is not backwards compatible with the op compliance specification version (1.1)}} + %0 = tosa.matmul %arg0, %arg1, %azp0, %bzp0 : (tensor<1x14x19xf8E4M3FN>, tensor<1x19x28xf8E4M3FN>, tensor<1xf8E4M3FN>, tensor<1xf8E4M3FN>) -> tensor<1x14x28xf32> + return %0 : tensor<1x14x28xf32> +} diff --git a/mlir/test/Dialect/Tosa/tosa-validation-version-1p1-valid.mlir b/mlir/test/Dialect/Tosa/tosa-validation-version-1p1-valid.mlir new file mode 100644 index 0000000..8164509 --- /dev/null +++ b/mlir/test/Dialect/Tosa/tosa-validation-version-1p1-valid.mlir @@ -0,0 +1,20 @@ +// RUN: mlir-opt %s -split-input-file -verify-diagnostics -tosa-attach-target="specification_version=1.1.draft profiles=pro_int,pro_fp extensions=int16,int4,bf16,fp8e4m3,fp8e5m2,fft,variable,controlflow,doubleround,inexactround" -tosa-validate="strict-op-spec-alignment" | FileCheck %s + +// ----- + +func.func @test_matmul_fp8_mixed_precision_operands(%arg0: tensor<1x14x19xf8E4M3FN>, %arg1: tensor<1x19x28xf8E5M2>) -> tensor<1x14x28xf16> { + %azp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E4M3FN>}> : () -> tensor<1xf8E4M3FN> + %bzp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E5M2>}> : () -> tensor<1xf8E5M2> + %0 = tosa.matmul %arg0, %arg1, %azp0, %bzp0 : (tensor<1x14x19xf8E4M3FN>, tensor<1x19x28xf8E5M2>, tensor<1xf8E4M3FN>, tensor<1xf8E5M2>) -> tensor<1x14x28xf16> + return %0 : tensor<1x14x28xf16> +} + +// ----- + +// CHECK-LABEL: test_matmul_fp8_input_fp32_acc_type +func.func @test_matmul_fp8_input_fp32_acc_type(%arg0: tensor<1x14x19xf8E4M3FN>, %arg1: tensor<1x19x28xf8E4M3FN>) -> tensor<1x14x28xf32> { + %azp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E4M3FN>}> : () -> tensor<1xf8E4M3FN> + %bzp0 = "tosa.const"() <{values = dense<0.0> : tensor<1xf8E4M3FN>}> : () -> tensor<1xf8E4M3FN> + %0 = tosa.matmul %arg0, %arg1, %azp0, %bzp0 : (tensor<1x14x19xf8E4M3FN>, tensor<1x19x28xf8E4M3FN>, tensor<1xf8E4M3FN>, tensor<1xf8E4M3FN>) -> tensor<1x14x28xf32> + return %0 : tensor<1x14x28xf32> +} diff --git a/mlir/test/Dialect/Vector/canonicalize/vector-step.mlir b/mlir/test/Dialect/Vector/canonicalize/vector-step.mlir new file mode 100644 index 0000000..023a0e5 --- /dev/null +++ b/mlir/test/Dialect/Vector/canonicalize/vector-step.mlir @@ -0,0 +1,311 @@ +// RUN: mlir-opt %s -canonicalize="test-convergence" -split-input-file | FileCheck %s + +///===----------------------------------------------===// +/// Tests of `StepCompareFolder` +///===----------------------------------------------===// + + +///===------------------------------------===// +/// Tests of `ugt` (unsigned greater than) +///===------------------------------------===// + +// CHECK-LABEL: @ugt_constant_3_lhs +// CHECK: %[[CST:.*]] = arith.constant dense<true> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ugt_constant_3_lhs() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // 3 > [0, 1, 2] => [true, true, true] => true for all indices => fold + %1 = arith.cmpi ugt, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ugt_constant_2_lhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ugt_constant_2_lhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // 2 > [0, 1, 2] => [true, true, false] => not same for all indices => don't fold + %1 = arith.cmpi ugt, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @ugt_constant_3_rhs +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ugt_constant_3_rhs() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] > 3 => [false, false, false] => false for all indices => fold + %1 = arith.cmpi ugt, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @ugt_constant_max_rhs +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ugt_constant_max_rhs() -> vector<3xi1> { + // The largest i64 possible: + %cst = arith.constant dense<0x7fffffffffffffff> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ugt, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + + +// ----- + +// CHECK-LABEL: @ugt_constant_2_rhs +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ugt_constant_2_rhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] > 2 => [false, false, false] => false for all indices => fold + %1 = arith.cmpi ugt, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ugt_constant_1_rhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ugt_constant_1_rhs() -> vector<3xi1> { + %cst = arith.constant dense<1> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] > 1 => [false, false, true] => not same for all indices => don't fold + %1 = arith.cmpi ugt, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +///===------------------------------------===// +/// Tests of `uge` (unsigned greater than or equal) +///===------------------------------------===// + + +// CHECK-LABEL: @uge_constant_2_lhs +// CHECK: %[[CST:.*]] = arith.constant dense<true> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @uge_constant_2_lhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // 2 >= [0, 1, 2] => [true, true, true] => true for all indices => fold + %1 = arith.cmpi uge, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_uge_constant_1_lhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_uge_constant_1_lhs() -> vector<3xi1> { + %cst = arith.constant dense<1> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // 1 >= [0, 1, 2] => [true, false, false] => not same for all indices => don't fold + %1 = arith.cmpi uge, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @uge_constant_3_rhs +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @uge_constant_3_rhs() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] >= 3 => [false, false, false] => false for all indices => fold + %1 = arith.cmpi uge, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_uge_constant_2_rhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_uge_constant_2_rhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] >= 2 => [false, false, true] => not same for all indices => don't fold + %1 = arith.cmpi uge, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + + +///===------------------------------------===// +/// Tests of `ult` (unsigned less than) +///===------------------------------------===// + + +// CHECK-LABEL: @ult_constant_2_lhs +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ult_constant_2_lhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // 2 < [0, 1, 2] => [false, false, false] => false for all indices => fold + %1 = arith.cmpi ult, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ult_constant_1_lhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ult_constant_1_lhs() -> vector<3xi1> { + %cst = arith.constant dense<1> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // 1 < [0, 1, 2] => [false, false, true] => not same for all indices => don't fold + %1 = arith.cmpi ult, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @ult_constant_3_rhs +// CHECK: %[[CST:.*]] = arith.constant dense<true> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ult_constant_3_rhs() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] < 3 => [true, true, true] => true for all indices => fold + %1 = arith.cmpi ult, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ult_constant_2_rhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ult_constant_2_rhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + // [0, 1, 2] < 2 => [true, true, false] => not same for all indices => don't fold + %1 = arith.cmpi ult, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +///===------------------------------------===// +/// Tests of `ule` (unsigned less than or equal) +///===------------------------------------===// + +// CHECK-LABEL: @ule_constant_3_lhs +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ule_constant_3_lhs() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ule, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ule_constant_2_lhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ule_constant_2_lhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ule, %cst, %0 : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @ule_constant_2_rhs +// CHECK: %[[CST:.*]] = arith.constant dense<true> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ule_constant_2_rhs() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ule, %0, %cst : vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ule_constant_1_rhs +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ule_constant_1_rhs() -> vector<3xi1> { + %cst = arith.constant dense<1> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ule, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +///===------------------------------------===// +/// Tests of `eq` (equal) +///===------------------------------------===// + +// CHECK-LABEL: @eq_constant_3 +// CHECK: %[[CST:.*]] = arith.constant dense<false> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @eq_constant_3() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi eq, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_eq_constant_2 +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_eq_constant_2() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi eq, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +///===------------------------------------===// +/// Tests of `ne` (not equal) +///===------------------------------------===// + +// CHECK-LABEL: @ne_constant_3 +// CHECK: %[[CST:.*]] = arith.constant dense<true> : vector<3xi1> +// CHECK: return %[[CST]] : vector<3xi1> +func.func @ne_constant_3() -> vector<3xi1> { + %cst = arith.constant dense<3> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ne, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + +// ----- + +// CHECK-LABEL: @negative_ne_constant_2 +// CHECK: %[[CMP:.*]] = arith.cmpi +// CHECK: return %[[CMP]] +func.func @negative_ne_constant_2() -> vector<3xi1> { + %cst = arith.constant dense<2> : vector<3xindex> + %0 = vector.step : vector<3xindex> + %1 = arith.cmpi ne, %0, %cst: vector<3xindex> + return %1 : vector<3xi1> +} + diff --git a/mlir/test/Dialect/Vector/vector-unroll-options.mlir b/mlir/test/Dialect/Vector/vector-unroll-options.mlir index 35db14e..e5a98b5 100644 --- a/mlir/test/Dialect/Vector/vector-unroll-options.mlir +++ b/mlir/test/Dialect/Vector/vector-unroll-options.mlir @@ -188,15 +188,38 @@ func.func @vector_fma(%a: vector<4x4xf32>, %b: vector<4x4xf32>, %c: vector<4x4xf // CHECK-LABEL: func @vector_fma // CHECK-COUNT-4: vector.fma %{{.+}}, %{{.+}}, %{{.+}} : vector<2x2xf32> -// TODO: We should be able to unroll this like the example above - this will require extending UnrollElementwisePattern. -func.func @negative_vector_fma_3d(%a: vector<3x2x2xf32>) -> vector<3x2x2xf32>{ +func.func @vector_fma_3d(%a: vector<3x2x2xf32>) -> vector<3x2x2xf32>{ %0 = vector.fma %a, %a, %a : vector<3x2x2xf32> return %0 : vector<3x2x2xf32> } -// CHECK-LABEL: func @negative_vector_fma_3d -// CHECK-NOT: vector.extract_strided_slice -// CHECK: %[[R0:.*]] = vector.fma %{{.+}} : vector<3x2x2xf32> -// CHECK: return +// CHECK-LABEL: func @vector_fma_3d +// CHECK-SAME: (%[[SRC:.*]]: vector<3x2x2xf32>) -> vector<3x2x2xf32> { +// CHECK: %[[CST:.*]] = arith.constant dense<0.000000e+00> : vector<3x2x2xf32> +// CHECK: %[[E_LHS_0:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [0, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_LHS_0:.*]] = vector.shape_cast %[[E_LHS_0]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_RHS_0:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [0, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_RHS_0:.*]] = vector.shape_cast %[[E_RHS_0]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_OUT_0:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [0, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_OUT_0:.*]] = vector.shape_cast %[[E_OUT_0]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[FMA0:.*]] = vector.fma %[[S_LHS_0]], %[[S_RHS_0]], %[[S_OUT_0]] : vector<2x2xf32> +// CHECK: %[[I0:.*]] = vector.insert_strided_slice %[[FMA0]], %[[CST]] {offsets = [0, 0, 0], strides = [1, 1]} : vector<2x2xf32> into vector<3x2x2xf32> +// CHECK: %[[E_LHS_1:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [1, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_LHS_1:.*]] = vector.shape_cast %[[E_LHS_1]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_RHS_1:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [1, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_RHS_1:.*]] = vector.shape_cast %[[E_RHS_1]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_OUT_1:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [1, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_OUT_1:.*]] = vector.shape_cast %[[E_OUT_1]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[FMA1:.*]] = vector.fma %[[S_LHS_1]], %[[S_RHS_1]], %[[S_OUT_1]] : vector<2x2xf32> +// CHECK: %[[I1:.*]] = vector.insert_strided_slice %[[FMA1]], %[[I0]] {offsets = [1, 0, 0], strides = [1, 1]} : vector<2x2xf32> into vector<3x2x2xf32> +// CHECK: %[[E_LHS_2:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [2, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_LHS_2:.*]] = vector.shape_cast %[[E_LHS_2]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_RHS_2:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [2, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_RHS_2:.*]] = vector.shape_cast %[[E_RHS_2]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_OUT_2:.*]] = vector.extract_strided_slice %[[SRC]] {offsets = [2, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<3x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_OUT_2:.*]] = vector.shape_cast %[[E_OUT_2]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[FMA2:.*]] = vector.fma %[[S_LHS_2]], %[[S_RHS_2]], %[[S_OUT_2]] : vector<2x2xf32> +// CHECK: %[[I2:.*]] = vector.insert_strided_slice %[[FMA2]], %[[I1]] {offsets = [2, 0, 0], strides = [1, 1]} : vector<2x2xf32> into vector<3x2x2xf32> +// CHECK: return %[[I2]] : vector<3x2x2xf32> func.func @vector_multi_reduction(%v : vector<4x6xf32>, %acc: vector<4xf32>) -> vector<4xf32> { %0 = vector.multi_reduction #vector.kind<add>, %v, %acc [1] : vector<4x6xf32> to vector<4xf32> @@ -440,3 +463,36 @@ func.func @vector_step() -> vector<32xindex> { // CHECK: %[[ADD3:.*]] = arith.addi %[[STEP]], %[[CST]] : vector<8xindex> // CHECK: %[[INS3:.*]] = vector.insert_strided_slice %[[ADD3]], %[[INS2]] {offsets = [24], strides = [1]} : vector<8xindex> into vector<32xindex> // CHECK: return %[[INS3]] : vector<32xindex> + + +func.func @elementwise_3D_to_2D(%v1: vector<2x2x2xf32>, %v2: vector<2x2x2xf32>) -> vector<2x2x2xf32> { + %0 = arith.addf %v1, %v2 : vector<2x2x2xf32> + return %0 : vector<2x2x2xf32> +} +// CHECK-LABEL: func @elementwise_3D_to_2D +// CHECK-SAME: (%[[ARG0:.*]]: vector<2x2x2xf32>, %[[ARG1:.*]]: vector<2x2x2xf32>) -> vector<2x2x2xf32> { +// CHECK: %[[CST:.*]] = arith.constant dense<0.000000e+00> : vector<2x2x2xf32> +// CHECK: %[[E_LHS_0:.*]] = vector.extract_strided_slice %[[ARG0]] {offsets = [0, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<2x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_LHS_0:.*]] = vector.shape_cast %[[E_LHS_0]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_RHS_0:.*]] = vector.extract_strided_slice %[[ARG1]] {offsets = [0, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<2x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_RHS_0:.*]] = vector.shape_cast %[[E_RHS_0]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[ADD0:.*]] = arith.addf %[[S_LHS_0]], %[[S_RHS_0]] : vector<2x2xf32> +// CHECK: %[[I0:.*]] = vector.insert_strided_slice %[[ADD0]], %[[CST]] {offsets = [0, 0, 0], strides = [1, 1]} : vector<2x2xf32> into vector<2x2x2xf32> +// CHECK: %[[E_LHS_1:.*]] = vector.extract_strided_slice %[[ARG0]] {offsets = [1, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<2x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_LHS_1:.*]] = vector.shape_cast %[[E_LHS_1]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[E_RHS_1:.*]] = vector.extract_strided_slice %[[ARG1]] {offsets = [1, 0, 0], sizes = [1, 2, 2], strides = [1, 1, 1]} : vector<2x2x2xf32> to vector<1x2x2xf32> +// CHECK: %[[S_RHS_1:.*]] = vector.shape_cast %[[E_RHS_1]] : vector<1x2x2xf32> to vector<2x2xf32> +// CHECK: %[[ADD1:.*]] = arith.addf %[[S_LHS_1]], %[[S_RHS_1]] : vector<2x2xf32> +// CHECK: %[[I1:.*]] = vector.insert_strided_slice %[[ADD1]], %[[I0]] {offsets = [1, 0, 0], strides = [1, 1]} : vector<2x2xf32> into vector<2x2x2xf32> +// CHECK: return %[[I1]] : vector<2x2x2xf32> + + +func.func @elementwise_4D_to_2D(%v1: vector<2x2x2x2xf32>, %v2: vector<2x2x2x2xf32>) -> vector<2x2x2x2xf32> { + %0 = arith.addf %v1, %v2 : vector<2x2x2x2xf32> + return %0 : vector<2x2x2x2xf32> +} + +// CHECK-LABEL: func @elementwise_4D_to_2D +// CHECK-COUNT-4: arith.addf %{{.*}}, %{{.*}} : vector<2x2xf32> +// CHECK-NOT: arith.addf +// CHECK: return diff --git a/mlir/test/Dialect/Vector/vector-warp-distribute.mlir b/mlir/test/Dialect/Vector/vector-warp-distribute.mlir index bb76392..401cdd29 100644 --- a/mlir/test/Dialect/Vector/vector-warp-distribute.mlir +++ b/mlir/test/Dialect/Vector/vector-warp-distribute.mlir @@ -1925,3 +1925,22 @@ func.func @warp_scf_if_distribute(%pred : i1) { // CHECK-PROP: "some_use"(%[[IF_YIELD_DIST]]) : (vector<1xf32>) -> () // CHECK-PROP: return // CHECK-PROP: } + +// ----- +func.func @dedup_unused_result(%laneid : index) -> (vector<1xf32>) { + %r:3 = gpu.warp_execute_on_lane_0(%laneid)[32] -> + (vector<1xf32>, vector<2xf32>, vector<1xf32>) { + %2 = "some_def"() : () -> (vector<32xf32>) + %3 = "some_def"() : () -> (vector<64xf32>) + gpu.yield %2, %3, %2 : vector<32xf32>, vector<64xf32>, vector<32xf32> + } + %r0 = "some_use"(%r#2, %r#2) : (vector<1xf32>, vector<1xf32>) -> (vector<1xf32>) + return %r0 : vector<1xf32> +} + +// CHECK-PROP: func @dedup_unused_result +// CHECK-PROP: %[[R:.*]] = gpu.warp_execute_on_lane_0(%arg0)[32] -> (vector<1xf32>) +// CHECK-PROP: %[[Y0:.*]] = "some_def"() : () -> vector<32xf32> +// CHECK-PROP: %[[Y1:.*]] = "some_def"() : () -> vector<64xf32> +// CHECK-PROP: gpu.yield %[[Y0]] : vector<32xf32> +// CHECK-PROP: "some_use"(%[[R]], %[[R]]) : (vector<1xf32>, vector<1xf32>) -> vector<1xf32> diff --git a/mlir/test/Dialect/WasmSSA/custom_parser/global.mlir b/mlir/test/Dialect/WasmSSA/custom_parser/global.mlir index b9b3420..a25abbd 100644 --- a/mlir/test/Dialect/WasmSSA/custom_parser/global.mlir +++ b/mlir/test/Dialect/WasmSSA/custom_parser/global.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt %s | FileCheck %s module { - wasmssa.import_global "from_js" from "env" as @global_0 nested : i32 + wasmssa.import_global "from_js" from "env" as @global_0 : i32 wasmssa.global @global_1 i32 : { %0 = wasmssa.const 10 : i32 @@ -21,7 +21,7 @@ module { } } -// CHECK-LABEL: wasmssa.import_global "from_js" from "env" as @global_0 nested : i32 +// CHECK-LABEL: wasmssa.import_global "from_js" from "env" as @global_0 : i32 // CHECK-LABEL: wasmssa.global @global_1 i32 : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 diff --git a/mlir/test/Dialect/WasmSSA/custom_parser/if.mlir b/mlir/test/Dialect/WasmSSA/custom_parser/if.mlir index 01068cb..cee3c69 100644 --- a/mlir/test/Dialect/WasmSSA/custom_parser/if.mlir +++ b/mlir/test/Dialect/WasmSSA/custom_parser/if.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt %s | FileCheck %s -// CHECK-LABEL: wasmssa.func nested @func_0( +// CHECK-LABEL: wasmssa.func @func_0( // CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 // CHECK: wasmssa.if %[[VAL_0]] : { @@ -12,7 +12,7 @@ // CHECK: }> ^bb1 // CHECK: ^bb1(%[[VAL_3:.*]]: f32): // CHECK: wasmssa.return %[[VAL_3]] : f32 -wasmssa.func nested @func_0(%arg0 : !wasmssa<local ref to i32>) -> i32 { +wasmssa.func @func_0(%arg0 : !wasmssa<local ref to i32>) -> i32 { %cond = wasmssa.local_get %arg0 : ref to i32 wasmssa.if %cond : { %c0 = wasmssa.const 0.5 : f32 @@ -25,7 +25,7 @@ wasmssa.func nested @func_0(%arg0 : !wasmssa<local ref to i32>) -> i32 { wasmssa.return %retVal : f32 } -// CHECK-LABEL: wasmssa.func nested @func_1( +// CHECK-LABEL: wasmssa.func @func_1( // CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 // CHECK: %[[VAL_1:.*]] = wasmssa.local of type i32 @@ -38,7 +38,7 @@ wasmssa.func nested @func_0(%arg0 : !wasmssa<local ref to i32>) -> i32 { // CHECK: ^bb1: // CHECK: %[[VAL_4:.*]] = wasmssa.local_get %[[VAL_1]] : ref to i32 // CHECK: wasmssa.return %[[VAL_4]] : i32 -wasmssa.func nested @func_1(%arg0 : !wasmssa<local ref to i32>) -> i32 { +wasmssa.func @func_1(%arg0 : !wasmssa<local ref to i32>) -> i32 { %cond = wasmssa.local_get %arg0 : ref to i32 %var = wasmssa.local of type i32 %zero = wasmssa.const 0 diff --git a/mlir/test/Dialect/WasmSSA/custom_parser/import.mlir b/mlir/test/Dialect/WasmSSA/custom_parser/import.mlir index 3cc0548..dc23229 100644 --- a/mlir/test/Dialect/WasmSSA/custom_parser/import.mlir +++ b/mlir/test/Dialect/WasmSSA/custom_parser/import.mlir @@ -5,13 +5,13 @@ module { wasmssa.import_func "bar" from "my_module" as @func_1 {sym_visibility = "nested", type = (i32) -> ()} wasmssa.import_table "table" from "my_module" as @table_0 {sym_visibility = "nested", type = !wasmssa<tabletype !wasmssa.funcref [2:]>} wasmssa.import_mem "mem" from "my_module" as @mem_0 {limits = !wasmssa<limit[2:]>, sym_visibility = "nested"} - wasmssa.import_global "glob" from "my_module" as @global_0 nested : i32 - wasmssa.import_global "glob_mut" from "my_other_module" as @global_1 mutable nested : i32 + wasmssa.import_global "glob" from "my_module" as @global_0 : i32 + wasmssa.import_global "glob_mut" from "my_other_module" as @global_1 mutable : i32 } // CHECK-LABEL: wasmssa.import_func "foo" from "my_module" as @func_0 {sym_visibility = "nested", type = (i32) -> ()} // CHECK: wasmssa.import_func "bar" from "my_module" as @func_1 {sym_visibility = "nested", type = (i32) -> ()} // CHECK: wasmssa.import_table "table" from "my_module" as @table_0 {sym_visibility = "nested", type = !wasmssa<tabletype !wasmssa.funcref [2:]>} // CHECK: wasmssa.import_mem "mem" from "my_module" as @mem_0 {limits = !wasmssa<limit[2:]>, sym_visibility = "nested"} -// CHECK: wasmssa.import_global "glob" from "my_module" as @global_0 nested : i32 -// CHECK: wasmssa.import_global "glob_mut" from "my_other_module" as @global_1 mutable nested : i32 +// CHECK: wasmssa.import_global "glob" from "my_module" as @global_0 : i32 +// CHECK: wasmssa.import_global "glob_mut" from "my_other_module" as @global_1 mutable : i32 diff --git a/mlir/test/Dialect/WasmSSA/custom_parser/local.mlir b/mlir/test/Dialect/WasmSSA/custom_parser/local.mlir index 3f6423f..f613ebf 100644 --- a/mlir/test/Dialect/WasmSSA/custom_parser/local.mlir +++ b/mlir/test/Dialect/WasmSSA/custom_parser/local.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt %s | FileCheck %s module { - wasmssa.func nested @func_0() -> f32 { + wasmssa.func @func_0() -> f32 { %0 = wasmssa.local of type f32 %1 = wasmssa.local of type f32 %2 = wasmssa.const 8.000000e+00 : f32 @@ -9,7 +9,7 @@ module { %4 = wasmssa.add %2 %3 : f32 wasmssa.return %4 : f32 } - wasmssa.func nested @func_1() -> i32 { + wasmssa.func @func_1() -> i32 { %0 = wasmssa.local of type i32 %1 = wasmssa.local of type i32 %2 = wasmssa.const 8 : i32 @@ -17,13 +17,13 @@ module { %4 = wasmssa.add %2 %3 : i32 wasmssa.return %4 : i32 } - wasmssa.func nested @func_2(%arg0: !wasmssa<local ref to i32>) -> i32 { + wasmssa.func @func_2(%arg0: !wasmssa<local ref to i32>) -> i32 { %0 = wasmssa.const 3 : i32 wasmssa.return %0 : i32 } } -// CHECK-LABEL: wasmssa.func nested @func_0() -> f32 { +// CHECK-LABEL: wasmssa.func @func_0() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.local of type f32 // CHECK: %[[VAL_1:.*]] = wasmssa.local of type f32 // CHECK: %[[VAL_2:.*]] = wasmssa.const 8.000000e+00 : f32 @@ -31,7 +31,7 @@ module { // CHECK: %[[VAL_4:.*]] = wasmssa.add %[[VAL_2]] %[[VAL_3]] : f32 // CHECK: wasmssa.return %[[VAL_4]] : f32 -// CHECK-LABEL: wasmssa.func nested @func_1() -> i32 { +// CHECK-LABEL: wasmssa.func @func_1() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.local of type i32 // CHECK: %[[VAL_1:.*]] = wasmssa.local of type i32 // CHECK: %[[VAL_2:.*]] = wasmssa.const 8 : i32 @@ -39,7 +39,7 @@ module { // CHECK: %[[VAL_4:.*]] = wasmssa.add %[[VAL_2]] %[[VAL_3]] : i32 // CHECK: wasmssa.return %[[VAL_4]] : i32 -// CHECK-LABEL: wasmssa.func nested @func_2( +// CHECK-LABEL: wasmssa.func @func_2( // CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 3 : i32 // CHECK: wasmssa.return %[[VAL_0]] : i32 diff --git a/mlir/test/Dialect/WasmSSA/custom_parser/memory.mlir b/mlir/test/Dialect/WasmSSA/custom_parser/memory.mlir index 47551db..ca6ebe0 100644 --- a/mlir/test/Dialect/WasmSSA/custom_parser/memory.mlir +++ b/mlir/test/Dialect/WasmSSA/custom_parser/memory.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt %s | FileCheck %s -// CHECK: wasmssa.memory @mem0 public !wasmssa<limit[0: 65536]> -wasmssa.memory @mem0 public !wasmssa<limit[0:65536]> - -// CHECK: wasmssa.memory @mem1 nested !wasmssa<limit[512:]> +// CHECK: wasmssa.memory @mem1 !wasmssa<limit[512:]> wasmssa.memory @mem1 !wasmssa<limit[512:]> + +// CHECK: wasmssa.memory exported @mem2 !wasmssa<limit[0: 65536]> +wasmssa.memory exported @mem2 !wasmssa<limit[0:65536]> diff --git a/mlir/test/Dialect/WasmSSA/custom_parser/table.mlir b/mlir/test/Dialect/WasmSSA/custom_parser/table.mlir index 5a874f4..ea630de 100644 --- a/mlir/test/Dialect/WasmSSA/custom_parser/table.mlir +++ b/mlir/test/Dialect/WasmSSA/custom_parser/table.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt %s | FileCheck %s -// CHECK: wasmssa.table @tab0 public !wasmssa<tabletype !wasmssa.externref [0: 65536]> -wasmssa.table @tab0 public !wasmssa<tabletype !wasmssa.externref [0:65536]> +// CHECK: wasmssa.table exported @tab0 !wasmssa<tabletype !wasmssa.externref [0: 65536]> +wasmssa.table exported @tab0 !wasmssa<tabletype !wasmssa.externref [0:65536]> -// CHECK: wasmssa.table @tab1 nested !wasmssa<tabletype !wasmssa.funcref [348:]> +// CHECK: wasmssa.table @tab1 !wasmssa<tabletype !wasmssa.funcref [348:]> wasmssa.table @tab1 !wasmssa<tabletype !wasmssa.funcref [348:]> diff --git a/mlir/test/Dialect/WasmSSA/extend-invalid.mlir b/mlir/test/Dialect/WasmSSA/extend-invalid.mlir index 8d78280..7687e5f 100644 --- a/mlir/test/Dialect/WasmSSA/extend-invalid.mlir +++ b/mlir/test/Dialect/WasmSSA/extend-invalid.mlir @@ -1,7 +1,7 @@ // RUN: mlir-opt %s -split-input-file -verify-diagnostics -wasmssa.func nested @extend_low_64() -> i32 { +wasmssa.func @extend_low_64() -> i32 { %0 = wasmssa.const 10 : i32 // expected-error@+1 {{extend op can only take 8, 16 or 32 bits. Got 64}} %1 = wasmssa.extend 64 low bits from %0: i32 @@ -10,7 +10,7 @@ wasmssa.func nested @extend_low_64() -> i32 { // ----- -wasmssa.func nested @extend_too_much() -> i32 { +wasmssa.func @extend_too_much() -> i32 { %0 = wasmssa.const 10 : i32 // expected-error@+1 {{trying to extend the 32 low bits from a 'i32' value is illegal}} %1 = wasmssa.extend 32 low bits from %0: i32 diff --git a/mlir/test/Dialect/WasmSSA/global-invalid.mlir b/mlir/test/Dialect/WasmSSA/global-invalid.mlir index b9cafd8..c5bc606 100644 --- a/mlir/test/Dialect/WasmSSA/global-invalid.mlir +++ b/mlir/test/Dialect/WasmSSA/global-invalid.mlir @@ -13,7 +13,7 @@ module { // ----- module { - wasmssa.import_global "glob" from "my_module" as @global_0 mutable nested : i32 + wasmssa.import_global "glob" from "my_module" as @global_0 mutable : i32 wasmssa.global @global_1 i32 : { // expected-error@+1 {{global.get op is considered constant if it's referring to a import.global symbol marked non-mutable}} %0 = wasmssa.global_get @global_0 : i32 @@ -30,3 +30,13 @@ module { wasmssa.return %0 : i32 } } + +// ----- + +module { + // expected-error@+1 {{expecting either `exported` or symbol name. got exproted}} + wasmssa.global exproted @global_1 i32 : { + %0 = wasmssa.const 17 : i32 + wasmssa.return %0 : i32 + } +} diff --git a/mlir/test/Dialect/WasmSSA/locals-invalid.mlir b/mlir/test/Dialect/WasmSSA/locals-invalid.mlir index 35c590b..eaad80e 100644 --- a/mlir/test/Dialect/WasmSSA/locals-invalid.mlir +++ b/mlir/test/Dialect/WasmSSA/locals-invalid.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt %s -split-input-file -verify-diagnostics -wasmssa.func nested @local_set_err(%arg0: !wasmssa<local ref to i32>) -> i64 { +wasmssa.func @local_set_err(%arg0: !wasmssa<local ref to i32>) -> i64 { %0 = wasmssa.const 3 : i64 // expected-error@+1 {{input type and result type of local.set do not match}} wasmssa.local_set %arg0 : ref to i32 to %0 : i64 @@ -9,7 +9,7 @@ wasmssa.func nested @local_set_err(%arg0: !wasmssa<local ref to i32>) -> i64 { // ----- -wasmssa.func nested @local_tee_err(%arg0: !wasmssa<local ref to i32>) -> i32 { +wasmssa.func @local_tee_err(%arg0: !wasmssa<local ref to i32>) -> i32 { %0 = wasmssa.const 3 : i64 // expected-error@+1 {{input type and output type of local.tee do not match}} %1 = wasmssa.local_tee %arg0 : ref to i32 to %0 : i64 diff --git a/mlir/test/Dialect/XeGPU/invalid.mlir b/mlir/test/Dialect/XeGPU/invalid.mlir index 228ef69d..ebbe3ce 100644 --- a/mlir/test/Dialect/XeGPU/invalid.mlir +++ b/mlir/test/Dialect/XeGPU/invalid.mlir @@ -858,7 +858,7 @@ func.func @load_mem_desc_mismatch_element_type(%arg0: !xegpu.mem_desc<16x64xf16> // ----- func.func @load_mem_desc_invalid_result_size(%arg0: !xegpu.mem_desc<16x64xf16>) { - // expected-error@+1 {{result shape must not exceed mem_desc shape}} + // expected-error@+1 {{data shape must not exceed mem_desc shape}} %data = xegpu.load_matrix %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> vector<32x16xf16> return } @@ -871,6 +871,14 @@ func.func @load_mem_desc_invalid_rank(%arg0: !xegpu.mem_desc<64xf16>) { } // ----- +func.func @load_mem_desc_invalid_attr2(%arg0: !xegpu.mem_desc<16x64xf16>) { + // expected-error@+1 {{subgroup_block_io are only allowed when result is a 1D VectorType.}} + %data2 = xegpu.load_matrix %arg0[8, 8] <{subgroup_block_io}>: !xegpu.mem_desc<16x64xf16> -> vector<16x16xf16> + return +} + + +// ----- func.func @store_mem_desc_mismatch_element_type(%arg0: !xegpu.mem_desc<16x64xf16>, %arg1: vector<16x16xf32>) { // expected-error@+1 {{failed to verify that all of {mem_desc, data} have same element type}} xegpu.store_matrix %arg1, %arg0[8, 8] : vector<16x16xf32>, !xegpu.mem_desc<16x64xf16> @@ -892,30 +900,16 @@ func.func @store_mem_desc_invalid_rank(%arg0: !xegpu.mem_desc<64xf16>, %arg1: ve } // ----- -func.func @mem_desc_subview_size_mismatch(%arg0: !xegpu.mem_desc<16x64xf16>) { - // expected-error@+1 {{result shape must not exceed source shape}} - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<32x16xf16> - return -} - -// ----- -func.func @mem_desc_subview_layout_mismatch(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride =[1, 16]>>) { - // expected-error@+1 {{result must inherit the source strides}} - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride =[1, 16]>> -> !xegpu.mem_desc<8x16xf16> - return -} - -// ----- -func.func @mem_desc_subview_element_type_mismatch(%arg0: !xegpu.mem_desc<16x64xf16>) { - // expected-error@+1 {{failed to verify that all of {src, res} have same element type}} - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<8x16xf32, #xegpu.mem_layout<stride =[64, 1]>> +func.func @store_mem_desc_invalid_attr2(%arg0: !xegpu.mem_desc<16x64xf16>, %data: vector<16x16xf16>) { + // expected-error@+1 {{subgroup_block_io are only allowed when result is a 1D VectorType.}} + xegpu.store_matrix %data, %arg0[8, 8] <{subgroup_block_io}>: vector<16x16xf16>, !xegpu.mem_desc<16x64xf16> return } // ----- -func.func @mem_desc_subview_rank_mismatch(%arg0: !xegpu.mem_desc<16x64xf16>) { - // expected-error@+1 {{result rank must not exceed source rank}} - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<4x8x16xf16> +func.func @store_mem_desc_invalid_attr2(%arg0: !xegpu.mem_desc<16x64xf16>, %data: vector<16x16xf16>) { + // expected-error@+1 {{subgroup_block_io are only allowed when result is a 1D VectorType.}} + xegpu.store_matrix %data, %arg0[8, 8] <{subgroup_block_io}>: vector<16x16xf16>, !xegpu.mem_desc<16x64xf16> return } diff --git a/mlir/test/Dialect/XeGPU/ops.mlir b/mlir/test/Dialect/XeGPU/ops.mlir index bb37902..0a10f68 100644 --- a/mlir/test/Dialect/XeGPU/ops.mlir +++ b/mlir/test/Dialect/XeGPU/ops.mlir @@ -825,53 +825,73 @@ gpu.func @create_mem_desc_with_stride() { gpu.return } -// CHECK: gpu.func @load_mem_desc([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16>) -gpu.func @load_mem_desc(%arg0: !xegpu.mem_desc<16x64xf16>) { +// CHECK: gpu.func @load_matrix([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16>) +gpu.func @load_matrix(%arg0: !xegpu.mem_desc<16x64xf16>) { // CHECK: xegpu.load_matrix [[ARG0]][8, 8] : !xegpu.mem_desc<16x64xf16> -> vector<8x16xf16> %data = xegpu.load_matrix %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> vector<8x16xf16> gpu.return } -// CHECK: gpu.func @load_mem_desc_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) -gpu.func @load_mem_desc_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) { +// CHECK: gpu.func @load_matrix_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) +gpu.func @load_matrix_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) { // CHECK: xegpu.load_matrix [[ARG0]][8, 8] : !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> -> vector<8x16xf16> %data = xegpu.load_matrix %arg0[8, 8]: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> -> vector<8x16xf16> gpu.return } +// CHECK: gpu.func @simt_load_matrix(%arg0: !xegpu.mem_desc<16x64xf16>) +gpu.func @simt_load_matrix(%arg0: !xegpu.mem_desc<16x64xf16>) { + // CHECK: xegpu.load_matrix [[ARG0]][8, 16] : !xegpu.mem_desc<16x64xf16> -> vector<1xf16> + %data = xegpu.load_matrix %arg0[8, 16]: !xegpu.mem_desc<16x64xf16> -> vector<1xf16> + gpu.return +} + +// CHECK: gpu.func @simt_load_matrix_subgroup_block_io(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>>) +gpu.func @simt_load_matrix_subgroup_block_io(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>>) { + // CHECK: xegpu.load_matrix [[ARG0]][8, 16] <{subgroup_block_io}>: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>> -> vector<8xf16> + %data = xegpu.load_matrix %arg0[8, 16] <{subgroup_block_io}>: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>> -> vector<8xf16> + gpu.return +} + +// CHECK: gpu.func @simt_load_matrix_vector(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) +gpu.func @simt_load_matrix_vector(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) { + // CHECK: xegpu.load_matrix [[ARG0]][8, 8] : !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> -> vector<8xf16> + %data = xegpu.load_matrix %arg0[8, 8] : !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> -> vector<8xf16> + gpu.return +} -// CHECK: gpu.func @store_mem_desc([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16>, [[ARG1:%.+]]: vector<16x16xf16>) -gpu.func @store_mem_desc(%arg0: !xegpu.mem_desc<16x64xf16>, %arg1: vector<16x16xf16>) { +// CHECK: gpu.func @store_matrix([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16>, [[ARG1:%.+]]: vector<16x16xf16>) +gpu.func @store_matrix(%arg0: !xegpu.mem_desc<16x64xf16>, %arg1: vector<16x16xf16>) { // CHECK: xegpu.store_matrix [[ARG1]], [[ARG0]][8, 8] : vector<16x16xf16>, !xegpu.mem_desc<16x64xf16> xegpu.store_matrix %arg1, %arg0[8, 8]: vector<16x16xf16>, !xegpu.mem_desc<16x64xf16> gpu.return } -// CHECK: gpu.func @store_mem_desc_with_stride([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>, [[ARG1:%.+]]: vector<16x16xf16>) -gpu.func @store_mem_desc_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>, %arg1: vector<16x16xf16>) { +// CHECK: gpu.func @store_matrix_with_stride([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>, [[ARG1:%.+]]: vector<16x16xf16>) +gpu.func @store_matrix_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>, %arg1: vector<16x16xf16>) { // CHECK: xegpu.store_matrix [[ARG1]], [[ARG0]][0, 8] : vector<16x16xf16>, !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> xegpu.store_matrix %arg1, %arg0[0, 8]: vector<16x16xf16>, !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> gpu.return } -// CHECK: gpu.func @mem_desc_subview([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16>) -gpu.func @mem_desc_subview(%arg0: !xegpu.mem_desc<16x64xf16>) { - //CHECK: xegpu.mem_desc_subview [[ARG0]][8, 8] : !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<8x16xf16, #xegpu.mem_layout<stride = [64, 1]>> - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<8x16xf16, #xegpu.mem_layout<stride = [64, 1]>> +// CHECK: gpu.func @simt_store_matrix(%arg0: !xegpu.mem_desc<16x64xf16>, %arg1: vector<1xf16>) { +gpu.func @simt_store_matrix(%arg0: !xegpu.mem_desc<16x64xf16>, %arg1: vector<1xf16>) { + // CHECK: xegpu.store_matrix [[ARG1]], [[ARG0]][8, 16] : vector<1xf16>, !xegpu.mem_desc<16x64xf16> + xegpu.store_matrix %arg1, %arg0[8, 16]: vector<1xf16>, !xegpu.mem_desc<16x64xf16> gpu.return } -// CHECK: gpu.func @mem_desc_subview_lower_rank([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16>) -gpu.func @mem_desc_subview_lower_rank(%arg0: !xegpu.mem_desc<16x64xf16>) { - //CHECK: xegpu.mem_desc_subview [[ARG0]][8, 8] : !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<16xf16, #xegpu.mem_layout<stride = [64, 1]>> - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16> -> !xegpu.mem_desc<16xf16, #xegpu.mem_layout<stride = [64, 1]>> +// CHECK: gpu.func @simt_store_matrix_subgroup_block_io(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>>, %arg1: vector<8xf16>) +gpu.func @simt_store_matrix_subgroup_block_io(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>>, %arg1: vector<8xf16>) { + // CHECK: xegpu.store_matrix [[ARG1]], [[ARG0]][8, 16] <{subgroup_block_io}>: vector<8xf16>, !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>> + xegpu.store_matrix %arg1, %arg0[8, 16] <{subgroup_block_io}>: vector<8xf16>, !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<block = [16, 16]>> gpu.return } -// CHECK: gpu.func @mem_desc_subview_with_stride([[ARG0:%.+]]: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) -gpu.func @mem_desc_subview_with_stride(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>) { - //CHECK: xegpu.mem_desc_subview [[ARG0]][8, 8] : !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> -> !xegpu.mem_desc<8x16xf16, #xegpu.mem_layout<stride = [1, 16]>> - %data = xegpu.mem_desc_subview %arg0[8, 8]: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> -> !xegpu.mem_desc<8x16xf16, #xegpu.mem_layout<stride = [1, 16]>> +// CHECK: gpu.func @simt_store_matrix_vector(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>, %arg1: vector<8xf16>) { +gpu.func @simt_store_matrix_vector(%arg0: !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>>, %arg1: vector<8xf16>) { + // CHECK: xegpu.store_matrix [[ARG1]], [[ARG0]][8, 8] : vector<8xf16>, !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> + xegpu.store_matrix %arg1, %arg0[8, 8] : vector<8xf16>, !xegpu.mem_desc<16x64xf16, #xegpu.mem_layout<stride = [1, 16]>> gpu.return } diff --git a/mlir/test/Target/LLVMIR/Import/function-attributes.ll b/mlir/test/Target/LLVMIR/Import/function-attributes.ll index cc3d799..00d09ba 100644 --- a/mlir/test/Target/LLVMIR/Import/function-attributes.ll +++ b/mlir/test/Target/LLVMIR/Import/function-attributes.ll @@ -393,6 +393,12 @@ declare void @alwaysinline_attribute() alwaysinline // ----- +; CHECK-LABEL: @inlinehint_attribute +; CHECK-SAME: attributes {inline_hint} +declare void @inlinehint_attribute() inlinehint + +// ----- + ; CHECK-LABEL: @optnone_attribute ; CHECK-SAME: attributes {no_inline, optimize_none} declare void @optnone_attribute() noinline optnone diff --git a/mlir/test/Target/LLVMIR/llvmir.mlir b/mlir/test/Target/LLVMIR/llvmir.mlir index 69814f2..cc243c8 100644 --- a/mlir/test/Target/LLVMIR/llvmir.mlir +++ b/mlir/test/Target/LLVMIR/llvmir.mlir @@ -2555,6 +2555,17 @@ llvm.func @always_inline() attributes { always_inline } { // ----- +// CHECK-LABEL: @inline_hint +// CHECK-SAME: #[[ATTRS:[0-9]+]] +llvm.func @inline_hint() attributes { inline_hint } { + llvm.return +} + +// CHECK: #[[ATTRS]] +// CHECK-SAME: inlinehint + +// ----- + // CHECK-LABEL: @optimize_none // CHECK-SAME: #[[ATTRS:[0-9]+]] llvm.func @optimize_none() attributes { no_inline, optimize_none } { diff --git a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir index 78e1e659..6cccfe4 100644 --- a/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir +++ b/mlir/test/Target/LLVMIR/nvvmir-invalid.mlir @@ -567,3 +567,25 @@ llvm.func @clusterlaunchcontrol_query_cancel_get_first_cta_id_invalid_return_typ %res = nvvm.clusterlaunchcontrol.query.cancel query = get_first_cta_id_x, %try_cancel_response : i1 llvm.return } + +// ----- + +// Test that ensures invalid row/col layouts for matrices A and B are not accepted +llvm.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32, %b0 : i32, %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32,i32,i32,i32)> { + // expected-error@+1 {{Only m8n8k4 with f16 supports other layouts.}} + %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3] + {layoutA = #nvvm.mma_layout<col>, layoutB = #nvvm.mma_layout<col>, + multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>, + intOverflowBehavior=#nvvm.mma_int_overflow<satfinite>, + shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)> + llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)> +} + +// ----- + +// Test for range validation - invalid range where lower == upper but not at extremes +func.func @invalid_range_equal_bounds() { + // expected-error @below {{invalid range attribute: Lower == Upper, but they aren't min (0) or max (4294967295) value! This is an invalid constant range.}} + %0 = nvvm.read.ptx.sreg.warpsize range <i32, 32, 32> : i32 + return +} diff --git a/mlir/test/Target/LLVMIR/nvvmir.mlir b/mlir/test/Target/LLVMIR/nvvmir.mlir index 00a479d..594ae48 100644 --- a/mlir/test/Target/LLVMIR/nvvmir.mlir +++ b/mlir/test/Target/LLVMIR/nvvmir.mlir @@ -152,6 +152,10 @@ llvm.func @nvvm_special_regs() -> i32 { %74 = nvvm.read.ptx.sreg.lanemask.ge : i32 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.gt %75 = nvvm.read.ptx.sreg.lanemask.gt : i32 + // CHECK: %76 = call range(i32 0, 0) i32 @llvm.nvvm.read.ptx.sreg.tid.x() + %76 = nvvm.read.ptx.sreg.tid.x range <i32, 0, 0> : i32 + // CHECK: %77 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() + %77 = nvvm.read.ptx.sreg.tid.x range <i32, 4294967295, 4294967295> : i32 llvm.return %1 : i32 } diff --git a/mlir/test/Target/LLVMIR/rocdl.mlir b/mlir/test/Target/LLVMIR/rocdl.mlir index fdd2c91..6536fac 100644 --- a/mlir/test/Target/LLVMIR/rocdl.mlir +++ b/mlir/test/Target/LLVMIR/rocdl.mlir @@ -276,6 +276,20 @@ llvm.func @rocdl.s.wait.expcnt() { llvm.return } +llvm.func @rocdl.s.wait.asynccnt() { + // CHECK-LABEL: rocdl.s.wait.asynccnt + // CHECK-NEXT: call void @llvm.amdgcn.s.wait.asynccnt(i16 0) + rocdl.s.wait.asynccnt 0 + llvm.return +} + +llvm.func @rocdl.s.wait.tensorcnt() { + // CHECK-LABEL: rocdl.s.wait.tensorcnt + // CHECK-NEXT: call void @llvm.amdgcn.s.wait.tensorcnt(i16 0) + rocdl.s.wait.tensorcnt 0 + llvm.return +} + llvm.func @rocdl.setprio() { // CHECK: call void @llvm.amdgcn.s.setprio(i16 0) rocdl.s.setprio 0 diff --git a/mlir/test/Target/Wasm/abs.mlir b/mlir/test/Target/Wasm/abs.mlir index 9c45ba7..fe3602a 100644 --- a/mlir/test/Target/Wasm/abs.mlir +++ b/mlir/test/Target/Wasm/abs.mlir @@ -12,12 +12,12 @@ ) */ -// CHECK-LABEL: wasmssa.func @abs_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @abs_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.abs %[[VAL_0]] : f32 // CHECK: wasmssa.return %[[VAL_1]] : f32 -// CHECK-LABEL: wasmssa.func @abs_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @abs_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.abs %[[VAL_0]] : f64 // CHECK: wasmssa.return %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/add_div.mlir b/mlir/test/Target/Wasm/add_div.mlir new file mode 100644 index 0000000..8a87c60 --- /dev/null +++ b/mlir/test/Target/Wasm/add_div.mlir @@ -0,0 +1,40 @@ +// RUN: yaml2obj %S/inputs/add_div.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: + (module $test.wasm + (type (;0;) (func (param i32) (result i32))) + (type (;1;) (func (param i32 i32) (result i32))) + (import "env" "twoTimes" (func $twoTimes (type 0))) + (func $add (type 1) (param i32 i32) (result i32) + local.get 0 + call $twoTimes + local.get 1 + call $twoTimes + i32.add + i32.const 2 + i32.div_s) + (memory (;0;) 2) + (global $__stack_pointer (mut i32) (i32.const 66560)) + (export "memory" (memory 0)) + (export "add" (func $add))) +*/ + +// CHECK-LABEL: wasmssa.import_func "twoTimes" from "env" as @func_0 {type = (i32) -> i32} + +// CHECK-LABEL: wasmssa.func exported @add( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>, +// CHECK-SAME: %[[ARG1:.*]]: !wasmssa<local ref to i32>) -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.call @func_0(%[[VAL_0]]) : (i32) -> i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.local_get %[[ARG1]] : ref to i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.call @func_0(%[[VAL_2]]) : (i32) -> i32 +// CHECK: %[[VAL_4:.*]] = wasmssa.add %[[VAL_1]] %[[VAL_3]] : i32 +// CHECK: %[[VAL_5:.*]] = wasmssa.const 2 : i32 +// CHECK: %[[VAL_6:.*]] = wasmssa.div_si %[[VAL_4]] %[[VAL_5]] : i32 +// CHECK: wasmssa.return %[[VAL_6]] : i32 +// CHECK: } +// CHECK: wasmssa.memory exported @memory !wasmssa<limit[2:]> + +// CHECK-LABEL: wasmssa.global @global_0 i32 mutable : { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 66560 : i32 +// CHECK: wasmssa.return %[[VAL_0]] : i32 diff --git a/mlir/test/Target/Wasm/and.mlir b/mlir/test/Target/Wasm/and.mlir index 4c0fea0..323d41a 100644 --- a/mlir/test/Target/Wasm/and.mlir +++ b/mlir/test/Target/Wasm/and.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @and_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @and_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.and %0 %1 : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @and_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @and_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.and %0 %1 : i64 diff --git a/mlir/test/Target/Wasm/block.mlir b/mlir/test/Target/Wasm/block.mlir new file mode 100644 index 0000000..c85fc1e --- /dev/null +++ b/mlir/test/Target/Wasm/block.mlir @@ -0,0 +1,16 @@ +// RUN: yaml2obj %S/inputs/block.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module +(func(export "i_am_a_block") +(block $i_am_a_block) +) +) +*/ + +// CHECK-LABEL: wasmssa.func exported @i_am_a_block() { +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.return diff --git a/mlir/test/Target/Wasm/block_complete_type.mlir b/mlir/test/Target/Wasm/block_complete_type.mlir new file mode 100644 index 0000000..67df198 --- /dev/null +++ b/mlir/test/Target/Wasm/block_complete_type.mlir @@ -0,0 +1,24 @@ +// RUN: yaml2obj %S/inputs/block_complete_type.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module + (type (;0;) (func (param i32) (result i32))) + (type (;1;) (func (result i32))) + (func (;0;) (type 1) (result i32) + i32.const 14 + block (param i32) (result i32) ;; label = @1 + i32.const 1 + i32.add + end)) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 14 : i32 +// CHECK: wasmssa.block(%[[VAL_0]]) : i32 : { +// CHECK: ^bb0(%[[VAL_1:.*]]: i32): +// CHECK: %[[VAL_2:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.add %[[VAL_1]] %[[VAL_2]] : i32 +// CHECK: wasmssa.block_return %[[VAL_3]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_4:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_4]] : i32 diff --git a/mlir/test/Target/Wasm/block_value_type.mlir b/mlir/test/Target/Wasm/block_value_type.mlir new file mode 100644 index 0000000..fa30f08 --- /dev/null +++ b/mlir/test/Target/Wasm/block_value_type.mlir @@ -0,0 +1,19 @@ +// RUN: yaml2obj %S/inputs/block_value_type.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module + (type (;0;) (func (result i32))) + (func (;0;) (type 0) (result i32) + block (result i32) ;; label = @1 + i32.const 17 + end)) +*/ + + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: wasmssa.block : { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 17 : i32 +// CHECK: wasmssa.block_return %[[VAL_0]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_1:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_1]] : i32 diff --git a/mlir/test/Target/Wasm/branch_if.mlir b/mlir/test/Target/Wasm/branch_if.mlir new file mode 100644 index 0000000..c91ff37 --- /dev/null +++ b/mlir/test/Target/Wasm/branch_if.mlir @@ -0,0 +1,29 @@ +// RUN: yaml2obj %S/inputs/branch_if.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module + (type $produce_i32 (func (result i32))) + (func (type $produce_i32) + (block $my_block (type $produce_i32) + i32.const 1 + i32.const 2 + br_if $my_block + i32.const 1 + i32.add + ) + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: wasmssa.block : { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 2 : i32 +// CHECK: wasmssa.branch_if %[[VAL_1]] to level 0 with args(%[[VAL_0]] : i32) else ^bb1 +// CHECK: ^bb1: +// CHECK: %[[VAL_2:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.add %[[VAL_0]] %[[VAL_2]] : i32 +// CHECK: wasmssa.block_return %[[VAL_3]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_4:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_4]] : i32 diff --git a/mlir/test/Target/Wasm/call.mlir b/mlir/test/Target/Wasm/call.mlir new file mode 100644 index 0000000..c0169aa --- /dev/null +++ b/mlir/test/Target/Wasm/call.mlir @@ -0,0 +1,17 @@ +// RUN: yaml2obj %S/inputs/call.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module +(func $forty_two (result i32) +i32.const 42) +(func(export "forty_two")(result i32) +call $forty_two)) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 42 : i32 +// CHECK: wasmssa.return %[[VAL_0]] : i32 + +// CHECK-LABEL: wasmssa.func exported @forty_two() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.call @func_0 : () -> i32 +// CHECK: wasmssa.return %[[VAL_0]] : i32 diff --git a/mlir/test/Target/Wasm/clz.mlir b/mlir/test/Target/Wasm/clz.mlir index 3e6641d..858c09d 100644 --- a/mlir/test/Target/Wasm/clz.mlir +++ b/mlir/test/Target/Wasm/clz.mlir @@ -14,12 +14,12 @@ ) */ -// CHECK-LABEL: wasmssa.func @clz_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @clz_i32() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.clz %[[VAL_0]] : i32 // CHECK: wasmssa.return %[[VAL_1]] : i32 -// CHECK-LABEL: wasmssa.func @clz_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @clz_i64() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.clz %[[VAL_0]] : i64 // CHECK: wasmssa.return %[[VAL_1]] : i64 diff --git a/mlir/test/Target/Wasm/comparison_ops.mlir b/mlir/test/Target/Wasm/comparison_ops.mlir new file mode 100644 index 0000000..91e3a6a --- /dev/null +++ b/mlir/test/Target/Wasm/comparison_ops.mlir @@ -0,0 +1,269 @@ +// RUN: yaml2obj %S/inputs/comparison_ops.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s +/* Source code used to create this test: +(module + (func $lt_si32 (result i32) + i32.const 12 + i32.const 50 + i32.lt_s + ) + (func $le_si32 (result i32) + i32.const 12 + i32.const 50 + i32.le_s + ) + (func $lt_ui32 (result i32) + i32.const 12 + i32.const 50 + i32.lt_u + ) + (func $le_ui32 (result i32) + i32.const 12 + i32.const 50 + i32.le_u + ) + (func $gt_si32 (result i32) + i32.const 12 + i32.const 50 + i32.gt_s + ) + (func $gt_ui32 (result i32) + i32.const 12 + i32.const 50 + i32.gt_u + ) + (func $ge_si32 (result i32) + i32.const 12 + i32.const 50 + i32.ge_s + ) + (func $ge_ui32 (result i32) + i32.const 12 + i32.const 50 + i32.ge_u + ) + (func $lt_si64 (result i32) + i64.const 12 + i64.const 50 + i64.lt_s + ) + (func $le_si64 (result i32) + i64.const 12 + i64.const 50 + i64.le_s + ) + (func $lt_ui64 (result i32) + i64.const 12 + i64.const 50 + i64.lt_u + ) + (func $le_ui64 (result i32) + i64.const 12 + i64.const 50 + i64.le_u + ) + (func $gt_si64 (result i32) + i64.const 12 + i64.const 50 + i64.gt_s + ) + (func $gt_ui64 (result i32) + i64.const 12 + i64.const 50 + i64.gt_u + ) + (func $ge_si64 (result i32) + i64.const 12 + i64.const 50 + i64.ge_s + ) + (func $ge_ui64 (result i32) + i64.const 12 + i64.const 50 + i64.ge_u + ) + (func $lt_f32 (result i32) + f32.const 5 + f32.const 14 + f32.lt + ) + (func $le_f32 (result i32) + f32.const 5 + f32.const 14 + f32.le + ) + (func $gt_f32 (result i32) + f32.const 5 + f32.const 14 + f32.gt + ) + (func $ge_f32 (result i32) + f32.const 5 + f32.const 14 + f32.ge + ) + (func $lt_f64 (result i32) + f64.const 5 + f64.const 14 + f64.lt + ) + (func $le_f64 (result i32) + f64.const 5 + f64.const 14 + f64.le + ) + (func $gt_f64 (result i32) + f64.const 5 + f64.const 14 + f64.gt + ) + (func $ge_f64 (result i32) + f64.const 5 + f64.const 14 + f64.ge + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.lt_si %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_1() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.le_si %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_2() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.lt_ui %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_3() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.le_ui %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_4() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.gt_si %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_5() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.gt_ui %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_6() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.ge_si %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_7() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.ge_ui %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_8() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.lt_si %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_9() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.le_si %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_10() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.lt_ui %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_11() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.le_ui %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_12() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.gt_si %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_13() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.gt_ui %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_14() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.ge_si %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_15() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.ge_ui %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_16() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 +// CHECK: %[[VAL_2:.*]] = wasmssa.lt %[[VAL_0]] %[[VAL_1]] : f32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_17() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 +// CHECK: %[[VAL_2:.*]] = wasmssa.le %[[VAL_0]] %[[VAL_1]] : f32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_18() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 +// CHECK: %[[VAL_2:.*]] = wasmssa.gt %[[VAL_0]] %[[VAL_1]] : f32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_19() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 +// CHECK: %[[VAL_2:.*]] = wasmssa.ge %[[VAL_0]] %[[VAL_1]] : f32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_20() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f64 +// CHECK: %[[VAL_2:.*]] = wasmssa.lt %[[VAL_0]] %[[VAL_1]] : f64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_21() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f64 +// CHECK: %[[VAL_2:.*]] = wasmssa.le %[[VAL_0]] %[[VAL_1]] : f64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_22() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f64 +// CHECK: %[[VAL_2:.*]] = wasmssa.gt %[[VAL_0]] %[[VAL_1]] : f64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_23() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f64 +// CHECK: %[[VAL_2:.*]] = wasmssa.ge %[[VAL_0]] %[[VAL_1]] : f64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 diff --git a/mlir/test/Target/Wasm/const.mlir b/mlir/test/Target/Wasm/const.mlir index aa9e76f..adb792a 100644 --- a/mlir/test/Target/Wasm/const.mlir +++ b/mlir/test/Target/Wasm/const.mlir @@ -16,22 +16,22 @@ ) */ -// CHECK-LABEL: wasmssa.func nested @func_0() -> i32 { +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1 : i32 // CHECK: wasmssa.return %[[VAL_0]] : i32 // CHECK: } -// CHECK-LABEL: wasmssa.func nested @func_1() -> i64 { +// CHECK-LABEL: wasmssa.func @func_1() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 3 : i64 // CHECK: wasmssa.return %[[VAL_0]] : i64 // CHECK: } -// CHECK-LABEL: wasmssa.func nested @func_2() -> f32 { +// CHECK-LABEL: wasmssa.func @func_2() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 4.000000e+00 : f32 // CHECK: wasmssa.return %[[VAL_0]] : f32 // CHECK: } -// CHECK-LABEL: wasmssa.func nested @func_3() -> f64 { +// CHECK-LABEL: wasmssa.func @func_3() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 9.000000e+00 : f64 // CHECK: wasmssa.return %[[VAL_0]] : f64 // CHECK: } diff --git a/mlir/test/Target/Wasm/convert.mlir b/mlir/test/Target/Wasm/convert.mlir new file mode 100644 index 0000000..ddc29a7 --- /dev/null +++ b/mlir/test/Target/Wasm/convert.mlir @@ -0,0 +1,85 @@ +// RUN: yaml2obj %S/inputs/convert.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to generate this test: +(module + (func (export "convert_i32_u_to_f32") (result f32) + i32.const 10 + f32.convert_i32_u + ) + + (func (export "convert_i32_s_to_f32") (result f32) + i32.const 42 + f32.convert_i32_s + ) + + (func (export "convert_i64_u_to_f32") (result f32) + i64.const 17 + f32.convert_i64_u + ) + + (func (export "convert_i64s_to_f32") (result f32) + i64.const 10 + f32.convert_i64_s + ) + + (func (export "convert_i32_u_to_f64") (result f64) + i32.const 10 + f64.convert_i32_u + ) + + (func (export "convert_i32_s_to_f64") (result f64) + i32.const 42 + f64.convert_i32_s + ) + + (func (export "convert_i64_u_to_f64") (result f64) + i64.const 17 + f64.convert_i64_u + ) + + (func (export "convert_i64s_to_f64") (result f64) + i64.const 10 + f64.convert_i64_s + ) +) +*/ + +// CHECK-LABEL: wasmssa.func exported @convert_i32_u_to_f32() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_u %[[VAL_0]] : i32 to f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func exported @convert_i32_s_to_f32() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 42 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_s %[[VAL_0]] : i32 to f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func exported @convert_i64_u_to_f32() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 17 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_u %[[VAL_0]] : i64 to f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func exported @convert_i64s_to_f32() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_s %[[VAL_0]] : i64 to f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func exported @convert_i32_u_to_f64() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_u %[[VAL_0]] : i32 to f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 + +// CHECK-LABEL: wasmssa.func exported @convert_i32_s_to_f64() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 42 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_s %[[VAL_0]] : i32 to f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 + +// CHECK-LABEL: wasmssa.func exported @convert_i64_u_to_f64() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 17 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_u %[[VAL_0]] : i64 to f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 + +// CHECK-LABEL: wasmssa.func exported @convert_i64s_to_f64() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.convert_s %[[VAL_0]] : i64 to f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/copysign.mlir b/mlir/test/Target/Wasm/copysign.mlir index 33d7a56..90c5b11 100644 --- a/mlir/test/Target/Wasm/copysign.mlir +++ b/mlir/test/Target/Wasm/copysign.mlir @@ -16,14 +16,14 @@ ) */ -// CHECK-LABEL: wasmssa.func @copysign_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @copysign_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.000000e+00 : f32 // CHECK: %[[VAL_2:.*]] = wasmssa.copysign %[[VAL_0]] %[[VAL_1]] : f32 // CHECK: wasmssa.return %[[VAL_2]] : f32 // CHECK: } -// CHECK-LABEL: wasmssa.func @copysign_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @copysign_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.000000e+00 : f64 // CHECK: %[[VAL_2:.*]] = wasmssa.copysign %[[VAL_0]] %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/ctz.mlir b/mlir/test/Target/Wasm/ctz.mlir index 6c0806f..9e7cc5e 100644 --- a/mlir/test/Target/Wasm/ctz.mlir +++ b/mlir/test/Target/Wasm/ctz.mlir @@ -14,12 +14,12 @@ ) */ -// CHECK-LABEL: wasmssa.func @ctz_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @ctz_i32() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.ctz %[[VAL_0]] : i32 // CHECK: wasmssa.return %[[VAL_1]] : i32 -// CHECK-LABEL: wasmssa.func @ctz_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @ctz_i64() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.ctz %[[VAL_0]] : i64 // CHECK: wasmssa.return %[[VAL_1]] : i64 diff --git a/mlir/test/Target/Wasm/demote.mlir b/mlir/test/Target/Wasm/demote.mlir new file mode 100644 index 0000000..3d2bc05 --- /dev/null +++ b/mlir/test/Target/Wasm/demote.mlir @@ -0,0 +1,15 @@ +// RUN: yaml2obj %S/inputs/demote.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module + (func $main (result f32) + f64.const 2.24 + f32.demote_f64 + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 2.240000e+00 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.demote %[[VAL_0]] : f64 to f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 diff --git a/mlir/test/Target/Wasm/div.mlir b/mlir/test/Target/Wasm/div.mlir index c91f780..4967d96 100644 --- a/mlir/test/Target/Wasm/div.mlir +++ b/mlir/test/Target/Wasm/div.mlir @@ -66,61 +66,61 @@ ) */ -// CHECK-LABEL: wasmssa.func @div_u_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @div_u_i32() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 2 : i32 // CHECK: %[[VAL_2:.*]] = wasmssa.div_ui %[[VAL_0]] %[[VAL_1]] : i32 // CHECK: wasmssa.return %[[VAL_2]] : i32 -// CHECK-LABEL: wasmssa.func @div_u_i32_zero() -> i32 { +// CHECK-LABEL: wasmssa.func exported @div_u_i32_zero() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 0 : i32 // CHECK: %[[VAL_2:.*]] = wasmssa.div_ui %[[VAL_0]] %[[VAL_1]] : i32 // CHECK: wasmssa.return %[[VAL_2]] : i32 -// CHECK-LABEL: wasmssa.func @div_s_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @div_s_i32() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 2 : i32 // CHECK: %[[VAL_2:.*]] = wasmssa.div_si %[[VAL_0]] %[[VAL_1]] : i32 // CHECK: wasmssa.return %[[VAL_2]] : i32 -// CHECK-LABEL: wasmssa.func @div_s_i32_zero() -> i32 { +// CHECK-LABEL: wasmssa.func exported @div_s_i32_zero() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 0 : i32 // CHECK: %[[VAL_2:.*]] = wasmssa.div_si %[[VAL_0]] %[[VAL_1]] : i32 // CHECK: wasmssa.return %[[VAL_2]] : i32 -// CHECK-LABEL: wasmssa.func @div_u_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @div_u_i64() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 2 : i64 // CHECK: %[[VAL_2:.*]] = wasmssa.div_ui %[[VAL_0]] %[[VAL_1]] : i64 // CHECK: wasmssa.return %[[VAL_2]] : i64 -// CHECK-LABEL: wasmssa.func @div_u_i64_zero() -> i64 { +// CHECK-LABEL: wasmssa.func exported @div_u_i64_zero() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 0 : i64 // CHECK: %[[VAL_2:.*]] = wasmssa.div_ui %[[VAL_0]] %[[VAL_1]] : i64 // CHECK: wasmssa.return %[[VAL_2]] : i64 -// CHECK-LABEL: wasmssa.func @div_s_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @div_s_i64() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 2 : i64 // CHECK: %[[VAL_2:.*]] = wasmssa.div_si %[[VAL_0]] %[[VAL_1]] : i64 // CHECK: wasmssa.return %[[VAL_2]] : i64 -// CHECK-LABEL: wasmssa.func @div_s_i64_zero() -> i64 { +// CHECK-LABEL: wasmssa.func exported @div_s_i64_zero() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 0 : i64 // CHECK: %[[VAL_2:.*]] = wasmssa.div_si %[[VAL_0]] %[[VAL_1]] : i64 // CHECK: wasmssa.return %[[VAL_2]] : i64 -// CHECK-LABEL: wasmssa.func @div_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @div_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 2.000000e+00 : f32 // CHECK: %[[VAL_2:.*]] = wasmssa.div %[[VAL_0]] %[[VAL_1]] : f32 // CHECK: wasmssa.return %[[VAL_2]] : f32 -// CHECK-LABEL: wasmssa.func @div_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @div_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 2.000000e+00 : f64 // CHECK: %[[VAL_2:.*]] = wasmssa.div %[[VAL_0]] %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/double_nested_loop.mlir b/mlir/test/Target/Wasm/double_nested_loop.mlir new file mode 100644 index 0000000..8b3e499 --- /dev/null +++ b/mlir/test/Target/Wasm/double_nested_loop.mlir @@ -0,0 +1,63 @@ +// RUN: yaml2obj %S/inputs/double_nested_loop.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* +(module + (func + ;; create a local variable and initialize it to 0 + (local $i i32) + (local $j i32) + + (loop $my_loop + + ;; add one to $i + local.get $i + i32.const 1 + i32.add + local.set $i + (loop $my_second_loop (result i32) + i32.const 1 + local.get $j + i32.const 12 + i32.add + local.tee $j + local.get $i + i32.gt_s + br_if $my_second_loop + ) + i32.const 10 + i32.lt_s + br_if $my_loop + ) + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() { +// CHECK: %[[VAL_0:.*]] = wasmssa.local of type i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.local of type i32 +// CHECK: wasmssa.loop : { +// CHECK: %[[VAL_2:.*]] = wasmssa.local_get %[[VAL_0]] : ref to i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_4:.*]] = wasmssa.add %[[VAL_2]] %[[VAL_3]] : i32 +// CHECK: wasmssa.local_set %[[VAL_0]] : ref to i32 to %[[VAL_4]] : i32 +// CHECK: wasmssa.loop : { +// CHECK: %[[VAL_5:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_6:.*]] = wasmssa.local_get %[[VAL_1]] : ref to i32 +// CHECK: %[[VAL_7:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_8:.*]] = wasmssa.add %[[VAL_6]] %[[VAL_7]] : i32 +// CHECK: %[[VAL_9:.*]] = wasmssa.local_tee %[[VAL_1]] : ref to i32 to %[[VAL_8]] : i32 +// CHECK: %[[VAL_10:.*]] = wasmssa.local_get %[[VAL_0]] : ref to i32 +// CHECK: %[[VAL_11:.*]] = wasmssa.gt_si %[[VAL_9]] %[[VAL_10]] : i32 -> i32 +// CHECK: wasmssa.branch_if %[[VAL_11]] to level 0 else ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.block_return %[[VAL_5]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_12:.*]]: i32): +// CHECK: %[[VAL_13:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_14:.*]] = wasmssa.lt_si %[[VAL_12]] %[[VAL_13]] : i32 -> i32 +// CHECK: wasmssa.branch_if %[[VAL_14]] to level 0 else ^bb2 +// CHECK: ^bb2: +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.return diff --git a/mlir/test/Target/Wasm/empty_blocks_list_and_stack.mlir b/mlir/test/Target/Wasm/empty_blocks_list_and_stack.mlir new file mode 100644 index 0000000..5c98f1a --- /dev/null +++ b/mlir/test/Target/Wasm/empty_blocks_list_and_stack.mlir @@ -0,0 +1,53 @@ +// RUN: yaml2obj %S/inputs/empty_blocks_list_and_stack.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module + (func (param $num i32) + (block $b1 + (block $b2 + (block $b3 + ) + ) + ) + ) + + (func (param $num i32) + (block $b1) + (block $b2) + (block $b3) + ) +) + +*/ + +// CHECK-LABEL: wasmssa.func @func_0( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) { +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.return + +// CHECK-LABEL: wasmssa.func @func_1( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) { +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block_return +// CHECK: }> ^bb2 +// CHECK: ^bb2: +// CHECK: wasmssa.block : { +// CHECK: wasmssa.block_return +// CHECK: }> ^bb3 +// CHECK: ^bb3: +// CHECK: wasmssa.return diff --git a/mlir/test/Target/Wasm/eq.mlir b/mlir/test/Target/Wasm/eq.mlir new file mode 100644 index 0000000..ba3ae2f --- /dev/null +++ b/mlir/test/Target/Wasm/eq.mlir @@ -0,0 +1,56 @@ +// RUN: yaml2obj %S/inputs/eq.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s +/* Source code used to create this test: +(module + (func $eq_i32 (result i32) + i32.const 12 + i32.const 50 + i32.eq + ) + + (func $eq_i64 (result i32) + i64.const 20 + i64.const 5 + i64.eq + ) + + (func $eq_f32 (result i32) + f32.const 5 + f32.const 14 + f32.eq + ) + + (func $eq_f64 (result i32) + f64.const 17 + f64.const 0 + f64.eq + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.eq %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 +// CHECK: } + +// CHECK-LABEL: wasmssa.func @func_1() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 20 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 5 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.eq %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 +// CHECK: } + +// CHECK-LABEL: wasmssa.func @func_2() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 +// CHECK: %[[VAL_2:.*]] = wasmssa.eq %[[VAL_0]] %[[VAL_1]] : f32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 +// CHECK: } + +// CHECK-LABEL: wasmssa.func @func_3() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1.700000e+01 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 0.000000e+00 : f64 +// CHECK: %[[VAL_2:.*]] = wasmssa.eq %[[VAL_0]] %[[VAL_1]] : f64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 +// CHECK: } diff --git a/mlir/test/Target/Wasm/eqz.mlir b/mlir/test/Target/Wasm/eqz.mlir new file mode 100644 index 0000000..55cf94a --- /dev/null +++ b/mlir/test/Target/Wasm/eqz.mlir @@ -0,0 +1,21 @@ +// RUN: yaml2obj %S/inputs/eqz.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s +/* Source code used to create this test: +(module + (func (export "eqz_i32") (result i32) + i32.const 13 + i32.eqz) + + (func (export "eqz_i64") (result i32) + i64.const 13 + i64.eqz) +) +*/ +// CHECK-LABEL: wasmssa.func exported @eqz_i32() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 13 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.eqz %[[VAL_0]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_1]] : i32 + +// CHECK-LABEL: wasmssa.func exported @eqz_i64() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 13 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.eqz %[[VAL_0]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_1]] : i32 diff --git a/mlir/test/Target/Wasm/extend.mlir b/mlir/test/Target/Wasm/extend.mlir new file mode 100644 index 0000000..5d4446a --- /dev/null +++ b/mlir/test/Target/Wasm/extend.mlir @@ -0,0 +1,69 @@ +// RUN: yaml2obj %S/inputs/extend.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module + (func $i32_s (result i64) + i32.const 10 + i64.extend_i32_s + ) + (func $i32_u (result i64) + i32.const 10 + i64.extend_i32_u + ) + (func $extend8_32 (result i32) + i32.const 10 + i32.extend8_s + ) + (func $extend16_32 (result i32) + i32.const 10 + i32.extend16_s + ) + (func $extend8_64 (result i64) + i64.const 10 + i64.extend8_s + ) + (func $extend16_64 (result i64) + i64.const 10 + i64.extend16_s + ) + (func $extend32_64 (result i64) + i64.const 10 + i64.extend32_s + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend_i32_s %[[VAL_0]] to i64 +// CHECK: wasmssa.return %[[VAL_1]] : i64 + +// CHECK-LABEL: wasmssa.func @func_1() -> i64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend_i32_u %[[VAL_0]] to i64 +// CHECK: wasmssa.return %[[VAL_1]] : i64 + +// CHECK-LABEL: wasmssa.func @func_2() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend 8 : ui32 low bits from %[[VAL_0]] : i32 +// CHECK: wasmssa.return %[[VAL_1]] : i32 + +// CHECK-LABEL: wasmssa.func @func_3() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend 16 : ui32 low bits from %[[VAL_0]] : i32 +// CHECK: wasmssa.return %[[VAL_1]] : i32 + +// CHECK-LABEL: wasmssa.func @func_4() -> i64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend 8 : ui32 low bits from %[[VAL_0]] : i64 +// CHECK: wasmssa.return %[[VAL_1]] : i64 + +// CHECK-LABEL: wasmssa.func @func_5() -> i64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend 16 : ui32 low bits from %[[VAL_0]] : i64 +// CHECK: wasmssa.return %[[VAL_1]] : i64 + +// CHECK-LABEL: wasmssa.func @func_6() -> i64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.extend 32 : ui32 low bits from %[[VAL_0]] : i64 +// CHECK: wasmssa.return %[[VAL_1]] : i64 diff --git a/mlir/test/Target/Wasm/global.mlir b/mlir/test/Target/Wasm/global.mlir index e72fe69..1e4fe44 100644 --- a/mlir/test/Target/Wasm/global.mlir +++ b/mlir/test/Target/Wasm/global.mlir @@ -29,9 +29,9 @@ i32.add ) */ -// CHECK-LABEL: wasmssa.import_global "from_js" from "env" as @global_0 nested : i32 +// CHECK-LABEL: wasmssa.import_global "from_js" from "env" as @global_0 : i32 -// CHECK-LABEL: wasmssa.func nested @func_0() -> i32 { +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.global_get @global_0 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.global_get @global_1 : i32 // CHECK: %[[VAL_2:.*]] = wasmssa.add %[[VAL_0]] %[[VAL_1]] : i32 @@ -41,26 +41,26 @@ i32.add // CHECK: %[[VAL_6:.*]] = wasmssa.add %[[VAL_2]] %[[VAL_5]] : i32 // CHECK: wasmssa.return %[[VAL_6]] : i32 -// CHECK-LABEL: wasmssa.global @global_1 i32 nested : { +// CHECK-LABEL: wasmssa.global @global_1 i32 : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: wasmssa.return %[[VAL_0]] : i32 -// CHECK-LABEL: wasmssa.global @global_2 i32 mutable nested : { +// CHECK-LABEL: wasmssa.global @global_2 i32 mutable : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: wasmssa.return %[[VAL_0]] : i32 -// CHECK-LABEL: wasmssa.global @global_3 i32 mutable nested : { +// CHECK-LABEL: wasmssa.global @global_3 i32 mutable : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: wasmssa.return %[[VAL_0]] : i32 -// CHECK-LABEL: wasmssa.global @global_4 i64 nested : { +// CHECK-LABEL: wasmssa.global @global_4 i64 : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 11 : i64 // CHECK: wasmssa.return %[[VAL_0]] : i64 -// CHECK-LABEL: wasmssa.global @global_5 f32 nested : { +// CHECK-LABEL: wasmssa.global @global_5 f32 : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.200000e+01 : f32 // CHECK: wasmssa.return %[[VAL_0]] : f32 -// CHECK-LABEL: wasmssa.global @global_6 f64 nested : { +// CHECK-LABEL: wasmssa.global @global_6 f64 : { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.300000e+01 : f64 // CHECK: wasmssa.return %[[VAL_0]] : f64 diff --git a/mlir/test/Target/Wasm/if.mlir b/mlir/test/Target/Wasm/if.mlir new file mode 100644 index 0000000..2d7bfbe --- /dev/null +++ b/mlir/test/Target/Wasm/if.mlir @@ -0,0 +1,112 @@ +// RUN: yaml2obj %S/inputs/if.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to create this test: +(module +(type $intMapper (func (param $input i32) (result i32))) +(func $if_else (type $intMapper) + local.get 0 + i32.const 1 + i32.and + if $isOdd (result i32) + local.get 0 + i32.const 3 + i32.mul + i32.const 1 + i32.add + else + local.get 0 + i32.const 1 + i32.shr_u + end +) + +(func $if_only (type $intMapper) + local.get 0 + local.get 0 + i32.const 1 + i32.and + if $isOdd (type $intMapper) + i32.const 1 + i32.add + end +) + +(func $if_if (type $intMapper) + local.get 0 + i32.ctz + if $isEven (result i32) + i32.const 2 + local.get 0 + i32.const 1 + i32.shr_u + i32.ctz + if $isMultipleOfFour (type $intMapper) + i32.const 2 + i32.add + end + else + i32.const 1 + end +) +) +*/ +// CHECK-LABEL: wasmssa.func @func_0( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.and %[[VAL_0]] %[[VAL_1]] : i32 +// CHECK: wasmssa.if %[[VAL_2]] : { +// CHECK: %[[VAL_3:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_4:.*]] = wasmssa.const 3 : i32 +// CHECK: %[[VAL_5:.*]] = wasmssa.mul %[[VAL_3]] %[[VAL_4]] : i32 +// CHECK: %[[VAL_6:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_7:.*]] = wasmssa.add %[[VAL_5]] %[[VAL_6]] : i32 +// CHECK: wasmssa.block_return %[[VAL_7]] : i32 +// CHECK: } "else "{ +// CHECK: %[[VAL_8:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_9:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_10:.*]] = wasmssa.shr_u %[[VAL_8]] by %[[VAL_9]] bits : i32 +// CHECK: wasmssa.block_return %[[VAL_10]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_11:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_11]] : i32 + +// CHECK-LABEL: wasmssa.func @func_1( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.and %[[VAL_1]] %[[VAL_2]] : i32 +// CHECK: wasmssa.if %[[VAL_3]](%[[VAL_0]]) : i32 : { +// CHECK: ^bb0(%[[VAL_4:.*]]: i32): +// CHECK: %[[VAL_5:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_6:.*]] = wasmssa.add %[[VAL_4]] %[[VAL_5]] : i32 +// CHECK: wasmssa.block_return %[[VAL_6]] : i32 +// CHECK: } > ^bb1 +// CHECK: ^bb1(%[[VAL_7:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_7]] : i32 + +// CHECK-LABEL: wasmssa.func @func_2( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.ctz %[[VAL_0]] : i32 +// CHECK: wasmssa.if %[[VAL_1]] : { +// CHECK: %[[VAL_2:.*]] = wasmssa.const 2 : i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.local_get %[[ARG0]] : ref to i32 +// CHECK: %[[VAL_4:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_5:.*]] = wasmssa.shr_u %[[VAL_3]] by %[[VAL_4]] bits : i32 +// CHECK: %[[VAL_6:.*]] = wasmssa.ctz %[[VAL_5]] : i32 +// CHECK: wasmssa.if %[[VAL_6]](%[[VAL_2]]) : i32 : { +// CHECK: ^bb0(%[[VAL_7:.*]]: i32): +// CHECK: %[[VAL_8:.*]] = wasmssa.const 2 : i32 +// CHECK: %[[VAL_9:.*]] = wasmssa.add %[[VAL_7]] %[[VAL_8]] : i32 +// CHECK: wasmssa.block_return %[[VAL_9]] : i32 +// CHECK: } > ^bb1 +// CHECK: ^bb1(%[[VAL_10:.*]]: i32): +// CHECK: wasmssa.block_return %[[VAL_10]] : i32 +// CHECK: } "else "{ +// CHECK: %[[VAL_11:.*]] = wasmssa.const 1 : i32 +// CHECK: wasmssa.block_return %[[VAL_11]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_12:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_12]] : i32 diff --git a/mlir/test/Target/Wasm/import.mlir b/mlir/test/Target/Wasm/import.mlir index 541dcf3..dcdfa52 100644 --- a/mlir/test/Target/Wasm/import.mlir +++ b/mlir/test/Target/Wasm/import.mlir @@ -11,9 +11,9 @@ ) */ -// CHECK-LABEL: wasmssa.import_func "foo" from "my_module" as @func_0 {sym_visibility = "nested", type = (i32) -> ()} -// CHECK: wasmssa.import_func "bar" from "my_module" as @func_1 {sym_visibility = "nested", type = (i32) -> ()} -// CHECK: wasmssa.import_table "table" from "my_module" as @table_0 {sym_visibility = "nested", type = !wasmssa<tabletype !wasmssa.funcref [2:]>} -// CHECK: wasmssa.import_mem "mem" from "my_module" as @mem_0 {limits = !wasmssa<limit[2:]>, sym_visibility = "nested"} -// CHECK: wasmssa.import_global "glob" from "my_module" as @global_0 nested : i32 -// CHECK: wasmssa.import_global "glob_mut" from "my_other_module" as @global_1 mutable nested : i32 +// CHECK-LABEL: wasmssa.import_func "foo" from "my_module" as @func_0 {type = (i32) -> ()} +// CHECK: wasmssa.import_func "bar" from "my_module" as @func_1 {type = (i32) -> ()} +// CHECK: wasmssa.import_table "table" from "my_module" as @table_0 {type = !wasmssa<tabletype !wasmssa.funcref [2:]>} +// CHECK: wasmssa.import_mem "mem" from "my_module" as @mem_0 {limits = !wasmssa<limit[2:]>} +// CHECK: wasmssa.import_global "glob" from "my_module" as @global_0 : i32 +// CHECK: wasmssa.import_global "glob_mut" from "my_other_module" as @global_1 mutable : i32 diff --git a/mlir/test/Target/Wasm/inputs/add_div.yaml.wasm b/mlir/test/Target/Wasm/inputs/add_div.yaml.wasm new file mode 100644 index 0000000..865c315 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/add_div.yaml.wasm @@ -0,0 +1,50 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: + - I32 + ReturnTypes: + - I32 + - Index: 1 + ParamTypes: + - I32 + - I32 + ReturnTypes: + - I32 + - Type: IMPORT + Imports: + - Module: env + Field: twoTimes + Kind: FUNCTION + SigIndex: 0 + - Type: FUNCTION + FunctionTypes: [ 1 ] + - Type: MEMORY + Memories: + - Minimum: 0x2 + - Type: GLOBAL + Globals: + - Index: 0 + Type: I32 + Mutable: true + InitExpr: + Opcode: I32_CONST + Value: 66560 + - Type: EXPORT + Exports: + - Name: memory + Kind: MEMORY + Index: 0 + - Name: add + Kind: FUNCTION + Index: 1 + - Type: CODE + Functions: + - Index: 1 + Locals: [] + Body: 20001000200110006A41026D0B +... diff --git a/mlir/test/Target/Wasm/inputs/block.yaml.wasm b/mlir/test/Target/Wasm/inputs/block.yaml.wasm new file mode 100644 index 0000000..dd5118a --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/block.yaml.wasm @@ -0,0 +1,22 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: [] + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: EXPORT + Exports: + - Name: i_am_a_block + Kind: FUNCTION + Index: 0 + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 02400B0B +... diff --git a/mlir/test/Target/Wasm/inputs/block_complete_type.yaml.wasm b/mlir/test/Target/Wasm/inputs/block_complete_type.yaml.wasm new file mode 100644 index 0000000..7a125bf --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/block_complete_type.yaml.wasm @@ -0,0 +1,23 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: + - I32 + ReturnTypes: + - I32 + - Index: 1 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 1 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410E020041016A0B0B +... diff --git a/mlir/test/Target/Wasm/inputs/block_value_type.yaml.wasm b/mlir/test/Target/Wasm/inputs/block_value_type.yaml.wasm new file mode 100644 index 0000000..4ba291d --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/block_value_type.yaml.wasm @@ -0,0 +1,18 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 027F41110B0B +... diff --git a/mlir/test/Target/Wasm/inputs/branch_if.yaml.wasm b/mlir/test/Target/Wasm/inputs/branch_if.yaml.wasm new file mode 100644 index 0000000..40536ed --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/branch_if.yaml.wasm @@ -0,0 +1,18 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 027F410141020D0041016A0B0B +... diff --git a/mlir/test/Target/Wasm/inputs/call.yaml.wasm b/mlir/test/Target/Wasm/inputs/call.yaml.wasm new file mode 100644 index 0000000..535a623 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/call.yaml.wasm @@ -0,0 +1,26 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0 ] + - Type: EXPORT + Exports: + - Name: forty_two + Kind: FUNCTION + Index: 1 + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 412A0B + - Index: 1 + Locals: [] + Body: 10000B +... diff --git a/mlir/test/Target/Wasm/inputs/comparison_ops.yaml.wasm b/mlir/test/Target/Wasm/inputs/comparison_ops.yaml.wasm new file mode 100644 index 0000000..cde9ee1 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/comparison_ops.yaml.wasm @@ -0,0 +1,88 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410C4132480B + - Index: 1 + Locals: [] + Body: 410C41324C0B + - Index: 2 + Locals: [] + Body: 410C4132490B + - Index: 3 + Locals: [] + Body: 410C41324D0B + - Index: 4 + Locals: [] + Body: 410C41324A0B + - Index: 5 + Locals: [] + Body: 410C41324B0B + - Index: 6 + Locals: [] + Body: 410C41324E0B + - Index: 7 + Locals: [] + Body: 410C41324F0B + - Index: 8 + Locals: [] + Body: 420C4232530B + - Index: 9 + Locals: [] + Body: 420C4232570B + - Index: 10 + Locals: [] + Body: 420C4232540B + - Index: 11 + Locals: [] + Body: 420C4232580B + - Index: 12 + Locals: [] + Body: 420C4232550B + - Index: 13 + Locals: [] + Body: 420C4232560B + - Index: 14 + Locals: [] + Body: 420C4232590B + - Index: 15 + Locals: [] + Body: 420C42325A0B + - Index: 16 + Locals: [] + Body: 430000A04043000060415D0B + - Index: 17 + Locals: [] + Body: 430000A04043000060415F0B + - Index: 18 + Locals: [] + Body: 430000A04043000060415E0B + - Index: 19 + Locals: [] + Body: 430000A0404300006041600B + - Index: 20 + Locals: [] + Body: 440000000000001440440000000000002C40630B + - Index: 21 + Locals: [] + Body: 440000000000001440440000000000002C40650B + - Index: 22 + Locals: [] + Body: 440000000000001440440000000000002C40640B + - Index: 23 + Locals: [] + Body: 440000000000001440440000000000002C40660B +... diff --git a/mlir/test/Target/Wasm/inputs/convert.yaml.wasm b/mlir/test/Target/Wasm/inputs/convert.yaml.wasm new file mode 100644 index 0000000..c346a75 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/convert.yaml.wasm @@ -0,0 +1,69 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - F32 + - Index: 1 + ParamTypes: [] + ReturnTypes: + - F64 + - Type: FUNCTION + FunctionTypes: [ 0, 0, 0, 0, 1, 1, 1, 1 ] + - Type: EXPORT + Exports: + - Name: convert_i32_u_to_f32 + Kind: FUNCTION + Index: 0 + - Name: convert_i32_s_to_f32 + Kind: FUNCTION + Index: 1 + - Name: convert_i64_u_to_f32 + Kind: FUNCTION + Index: 2 + - Name: convert_i64s_to_f32 + Kind: FUNCTION + Index: 3 + - Name: convert_i32_u_to_f64 + Kind: FUNCTION + Index: 4 + - Name: convert_i32_s_to_f64 + Kind: FUNCTION + Index: 5 + - Name: convert_i64_u_to_f64 + Kind: FUNCTION + Index: 6 + - Name: convert_i64s_to_f64 + Kind: FUNCTION + Index: 7 + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410AB30B + - Index: 1 + Locals: [] + Body: 412AB20B + - Index: 2 + Locals: [] + Body: 4211B50B + - Index: 3 + Locals: [] + Body: 420AB40B + - Index: 4 + Locals: [] + Body: 410AB80B + - Index: 5 + Locals: [] + Body: 412AB70B + - Index: 6 + Locals: [] + Body: 4211BA0B + - Index: 7 + Locals: [] + Body: 420AB90B +... diff --git a/mlir/test/Target/Wasm/inputs/demote.yaml.wasm b/mlir/test/Target/Wasm/inputs/demote.yaml.wasm new file mode 100644 index 0000000..3997045 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/demote.yaml.wasm @@ -0,0 +1,18 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - F32 + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 44EC51B81E85EB0140B60B +... diff --git a/mlir/test/Target/Wasm/inputs/double_nested_loop.yaml.wasm b/mlir/test/Target/Wasm/inputs/double_nested_loop.yaml.wasm new file mode 100644 index 0000000..41a2944 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/double_nested_loop.yaml.wasm @@ -0,0 +1,19 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: [] + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: + - Type: I32 + Count: 2 + Body: 0340200041016A2100037F41012001410C6A220120004A0D000B410A480D000B0B +... diff --git a/mlir/test/Target/Wasm/inputs/empty_blocks_list_and_stack.yaml.wasm b/mlir/test/Target/Wasm/inputs/empty_blocks_list_and_stack.yaml.wasm new file mode 100644 index 0000000..3171409 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/empty_blocks_list_and_stack.yaml.wasm @@ -0,0 +1,21 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: + - I32 + ReturnTypes: [] + - Type: FUNCTION + FunctionTypes: [ 0, 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 0240024002400B0B0B0B + - Index: 1 + Locals: [] + Body: 02400B02400B02400B0B +... diff --git a/mlir/test/Target/Wasm/inputs/eq.yaml.wasm b/mlir/test/Target/Wasm/inputs/eq.yaml.wasm new file mode 100644 index 0000000..1998369 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/eq.yaml.wasm @@ -0,0 +1,27 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0, 0, 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410C4132460B + - Index: 1 + Locals: [] + Body: 42144205510B + - Index: 2 + Locals: [] + Body: 430000A04043000060415B0B + - Index: 3 + Locals: [] + Body: 440000000000003140440000000000000000610B +... diff --git a/mlir/test/Target/Wasm/inputs/eqz.yaml.wasm b/mlir/test/Target/Wasm/inputs/eqz.yaml.wasm new file mode 100644 index 0000000..894ac50 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/eqz.yaml.wasm @@ -0,0 +1,29 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0 ] + - Type: EXPORT + Exports: + - Name: eqz_i32 + Kind: FUNCTION + Index: 0 + - Name: eqz_i64 + Kind: FUNCTION + Index: 1 + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410D450B + - Index: 1 + Locals: [] + Body: 420D500B +... diff --git a/mlir/test/Target/Wasm/inputs/extend.yaml.wasm b/mlir/test/Target/Wasm/inputs/extend.yaml.wasm new file mode 100644 index 0000000..7e872ba --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/extend.yaml.wasm @@ -0,0 +1,40 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I64 + - Index: 1 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0, 1, 1, 0, 0, 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410AAC0B + - Index: 1 + Locals: [] + Body: 410AAD0B + - Index: 2 + Locals: [] + Body: 410AC00B + - Index: 3 + Locals: [] + Body: 410AC10B + - Index: 4 + Locals: [] + Body: 420AC20B + - Index: 5 + Locals: [] + Body: 420AC30B + - Index: 6 + Locals: [] + Body: 420AC40B +... diff --git a/mlir/test/Target/Wasm/inputs/if.yaml.wasm b/mlir/test/Target/Wasm/inputs/if.yaml.wasm new file mode 100644 index 0000000..ccc38f6 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/if.yaml.wasm @@ -0,0 +1,25 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: + - I32 + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0, 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 2000410171047F200041036C41016A0520004101760B0B + - Index: 1 + Locals: [] + Body: 20002000410171040041016A0B0B + - Index: 2 + Locals: [] + Body: 200068047F4102200041017668040041026A0B0541010B0B +... diff --git a/mlir/test/Target/Wasm/inputs/loop.yaml.wasm b/mlir/test/Target/Wasm/inputs/loop.yaml.wasm new file mode 100644 index 0000000..9d33894 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/loop.yaml.wasm @@ -0,0 +1,17 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: [] + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 03400B0B +... diff --git a/mlir/test/Target/Wasm/inputs/loop_with_inst.yaml.wasm b/mlir/test/Target/Wasm/inputs/loop_with_inst.yaml.wasm new file mode 100644 index 0000000..4b8cc54 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/loop_with_inst.yaml.wasm @@ -0,0 +1,20 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: + - Type: I32 + Count: 1 + Body: 037F200041016A21002000410A480B0B +... diff --git a/mlir/test/Target/Wasm/inputs/ne.yaml.wasm b/mlir/test/Target/Wasm/inputs/ne.yaml.wasm new file mode 100644 index 0000000..0167519 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/ne.yaml.wasm @@ -0,0 +1,27 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0, 0, 0, 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410C4132470B + - Index: 1 + Locals: [] + Body: 42144205520B + - Index: 2 + Locals: [] + Body: 430000A04043000060415C0B + - Index: 3 + Locals: [] + Body: 440000000000003140440000000000000000620B +... diff --git a/mlir/test/Target/Wasm/inputs/promote.yaml.wasm b/mlir/test/Target/Wasm/inputs/promote.yaml.wasm new file mode 100644 index 0000000..d38603e --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/promote.yaml.wasm @@ -0,0 +1,18 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - F64 + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 4300002841BB0B +... diff --git a/mlir/test/Target/Wasm/inputs/reinterpret.yaml.wasm b/mlir/test/Target/Wasm/inputs/reinterpret.yaml.wasm new file mode 100644 index 0000000..c01c1b1 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/reinterpret.yaml.wasm @@ -0,0 +1,53 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - I32 + - Index: 1 + ParamTypes: [] + ReturnTypes: + - I64 + - Index: 2 + ParamTypes: [] + ReturnTypes: + - F32 + - Index: 3 + ParamTypes: [] + ReturnTypes: + - F64 + - Type: FUNCTION + FunctionTypes: [ 0, 1, 2, 3 ] + - Type: EXPORT + Exports: + - Name: i32.reinterpret_f32 + Kind: FUNCTION + Index: 0 + - Name: i64.reinterpret_f64 + Kind: FUNCTION + Index: 1 + - Name: f32.reinterpret_i32 + Kind: FUNCTION + Index: 2 + - Name: f64.reinterpret_i64 + Kind: FUNCTION + Index: 3 + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 43000080BFBC0B + - Index: 1 + Locals: [] + Body: 44000000000000F0BFBD0B + - Index: 2 + Locals: [] + Body: 417FBE0B + - Index: 3 + Locals: [] + Body: 427FBF0B +... diff --git a/mlir/test/Target/Wasm/inputs/rounding.yaml.wasm b/mlir/test/Target/Wasm/inputs/rounding.yaml.wasm new file mode 100644 index 0000000..c6e8bf6 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/rounding.yaml.wasm @@ -0,0 +1,37 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: [] + ReturnTypes: + - F64 + - Index: 1 + ParamTypes: [] + ReturnTypes: + - F32 + - Type: FUNCTION + FunctionTypes: [ 0, 1, 0, 1, 0, 1 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 4433333333333328C09B0B + - Index: 1 + Locals: [] + Body: 43A01ACF3F8D0B + - Index: 2 + Locals: [] + Body: 4433333333333328C09C0B + - Index: 3 + Locals: [] + Body: 43A01ACF3F8E0B + - Index: 4 + Locals: [] + Body: 4433333333333328C09D0B + - Index: 5 + Locals: [] + Body: 43A01ACF3F8F0B +... diff --git a/mlir/test/Target/Wasm/inputs/wrap.yaml.wasm b/mlir/test/Target/Wasm/inputs/wrap.yaml.wasm new file mode 100644 index 0000000..51c0b02 --- /dev/null +++ b/mlir/test/Target/Wasm/inputs/wrap.yaml.wasm @@ -0,0 +1,24 @@ +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: + - I64 + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 0 ] + - Type: EXPORT + Exports: + - Name: i64_wrap + Kind: FUNCTION + Index: 0 + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 2000A70B +... diff --git a/mlir/test/Target/Wasm/invalid_block_type_index.yaml b/mlir/test/Target/Wasm/invalid_block_type_index.yaml new file mode 100644 index 0000000..5b83e2e --- /dev/null +++ b/mlir/test/Target/Wasm/invalid_block_type_index.yaml @@ -0,0 +1,28 @@ + +# RUN: yaml2obj %s | not mlir-translate --import-wasm -o - 2>&1 | FileCheck %s + +# CHECK: type index references nonexistent type (2) + +--- !WASM +FileHeader: + Version: 0x1 +Sections: + - Type: TYPE + Signatures: + - Index: 0 + ParamTypes: + - I32 + ReturnTypes: + - I32 + - Index: 1 + ParamTypes: [] + ReturnTypes: + - I32 + - Type: FUNCTION + FunctionTypes: [ 1 ] + - Type: CODE + Functions: + - Index: 0 + Locals: [] + Body: 410E020241016A0B0B +# -----------------------------^^ Invalid type ID diff --git a/mlir/test/Target/Wasm/local.mlir b/mlir/test/Target/Wasm/local.mlir index 32f5900..9844f9c 100644 --- a/mlir/test/Target/Wasm/local.mlir +++ b/mlir/test/Target/Wasm/local.mlir @@ -29,7 +29,7 @@ ) */ -// CHECK-LABEL: wasmssa.func nested @func_0() -> f32 { +// CHECK-LABEL: wasmssa.func @func_0() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.local of type f32 // CHECK: %[[VAL_1:.*]] = wasmssa.local of type f32 // CHECK: %[[VAL_2:.*]] = wasmssa.const 8.000000e+00 : f32 @@ -40,7 +40,7 @@ // CHECK: %[[VAL_6:.*]] = wasmssa.add %[[VAL_3]] %[[VAL_5]] : f32 // CHECK: wasmssa.return %[[VAL_6]] : f32 -// CHECK-LABEL: wasmssa.func nested @func_1() -> i32 { +// CHECK-LABEL: wasmssa.func @func_1() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.local of type i32 // CHECK: %[[VAL_1:.*]] = wasmssa.local of type i32 // CHECK: %[[VAL_2:.*]] = wasmssa.const 8 : i32 @@ -51,7 +51,7 @@ // CHECK: %[[VAL_6:.*]] = wasmssa.add %[[VAL_3]] %[[VAL_5]] : i32 // CHECK: wasmssa.return %[[VAL_6]] : i32 -// CHECK-LABEL: wasmssa.func nested @func_2( +// CHECK-LABEL: wasmssa.func @func_2( // CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i32>) -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 3 : i32 // CHECK: wasmssa.local_set %[[ARG0]] : ref to i32 to %[[VAL_0]] : i32 diff --git a/mlir/test/Target/Wasm/loop.mlir b/mlir/test/Target/Wasm/loop.mlir new file mode 100644 index 0000000..29ad502 --- /dev/null +++ b/mlir/test/Target/Wasm/loop.mlir @@ -0,0 +1,17 @@ +// RUN: yaml2obj %S/inputs/loop.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* IR generated from: +(module + (func + (loop $my_loop + ) + ) +)*/ + +// CHECK-LABEL: wasmssa.func @func_0() { +// CHECK: wasmssa.loop : { +// CHECK: wasmssa.block_return +// CHECK: }> ^bb1 +// CHECK: ^bb1: +// CHECK: wasmssa.return +// CHECK: } diff --git a/mlir/test/Target/Wasm/loop_with_inst.mlir b/mlir/test/Target/Wasm/loop_with_inst.mlir new file mode 100644 index 0000000..311d007 --- /dev/null +++ b/mlir/test/Target/Wasm/loop_with_inst.mlir @@ -0,0 +1,33 @@ +// RUN: yaml2obj %S/inputs/loop_with_inst.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Code used to create this test: + +(module + (func (result i32) + (local $i i32) + (loop $my_loop (result i32) + local.get $i + i32.const 1 + i32.add + local.set $i + local.get $i + i32.const 10 + i32.lt_s + ) + ) +)*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.local of type i32 +// CHECK: wasmssa.loop : { +// CHECK: %[[VAL_1:.*]] = wasmssa.local_get %[[VAL_0]] : ref to i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.const 1 : i32 +// CHECK: %[[VAL_3:.*]] = wasmssa.add %[[VAL_1]] %[[VAL_2]] : i32 +// CHECK: wasmssa.local_set %[[VAL_0]] : ref to i32 to %[[VAL_3]] : i32 +// CHECK: %[[VAL_4:.*]] = wasmssa.local_get %[[VAL_0]] : ref to i32 +// CHECK: %[[VAL_5:.*]] = wasmssa.const 10 : i32 +// CHECK: %[[VAL_6:.*]] = wasmssa.lt_si %[[VAL_4]] %[[VAL_5]] : i32 -> i32 +// CHECK: wasmssa.block_return %[[VAL_6]] : i32 +// CHECK: }> ^bb1 +// CHECK: ^bb1(%[[VAL_7:.*]]: i32): +// CHECK: wasmssa.return %[[VAL_7]] : i32 diff --git a/mlir/test/Target/Wasm/max.mlir b/mlir/test/Target/Wasm/max.mlir index 4ef2042..9160bde 100644 --- a/mlir/test/Target/Wasm/max.mlir +++ b/mlir/test/Target/Wasm/max.mlir @@ -16,14 +16,14 @@ ) */ -// CHECK-LABEL: wasmssa.func @min_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @min_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.000000e+00 : f32 // CHECK: %[[VAL_2:.*]] = wasmssa.max %[[VAL_0]] %[[VAL_1]] : f32 // CHECK: wasmssa.return %[[VAL_2]] : f32 -// CHECK-LABEL: wasmssa.func @min_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @min_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.000000e+00 : f64 // CHECK: %[[VAL_2:.*]] = wasmssa.max %[[VAL_0]] %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/memory_min_eq_max.mlir b/mlir/test/Target/Wasm/memory_min_eq_max.mlir index 2ba5ab5..ea8f719 100644 --- a/mlir/test/Target/Wasm/memory_min_eq_max.mlir +++ b/mlir/test/Target/Wasm/memory_min_eq_max.mlir @@ -4,4 +4,4 @@ (module (memory 0 0)) */ -// CHECK-LABEL: wasmssa.memory @mem_0 nested !wasmssa<limit[0: 0]> +// CHECK-LABEL: wasmssa.memory @mem_0 !wasmssa<limit[0: 0]> diff --git a/mlir/test/Target/Wasm/memory_min_max.mlir b/mlir/test/Target/Wasm/memory_min_max.mlir index ebf6418..88782ec 100644 --- a/mlir/test/Target/Wasm/memory_min_max.mlir +++ b/mlir/test/Target/Wasm/memory_min_max.mlir @@ -4,4 +4,4 @@ (module (memory 0 65536)) */ -// CHECK-LABEL: wasmssa.memory @mem_0 nested !wasmssa<limit[0: 65536]> +// CHECK-LABEL: wasmssa.memory @mem_0 !wasmssa<limit[0: 65536]> diff --git a/mlir/test/Target/Wasm/memory_min_no_max.mlir b/mlir/test/Target/Wasm/memory_min_no_max.mlir index 8d88786..c10c5cc 100644 --- a/mlir/test/Target/Wasm/memory_min_no_max.mlir +++ b/mlir/test/Target/Wasm/memory_min_no_max.mlir @@ -4,4 +4,4 @@ (module (memory 1)) */ -// CHECK-LABEL: wasmssa.memory @mem_0 nested !wasmssa<limit[1:]> +// CHECK-LABEL: wasmssa.memory @mem_0 !wasmssa<limit[1:]> diff --git a/mlir/test/Target/Wasm/min.mlir b/mlir/test/Target/Wasm/min.mlir index 1058c7d..2372bcc 100644 --- a/mlir/test/Target/Wasm/min.mlir +++ b/mlir/test/Target/Wasm/min.mlir @@ -16,13 +16,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @min_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @min_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.000000e+00 : f32 // CHECK: %[[VAL_2:.*]] = wasmssa.min %[[VAL_0]] %[[VAL_1]] : f32 // CHECK: wasmssa.return %[[VAL_2]] : f32 -// CHECK-LABEL: wasmssa.func @min_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @min_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.000000e+00 : f64 // CHECK: %[[VAL_2:.*]] = wasmssa.min %[[VAL_0]] %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/ne.mlir b/mlir/test/Target/Wasm/ne.mlir new file mode 100644 index 0000000..331df75 --- /dev/null +++ b/mlir/test/Target/Wasm/ne.mlir @@ -0,0 +1,52 @@ +// RUN: yaml2obj %S/inputs/ne.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s +/* Source code used to create this test: +(module + (func $ne_i32 (result i32) + i32.const 12 + i32.const 50 + i32.ne + ) + + (func $ne_i64 (result i32) + i64.const 20 + i64.const 5 + i64.ne + ) + + (func $ne_f32 (result i32) + f32.const 5 + f32.const 14 + f32.ne + ) + + (func $ne_f64 (result i32) + f64.const 17 + f64.const 0 + f64.ne + ) +) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 +// CHECK: %[[VAL_2:.*]] = wasmssa.ne %[[VAL_0]] %[[VAL_1]] : i32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_1() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 20 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 5 : i64 +// CHECK: %[[VAL_2:.*]] = wasmssa.ne %[[VAL_0]] %[[VAL_1]] : i64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_2() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 +// CHECK: %[[VAL_2:.*]] = wasmssa.ne %[[VAL_0]] %[[VAL_1]] : f32 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 + +// CHECK-LABEL: wasmssa.func @func_3() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1.700000e+01 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.const 0.000000e+00 : f64 +// CHECK: %[[VAL_2:.*]] = wasmssa.ne %[[VAL_0]] %[[VAL_1]] : f64 -> i32 +// CHECK: wasmssa.return %[[VAL_2]] : i32 diff --git a/mlir/test/Target/Wasm/neg.mlir b/mlir/test/Target/Wasm/neg.mlir index 5811ab50..dae8ee5 100644 --- a/mlir/test/Target/Wasm/neg.mlir +++ b/mlir/test/Target/Wasm/neg.mlir @@ -12,12 +12,12 @@ ) */ -// CHECK-LABEL: wasmssa.func @neg_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @neg_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.neg %[[VAL_0]] : f32 // CHECK: wasmssa.return %[[VAL_1]] : f32 -// CHECK-LABEL: wasmssa.func @neg_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @neg_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.neg %[[VAL_0]] : f64 // CHECK: wasmssa.return %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/or.mlir b/mlir/test/Target/Wasm/or.mlir index 521f2ba..be0b3d7 100644 --- a/mlir/test/Target/Wasm/or.mlir +++ b/mlir/test/Target/Wasm/or.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @or_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @or_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.or %0 %1 : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @or_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @or_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.or %0 %1 : i64 diff --git a/mlir/test/Target/Wasm/popcnt.mlir b/mlir/test/Target/Wasm/popcnt.mlir index 235333a..bfaa8eb 100644 --- a/mlir/test/Target/Wasm/popcnt.mlir +++ b/mlir/test/Target/Wasm/popcnt.mlir @@ -14,12 +14,12 @@ ) */ -// CHECK-LABEL: wasmssa.func @popcnt_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @popcnt_i32() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.popcnt %[[VAL_0]] : i32 // CHECK: wasmssa.return %[[VAL_1]] : i32 -// CHECK-LABEL: wasmssa.func @popcnt_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @popcnt_i64() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 10 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.popcnt %[[VAL_0]] : i64 // CHECK: wasmssa.return %[[VAL_1]] : i64 diff --git a/mlir/test/Target/Wasm/promote.mlir b/mlir/test/Target/Wasm/promote.mlir new file mode 100644 index 0000000..44c31b6 --- /dev/null +++ b/mlir/test/Target/Wasm/promote.mlir @@ -0,0 +1,14 @@ +// RUN: yaml2obj %S/inputs/promote.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* Source code used to generate this test: +(module + (func $main (result f64) + f32.const 10.5 + f64.promote_f32 + ) +)*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1.050000e+01 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.promote %[[VAL_0]] : f32 to f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/reinterpret.mlir b/mlir/test/Target/Wasm/reinterpret.mlir new file mode 100644 index 0000000..574d13f --- /dev/null +++ b/mlir/test/Target/Wasm/reinterpret.mlir @@ -0,0 +1,46 @@ +// RUN: yaml2obj %S/inputs/reinterpret.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s + +/* +Test generated from: +(module + (func (export "i32.reinterpret_f32") (result i32) + f32.const -1 + i32.reinterpret_f32 + ) + + (func (export "i64.reinterpret_f64") (result i64) + f64.const -1 + i64.reinterpret_f64 + ) + + (func (export "f32.reinterpret_i32") (result f32) + i32.const -1 + f32.reinterpret_i32 + ) + + (func (export "f64.reinterpret_i64") (result f64) + i64.const -1 + f64.reinterpret_i64 + ) +) +*/ + +// CHECK-LABEL: wasmssa.func exported @i32.reinterpret_f32() -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1.000000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.reinterpret %[[VAL_0]] : f32 as i32 +// CHECK: wasmssa.return %[[VAL_1]] : i32 + +// CHECK-LABEL: wasmssa.func exported @i64.reinterpret_f64() -> i64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1.000000e+00 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.reinterpret %[[VAL_0]] : f64 as i64 +// CHECK: wasmssa.return %[[VAL_1]] : i64 + +// CHECK-LABEL: wasmssa.func exported @f32.reinterpret_i32() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1 : i32 +// CHECK: %[[VAL_1:.*]] = wasmssa.reinterpret %[[VAL_0]] : i32 as f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func exported @f64.reinterpret_i64() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1 : i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.reinterpret %[[VAL_0]] : i64 as f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/rem.mlir b/mlir/test/Target/Wasm/rem.mlir index b19b8d9..16c9c78 100644 --- a/mlir/test/Target/Wasm/rem.mlir +++ b/mlir/test/Target/Wasm/rem.mlir @@ -24,28 +24,28 @@ ) */ -// CHECK-LABEL: wasmssa.func @rem_u_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @rem_u_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.rem_ui %0 %1 : i32 // CHECK: wasmssa.return %2 : i32 // CHECK: } -// CHECK-LABEL: wasmssa.func @rem_u_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @rem_u_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.rem_ui %0 %1 : i64 // CHECK: wasmssa.return %2 : i64 // CHECK: } -// CHECK-LABEL: wasmssa.func @rem_s_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @rem_s_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.rem_si %0 %1 : i32 // CHECK: wasmssa.return %2 : i32 // CHECK: } -// CHECK-LABEL: wasmssa.func @rem_s_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @rem_s_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.rem_si %0 %1 : i64 diff --git a/mlir/test/Target/Wasm/rotl.mlir b/mlir/test/Target/Wasm/rotl.mlir index ec573554..4c2e5af 100644 --- a/mlir/test/Target/Wasm/rotl.mlir +++ b/mlir/test/Target/Wasm/rotl.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @rotl_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @rotl_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.rotl %0 by %1 bits : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @rotl_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @rotl_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.rotl %0 by %1 bits : i64 diff --git a/mlir/test/Target/Wasm/rotr.mlir b/mlir/test/Target/Wasm/rotr.mlir index 5618b43..ec403d0 100644 --- a/mlir/test/Target/Wasm/rotr.mlir +++ b/mlir/test/Target/Wasm/rotr.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @rotr_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @rotr_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.rotr %0 by %1 bits : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @rotr_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @rotr_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.rotr %0 by %1 bits : i64 diff --git a/mlir/test/Target/Wasm/rounding.mlir b/mlir/test/Target/Wasm/rounding.mlir new file mode 100644 index 0000000..947637e --- /dev/null +++ b/mlir/test/Target/Wasm/rounding.mlir @@ -0,0 +1,50 @@ +// RUN: yaml2obj %S/inputs/rounding.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s +/* Source code used to create this test: +(module + (func $ceil_f64 (result f64) + f64.const -12.1 + f64.ceil + ) + (func $ceil_f32 (result f32) + f32.const 1.618 + f32.ceil + ) + (func $floor_f64 (result f64) + f64.const -12.1 + f64.floor + ) + (func $floor_f32 (result f32) + f32.const 1.618 + f32.floor + ) +*/ + +// CHECK-LABEL: wasmssa.func @func_0() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1.210000e+01 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.ceil %[[VAL_0]] : f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 + +// CHECK-LABEL: wasmssa.func @func_1() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1.618000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.ceil %[[VAL_0]] : f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func @func_2() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1.210000e+01 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.floor %[[VAL_0]] : f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 + +// CHECK-LABEL: wasmssa.func @func_3() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1.618000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.floor %[[VAL_0]] : f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 + +// CHECK-LABEL: wasmssa.func @func_4() -> f64 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const -1.210000e+01 : f64 +// CHECK: %[[VAL_1:.*]] = wasmssa.trunc %[[VAL_0]] : f64 +// CHECK: wasmssa.return %[[VAL_1]] : f64 + +// CHECK-LABEL: wasmssa.func @func_5() -> f32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.const 1.618000e+00 : f32 +// CHECK: %[[VAL_1:.*]] = wasmssa.trunc %[[VAL_0]] : f32 +// CHECK: wasmssa.return %[[VAL_1]] : f32 diff --git a/mlir/test/Target/Wasm/shl.mlir b/mlir/test/Target/Wasm/shl.mlir index f2bdd57..1363112 100644 --- a/mlir/test/Target/Wasm/shl.mlir +++ b/mlir/test/Target/Wasm/shl.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @shl_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @shl_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.shl %0 by %1 bits : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @shl_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @shl_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.shl %0 by %1 bits : i64 diff --git a/mlir/test/Target/Wasm/shr_s.mlir b/mlir/test/Target/Wasm/shr_s.mlir index 247d9be..da1a38f 100644 --- a/mlir/test/Target/Wasm/shr_s.mlir +++ b/mlir/test/Target/Wasm/shr_s.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @shr_s_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @shr_s_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.shr_s %0 by %1 bits : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @shr_s_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @shr_s_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.shr_s %0 by %1 bits : i64 diff --git a/mlir/test/Target/Wasm/shr_u.mlir b/mlir/test/Target/Wasm/shr_u.mlir index 9a79eed..2991c2a 100644 --- a/mlir/test/Target/Wasm/shr_u.mlir +++ b/mlir/test/Target/Wasm/shr_u.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @shr_u_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @shr_u_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.shr_u %0 by %1 bits : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @shr_u_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @shr_u_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.shr_u %0 by %1 bits : i64 diff --git a/mlir/test/Target/Wasm/sqrt.mlir b/mlir/test/Target/Wasm/sqrt.mlir index 77444ad..6b968d6 100644 --- a/mlir/test/Target/Wasm/sqrt.mlir +++ b/mlir/test/Target/Wasm/sqrt.mlir @@ -12,12 +12,12 @@ ) */ -// CHECK-LABEL: wasmssa.func @sqrt_f32() -> f32 { +// CHECK-LABEL: wasmssa.func exported @sqrt_f32() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.sqrt %[[VAL_0]] : f32 // CHECK: wasmssa.return %[[VAL_1]] : f32 -// CHECK-LABEL: wasmssa.func @sqrt_f64() -> f64 { +// CHECK-LABEL: wasmssa.func exported @sqrt_f64() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.000000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.sqrt %[[VAL_0]] : f64 // CHECK: wasmssa.return %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/sub.mlir b/mlir/test/Target/Wasm/sub.mlir index b9c6caf..5b242f4 100644 --- a/mlir/test/Target/Wasm/sub.mlir +++ b/mlir/test/Target/Wasm/sub.mlir @@ -27,25 +27,25 @@ ) */ -// CHECK-LABEL: wasmssa.func nested @func_0() -> i32 { +// CHECK-LABEL: wasmssa.func @func_0() -> i32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 12 : i32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 50 : i32 // CHECK: %[[VAL_2:.*]] = wasmssa.sub %[[VAL_0]] %[[VAL_1]] : i32 // CHECK: wasmssa.return %[[VAL_2]] : i32 -// CHECK-LABEL: wasmssa.func nested @func_1() -> i64 { +// CHECK-LABEL: wasmssa.func @func_1() -> i64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 20 : i64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 5 : i64 // CHECK: %[[VAL_2:.*]] = wasmssa.sub %[[VAL_0]] %[[VAL_1]] : i64 // CHECK: wasmssa.return %[[VAL_2]] : i64 -// CHECK-LABEL: wasmssa.func nested @func_2() -> f32 { +// CHECK-LABEL: wasmssa.func @func_2() -> f32 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 5.000000e+00 : f32 // CHECK: %[[VAL_1:.*]] = wasmssa.const 1.400000e+01 : f32 // CHECK: %[[VAL_2:.*]] = wasmssa.sub %[[VAL_0]] %[[VAL_1]] : f32 // CHECK: wasmssa.return %[[VAL_2]] : f32 -// CHECK-LABEL: wasmssa.func nested @func_3() -> f64 { +// CHECK-LABEL: wasmssa.func @func_3() -> f64 { // CHECK: %[[VAL_0:.*]] = wasmssa.const 1.700000e+01 : f64 // CHECK: %[[VAL_1:.*]] = wasmssa.const 0.000000e+00 : f64 // CHECK: %[[VAL_2:.*]] = wasmssa.sub %[[VAL_0]] %[[VAL_1]] : f64 diff --git a/mlir/test/Target/Wasm/wrap.mlir b/mlir/test/Target/Wasm/wrap.mlir new file mode 100644 index 0000000..1266758 --- /dev/null +++ b/mlir/test/Target/Wasm/wrap.mlir @@ -0,0 +1,15 @@ +// RUN: yaml2obj %S/inputs/wrap.yaml.wasm -o - | mlir-translate --import-wasm | FileCheck %s +/* Source code used to create this test: +(module + (func (export "i64_wrap") (param $in i64) (result i32) + local.get $in + i32.wrap_i64 + ) +) +*/ + +// CHECK-LABEL: wasmssa.func exported @i64_wrap( +// CHECK-SAME: %[[ARG0:.*]]: !wasmssa<local ref to i64>) -> i32 { +// CHECK: %[[VAL_0:.*]] = wasmssa.local_get %[[ARG0]] : ref to i64 +// CHECK: %[[VAL_1:.*]] = wasmssa.wrap %[[VAL_0]] : i64 to i32 +// CHECK: wasmssa.return %[[VAL_1]] : i32 diff --git a/mlir/test/Target/Wasm/xor.mlir b/mlir/test/Target/Wasm/xor.mlir index 94691de..56407db 100644 --- a/mlir/test/Target/Wasm/xor.mlir +++ b/mlir/test/Target/Wasm/xor.mlir @@ -14,13 +14,13 @@ ) */ -// CHECK-LABEL: wasmssa.func @xor_i32() -> i32 { +// CHECK-LABEL: wasmssa.func exported @xor_i32() -> i32 { // CHECK: %0 = wasmssa.const 10 : i32 // CHECK: %1 = wasmssa.const 3 : i32 // CHECK: %2 = wasmssa.xor %0 %1 : i32 // CHECK: wasmssa.return %2 : i32 -// CHECK-LABEL: wasmssa.func @xor_i64() -> i64 { +// CHECK-LABEL: wasmssa.func exported @xor_i64() -> i64 { // CHECK: %0 = wasmssa.const 10 : i64 // CHECK: %1 = wasmssa.const 3 : i64 // CHECK: %2 = wasmssa.xor %0 %1 : i64 diff --git a/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp b/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp index 1e2d4a7..4069a74 100644 --- a/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp +++ b/mlir/test/lib/Dialect/Bufferization/TestOneShotModuleBufferize.cpp @@ -11,11 +11,25 @@ #include "mlir/Dialect/Bufferization/Transforms/Bufferize.h" #include "mlir/Dialect/Bufferization/Transforms/OneShotModuleBufferize.h" #include "mlir/Dialect/Bufferization/Transforms/Transforms.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Pass/Pass.h" +#include "TestAttributes.h" // TestTensorEncodingAttr, TestMemRefLayoutAttr +#include "TestDialect.h" + using namespace mlir; namespace { +MemRefLayoutAttrInterface +getMemRefLayoutForTensorEncoding(RankedTensorType tensorType) { + if (auto encoding = dyn_cast_if_present<test::TestTensorEncodingAttr>( + tensorType.getEncoding())) { + return cast<MemRefLayoutAttrInterface>(test::TestMemRefLayoutAttr::get( + tensorType.getContext(), encoding.getDummy())); + } + return {}; +} + struct TestOneShotModuleBufferizePass : public PassWrapper<TestOneShotModuleBufferizePass, OperationPass<>> { MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TestOneShotModuleBufferizePass) @@ -25,6 +39,7 @@ struct TestOneShotModuleBufferizePass : PassWrapper(pass) {} void getDependentDialects(DialectRegistry ®istry) const override { + registry.insert<test::TestDialect>(); registry.insert<bufferization::BufferizationDialect>(); } StringRef getArgument() const final { @@ -41,6 +56,17 @@ struct TestOneShotModuleBufferizePass bufferization::OneShotBufferizationOptions opt; opt.bufferizeFunctionBoundaries = true; + opt.functionArgTypeConverterFn = + [&](bufferization::TensorLikeType tensor, Attribute memSpace, + func::FuncOp, const bufferization::BufferizationOptions &) { + assert(isa<RankedTensorType>(tensor) && "tests only builtin tensors"); + auto tensorType = cast<RankedTensorType>(tensor); + auto layout = getMemRefLayoutForTensorEncoding(tensorType); + return cast<bufferization::BufferLikeType>( + MemRefType::get(tensorType.getShape(), + tensorType.getElementType(), layout, memSpace)); + }; + bufferization::BufferizationState bufferizationState; if (failed(bufferization::runOneShotModuleBufferize(getOperation(), opt, diff --git a/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp b/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp index 727c84c..8c5c8e8 100644 --- a/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp +++ b/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp @@ -276,10 +276,8 @@ void TestLinalgTransforms::runOnOperation() { Operation *consumer = opOperand->getOwner(); // If we have a pack/unpack consumer and a producer that has multiple // uses, do not apply the folding patterns. - if (isa<linalg::PackOp, linalg::UnPackOp>(consumer) && - isa<TilingInterface>(producer) && !producer->hasOneUse()) - return false; - return true; + return !(isa<linalg::PackOp, linalg::UnPackOp>(consumer) && + isa<TilingInterface>(producer) && !producer->hasOneUse()); }; applyFoldIntoPackAndUnpackPatterns(rootOp, controlFn); } diff --git a/mlir/test/lib/Dialect/Test/TestAttrDefs.td b/mlir/test/lib/Dialect/Test/TestAttrDefs.td index 5685004..9e7e4f8 100644 --- a/mlir/test/lib/Dialect/Test/TestAttrDefs.td +++ b/mlir/test/lib/Dialect/Test/TestAttrDefs.td @@ -22,6 +22,7 @@ include "mlir/IR/AttrTypeBase.td" include "mlir/IR/BuiltinAttributeInterfaces.td" include "mlir/IR/EnumAttr.td" include "mlir/IR/OpAsmInterface.td" +include "mlir/IR/TensorEncoding.td" // All of the attributes will extend this class. class Test_Attr<string name, list<Trait> traits = []> @@ -439,4 +440,20 @@ def TestCustomStorageCtorAttr : Test_Attr<"TestCustomStorageCtorAttr"> { let hasStorageCustomConstructor = 1; } +def TestTensorEncodingAttr : Test_Attr<"TestTensorEncoding", + [DeclareAttrInterfaceMethods<VerifiableTensorEncoding>]> { + let mnemonic = "tensor_encoding"; + + let parameters = (ins "mlir::StringAttr":$dummy); + let assemblyFormat = "`<` $dummy `>`"; +} + +def TestMemRefLayoutAttr : Test_Attr<"TestMemRefLayout", + [DeclareAttrInterfaceMethods<MemRefLayoutAttrInterface>]> { + let mnemonic = "memref_layout"; + + let parameters = (ins "mlir::StringAttr":$dummy); + let assemblyFormat = "`<` $dummy `>`"; +} + #endif // TEST_ATTRDEFS diff --git a/mlir/test/lib/Dialect/Test/TestAttributes.cpp b/mlir/test/lib/Dialect/Test/TestAttributes.cpp index fe1e916..9db7b01 100644 --- a/mlir/test/lib/Dialect/Test/TestAttributes.cpp +++ b/mlir/test/lib/Dialect/Test/TestAttributes.cpp @@ -542,6 +542,24 @@ test::detail::TestCustomStorageCtorAttrAttrStorage::construct( } //===----------------------------------------------------------------------===// +// TestTensorEncodingAttr +//===----------------------------------------------------------------------===// + +::llvm::LogicalResult TestTensorEncodingAttr::verifyEncoding( + mlir::ArrayRef<int64_t> shape, mlir::Type elementType, + llvm::function_ref<::mlir::InFlightDiagnostic()> emitError) const { + return mlir::success(); +} + +//===----------------------------------------------------------------------===// +// TestMemRefLayoutAttr +//===----------------------------------------------------------------------===// + +mlir::AffineMap TestMemRefLayoutAttr::getAffineMap() const { + return mlir::AffineMap::getMultiDimIdentityMap(1, getContext()); +} + +//===----------------------------------------------------------------------===// // TestDialect //===----------------------------------------------------------------------===// diff --git a/mlir/test/lib/Dialect/Test/TestAttributes.h b/mlir/test/lib/Dialect/Test/TestAttributes.h index 778d84fa..0ad5ab6 100644 --- a/mlir/test/lib/Dialect/Test/TestAttributes.h +++ b/mlir/test/lib/Dialect/Test/TestAttributes.h @@ -24,6 +24,7 @@ #include "mlir/IR/Dialect.h" #include "mlir/IR/DialectImplementation.h" #include "mlir/IR/DialectResourceBlobManager.h" +#include "mlir/IR/TensorEncoding.h" // generated files require above includes to come first #include "TestAttrInterfaces.h.inc" diff --git a/mlir/test/lib/Dialect/Test/TestDialect.h b/mlir/test/lib/Dialect/Test/TestDialect.h index f2adca6..bcf3b55d 100644 --- a/mlir/test/lib/Dialect/Test/TestDialect.h +++ b/mlir/test/lib/Dialect/Test/TestDialect.h @@ -18,6 +18,7 @@ #include "TestInterfaces.h" #include "TestTypes.h" #include "mlir/Bytecode/BytecodeImplementation.h" +#include "mlir/Dialect/Bufferization/IR/Bufferization.h" #include "mlir/Dialect/DLTI/DLTI.h" #include "mlir/Dialect/DLTI/Traits.h" #include "mlir/Dialect/Func/IR/FuncOps.h" diff --git a/mlir/test/lib/Dialect/Test/TestDialect.td b/mlir/test/lib/Dialect/Test/TestDialect.td index 2b5491f..37a263f 100644 --- a/mlir/test/lib/Dialect/Test/TestDialect.td +++ b/mlir/test/lib/Dialect/Test/TestDialect.td @@ -24,7 +24,10 @@ def Test_Dialect : Dialect { let useDefaultTypePrinterParser = 0; let useDefaultAttributePrinterParser = 1; let isExtensible = 1; - let dependentDialects = ["::mlir::DLTIDialect"]; + let dependentDialects = [ + "::mlir::DLTIDialect", + "::mlir::bufferization::BufferizationDialect" + ]; let discardableAttrs = (ins "mlir::IntegerAttr":$discardable_attr_key, "SimpleAAttr":$other_discardable_attr_key diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp index 53055fe..b211e24 100644 --- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp +++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp @@ -1425,6 +1425,39 @@ TestMultiSlotAlloca::handleDestructuringComplete( return createNewMultiAllocaWithoutSlot(slot, builder, *this); } +namespace { +/// Returns test dialect's memref layout for test dialect's tensor encoding when +/// applicable. +MemRefLayoutAttrInterface +getMemRefLayoutForTensorEncoding(RankedTensorType tensorType) { + if (auto encoding = + dyn_cast<test::TestTensorEncodingAttr>(tensorType.getEncoding())) { + return cast<MemRefLayoutAttrInterface>(test::TestMemRefLayoutAttr::get( + tensorType.getContext(), encoding.getDummy())); + } + return {}; +} + +/// Auxiliary bufferization function for test and builtin tensors. +bufferization::BufferLikeType +convertTensorToBuffer(mlir::Operation *op, + const bufferization::BufferizationOptions &options, + bufferization::TensorLikeType tensorLike) { + auto buffer = + *tensorLike.getBufferType(options, [&]() { return op->emitError(); }); + if (auto memref = dyn_cast<MemRefType>(buffer)) { + // Note: For the sake of testing, we want to ensure that encoding -> layout + // bufferization happens. This is currently achieved manually. + auto layout = + getMemRefLayoutForTensorEncoding(cast<RankedTensorType>(tensorLike)); + return cast<bufferization::BufferLikeType>( + MemRefType::get(memref.getShape(), memref.getElementType(), layout, + memref.getMemorySpace())); + } + return buffer; +} +} // namespace + ::mlir::LogicalResult test::TestDummyTensorOp::bufferize( ::mlir::RewriterBase &rewriter, const ::mlir::bufferization::BufferizationOptions &options, @@ -1435,8 +1468,8 @@ TestMultiSlotAlloca::handleDestructuringComplete( return failure(); const auto outType = getOutput().getType(); - const auto bufferizedOutType = test::TestMemrefType::get( - getContext(), outType.getShape(), outType.getElementType(), nullptr); + const auto bufferizedOutType = + convertTensorToBuffer(getOperation(), options, outType); // replace op with memref analogy auto dummyMemrefOp = test::TestDummyMemrefOp::create( rewriter, getLoc(), bufferizedOutType, *buffer); @@ -1470,13 +1503,12 @@ TestMultiSlotAlloca::handleDestructuringComplete( mlir::FailureOr<mlir::bufferization::BufferLikeType> test::TestCreateTensorOp::getBufferType( - mlir::Value value, const mlir::bufferization::BufferizationOptions &, + mlir::Value value, const mlir::bufferization::BufferizationOptions &options, const mlir::bufferization::BufferizationState &, llvm::SmallVector<::mlir::Value> &) { - const auto type = dyn_cast<test::TestTensorType>(value.getType()); + const auto type = dyn_cast<bufferization::TensorLikeType>(value.getType()); if (type == nullptr) return failure(); - return cast<mlir::bufferization::BufferLikeType>(test::TestMemrefType::get( - getContext(), type.getShape(), type.getElementType(), nullptr)); + return convertTensorToBuffer(getOperation(), options, type); } diff --git a/mlir/test/lib/Dialect/Test/TestOps.td b/mlir/test/lib/Dialect/Test/TestOps.td index 6329d61..05a33cf 100644 --- a/mlir/test/lib/Dialect/Test/TestOps.td +++ b/mlir/test/lib/Dialect/Test/TestOps.td @@ -32,6 +32,7 @@ include "mlir/Interfaces/MemorySlotInterfaces.td" include "mlir/Interfaces/SideEffectInterfaces.td" include "mlir/Interfaces/ValueBoundsOpInterface.td" include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.td" +include "mlir/Dialect/Bufferization/IR/BufferizationTypeInterfaces.td" // Include the attribute definitions. include "TestAttrDefs.td" @@ -2335,7 +2336,7 @@ def SideEffectWithRegionOp : TEST_Op<"side_effect_with_region_op", } //===----------------------------------------------------------------------===// -// Copy Operation Test +// Copy Operation Test //===----------------------------------------------------------------------===// def CopyOp : TEST_Op<"copy", []> { @@ -3676,10 +3677,10 @@ def TestDummyTensorOp : TEST_Op<"dummy_tensor_op", ["bufferize", "bufferizesToMemoryRead", "bufferizesToMemoryWrite", "getAliasingValues"]>]> { let arguments = (ins - Arg<TestTensorType>:$input + Arg<Bufferization_TensorLikeTypeInterface>:$input ); let results = (outs - Arg<TestTensorType>:$output + Arg<Bufferization_TensorLikeTypeInterface>:$output ); let extraClassDefinition = [{ @@ -3701,10 +3702,10 @@ def TestDummyTensorOp : TEST_Op<"dummy_tensor_op", def TestDummyMemrefOp : TEST_Op<"dummy_memref_op", []> { let arguments = (ins - Arg<TestMemrefType>:$input + Arg<Bufferization_BufferLikeTypeInterface>:$input ); let results = (outs - Arg<TestMemrefType>:$output + Arg<Bufferization_BufferLikeTypeInterface>:$output ); } @@ -3714,7 +3715,7 @@ def TestCreateTensorOp : TEST_Op<"create_tensor_op", "bufferizesToMemoryWrite", "getAliasingValues", "bufferizesToAllocation"]>]> { let arguments = (ins); - let results = (outs Arg<TestTensorType>:$output); + let results = (outs Arg<Bufferization_TensorLikeTypeInterface>:$output); let extraClassDefinition = [{ bool test::TestCreateTensorOp::bufferizesToMemoryRead(::mlir::OpOperand&, const ::mlir::bufferization::AnalysisState&) { @@ -3738,7 +3739,7 @@ def TestCreateTensorOp : TEST_Op<"create_tensor_op", def TestCreateMemrefOp : TEST_Op<"create_memref_op"> { let arguments = (ins); - let results = (outs Arg<TestMemrefType>:$output); + let results = (outs Arg<Bufferization_BufferLikeTypeInterface>:$output); } //===----------------------------------------------------------------------===// diff --git a/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp b/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp index 97fc699..496f18b 100644 --- a/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp +++ b/mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp @@ -938,10 +938,10 @@ public: // These are automatically generated by ODS but are not used as the Transform // dialect uses a different dispatch mechanism to support dialect extensions. -LLVM_ATTRIBUTE_UNUSED static OptionalParseResult +[[maybe_unused]] static OptionalParseResult generatedTypeParser(AsmParser &parser, StringRef *mnemonic, Type &value); -LLVM_ATTRIBUTE_UNUSED static LogicalResult -generatedTypePrinter(Type def, AsmPrinter &printer); +[[maybe_unused]] static LogicalResult generatedTypePrinter(Type def, + AsmPrinter &printer); #define GET_TYPEDEF_CLASSES #include "TestTransformDialectExtensionTypes.cpp.inc" diff --git a/mlir/test/mlir-pdll/CodeGen/CPP/general.pdll b/mlir/test/mlir-pdll/CodeGen/CPP/general.pdll index 4e869e5..4be30d8 100644 --- a/mlir/test/mlir-pdll/CodeGen/CPP/general.pdll +++ b/mlir/test/mlir-pdll/CodeGen/CPP/general.pdll @@ -28,7 +28,7 @@ // CHECK: operation "test.op3" // CHECK: )mlir", context), std::forward<ConfigsT>(configs)...) -// CHECK: static void LLVM_ATTRIBUTE_UNUSED populateGeneratedPDLLPatterns(::mlir::RewritePatternSet &patterns, ConfigsT &&...configs) { +// CHECK{LITERAL}: [[maybe_unused]] static void populateGeneratedPDLLPatterns(::mlir::RewritePatternSet &patterns, ConfigsT &&...configs) { // CHECK-NEXT: patterns.add<GeneratedPDLLPattern0>(patterns.getContext(), configs...); // CHECK-NEXT: patterns.add<NamedPattern>(patterns.getContext(), configs...); // CHECK-NEXT: patterns.add<GeneratedPDLLPattern1>(patterns.getContext(), configs...); diff --git a/mlir/test/mlir-tblgen/cpp-class-comments.td b/mlir/test/mlir-tblgen/cpp-class-comments.td index a896888..9dcf975 100644 --- a/mlir/test/mlir-tblgen/cpp-class-comments.td +++ b/mlir/test/mlir-tblgen/cpp-class-comments.td @@ -96,17 +96,14 @@ def EncodingTrait : AttrInterface<"EncodingTrait"> { }]; let methods = [ ]; -// ATTR-INTERFACE: namespace mlir -// ATTR-INTERFACE-NEXT: namespace a -// ATTR-INTERFACE-NEXT: namespace traits +// ATTR-INTERFACE: namespace mlir::a::traits { // ATTR-INTERFACE-NEXT: /// Common trait for all layouts. // ATTR-INTERFACE-NEXT: class EncodingTrait; } def SimpleEncodingTrait : AttrInterface<"SimpleEncodingTrait"> { let cppNamespace = "a::traits"; -// ATTR-INTERFACE: namespace a { -// ATTR-INTERFACE-NEXT: namespace traits { +// ATTR-INTERFACE: namespace a::traits { // ATTR-INTERFACE-NEXT: class SimpleEncodingTrait; } @@ -116,8 +113,7 @@ def SimpleOpInterface : OpInterface<"SimpleOpInterface"> { Simple Op Interface description }]; -// OP-INTERFACE: namespace a { -// OP-INTERFACE-NEXT: namespace traits { +// OP-INTERFACE: namespace a::traits { // OP-INTERFACE-NEXT: /// Simple Op Interface description // OP-INTERFACE-NEXT: class SimpleOpInterface; } diff --git a/mlir/test/python/dialects/gpu/dialect.py b/mlir/test/python/dialects/gpu/dialect.py index 26ee9f3..66c4018 100644 --- a/mlir/test/python/dialects/gpu/dialect.py +++ b/mlir/test/python/dialects/gpu/dialect.py @@ -1,6 +1,7 @@ # RUN: %PYTHON %s | FileCheck %s from mlir.ir import * +import mlir.ir as ir import mlir.dialects.gpu as gpu import mlir.dialects.gpu.passes from mlir.passmanager import * @@ -64,3 +65,95 @@ def testObjectAttr(): # CHECK: #gpu.object<#nvvm.target, kernels = <[#gpu.kernel_metadata<"kernel", () -> ()>]>, "BC\C0\DE5\14\00\00\05\00\00\00b\0C0$MY\BEf"> print(o) assert o.kernels == kernelTable + + +# CHECK-LABEL: testGPUFuncOp +@run +def testGPUFuncOp(): + assert gpu.GPUFuncOp.__doc__ is not None + module = Module.create() + with InsertionPoint(module.body): + gpu_module_name = StringAttr.get("gpu_module") + gpumodule = gpu.GPUModuleOp(gpu_module_name) + block = gpumodule.bodyRegion.blocks.append() + + def builder(func: gpu.GPUFuncOp) -> None: + gpu.GlobalIdOp(gpu.Dimension.x) + gpu.ReturnOp([]) + + with InsertionPoint(block): + name = StringAttr.get("kernel0") + func_type = ir.FunctionType.get(inputs=[], results=[]) + type_attr = TypeAttr.get(func_type) + func = gpu.GPUFuncOp(type_attr, name) + func.attributes["sym_name"] = name + func.attributes["gpu.kernel"] = UnitAttr.get() + + try: + func.entry_block + assert False, "Expected RuntimeError" + except RuntimeError as e: + assert ( + str(e) + == "Entry block does not exist for kernel0. Do you need to call the add_entry_block() method on this GPUFuncOp?" + ) + + block = func.add_entry_block() + with InsertionPoint(block): + builder(func) + + try: + func.add_entry_block() + assert False, "Expected RuntimeError" + except RuntimeError as e: + assert str(e) == "Entry block already exists for kernel0" + + func = gpu.GPUFuncOp( + func_type, + sym_name="kernel1", + kernel=True, + body_builder=builder, + known_block_size=[1, 2, 3], + known_grid_size=DenseI32ArrayAttr.get([4, 5, 6]), + ) + + assert func.name.value == "kernel1" + assert func.function_type.value == func_type + assert func.arg_attrs == None + assert func.res_attrs == None + assert func.arguments == [] + assert func.entry_block == func.body.blocks[0] + assert func.is_kernel + assert func.known_block_size == DenseI32ArrayAttr.get( + [1, 2, 3] + ), func.known_block_size + assert func.known_grid_size == DenseI32ArrayAttr.get( + [4, 5, 6] + ), func.known_grid_size + + func = gpu.GPUFuncOp( + func_type, + sym_name="non_kernel_func", + body_builder=builder, + ) + assert not func.is_kernel + assert func.known_block_size is None + assert func.known_grid_size is None + + print(module) + + # CHECK: gpu.module @gpu_module + # CHECK: gpu.func @kernel0() kernel { + # CHECK: %[[VAL_0:.*]] = gpu.global_id x + # CHECK: gpu.return + # CHECK: } + # CHECK: gpu.func @kernel1() kernel attributes + # CHECK-SAME: known_block_size = array<i32: 1, 2, 3> + # CHECK-SAME: known_grid_size = array<i32: 4, 5, 6> + # CHECK: %[[VAL_0:.*]] = gpu.global_id x + # CHECK: gpu.return + # CHECK: } + # CHECK: gpu.func @non_kernel_func() { + # CHECK: %[[VAL_0:.*]] = gpu.global_id x + # CHECK: gpu.return + # CHECK: } diff --git a/mlir/test/python/dialects/openacc.py b/mlir/test/python/dialects/openacc.py new file mode 100644 index 0000000..8f2142a --- /dev/null +++ b/mlir/test/python/dialects/openacc.py @@ -0,0 +1,171 @@ +# RUN: %PYTHON %s | FileCheck %s +from unittest import result +from mlir.ir import ( + Context, + FunctionType, + Location, + Module, + InsertionPoint, + IntegerType, + IndexType, + MemRefType, + F32Type, + Block, + ArrayAttr, + Attribute, + UnitAttr, + StringAttr, + DenseI32ArrayAttr, + ShapedType, +) +from mlir.dialects import openacc, func, arith, memref +from mlir.extras import types + + +def run(f): + print("\n// TEST:", f.__name__) + with Context(), Location.unknown(): + f() + return f + + +@run +def testParallelMemcpy(): + module = Module.create() + + dynamic = ShapedType.get_dynamic_size() + memref_f32_1d_any = MemRefType.get([dynamic], types.f32()) + + with InsertionPoint(module.body): + function_type = FunctionType.get( + [memref_f32_1d_any, memref_f32_1d_any, types.i64()], [] + ) + f = func.FuncOp( + type=function_type, + name="memcpy_idiom", + ) + f.attributes["sym_visibility"] = StringAttr.get("public") + + with InsertionPoint(f.add_entry_block()): + c1024 = arith.ConstantOp(types.i32(), 1024) + c128 = arith.ConstantOp(types.i32(), 128) + + arg0, arg1, arg2 = f.arguments + + copied = openacc.copyin( + acc_var=arg0.type, + var=arg0, + var_type=types.f32(), + bounds=[], + async_operands=[], + implicit=False, + structured=True, + ) + created = openacc.create_( + acc_var=arg1.type, + var=arg1, + var_type=types.f32(), + bounds=[], + async_operands=[], + implicit=False, + structured=True, + ) + + parallel_op = openacc.ParallelOp( + asyncOperands=[], + waitOperands=[], + numGangs=[c1024], + numWorkers=[], + vectorLength=[c128], + reductionOperands=[], + privateOperands=[], + firstprivateOperands=[], + dataClauseOperands=[], + ) + + # Set required device_type and segment attributes to satisfy verifier + acc_device_none = ArrayAttr.get([Attribute.parse("#acc.device_type<none>")]) + parallel_op.numGangsDeviceType = acc_device_none + parallel_op.numGangsSegments = DenseI32ArrayAttr.get([1]) + parallel_op.vectorLengthDeviceType = acc_device_none + + parallel_block = Block.create_at_start(parent=parallel_op.region, arg_types=[]) + + with InsertionPoint(parallel_block): + c0 = arith.ConstantOp(types.i64(), 0) + c1 = arith.ConstantOp(types.i64(), 1) + + loop_op = openacc.LoopOp( + results_=[], + lowerbound=[c0], + upperbound=[f.arguments[2]], + step=[c1], + gangOperands=[], + workerNumOperands=[], + vectorOperands=[], + tileOperands=[], + cacheOperands=[], + privateOperands=[], + reductionOperands=[], + firstprivateOperands=[], + ) + + # Set loop attributes: gang and independent on device_type<none> + acc_device_none = ArrayAttr.get([Attribute.parse("#acc.device_type<none>")]) + loop_op.gang = acc_device_none + loop_op.independent = acc_device_none + + loop_block = Block.create_at_start( + parent=loop_op.region, arg_types=[types.i64()] + ) + + with InsertionPoint(loop_block): + idx = arith.index_cast(out=IndexType.get(), in_=loop_block.arguments[0]) + val = memref.load(memref=copied, indices=[idx]) + memref.store(value=val, memref=created, indices=[idx]) + openacc.YieldOp([]) + + openacc.YieldOp([]) + + deleted = openacc.delete( + acc_var=copied, + bounds=[], + async_operands=[], + implicit=False, + structured=True, + ) + copied = openacc.copyout( + acc_var=created, + var=arg1, + var_type=types.f32(), + bounds=[], + async_operands=[], + implicit=False, + structured=True, + ) + func.ReturnOp([]) + + print(module) + + # CHECK: TEST: testParallelMemcpy + # CHECK-LABEL: func.func public @memcpy_idiom( + # CHECK-SAME: %[[ARG0:.*]]: memref<?xf32>, %[[ARG1:.*]]: memref<?xf32>, %[[ARG2:.*]]: i64) { + # CHECK: %[[CONSTANT_0:.*]] = arith.constant 1024 : i32 + # CHECK: %[[CONSTANT_1:.*]] = arith.constant 128 : i32 + # CHECK: %[[COPYIN_0:.*]] = acc.copyin varPtr(%[[ARG0]] : memref<?xf32>) -> memref<?xf32> + # CHECK: %[[CREATE_0:.*]] = acc.create varPtr(%[[ARG1]] : memref<?xf32>) -> memref<?xf32> + # CHECK: acc.parallel num_gangs({%[[CONSTANT_0]] : i32}) vector_length(%[[CONSTANT_1]] : i32) { + # CHECK: %[[CONSTANT_2:.*]] = arith.constant 0 : i64 + # CHECK: %[[CONSTANT_3:.*]] = arith.constant 1 : i64 + # CHECK: acc.loop gang control(%[[VAL_0:.*]] : i64) = (%[[CONSTANT_2]] : i64) to (%[[ARG2]] : i64) step (%[[CONSTANT_3]] : i64) { + # CHECK: %[[INDEX_CAST_0:.*]] = arith.index_cast %[[VAL_0]] : i64 to index + # CHECK: %[[LOAD_0:.*]] = memref.load %[[COPYIN_0]]{{\[}}%[[INDEX_CAST_0]]] : memref<?xf32> + # CHECK: memref.store %[[LOAD_0]], %[[CREATE_0]]{{\[}}%[[INDEX_CAST_0]]] : memref<?xf32> + # CHECK: acc.yield + # CHECK: } attributes {independent = [#acc.device_type<none>]} + # CHECK: acc.yield + # CHECK: } + # CHECK: acc.delete accPtr(%[[COPYIN_0]] : memref<?xf32>) + # CHECK: acc.copyout accPtr(%[[CREATE_0]] : memref<?xf32>) to varPtr(%[[ARG1]] : memref<?xf32>) + # CHECK: return + # CHECK: } diff --git a/mlir/test/python/ir/operation.py b/mlir/test/python/ir/operation.py index cb4cfc8c..1d4ede1 100644 --- a/mlir/test/python/ir/operation.py +++ b/mlir/test/python/ir/operation.py @@ -569,12 +569,30 @@ def testOperationAttributes(): # CHECK: Attribute value b'text' print(f"Attribute value {sattr.value_bytes}") + # Python dict-style iteration # We don't know in which order the attributes are stored. - # CHECK-DAG: NamedAttribute(dependent="text") - # CHECK-DAG: NamedAttribute(other.attribute=3.000000e+00 : f64) - # CHECK-DAG: NamedAttribute(some.attribute=1 : i8) - for attr in op.attributes: - print(str(attr)) + # CHECK-DAG: dependent + # CHECK-DAG: other.attribute + # CHECK-DAG: some.attribute + for name in op.attributes: + print(name) + + # Basic dict-like introspection + # CHECK: True + print("some.attribute" in op.attributes) + # CHECK: False + print("missing" in op.attributes) + # CHECK: Keys: ['dependent', 'other.attribute', 'some.attribute'] + print("Keys:", sorted(op.attributes.keys())) + # CHECK: Values count 3 + print("Values count", len(op.attributes.values())) + # CHECK: Items count 3 + print("Items count", len(op.attributes.items())) + + # Dict() conversion test + d = {k: v.value for k, v in dict(op.attributes).items()} + # CHECK: Dict mapping {'dependent': 'text', 'other.attribute': 3.0, 'some.attribute': 1} + print("Dict mapping", d) # Check that exceptions are raised as expected. try: |