Skip to content

Commit dce13b4

Browse files
MaskRayMaskRay
committed
[mlir,test] Convert text files from CRLF to LF
1 parent f34a520 commit dce13b4

File tree

4 files changed

+98
-98
lines changed

4 files changed

+98
-98
lines changed

‎mlir/test/Conversion/GPUCommon/transfer_write.mlir

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,13 @@
1-
// RUN: mlir-opt %s --gpu-to-llvm | FileCheck %s
2-
3-
func.func @warp_extract(%arg0: index, %arg1: memref<1024x1024xf32>, %arg2: index, %arg3: vector<1xf32>) {
4-
%c0 = arith.constant 0 : index
5-
vector.warp_execute_on_lane_0(%arg0)[32] {
6-
// CHECK:%[[val:[0-9]+]] = llvm.extractelement
7-
// CHECK:%[[base:[0-9]+]] = llvm.extractvalue
8-
// CHECK:%[[ptr:[0-9]+]] = llvm.getelementptr %[[base]]
9-
// CHECK:llvm.store %[[val]], %[[ptr]]
10-
vector.transfer_write %arg3, %arg1[%c0, %c0] {in_bounds = [true]} : vector<1xf32>, memref<1024x1024xf32>
11-
}
12-
return
13-
}
1+
// RUN: mlir-opt %s --gpu-to-llvm | FileCheck %s
2+
3+
func.func @warp_extract(%arg0: index, %arg1: memref<1024x1024xf32>, %arg2: index, %arg3: vector<1xf32>) {
4+
%c0 = arith.constant 0 : index
5+
vector.warp_execute_on_lane_0(%arg0)[32] {
6+
// CHECK:%[[val:[0-9]+]] = llvm.extractelement
7+
// CHECK:%[[base:[0-9]+]] = llvm.extractvalue
8+
// CHECK:%[[ptr:[0-9]+]] = llvm.getelementptr %[[base]]
9+
// CHECK:llvm.store %[[val]], %[[ptr]]
10+
vector.transfer_write %arg3, %arg1[%c0, %c0] {in_bounds = [true]} : vector<1xf32>, memref<1024x1024xf32>
11+
}
12+
return
13+
}

‎mlir/test/Dialect/Shape/arg_with_shape.mlir

