[Mlir-commits] [mlir] dce13b4 - [mlir, test] Convert text files from CRLF to LF
Fangrui Song
llvmlistbot at llvm.org
Fri May 3 10:27:47 PDT 2024
Author: Fangrui Song
Date: 2024-05-03T10:27:43-07:00
New Revision: dce13b42177565793a2031eb6f395f9508c701e0
URL: https://github.com/llvm/llvm-project/commit/dce13b42177565793a2031eb6f395f9508c701e0
DIFF: https://github.com/llvm/llvm-project/commit/dce13b42177565793a2031eb6f395f9508c701e0.diff
LOG: [mlir,test] Convert text files from CRLF to LF
Added:
Modified:
mlir/test/Conversion/GPUCommon/transfer_write.mlir
mlir/test/Dialect/Shape/arg_with_shape.mlir
mlir/test/Target/LLVMIR/Import/metadata-linker-options.ll
mlir/test/mlir-vulkan-runner/addf_if.mlir
Removed:
################################################################################
diff --git a/mlir/test/Conversion/GPUCommon/transfer_write.mlir b/mlir/test/Conversion/GPUCommon/transfer_write.mlir
index cba85915b49e43..cd62b7b13fa9ae 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 089c5031fa5531..fd46461417db0f 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 8702415c298862..d936632031a5ca 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 fbd1fae6d0b596..5638bd44682de6 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>)
+}
+
More information about the Mlir-commits
mailing list