diff options
author | Fangrui Song <i@maskray.me> | 2024-05-03 10:27:42 -0700 |
---|---|---|
committer | Fangrui Song <i@maskray.me> | 2024-05-03 10:27:43 -0700 |
commit | dce13b42177565793a2031eb6f395f9508c701e0 (patch) | |
tree | 61bbcd47a217a7ec89c9506cbfa47514b5eda584 | |
parent | f34a5205aa481a6d9a15054bcc5f7b9875906a17 (diff) |
[mlir,test] Convert text files from CRLF to LF
-rw-r--r-- | mlir/test/Conversion/GPUCommon/transfer_write.mlir | 26 | ||||
-rw-r--r-- | mlir/test/Dialect/Shape/arg_with_shape.mlir | 32 | ||||
-rw-r--r-- | mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll | 30 | ||||
-rw-r--r-- | mlir/test/mlir-vulkan-runner/addf_if.mlir | 108 |
4 files changed, 98 insertions, 98 deletions
diff --git a/mlir/test/Conversion/GPUCommon/transfer_write.mlir b/mlir/test/Conversion/GPUCommon/transfer_write.mlir index cba85915b49e..cd62b7b13fa9 100644 --- a/mlir/test/Conversion/GPUCommon/transfer_write.mlir +++ b/mlir/test/Conversion/GPUCommon/transfer_write.mlir @@ -1,13 +1,13 @@ -// RUN: mlir-opt %s --gpu-to-llvm | FileCheck %s
-
- func.func @warp_extract(%arg0: index, %arg1: memref<1024x1024xf32>, %arg2: index, %arg3: vector<1xf32>) {
- %c0 = arith.constant 0 : index
- vector.warp_execute_on_lane_0(%arg0)[32] {
- // CHECK:%[[val:[0-9]+]] = llvm.extractelement
- // CHECK:%[[base:[0-9]+]] = llvm.extractvalue
- // CHECK:%[[ptr:[0-9]+]] = llvm.getelementptr %[[base]]
- // CHECK:llvm.store %[[val]], %[[ptr]]
- vector.transfer_write %arg3, %arg1[%c0, %c0] {in_bounds = [true]} : vector<1xf32>, memref<1024x1024xf32>
- }
- return
- }
+// RUN: mlir-opt %s --gpu-to-llvm | FileCheck %s + + func.func @warp_extract(%arg0: index, %arg1: memref<1024x1024xf32>, %arg2: index, %arg3: vector<1xf32>) { + %c0 = arith.constant 0 : index + vector.warp_execute_on_lane_0(%arg0)[32] { + // CHECK:%[[val:[0-9]+]] = llvm.extractelement + // CHECK:%[[base:[0-9]+]] = llvm.extractvalue + // CHECK:%[[ptr:[0-9]+]] = llvm.getelementptr %[[base]] + // CHECK:llvm.store %[[val]], %[[ptr]] + vector.transfer_write %arg3, %arg1[%c0, %c0] {in_bounds = [true]} : vector<1xf32>, memref<1024x1024xf32> + } + return + } diff --git a/mlir/test/Dialect/Shape/arg_with_shape.mlir b/mlir/test/Dialect/Shape/arg_with_shape.mlir index 089c5031fa55..fd46461417db 100644 --- a/mlir/test/Dialect/Shape/arg_with_shape.mlir +++ b/mlir/test/Dialect/Shape/arg_with_shape.mlir @@ -1,16 +1,16 @@ -// RUN: mlir-opt -outline-shape-computation -split-input-file %s 2>%t | FileCheck %s
-
-func.func @func1(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape {
- %0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape
- %1 = shape.shape_of %arg1 : !shape.value_shape -> !shape.shape
- %2 = shape.meet %0, %1 : !shape.shape, !shape.shape -> !shape.shape
- return %2 : !shape.shape
-}
-// Make sure with_shape used by call not crash.
-// CHECK-LABEL:func.func @func
-func.func @func(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape {
- %0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape
- %1 = shape.with_shape %arg1, %0 : !shape.value_shape, !shape.shape
- %2 = call @func1(%arg0, %1) : (!shape.value_shape, !shape.value_shape) -> !shape.shape
- return %2 : !shape.shape
-}
+// RUN: mlir-opt -outline-shape-computation -split-input-file %s 2>%t | FileCheck %s + +func.func @func1(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape { + %0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape + %1 = shape.shape_of %arg1 : !shape.value_shape -> !shape.shape + %2 = shape.meet %0, %1 : !shape.shape, !shape.shape -> !shape.shape + return %2 : !shape.shape +} +// Make sure with_shape used by call not crash. +// CHECK-LABEL:func.func @func +func.func @func(%arg0: !shape.value_shape, %arg1: !shape.value_shape) -> !shape.shape { + %0 = shape.shape_of %arg0 : !shape.value_shape -> !shape.shape + %1 = shape.with_shape %arg1, %0 : !shape.value_shape, !shape.shape + %2 = call @func1(%arg0, %1) : (!shape.value_shape, !shape.value_shape) -> !shape.shape + return %2 : !shape.shape +} diff --git a/mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll b/mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll index 8702415c2988..d936632031a5 100644 --- a/mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll +++ b/mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll @@ -1,15 +1,15 @@ -; RUN: mlir-translate -import-llvm -split-input-file %s | FileCheck %s
-
-; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"]
-!llvm.linker.options = !{!0}
-!0 = !{!"DEFAULTLIB:", !"libcmt"}
-
-; // -----
-
-!llvm.linker.options = !{!0, !1, !2}
-; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"]
-!0 = !{!"DEFAULTLIB:", !"libcmt"}
-; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmtd"]
-!1 = !{!"DEFAULTLIB:", !"libcmtd"}
-; CHECK: llvm.linker_options ["-lm"]
-!2 = !{!"-lm"}
+; RUN: mlir-translate -import-llvm -split-input-file %s | FileCheck %s + +; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"] +!llvm.linker.options = !{!0} +!0 = !{!"DEFAULTLIB:", !"libcmt"} + +; // ----- + +!llvm.linker.options = !{!0, !1, !2} +; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmt"] +!0 = !{!"DEFAULTLIB:", !"libcmt"} +; CHECK: llvm.linker_options ["DEFAULTLIB:", "libcmtd"] +!1 = !{!"DEFAULTLIB:", !"libcmtd"} +; CHECK: llvm.linker_options ["-lm"] +!2 = !{!"-lm"} diff --git a/mlir/test/mlir-vulkan-runner/addf_if.mlir b/mlir/test/mlir-vulkan-runner/addf_if.mlir index fbd1fae6d0b5..5638bd44682d 100644 --- a/mlir/test/mlir-vulkan-runner/addf_if.mlir +++ b/mlir/test/mlir-vulkan-runner/addf_if.mlir @@ -1,54 +1,54 @@ -// RUN: mlir-vulkan-runner %s --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s
-
-// CHECK: [3.3, 3.3, 3.3, 3.3, 0, 0, 0, 0]
-module attributes {
- gpu.container_module,
- spirv.target_env = #spirv.target_env<
- #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
-} {
- gpu.module @kernels {
- gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
- kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
- %0 = gpu.block_id x
- %limit = arith.constant 4 : index
- %cond = arith.cmpi slt, %0, %limit : index
- scf.if %cond {
- %1 = memref.load %arg0[%0] : memref<8xf32>
- %2 = memref.load %arg1[%0] : memref<8xf32>
- %3 = arith.addf %1, %2 : f32
- memref.store %3, %arg2[%0] : memref<8xf32>
- }
- gpu.return
- }
- }
-
- func.func @main() {
- %arg0 = memref.alloc() : memref<8xf32>
- %arg1 = memref.alloc() : memref<8xf32>
- %arg2 = memref.alloc() : memref<8xf32>
- %0 = arith.constant 0 : i32
- %1 = arith.constant 1 : i32
- %2 = arith.constant 2 : i32
- %value0 = arith.constant 0.0 : f32
- %value1 = arith.constant 1.1 : f32
- %value2 = arith.constant 2.2 : f32
- %arg3 = memref.cast %arg0 : memref<8xf32> to memref<?xf32>
- %arg4 = memref.cast %arg1 : memref<8xf32> to memref<?xf32>
- %arg5 = memref.cast %arg2 : memref<8xf32> to memref<?xf32>
- call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> ()
- call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> ()
- call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> ()
-
- %cst1 = arith.constant 1 : index
- %cst8 = arith.constant 8 : index
- gpu.launch_func @kernels::@kernel_add
- blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1)
- args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
- %arg6 = memref.cast %arg5 : memref<?xf32> to memref<*xf32>
- call @printMemrefF32(%arg6) : (memref<*xf32>) -> ()
- return
- }
- func.func private @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32)
- func.func private @printMemrefF32(%ptr : memref<*xf32>)
-}
-
+// RUN: mlir-vulkan-runner %s --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s + +// CHECK: [3.3, 3.3, 3.3, 3.3, 0, 0, 0, 0] +module attributes { + gpu.container_module, + spirv.target_env = #spirv.target_env< + #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>> +} { + gpu.module @kernels { + gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>) + kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} { + %0 = gpu.block_id x + %limit = arith.constant 4 : index + %cond = arith.cmpi slt, %0, %limit : index + scf.if %cond { + %1 = memref.load %arg0[%0] : memref<8xf32> + %2 = memref.load %arg1[%0] : memref<8xf32> + %3 = arith.addf %1, %2 : f32 + memref.store %3, %arg2[%0] : memref<8xf32> + } + gpu.return + } + } + + func.func @main() { + %arg0 = memref.alloc() : memref<8xf32> + %arg1 = memref.alloc() : memref<8xf32> + %arg2 = memref.alloc() : memref<8xf32> + %0 = arith.constant 0 : i32 + %1 = arith.constant 1 : i32 + %2 = arith.constant 2 : i32 + %value0 = arith.constant 0.0 : f32 + %value1 = arith.constant 1.1 : f32 + %value2 = arith.constant 2.2 : f32 + %arg3 = memref.cast %arg0 : memref<8xf32> to memref<?xf32> + %arg4 = memref.cast %arg1 : memref<8xf32> to memref<?xf32> + %arg5 = memref.cast %arg2 : memref<8xf32> to memref<?xf32> + call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> () + call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> () + call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> () + + %cst1 = arith.constant 1 : index + %cst8 = arith.constant 8 : index + gpu.launch_func @kernels::@kernel_add + blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1) + args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>) + %arg6 = memref.cast %arg5 : memref<?xf32> to memref<*xf32> + call @printMemrefF32(%arg6) : (memref<*xf32>) -> () + return + } + func.func private @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32) + func.func private @printMemrefF32(%ptr : memref<*xf32>) +} + |