Lines changed: 16 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,16 +1,16 @@
1-
// RUN: mlir-opt -outline-shape-computation -split-input-file %s 2>%t | FileCheck %s
2-
3-
func.func @func1(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape {
4-
%0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape
5-
%1 = shape.shape_of %arg1 : !shape.value_shape -> !shape.shape
6-
%2 = shape.meet %0, %1 : !shape.shape, !shape.shape -> !shape.shape
7-
return %2 : !shape.shape
8-
}
9-
// Make sure with_shape used by call not crash.
10-
// CHECK-LABEL:func.func @func
11-
func.func @func(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape {
12-
%0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape
13-
%1 = shape.with_shape %arg1, %0 : !shape.value_shape, !shape.shape
14-
%2 = call @func1(%arg0, %1) : (!shape.value_shape, !shape.value_shape) -> !shape.shape
15-
return %2 : !shape.shape
16-
}
1+
// RUN: mlir-opt -outline-shape-computation -split-input-file %s 2>%t | FileCheck %s
2+
3+
func.func @func1(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape {
4+
%0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape
5+
%1 = shape.shape_of %arg1 : !shape.value_shape -> !shape.shape
6+
%2 = shape.meet %0, %1 : !shape.shape, !shape.shape -> !shape.shape
7+
return %2 : !shape.shape
8+
}
9+
// Make sure with_shape used by call not crash.
10+
// CHECK-LABEL:func.func @func
11+
func.func @func(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape {
12+
%0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape
13+
%1 = shape.with_shape %arg1, %0 : !shape.value_shape, !shape.shape
14+
%2 = call @func1(%arg0, %1) : (!shape.value_shape, !shape.value_shape) -> !shape.shape
15+
return %2 : !shape.shape
16+
}

‎mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll

Lines changed: 15 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,15 @@
1-
; RUN: mlir-translate -import-llvm -split-input-file %s | FileCheck %s
2-
3-
; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"]
4-
!llvm.linker.options = !{!0}
5-
!0 = !{!"DEFAULTLIB:", !"libcmt"}
6-
7-
; // -----
8-
9-
!llvm.linker.options = !{!0, !1, !2}
10-
; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"]
11-
!0 = !{!"DEFAULTLIB:", !"libcmt"}
12-
; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmtd"]
13-
!1 = !{!"DEFAULTLIB:", !"libcmtd"}
14-
; CHECK: llvm.linker_options ["-lm"]
15-
!2 = !{!"-lm"}
1+
; RUN: mlir-translate -import-llvm -split-input-file %s | FileCheck %s
2+
3+
; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"]
4+
!llvm.linker.options = !{!0}
5+
!0 = !{!"DEFAULTLIB:", !"libcmt"}
6+
7+
; // -----
8+
9+
!llvm.linker.options = !{!0, !1, !2}
10+
; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"]
11+
!0 = !{!"DEFAULTLIB:", !"libcmt"}
12+
; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmtd"]
13+
!1 = !{!"DEFAULTLIB:", !"libcmtd"}
14+
; CHECK: llvm.linker_options ["-lm"]
15+
!2 = !{!"-lm"}

‎mlir/test/mlir-vulkan-runner/addf_if.mlir

Lines changed: 54 additions & 54 deletions
Original file line numberDiff line numberDiff line change
@@ -1,54 +1,54 @@
1-
// RUN: mlir-vulkan-runner %s --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
2-
3-
// CHECK: [3.3, 3.3, 3.3, 3.3, 0, 0, 0, 0]
4-
module attributes {
5-
gpu.container_module,
6-
spirv.target_env = #spirv.target_env<
7-
#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
8-
} {
9-
gpu.module @kernels {
10-
gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
11-
kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
12-
%0 = gpu.block_id x
13-
%limit = arith.constant 4 : index
14-
%cond = arith.cmpi slt, %0, %limit : index
15-
scf.if %cond {
16-
%1 = memref.load %arg0[%0] : memref<8xf32>
17-
%2 = memref.load %arg1[%0] : memref<8xf32>
18-
%3 = arith.addf %1, %2 : f32
19-
memref.store %3, %arg2[%0] : memref<8xf32>
20-
}
21-
gpu.return
22-
}
23-
}
24-
25-
func.func @main() {
26-
%arg0 = memref.alloc() : memref<8xf32>
27-
%arg1 = memref.alloc() : memref<8xf32>
28-
%arg2 = memref.alloc() : memref<8xf32>
29-
%0 = arith.constant 0 : i32
30-
%1 = arith.constant 1 : i32
31-
%2 = arith.constant 2 : i32
32-
%value0 = arith.constant 0.0 : f32
33-
%value1 = arith.constant 1.1 : f32
34-
%value2 = arith.constant 2.2 : f32
35-
%arg3 = memref.cast %arg0 : memref<8xf32> to memref<?xf32>
36-
%arg4 = memref.cast %arg1 : memref<8xf32> to memref<?xf32>
37-
%arg5 = memref.cast %arg2 : memref<8xf32> to memref<?xf32>
38-
call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> ()
39-
call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> ()
40-
call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> ()
41-
42-
%cst1 = arith.constant 1 : index
43-
%cst8 = arith.constant 8 : index
44-
gpu.launch_func @kernels::@kernel_add
45-
blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1)
46-
args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
47-
%arg6 = memref.cast %arg5 : memref<?xf32> to memref<*xf32>
48-
call @printMemrefF32(%arg6) : (memref<*xf32>) -> ()
49-
return
50-
}
51-
func.func private @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32)
52-
func.func private @printMemrefF32(%ptr : memref<*xf32>)
53-
}
54-
1+
// RUN: mlir-vulkan-runner %s --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
2+
3+
// CHECK: [3.3, 3.3, 3.3, 3.3, 0, 0, 0, 0]
4+
module attributes {
5+
gpu.container_module,
6+
spirv.target_env = #spirv.target_env<
7+
#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
8+
} {
9+
gpu.module @kernels {
10+
gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
11+
kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
12+
%0 = gpu.block_id x
13+
%limit = arith.constant 4 : index
14+
%cond = arith.cmpi slt, %0, %limit : index
15+
scf.if %cond {
16+
%1 = memref.load %arg0[%0] : memref<8xf32>
17+
%2 = memref.load %arg1[%0] : memref<8xf32>
18+
%3 = arith.addf %1, %2 : f32
19+
memref.store %3, %arg2[%0] : memref<8xf32>
20+
}
21+
gpu.return
22+
}
23+
}
24+
25+
func.func @main() {
26+
%arg0 = memref.alloc() : memref<8xf32>
27+
%arg1 = memref.alloc() : memref<8xf32>
28+
%arg2 = memref.alloc() : memref<8xf32>
29+
%0 = arith.constant 0 : i32
30+
%1 = arith.constant 1 : i32
31+
%2 = arith.constant 2 : i32
32+
%value0 = arith.constant 0.0 : f32
33+
%value1 = arith.constant 1.1 : f32
34+
%value2 = arith.constant 2.2 : f32
35+
%arg3 = memref.cast %arg0 : memref<8xf32> to memref<?xf32>
36+
%arg4 = memref.cast %arg1 : memref<8xf32> to memref<?xf32>
37+
%arg5 = memref.cast %arg2 : memref<8xf32> to memref<?xf32>
38+
call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> ()
39+
call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> ()
40+
call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> ()
41+
42+
%cst1 = arith.constant 1 : index
43+
%cst8 = arith.constant 8 : index
44+
gpu.launch_func @kernels::@kernel_add
45+
blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1)
46+
args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
47+
%arg6 = memref.cast %arg5 : memref<?xf32> to memref<*xf32>
48+
call @printMemrefF32(%arg6) : (memref<*xf32>) -> ()
49+
return
50+
}
51+
func.func private @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32)
52+
func.func private @printMemrefF32(%ptr : memref<*xf32>)
53+
}
54+

0 commit comments

Comments
 (0)