OSSForks/tensorflow
1
0
Fork 0
mirror of https://github.com/zebrajr/tensorflow.git synced 2025-12-06 12:20:11 +01:00
Code Issues Actions 5 Packages Projects Releases Wiki Activity

Integrate StableHLO at openxla/stablehlo@8817ff1d

Browse Source 
PiperOrigin-RevId: 652528759
This commit is contained in:
Kevin Gleason 2024-07-15 10:37:14 -07:00 committed by TensorFlower Gardener
parent f1703df6aa
commit 119f5328a1
14 changed files with 566 additions and 332 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

163
tensorflow/compiler/mlir/tf2xla/tests/legalize-tf.mlir
View File

@ -1,5 +1,5 @@
// RUN: tf-opt "-xla-legalize-tf=legalize-chlo=false" -split-input-file %s | FILECHECK_OPTS="" FileCheck %s
// RUN: tf-opt "-xla-legalize-tf=legalize-chlo=true" -split-input-file -verify-diagnostics %s | FileCheck %s --check-prefix CHLO --dump-input-filter=all
// RUN: tf-opt "-xla-legalize-tf=legalize-chlo=true" -split-input-file -verify-diagnostics %s | FileCheck %s --check-prefix CHLO
// This test runs twice:
// 1. Through FILECHECK_OPTS="" FileCheck with chlo legalization disabled since verifying
// that the chlo ops emit produces more useful tests.
@ -2225,19 +2225,16 @@ func.func @abs_dynamic(%arg0: tensor<?xf32>) -> tensor<?xf32> {
// CHLO-LABEL: @acos
func.func @acos(%arg0: tensor<2xf32>) -> tensor<2xf32> {
// CHECK: chlo.acos %arg0 : tensor<2xf32>
// CHLO: %[[VAL_1:.*]] = mhlo.compare NE, {{.*}}
// CHLO: %[[VAL_3:.*]] = mhlo.constant dense<2.000000e+00>
// CHLO: %[[VAL_4:.*]] = mhlo.constant dense<1.000000e+00>
// CHLO: %[[VAL_5:.*]] = mhlo.multiply %arg0, %arg0
// CHLO: %[[VAL_6:.*]] = mhlo.subtract %[[VAL_4]], %[[VAL_5]]
// CHLO: %[[VAL_7:.*]] = mhlo.sqrt %[[VAL_6]]
// CHLO: %[[VAL_8:.*]] = mhlo.constant dense<1.000000e+00>
// CHLO: %[[VAL_9:.*]] = mhlo.add %[[VAL_8]], %arg0
// CHLO: %[[VAL_10:.*]] = mhlo.atan2 %[[VAL_7]], %[[VAL_9]]
// CHLO: %[[VAL_11:.*]] = mhlo.multiply %[[VAL_3]], %[[VAL_10]]
// CHLO: %[[VAL_12:.*]] = mhlo.constant dense<3.14159274>
// CHLO: %[[VAL_13:.*]] = mhlo.select %[[VAL_1]], %[[VAL_11]], %[[VAL_12]]
// CHLO: return %[[VAL_13]] : tensor<2xf32>
// CHLO: %[[VAL_0:.*]]: tensor<2xf32>) -> tensor<2xf32> {
// CHLO: %[[VAL_1:.*]] = mhlo.constant dense<2.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_2:.*]] = mhlo.constant dense<1.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_3:.*]] = mhlo.subtract %[[VAL_2]], %[[VAL_0]] : tensor<2xf32>
// CHLO: %[[VAL_4:.*]] = mhlo.add %[[VAL_2]], %[[VAL_0]] : tensor<2xf32>
// CHLO: %[[VAL_5:.*]] = mhlo.multiply %[[VAL_3]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_6:.*]] = mhlo.sqrt %[[VAL_5]] : tensor<2xf32>
// CHLO: %[[VAL_7:.*]] = mhlo.atan2 %[[VAL_6]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_8:.*]] = mhlo.multiply %[[VAL_1]], %[[VAL_7]] : tensor<2xf32>
// CHLO: return %[[VAL_8]] : tensor<2xf32>
%0 = "tf.Acos"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
func.return %0 : tensor<2xf32>
}
@ -2247,7 +2244,143 @@ func.func @acos(%arg0: tensor<2xf32>) -> tensor<2xf32> {
// CHECK-LABEL: @acos_complex
// CHLO-LABEL: @acos_complex
func.func @acos_complex(%arg0: tensor<2xcomplex<f32>>) -> tensor<2xcomplex<f32>> {
// CHLO: tf.Acos
// CHECK: chlo.acos
// CHLO: %[[VAL_0:.*]]: tensor<2xcomplex<f32>>) -> tensor<2xcomplex<f32>> {
// CHLO: %[[VAL_1:.*]] = mhlo.real %[[VAL_0]] : (tensor<2xcomplex<f32>>) -> tensor<2xf32>
// CHLO: %[[VAL_2:.*]] = mhlo.abs %[[VAL_1]] : tensor<2xf32>
// CHLO: %[[VAL_3:.*]] = mhlo.imag %[[VAL_0]] : (tensor<2xcomplex<f32>>) -> tensor<2xf32>
// CHLO: %[[VAL_4:.*]] = mhlo.abs %[[VAL_3]] : tensor<2xf32>
// CHLO: %[[VAL_5:.*]] = mhlo.maximum %[[VAL_2]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_6:.*]] = mhlo.constant dense<3.40282347E+38> : tensor<2xf32>
// CHLO: %[[VAL_7:.*]] = mhlo.sqrt %[[VAL_6]] : tensor<2xf32>
// CHLO: %[[VAL_8:.*]] = mhlo.constant dense<8.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_9:.*]] = mhlo.divide %[[VAL_7]], %[[VAL_8]] : tensor<2xf32>
// CHLO: %[[VAL_10:.*]] = mhlo.compare GE, %[[VAL_5]], %[[VAL_9]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_11:.*]] = mhlo.constant dense<1.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_12:.*]] = mhlo.compare LE, %[[VAL_2]], %[[VAL_11]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_13:.*]] = mhlo.constant dense<5.000000e-01> : tensor<2xf32>
// CHLO: %[[VAL_14:.*]] = mhlo.add %[[VAL_2]], %[[VAL_11]] : tensor<2xf32>
// CHLO: %[[VAL_15:.*]] = mhlo.abs %[[VAL_14]] : tensor<2xf32>
// CHLO: %[[VAL_16:.*]] = mhlo.maximum %[[VAL_15]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_17:.*]] = mhlo.minimum %[[VAL_15]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_18:.*]] = mhlo.compare EQ, %[[VAL_16]], %[[VAL_17]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_19:.*]] = mhlo.constant dense<2.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_20:.*]] = mhlo.sqrt %[[VAL_19]] : tensor<2xf32>
// CHLO: %[[VAL_21:.*]] = mhlo.multiply %[[VAL_20]], %[[VAL_16]] : tensor<2xf32>
// CHLO: %[[VAL_22:.*]] = mhlo.divide %[[VAL_17]], %[[VAL_16]] : tensor<2xf32>
// CHLO: %[[VAL_23:.*]] = mhlo.multiply %[[VAL_22]], %[[VAL_22]] : tensor<2xf32>
// CHLO: %[[VAL_24:.*]] = mhlo.add %[[VAL_11]], %[[VAL_23]] : tensor<2xf32>
// CHLO: %[[VAL_25:.*]] = mhlo.sqrt %[[VAL_24]] : tensor<2xf32>
// CHLO: %[[VAL_26:.*]] = mhlo.compare EQ, %[[VAL_25]], %[[VAL_11]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_27:.*]] = mhlo.constant dense<0.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_28:.*]] = mhlo.compare GT, %[[VAL_23]], %[[VAL_27]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_29:.*]] = mhlo.and %[[VAL_26]], %[[VAL_28]] : tensor<2xi1>
// CHLO: %[[VAL_30:.*]] = mhlo.multiply %[[VAL_16]], %[[VAL_23]] : tensor<2xf32>
// CHLO: %[[VAL_31:.*]] = mhlo.divide %[[VAL_30]], %[[VAL_19]] : tensor<2xf32>
// CHLO: %[[VAL_32:.*]] = mhlo.add %[[VAL_16]], %[[VAL_31]] : tensor<2xf32>
// CHLO: %[[VAL_33:.*]] = mhlo.multiply %[[VAL_16]], %[[VAL_25]] : tensor<2xf32>
// CHLO: %[[VAL_34:.*]] = mhlo.select %[[VAL_29]], %[[VAL_32]], %[[VAL_33]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_35:.*]] = mhlo.select %[[VAL_18]], %[[VAL_21]], %[[VAL_34]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_36:.*]] = mhlo.subtract %[[VAL_2]], %[[VAL_11]] : tensor<2xf32>
// CHLO: %[[VAL_37:.*]] = mhlo.abs %[[VAL_36]] : tensor<2xf32>
// CHLO: %[[VAL_38:.*]] = mhlo.maximum %[[VAL_37]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_39:.*]] = mhlo.minimum %[[VAL_37]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_40:.*]] = mhlo.compare EQ, %[[VAL_38]], %[[VAL_39]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_41:.*]] = mhlo.multiply %[[VAL_20]], %[[VAL_38]] : tensor<2xf32>
// CHLO: %[[VAL_42:.*]] = mhlo.divide %[[VAL_39]], %[[VAL_38]] : tensor<2xf32>
// CHLO: %[[VAL_43:.*]] = mhlo.multiply %[[VAL_42]], %[[VAL_42]] : tensor<2xf32>
// CHLO: %[[VAL_44:.*]] = mhlo.add %[[VAL_11]], %[[VAL_43]] : tensor<2xf32>
// CHLO: %[[VAL_45:.*]] = mhlo.sqrt %[[VAL_44]] : tensor<2xf32>
// CHLO: %[[VAL_46:.*]] = mhlo.compare EQ, %[[VAL_45]], %[[VAL_11]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_47:.*]] = mhlo.compare GT, %[[VAL_43]], %[[VAL_27]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_48:.*]] = mhlo.and %[[VAL_46]], %[[VAL_47]] : tensor<2xi1>
// CHLO: %[[VAL_49:.*]] = mhlo.multiply %[[VAL_38]], %[[VAL_43]] : tensor<2xf32>
// CHLO: %[[VAL_50:.*]] = mhlo.divide %[[VAL_49]], %[[VAL_19]] : tensor<2xf32>
// CHLO: %[[VAL_51:.*]] = mhlo.add %[[VAL_38]], %[[VAL_50]] : tensor<2xf32>
// CHLO: %[[VAL_52:.*]] = mhlo.multiply %[[VAL_38]], %[[VAL_45]] : tensor<2xf32>
// CHLO: %[[VAL_53:.*]] = mhlo.select %[[VAL_48]], %[[VAL_51]], %[[VAL_52]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_54:.*]] = mhlo.select %[[VAL_40]], %[[VAL_41]], %[[VAL_53]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_55:.*]] = mhlo.add %[[VAL_35]], %[[VAL_54]] : tensor<2xf32>
// CHLO: %[[VAL_56:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_55]] : tensor<2xf32>
// CHLO: %[[VAL_57:.*]] = mhlo.add %[[VAL_56]], %[[VAL_2]] : tensor<2xf32>
// CHLO: %[[VAL_58:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_57]] : tensor<2xf32>
// CHLO: %[[VAL_59:.*]] = mhlo.multiply %[[VAL_4]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_60:.*]] = mhlo.add %[[VAL_35]], %[[VAL_14]] : tensor<2xf32>
// CHLO: %[[VAL_61:.*]] = mhlo.divide %[[VAL_59]], %[[VAL_60]] : tensor<2xf32>
// CHLO: %[[VAL_62:.*]] = mhlo.subtract %[[VAL_54]], %[[VAL_36]] : tensor<2xf32>
// CHLO: %[[VAL_63:.*]] = mhlo.add %[[VAL_61]], %[[VAL_62]] : tensor<2xf32>
// CHLO: %[[VAL_64:.*]] = mhlo.multiply %[[VAL_58]], %[[VAL_63]] : tensor<2xf32>
// CHLO: %[[VAL_65:.*]] = mhlo.sqrt %[[VAL_64]] : tensor<2xf32>
// CHLO: %[[VAL_66:.*]] = mhlo.divide %[[VAL_58]], %[[VAL_60]] : tensor<2xf32>
// CHLO: %[[VAL_67:.*]] = mhlo.add %[[VAL_54]], %[[VAL_36]] : tensor<2xf32>
// CHLO: %[[VAL_68:.*]] = mhlo.divide %[[VAL_58]], %[[VAL_67]] : tensor<2xf32>
// CHLO: %[[VAL_69:.*]] = mhlo.add %[[VAL_66]], %[[VAL_68]] : tensor<2xf32>
// CHLO: %[[VAL_70:.*]] = mhlo.sqrt %[[VAL_69]] : tensor<2xf32>
// CHLO: %[[VAL_71:.*]] = mhlo.multiply %[[VAL_4]], %[[VAL_70]] : tensor<2xf32>
// CHLO: %[[VAL_72:.*]] = mhlo.select %[[VAL_12]], %[[VAL_65]], %[[VAL_71]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_73:.*]] = mhlo.select %[[VAL_10]], %[[VAL_4]], %[[VAL_72]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_74:.*]] = mhlo.atan2 %[[VAL_73]], %[[VAL_1]] : tensor<2xf32>
// CHLO: %[[VAL_75:.*]] = mhlo.compare LT, %[[VAL_3]], %[[VAL_27]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_76:.*]] = mhlo.constant dense<9.99999995E+11> : tensor<2xf32>
// CHLO: %[[VAL_77:.*]] = mhlo.multiply %[[VAL_9]], %[[VAL_76]] : tensor<2xf32>
// CHLO: %[[VAL_78:.*]] = mhlo.compare LT, %[[VAL_2]], %[[VAL_77]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_79:.*]] = mhlo.constant dense<9.99999997E-7> : tensor<2xf32>
// CHLO: %[[VAL_80:.*]] = mhlo.multiply %[[VAL_9]], %[[VAL_79]] : tensor<2xf32>
// CHLO: %[[VAL_81:.*]] = mhlo.constant dense<1.000000e+02> : tensor<2xf32>
// CHLO: %[[VAL_82:.*]] = mhlo.multiply %[[VAL_9]], %[[VAL_81]] : tensor<2xf32>
// CHLO: %[[VAL_83:.*]] = mhlo.select %[[VAL_78]], %[[VAL_80]], %[[VAL_82]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_84:.*]] = mhlo.compare GE, %[[VAL_4]], %[[VAL_83]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_85:.*]] = mhlo.select %[[VAL_84]], %[[VAL_4]], %[[VAL_2]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_86:.*]] = mhlo.select %[[VAL_84]], %[[VAL_83]], %[[VAL_9]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_87:.*]] = mhlo.compare GE, %[[VAL_85]], %[[VAL_86]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_88:.*]] = mhlo.log %[[VAL_19]] : tensor<2xf32>
// CHLO: %[[VAL_89:.*]] = mhlo.log %[[VAL_85]] : tensor<2xf32>
// CHLO: %[[VAL_90:.*]] = mhlo.add %[[VAL_88]], %[[VAL_89]] : tensor<2xf32>
// CHLO: %[[VAL_91:.*]] = mhlo.constant dense<0x7F800000> : tensor<2xf32>
// CHLO: %[[VAL_92:.*]] = mhlo.compare EQ, %[[VAL_4]], %[[VAL_91]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_93:.*]] = mhlo.not %[[VAL_92]] : tensor<2xi1>
// CHLO: %[[VAL_94:.*]] = mhlo.and %[[VAL_84]], %[[VAL_93]] : tensor<2xi1>
// CHLO: %[[VAL_95:.*]] = mhlo.divide %[[VAL_2]], %[[VAL_4]] : tensor<2xf32>
// CHLO: %[[VAL_96:.*]] = mhlo.select %[[VAL_94]], %[[VAL_95]], %[[VAL_27]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_97:.*]] = mhlo.multiply %[[VAL_96]], %[[VAL_96]] : tensor<2xf32>
// CHLO: %[[VAL_98:.*]] = mhlo.log_plus_one %[[VAL_97]] : tensor<2xf32>
// CHLO: %[[VAL_99:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_98]] : tensor<2xf32>
// CHLO: %[[VAL_100:.*]] = mhlo.add %[[VAL_90]], %[[VAL_99]] : tensor<2xf32>
// CHLO: %[[VAL_101:.*]] = mhlo.constant dense<1.17549435E-38> : tensor<2xf32>
// CHLO: %[[VAL_102:.*]] = mhlo.sqrt %[[VAL_101]] : tensor<2xf32>
// CHLO: %[[VAL_103:.*]] = mhlo.constant dense<4.000000e+00> : tensor<2xf32>
// CHLO: %[[VAL_104:.*]] = mhlo.multiply %[[VAL_102]], %[[VAL_103]] : tensor<2xf32>
// CHLO: %[[VAL_105:.*]] = mhlo.compare LT, %[[VAL_4]], %[[VAL_104]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_106:.*]] = mhlo.compare LT, %[[VAL_2]], %[[VAL_11]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_107:.*]] = mhlo.and %[[VAL_105]], %[[VAL_106]] : tensor<2xi1>
// CHLO: %[[VAL_108:.*]] = mhlo.multiply %[[VAL_14]], %[[VAL_36]] : tensor<2xf32>
// CHLO: %[[VAL_109:.*]] = mhlo.add %[[VAL_56]], %[[VAL_11]] : tensor<2xf32>
// CHLO: %[[VAL_110:.*]] = mhlo.divide %[[VAL_108]], %[[VAL_109]] : tensor<2xf32>
// CHLO: %[[VAL_111:.*]] = mhlo.negate %[[VAL_110]] : tensor<2xf32>
// CHLO: %[[VAL_112:.*]] = mhlo.compare GE, %[[VAL_2]], %[[VAL_11]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_113:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_59]] : tensor<2xf32>
// CHLO: %[[VAL_114:.*]] = mhlo.divide %[[VAL_113]], %[[VAL_60]] : tensor<2xf32>
// CHLO: %[[VAL_115:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_67]] : tensor<2xf32>
// CHLO: %[[VAL_116:.*]] = mhlo.add %[[VAL_114]], %[[VAL_115]] : tensor<2xf32>
// CHLO: %[[VAL_117:.*]] = mhlo.constant dense<1.500000e+00> : tensor<2xf32>
// CHLO: %[[VAL_118:.*]] = mhlo.compare LE, %[[VAL_56]], %[[VAL_117]] : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// CHLO: %[[VAL_119:.*]] = mhlo.divide %[[VAL_113]], %[[VAL_62]] : tensor<2xf32>
// CHLO: %[[VAL_120:.*]] = mhlo.add %[[VAL_114]], %[[VAL_119]] : tensor<2xf32>
// CHLO: %[[VAL_121:.*]] = mhlo.subtract %[[VAL_56]], %[[VAL_11]] : tensor<2xf32>
// CHLO: %[[VAL_122:.*]] = mhlo.select %[[VAL_118]], %[[VAL_120]], %[[VAL_121]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_123:.*]] = mhlo.select %[[VAL_112]], %[[VAL_116]], %[[VAL_122]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_124:.*]] = mhlo.select %[[VAL_107]], %[[VAL_111]], %[[VAL_123]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_125:.*]] = mhlo.multiply %[[VAL_124]], %[[VAL_109]] : tensor<2xf32>
// CHLO: %[[VAL_126:.*]] = mhlo.sqrt %[[VAL_125]] : tensor<2xf32>
// CHLO: %[[VAL_127:.*]] = mhlo.divide %[[VAL_4]], %[[VAL_126]] : tensor<2xf32>
// CHLO: %[[VAL_128:.*]] = mhlo.add %[[VAL_124]], %[[VAL_126]] : tensor<2xf32>
// CHLO: %[[VAL_129:.*]] = mhlo.log_plus_one %[[VAL_128]] : tensor<2xf32>
// CHLO: %[[VAL_130:.*]] = mhlo.select %[[VAL_107]], %[[VAL_127]], %[[VAL_129]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_131:.*]] = mhlo.select %[[VAL_87]], %[[VAL_100]], %[[VAL_130]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_132:.*]] = mhlo.negate %[[VAL_131]] : tensor<2xf32>
// CHLO: %[[VAL_133:.*]] = mhlo.select %[[VAL_75]], %[[VAL_131]], %[[VAL_132]] : tensor<2xi1>, tensor<2xf32>
// CHLO: %[[VAL_134:.*]] = mhlo.complex %[[VAL_74]], %[[VAL_133]] : tensor<2xcomplex<f32>>
// CHLO: return %[[VAL_134]] : tensor<2xcomplex<f32>>
%0 = "tf.Acos"(%arg0) : (tensor<2xcomplex<f32>>) -> tensor<2xcomplex<f32>>
func.return %0 : tensor<2xcomplex<f32>>
}

123
third_party/stablehlo/temporary.patch vendored
View File

@ -1,99 +1,32 @@
diff --ruN a/stablehlo/CMakeLists.txt b/stablehlo/CMakeLists.txt
--- stablehlo/CMakeLists.txt
+++ stablehlo/CMakeLists.txt
@@ -53,16 +53,24 @@
#-------------------------------------------------------------------------------
# Project setup and globals
#-------------------------------------------------------------------------------
+
+# There are 3 build modes, one will be set to ON
+# - Standalone: Build MLIR as a part of StableHLO, requires registering LLVM globals
+# - External: StableHLO built as an external LLVM project (XLA/MHLO uses this)
+# - Embedded: StableHLO built as a part of another MLIR project (torch-mlir uses this)
+#
+# If building as part of another project, let it handle the MLIR dependency.
+# The dependent project might use a bundled version of MLIR instead of installing.
set(STABLEHLO_EXTERNAL_PROJECT_BUILD OFF)
-
-if(NOT (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) AND NOT MLIR_BINARY_DIR)
- # Building as part of LLVM via the external project mechanism.
- set(STABLEHLO_EXTERNAL_PROJECT_BUILD ON)
-else()
- # Building standalone.
- project(stablehlo LANGUAGES CXX C)
- set(CMAKE_C_STANDARD 11)
- set(CMAKE_CXX_STANDARD 17)
+set(STABLEHLO_STANDALONE_BUILD OFF)
+# STABLEHLO_BUILD_EMBEDDED declared in above options
+
+if (NOT STABLEHLO_BUILD_EMBEDDED)
+ if ((CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) OR MLIR_BINARY_DIR)
+ set(STABLEHLO_STANDALONE_BUILD ON)
+ else()
+ set(STABLEHLO_EXTERNAL_PROJECT_BUILD ON)
+ endif()
endif()
diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.cpp b/stablehlo/stablehlo/dialect/StablehloOps.cpp
--- stablehlo/stablehlo/dialect/StablehloOps.cpp
+++ stablehlo/stablehlo/dialect/StablehloOps.cpp
@@ -518,6 +518,16 @@
return StringAttr::get(getContext(), "");
}
#-------------------------------------------------------------------------------
@@ -74,9 +82,6 @@
set(LLVM_ENABLE_PEDANTIC ON)
endif()
+// Returns if the backend config is unset, or if empty dict / string attribute.
+bool CustomCallOp::hasEmptyBackendConfig() {
+ if (!getBackendConfig().has_value()) return true;
+ Attribute backendConfig = getBackendConfigOrDefault();
+ if (auto strAttr = dyn_cast<StringAttr>(backendConfig)) {
+ return strAttr.empty();
+ }
+ return cast<DictionaryAttr>(backendConfig).empty();
+}
+
//===----------------------------------------------------------------------===//
// CholeskyOp
//===----------------------------------------------------------------------===//
diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.td b/stablehlo/stablehlo/dialect/StablehloOps.td
--- stablehlo/stablehlo/dialect/StablehloOps.td
+++ stablehlo/stablehlo/dialect/StablehloOps.td
@@ -2389,6 +2389,7 @@
-# Find MLIR to install if we are building standalone. If building as part of
-# another project, let it handle the MLIR dependency. The dependent project
-# might use a bundled version of MLIR instead of installing, for instance.
if(STABLEHLO_EXTERNAL_PROJECT_BUILD)
message(STATUS "Building StableHLO as an external LLVM project")
set(MLIR_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../mlir ) # --src-root
@@ -88,17 +93,35 @@
let extraClassDeclaration = commonClassDeclaration # [{
mlir::Attribute getBackendConfigOrDefault();
+ bool hasEmptyBackendConfig();
}];
}
set(BACKEND_PACKAGE_STRING "${PACKAGE_STRING}")
list(APPEND CMAKE_MODULE_PATH "${MLIR_MAIN_SRC_DIR}/cmake/modules")
-elseif(NOT STABLEHLO_BUILD_EMBEDDED)
- message(STATUS "Building StableHLO with an installed MLIR")
+endif()
+
+if(STABLEHLO_STANDALONE_BUILD)
+ message("Building StableHLO as a standalone project.")
+ project(stablehlo LANGUAGES CXX C)
+ set(CMAKE_C_STANDARD 11)
+ set(CMAKE_CXX_STANDARD 17)
+
find_package(MLIR REQUIRED CONFIG)
message(STATUS "Using MLIRConfig.cmake in: ${MLIR_DIR}")
message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}")
+
set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin)
set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib)
list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}")
list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}")
-else()
+
+ include(TableGen)
+ include(AddLLVM)
+ include(AddMLIR)
+ include(HandleLLVMOptions)
+endif()
+
+if(STABLEHLO_BUILD_EMBEDDED)
message(STATUS "Building StableHLO embedded in another project")
+ include(TableGen)
+ include(AddLLVM)
+ include(AddMLIR)
+ include(HandleLLVMOptions)
endif()
# Add the CMake modules specific to StableHLO
@@ -134,10 +157,7 @@
endif()
endif()
-include(TableGen)
-include(AddLLVM)
-include(AddMLIR)
-include(HandleLLVMOptions)
+#TODO: Where should these be?
include_directories(${LLVM_INCLUDE_DIRS})
include_directories(${MLIR_INCLUDE_DIRS})
include_directories(${CMAKE_CURRENT_SOURCE_DIR})

4
third_party/stablehlo/workspace.bzl vendored
View File

@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls")
def repo():
# LINT.IfChange
STABLEHLO_COMMIT = "1b08c4c0e8c893d202bd33c2735562fa0ccc849f"
STABLEHLO_SHA256 = "271fc32368df969421400765c0b383b7452ba4a4e2fea8fc02fd5cf276c850ec"
STABLEHLO_COMMIT = "8817ff1d7105b027f1f50e671a52121c7cf51ea7"
STABLEHLO_SHA256 = "37f97c0954afcc8cafeb66e6c031411c8789aea307d154218e77b6e253290264"
# LINT.ThenChange(Google-internal path)
tf_http_archive(

123
third_party/xla/third_party/stablehlo/temporary.patch vendored
View File

@ -1,99 +1,32 @@
diff --ruN a/stablehlo/CMakeLists.txt b/stablehlo/CMakeLists.txt
--- stablehlo/CMakeLists.txt
+++ stablehlo/CMakeLists.txt
@@ -53,16 +53,24 @@
#-------------------------------------------------------------------------------
# Project setup and globals
#-------------------------------------------------------------------------------
+
+# There are 3 build modes, one will be set to ON
+# - Standalone: Build MLIR as a part of StableHLO, requires registering LLVM globals
+# - External: StableHLO built as an external LLVM project (XLA/MHLO uses this)
+# - Embedded: StableHLO built as a part of another MLIR project (torch-mlir uses this)
+#
+# If building as part of another project, let it handle the MLIR dependency.
+# The dependent project might use a bundled version of MLIR instead of installing.
set(STABLEHLO_EXTERNAL_PROJECT_BUILD OFF)
-
-if(NOT (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) AND NOT MLIR_BINARY_DIR)
- # Building as part of LLVM via the external project mechanism.
- set(STABLEHLO_EXTERNAL_PROJECT_BUILD ON)
-else()
- # Building standalone.
- project(stablehlo LANGUAGES CXX C)
- set(CMAKE_C_STANDARD 11)
- set(CMAKE_CXX_STANDARD 17)
+set(STABLEHLO_STANDALONE_BUILD OFF)
+# STABLEHLO_BUILD_EMBEDDED declared in above options
+
+if (NOT STABLEHLO_BUILD_EMBEDDED)
+ if ((CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR) OR MLIR_BINARY_DIR)
+ set(STABLEHLO_STANDALONE_BUILD ON)
+ else()
+ set(STABLEHLO_EXTERNAL_PROJECT_BUILD ON)
+ endif()
endif()
diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.cpp b/stablehlo/stablehlo/dialect/StablehloOps.cpp
--- stablehlo/stablehlo/dialect/StablehloOps.cpp
+++ stablehlo/stablehlo/dialect/StablehloOps.cpp
@@ -518,6 +518,16 @@
return StringAttr::get(getContext(), "");
}
#-------------------------------------------------------------------------------
@@ -74,9 +82,6 @@
set(LLVM_ENABLE_PEDANTIC ON)
endif()
+// Returns if the backend config is unset, or if empty dict / string attribute.
+bool CustomCallOp::hasEmptyBackendConfig() {
+ if (!getBackendConfig().has_value()) return true;
+ Attribute backendConfig = getBackendConfigOrDefault();
+ if (auto strAttr = dyn_cast<StringAttr>(backendConfig)) {
+ return strAttr.empty();
+ }
+ return cast<DictionaryAttr>(backendConfig).empty();
+}
+
//===----------------------------------------------------------------------===//
// CholeskyOp
//===----------------------------------------------------------------------===//
diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.td b/stablehlo/stablehlo/dialect/StablehloOps.td
--- stablehlo/stablehlo/dialect/StablehloOps.td
+++ stablehlo/stablehlo/dialect/StablehloOps.td
@@ -2389,6 +2389,7 @@
-# Find MLIR to install if we are building standalone. If building as part of
-# another project, let it handle the MLIR dependency. The dependent project
-# might use a bundled version of MLIR instead of installing, for instance.
if(STABLEHLO_EXTERNAL_PROJECT_BUILD)
message(STATUS "Building StableHLO as an external LLVM project")
set(MLIR_MAIN_SRC_DIR ${LLVM_MAIN_SRC_DIR}/../mlir ) # --src-root
@@ -88,17 +93,35 @@
let extraClassDeclaration = commonClassDeclaration # [{
mlir::Attribute getBackendConfigOrDefault();
+ bool hasEmptyBackendConfig();
}];
}
set(BACKEND_PACKAGE_STRING "${PACKAGE_STRING}")
list(APPEND CMAKE_MODULE_PATH "${MLIR_MAIN_SRC_DIR}/cmake/modules")
-elseif(NOT STABLEHLO_BUILD_EMBEDDED)
- message(STATUS "Building StableHLO with an installed MLIR")
+endif()
+
+if(STABLEHLO_STANDALONE_BUILD)
+ message("Building StableHLO as a standalone project.")
+ project(stablehlo LANGUAGES CXX C)
+ set(CMAKE_C_STANDARD 11)
+ set(CMAKE_CXX_STANDARD 17)
+
find_package(MLIR REQUIRED CONFIG)
message(STATUS "Using MLIRConfig.cmake in: ${MLIR_DIR}")
message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}")
+
set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin)
set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib)
list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}")
list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}")
-else()
+
+ include(TableGen)
+ include(AddLLVM)
+ include(AddMLIR)
+ include(HandleLLVMOptions)
+endif()
+
+if(STABLEHLO_BUILD_EMBEDDED)
message(STATUS "Building StableHLO embedded in another project")
+ include(TableGen)
+ include(AddLLVM)
+ include(AddMLIR)
+ include(HandleLLVMOptions)
endif()
# Add the CMake modules specific to StableHLO
@@ -134,10 +157,7 @@
endif()
endif()
-include(TableGen)
-include(AddLLVM)
-include(AddMLIR)
-include(HandleLLVMOptions)
+#TODO: Where should these be?
include_directories(${LLVM_INCLUDE_DIRS})
include_directories(${MLIR_INCLUDE_DIRS})
include_directories(${CMAKE_CURRENT_SOURCE_DIR})

4
third_party/xla/third_party/stablehlo/workspace.bzl vendored
View File

@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls")
def repo():
# LINT.IfChange
STABLEHLO_COMMIT = "1b08c4c0e8c893d202bd33c2735562fa0ccc849f"
STABLEHLO_SHA256 = "271fc32368df969421400765c0b383b7452ba4a4e2fea8fc02fd5cf276c850ec"
STABLEHLO_COMMIT = "8817ff1d7105b027f1f50e671a52121c7cf51ea7"
STABLEHLO_SHA256 = "37f97c0954afcc8cafeb66e6c031411c8789aea307d154218e77b6e253290264"
# LINT.ThenChange(Google-internal path)
tf_http_archive(

8
third_party/xla/xla/mlir_hlo/mhlo/transforms/hlo_legalize_to_stablehlo/hlo_legalize_to_stablehlo.cc vendored
View File

@ -135,14 +135,6 @@ std::optional<int64_t> getPublicFeaturesNotInStablehlo(HloOpTy hloOp) {
// Version 1: Initial version for TanOp.
return 1;
}
// StableHLO CustomCall doesn't support API_VERSION_TYPED_FFI yet.
// Proposal: https://github.com/openxla/stablehlo/issues/637.
if constexpr (std::is_same<HloOpTy, mhlo::CustomCallOp>::value) {
// Version 1: Initial version for TYPED_FFI
if (hloOp.getApiVersion() ==
mhlo::CustomCallApiVersion::API_VERSION_TYPED_FFI)
return 1;
}
// StableHLO doesn't support TopK yet.
// Proposal: https://github.com/openxla/stablehlo/pull/1593
if constexpr (std::is_same<HloOpTy, mhlo::TopKOp>::value) {

6
third_party/xla/xla/mlir_hlo/mhlo/transforms/stablehlo_legalize_to_hlo/stablehlo_legalize_to_hlo.cc vendored
View File

@ -231,7 +231,7 @@ LogicalResult rewriteCustomCallAsMhloOp(stablehlo::CustomCallOp stablehloOp,
name == "mhlo.version";
};
if (!llvm::all_of(stablehloOp->getAttrs(), isSupportedAttrName) ||
!stablehloOp.getBackendConfig().empty()) {
!stablehloOp.hasEmptyBackendConfig()) {
return failure();
}
@ -294,8 +294,8 @@ LogicalResult fixupMhloBackendConfig(stablehlo::CustomCallOp stablehloOp,
mhlo::CustomCallOp hloOp) {
auto stablehloBackendConfig = stablehloOp->getAttr("mhlo.backend_config");
if (stablehloBackendConfig) {
if (auto oldHloBackendConfig =
mlir::dyn_cast_or_null<StringAttr>(hloOp.getBackendConfigAttr())) {
if (auto oldHloBackendConfig = mlir::dyn_cast<StringAttr>(
stablehloOp.getBackendConfigOrDefault())) {
if (!oldHloBackendConfig.empty()) return failure();
} else {
return failure();

6
third_party/xla/xla/mlir_hlo/stablehlo_ext/IR/stablehlo_ops.cpp vendored
View File

@ -51,7 +51,7 @@ LogicalResult DynamicReduceWindowOpAdaptor::verify() {
return op_.emitError()
<< attr.getName() << " is not a supported attribute";
}
if (!op_.getBackendConfig().empty())
if (!op_.hasEmptyBackendConfig())
return op_.emitError() << "expects an empty backend_config";
if (op_.getCallTargetName() != "stablehlo.dynamic_reduce_window")
return op_.emitError() << "expects @stablehlo.dynamic_reduce_window";
@ -296,7 +296,7 @@ LogicalResult DynamicRngBitGeneratorOpAdaptor::verify() {
return op_.emitError()
<< attr.getName() << " is not a supported attribute";
}
if (!op_.getBackendConfig().empty())
if (!op_.hasEmptyBackendConfig())
return op_.emitError() << "expects an empty backend_config";
if (op_.getCallTargetName() != "stablehlo.dynamic_rng_bit_generator")
return op_.emitError() << "expects @stablehlo.dynamic_rng_bit_generator";
@ -407,7 +407,7 @@ LogicalResult DynamicTopKOpAdaptor::verify() {
return op_.emitError()
<< attr.getName() << " is not a supported attribute";
}
if (!op_.getBackendConfig().empty())
if (!op_.hasEmptyBackendConfig())
return op_.emitError() << "expects an empty backend_config";
if (op_.getCallTargetName() != "stablehlo.dynamic_top_k")
return op_.emitError() << "expects @stablehlo.dynamic_top_k";

2
third_party/xla/xla/mlir_hlo/stablehlo_ext/transforms/chlo_recompose_ops.cpp vendored
View File

@ -80,7 +80,7 @@ LogicalResult recomposeChloOpFromCustomCall(stablehlo::CustomCallOp op,
name == "mhlo.version";
};
if (!llvm::all_of(op->getAttrs(), isSupportedAttrName) ||
!op.getBackendConfig().empty()) {
!op.hasEmptyBackendConfig()) {
return rewriter.notifyMatchFailure(
op, "CHLO Recompose custom call did not have required attributes.");
}

416
third_party/xla/xla/mlir_hlo/tests/Dialect/chlo/chlo_legalize_to_mhlo.mlir vendored
View File

@ -424,66 +424,183 @@ func.func @asinh_f32(%arg : tensor<f32>) -> tensor<f32> {
// -----
// CHECK-LABEL: @asinh_f64
// CHECK-SAME: %[[ARG:.*]]: tensor<f64>
// CHECK-LABEL: @asinh_f64
// CHECK-SAME: %[[VAL_0:.*]]: tensor<f64>) -> tensor<f64> {
func.func @asinh_f64(%arg : tensor<f64>) -> tensor<f64> {
// CHECK: %[[TMP_0:.*]] = mhlo.sign %[[ARG]]
// CHECK: %[[TMP_1:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_2:.*]] = mhlo.constant dense<1.797{{.*}}E+308>
// CHECK: %[[TMP_3:.*]] = mhlo.sqrt %[[TMP_2]]
// CHECK: %[[TMP_4:.*]] = mhlo.compare GE, %[[TMP_1]], %[[TMP_3]]
// CHECK: %[[TMP_5:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_6:.*]] = mhlo.log %[[TMP_5]]
// CHECK: %[[TMP_7:.*]] = mhlo.constant dense<2.000{{.*}}e+00>
// CHECK: %[[TMP_8:.*]] = mhlo.log %[[TMP_7]]
// CHECK: %[[TMP_9:.*]] = mhlo.add %[[TMP_6]], %[[TMP_8]]
// CHECK: %[[TMP_10:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_11:.*]] = mhlo.constant dense<1.000{{.*}}e+00>
// CHECK: %[[TMP_12:.*]] = mhlo.compare LE, %[[TMP_10]], %[[TMP_11]]
// CHECK: %[[TMP_13:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_14:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_15:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_22:.*]] = mhlo.constant dense<1.000{{.*}}e+00>
// CHECK: %[[TMP_16:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_17:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_18:.*]] = mhlo.multiply %[[TMP_16]], %[[TMP_17]]
// CHECK: %[[TMP_19:.*]] = mhlo.constant dense<1.000{{.*}}e+00>
// CHECK: %[[TMP_20:.*]] = mhlo.add %[[TMP_18]], %[[TMP_19]]
// CHECK: %[[TMP_21:.*]] = mhlo.sqrt %[[TMP_20]]
// CHECK: %[[TMP_23:.*]] = mhlo.add %[[TMP_22]], %[[TMP_21]]
// CHECK: %[[TMP_24:.*]] = mhlo.divide %[[TMP_15]], %[[TMP_23]]
// CHECK: %[[TMP_25:.*]] = mhlo.multiply %[[TMP_14]], %[[TMP_24]]
// CHECK: %[[TMP_26:.*]] = mhlo.add %[[TMP_13]], %[[TMP_25]]
// CHECK: %[[TMP_27:.*]] = mhlo.log_plus_one %[[TMP_26]]
// CHECK: %[[TMP_28:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_29:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_30:.*]] = mhlo.abs %[[ARG]]
// CHECK: %[[TMP_31:.*]] = mhlo.multiply %[[TMP_29]], %[[TMP_30]]
// CHECK: %[[TMP_32:.*]] = mhlo.constant dense<1.000{{.*}}e+00>
// CHECK: %[[TMP_33:.*]] = mhlo.add %[[TMP_31]], %[[TMP_32]]
// CHECK: %[[TMP_34:.*]] = mhlo.sqrt %[[TMP_33]]
// CHECK: %[[TMP_35:.*]] = mhlo.add %[[TMP_28]], %[[TMP_34]]
// CHECK: %[[TMP_36:.*]] = mhlo.log %[[TMP_35]]
// CHECK: %[[TMP_37:.*]] = mhlo.select %[[TMP_12]], %[[TMP_27]], %[[TMP_36]]
// CHECK: %[[TMP_38:.*]] = mhlo.select %[[TMP_4]], %[[TMP_9]], %[[TMP_37]]
// CHECK: %[[RES:.*]] = mhlo.multiply %[[TMP_0]], %[[TMP_38]]
// CHECK: return %[[RES]]
// CHECK: %[[VAL_1:.*]] = mhlo.sign %[[VAL_0]] : tensor<f64>
// CHECK: %[[VAL_2:.*]] = mhlo.abs %[[VAL_0]] : tensor<f64>
// CHECK: %[[VAL_3:.*]] = mhlo.constant dense<1.7976931348623157E+308> : tensor<f64>
// CHECK: %[[VAL_4:.*]] = mhlo.sqrt %[[VAL_3]] : tensor<f64>
// CHECK: %[[VAL_5:.*]] = mhlo.compare GE, %[[VAL_2]], %[[VAL_4]] : (tensor<f64>, tensor<f64>) -> tensor<i1>
// CHECK: %[[VAL_6:.*]] = mhlo.constant dense<2.000000e+00> : tensor<f64>
// CHECK: %[[VAL_7:.*]] = mhlo.log %[[VAL_6]] : tensor<f64>
// CHECK: %[[VAL_8:.*]] = mhlo.log %[[VAL_2]] : tensor<f64>
// CHECK: %[[VAL_9:.*]] = mhlo.add %[[VAL_7]], %[[VAL_8]] : tensor<f64>
// CHECK: %[[VAL_10:.*]] = mhlo.multiply %[[VAL_2]], %[[VAL_2]] : tensor<f64>
// CHECK: %[[VAL_11:.*]] = mhlo.constant dense<1.000000e+00> : tensor<f64>
// CHECK: %[[VAL_12:.*]] = mhlo.add %[[VAL_11]], %[[VAL_10]] : tensor<f64>
// CHECK: %[[VAL_13:.*]] = mhlo.sqrt %[[VAL_12]] : tensor<f64>
// CHECK: %[[VAL_14:.*]] = mhlo.add %[[VAL_11]], %[[VAL_13]] : tensor<f64>
// CHECK: %[[VAL_15:.*]] = mhlo.divide %[[VAL_10]], %[[VAL_14]] : tensor<f64>
// CHECK: %[[VAL_16:.*]] = mhlo.add %[[VAL_2]], %[[VAL_15]] : tensor<f64>
// CHECK: %[[VAL_17:.*]] = mhlo.log_plus_one %[[VAL_16]] : tensor<f64>
// CHECK: %[[VAL_18:.*]] = mhlo.select %[[VAL_5]], %[[VAL_9]], %[[VAL_17]] : tensor<i1>, tensor<f64>
// CHECK: %[[VAL_19:.*]] = mhlo.multiply %[[VAL_1]], %[[VAL_18]] : tensor<f64>
// CHECK: return %[[VAL_19]] : tensor<f64>
// CHECK: }
%result = "chlo.asinh"(%arg) : (tensor<f64>) -> tensor<f64>
func.return %result : tensor<f64>
}
// -----
// CHECK-LABEL: @asinh_complex_f32
// CHECK-SAME: %[[ARG:.*]]: tensor<complex<f32>>
// CHECK-LABEL: func.func @asinh_complex_f32(
// CHECK-SAME: %[[VAL_0:.*]]: tensor<complex<f32>>) -> tensor<complex<f32>> {
func.func @asinh_complex_f32(%arg : tensor<complex<f32>>) -> tensor<complex<f32>> {
// CHECK: %[[TMP_0:.*]] = mhlo.multiply %[[ARG]], %[[ARG]]
// CHECK: %[[TMP_1:.*]] = mhlo.constant dense<(1.000000e+00,0.000000e+00)>
// CHECK: %[[TMP_2:.*]] = mhlo.add %[[TMP_0]], %[[TMP_1]]
// CHECK: %[[TMP_3:.*]] = mhlo.sqrt %[[TMP_2]]
// CHECK: %[[TMP_4:.*]] = mhlo.add %[[ARG]], %[[TMP_3]]
// CHECK: %[[TMP_5:.*]] = mhlo.log %[[TMP_4]]
// CHECK: return %[[TMP_5]]
// CHECK: %[[VAL_1:.*]] = mhlo.imag %[[VAL_0]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_2:.*]] = mhlo.negate %[[VAL_1]] : tensor<f32>
// CHECK: %[[VAL_3:.*]] = mhlo.real %[[VAL_0]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_4:.*]] = mhlo.complex %[[VAL_2]], %[[VAL_3]] : tensor<complex<f32>>
// CHECK: %[[VAL_5:.*]] = mhlo.real %[[VAL_4]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_6:.*]] = mhlo.abs %[[VAL_5]] : tensor<f32>
// CHECK: %[[VAL_7:.*]] = mhlo.imag %[[VAL_4]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_8:.*]] = mhlo.abs %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_9:.*]] = mhlo.maximum %[[VAL_6]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_10:.*]] = mhlo.constant dense<3.40282347E+38> : tensor<f32>
// CHECK: %[[VAL_11:.*]] = mhlo.sqrt %[[VAL_10]] : tensor<f32>
// CHECK: %[[VAL_12:.*]] = mhlo.constant dense<8.000000e+00> : tensor<f32>
// CHECK: %[[VAL_13:.*]] = mhlo.divide %[[VAL_11]], %[[VAL_12]] : tensor<f32>
// CHECK: %[[VAL_14:.*]] = mhlo.compare GE, %[[VAL_9]], %[[VAL_13]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_15:.*]] = mhlo.constant dense<1.000000e+00> : tensor<f32>
// CHECK: %[[VAL_16:.*]] = mhlo.compare LE, %[[VAL_6]], %[[VAL_15]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_17:.*]] = mhlo.constant dense<5.000000e-01> : tensor<f32>
// CHECK: %[[VAL_18:.*]] = mhlo.add %[[VAL_6]], %[[VAL_15]] : tensor<f32>
// CHECK: %[[VAL_19:.*]] = mhlo.abs %[[VAL_18]] : tensor<f32>
// CHECK: %[[VAL_20:.*]] = mhlo.maximum %[[VAL_19]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_21:.*]] = mhlo.minimum %[[VAL_19]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_22:.*]] = mhlo.compare EQ, %[[VAL_20]], %[[VAL_21]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_23:.*]] = mhlo.constant dense<2.000000e+00> : tensor<f32>
// CHECK: %[[VAL_24:.*]] = mhlo.sqrt %[[VAL_23]] : tensor<f32>
// CHECK: %[[VAL_25:.*]] = mhlo.multiply %[[VAL_24]], %[[VAL_20]] : tensor<f32>
// CHECK: %[[VAL_26:.*]] = mhlo.divide %[[VAL_21]], %[[VAL_20]] : tensor<f32>
// CHECK: %[[VAL_27:.*]] = mhlo.multiply %[[VAL_26]], %[[VAL_26]] : tensor<f32>
// CHECK: %[[VAL_28:.*]] = mhlo.add %[[VAL_15]], %[[VAL_27]] : tensor<f32>
// CHECK: %[[VAL_29:.*]] = mhlo.sqrt %[[VAL_28]] : tensor<f32>
// CHECK: %[[VAL_30:.*]] = mhlo.compare EQ, %[[VAL_29]], %[[VAL_15]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_31:.*]] = mhlo.constant dense<0.000000e+00> : tensor<f32>
// CHECK: %[[VAL_32:.*]] = mhlo.compare GT, %[[VAL_27]], %[[VAL_31]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_33:.*]] = mhlo.and %[[VAL_30]], %[[VAL_32]] : tensor<i1>
// CHECK: %[[VAL_34:.*]] = mhlo.multiply %[[VAL_20]], %[[VAL_27]] : tensor<f32>
// CHECK: %[[VAL_35:.*]] = mhlo.divide %[[VAL_34]], %[[VAL_23]] : tensor<f32>
// CHECK: %[[VAL_36:.*]] = mhlo.add %[[VAL_20]], %[[VAL_35]] : tensor<f32>
// CHECK: %[[VAL_37:.*]] = mhlo.multiply %[[VAL_20]], %[[VAL_29]] : tensor<f32>
// CHECK: %[[VAL_38:.*]] = mhlo.select %[[VAL_33]], %[[VAL_36]], %[[VAL_37]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_39:.*]] = mhlo.select %[[VAL_22]], %[[VAL_25]], %[[VAL_38]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_40:.*]] = mhlo.subtract %[[VAL_6]], %[[VAL_15]] : tensor<f32>
// CHECK: %[[VAL_41:.*]] = mhlo.abs %[[VAL_40]] : tensor<f32>
// CHECK: %[[VAL_42:.*]] = mhlo.maximum %[[VAL_41]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_43:.*]] = mhlo.minimum %[[VAL_41]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_44:.*]] = mhlo.compare EQ, %[[VAL_42]], %[[VAL_43]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_45:.*]] = mhlo.multiply %[[VAL_24]], %[[VAL_42]] : tensor<f32>
// CHECK: %[[VAL_46:.*]] = mhlo.divide %[[VAL_43]], %[[VAL_42]] : tensor<f32>
// CHECK: %[[VAL_47:.*]] = mhlo.multiply %[[VAL_46]], %[[VAL_46]] : tensor<f32>
// CHECK: %[[VAL_48:.*]] = mhlo.add %[[VAL_15]], %[[VAL_47]] : tensor<f32>
// CHECK: %[[VAL_49:.*]] = mhlo.sqrt %[[VAL_48]] : tensor<f32>
// CHECK: %[[VAL_50:.*]] = mhlo.compare EQ, %[[VAL_49]], %[[VAL_15]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_51:.*]] = mhlo.compare GT, %[[VAL_47]], %[[VAL_31]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_52:.*]] = mhlo.and %[[VAL_50]], %[[VAL_51]] : tensor<i1>
// CHECK: %[[VAL_53:.*]] = mhlo.multiply %[[VAL_42]], %[[VAL_47]] : tensor<f32>
// CHECK: %[[VAL_54:.*]] = mhlo.divide %[[VAL_53]], %[[VAL_23]] : tensor<f32>
// CHECK: %[[VAL_55:.*]] = mhlo.add %[[VAL_42]], %[[VAL_54]] : tensor<f32>
// CHECK: %[[VAL_56:.*]] = mhlo.multiply %[[VAL_42]], %[[VAL_49]] : tensor<f32>
// CHECK: %[[VAL_57:.*]] = mhlo.select %[[VAL_52]], %[[VAL_55]], %[[VAL_56]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_58:.*]] = mhlo.select %[[VAL_44]], %[[VAL_45]], %[[VAL_57]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_59:.*]] = mhlo.add %[[VAL_39]], %[[VAL_58]] : tensor<f32>
// CHECK: %[[VAL_60:.*]] = mhlo.multiply %[[VAL_17]], %[[VAL_59]] : tensor<f32>
// CHECK: %[[VAL_61:.*]] = mhlo.add %[[VAL_60]], %[[VAL_6]] : tensor<f32>
// CHECK: %[[VAL_62:.*]] = mhlo.multiply %[[VAL_17]], %[[VAL_61]] : tensor<f32>
// CHECK: %[[VAL_63:.*]] = mhlo.multiply %[[VAL_8]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_64:.*]] = mhlo.add %[[VAL_39]], %[[VAL_18]] : tensor<f32>
// CHECK: %[[VAL_65:.*]] = mhlo.divide %[[VAL_63]], %[[VAL_64]] : tensor<f32>
// CHECK: %[[VAL_66:.*]] = mhlo.subtract %[[VAL_58]], %[[VAL_40]] : tensor<f32>
// CHECK: %[[VAL_67:.*]] = mhlo.add %[[VAL_65]], %[[VAL_66]] : tensor<f32>
// CHECK: %[[VAL_68:.*]] = mhlo.multiply %[[VAL_62]], %[[VAL_67]] : tensor<f32>
// CHECK: %[[VAL_69:.*]] = mhlo.sqrt %[[VAL_68]] : tensor<f32>
// CHECK: %[[VAL_70:.*]] = mhlo.divide %[[VAL_62]], %[[VAL_64]] : tensor<f32>
// CHECK: %[[VAL_71:.*]] = mhlo.add %[[VAL_58]], %[[VAL_40]] : tensor<f32>
// CHECK: %[[VAL_72:.*]] = mhlo.divide %[[VAL_62]], %[[VAL_71]] : tensor<f32>
// CHECK: %[[VAL_73:.*]] = mhlo.add %[[VAL_70]], %[[VAL_72]] : tensor<f32>
// CHECK: %[[VAL_74:.*]] = mhlo.sqrt %[[VAL_73]] : tensor<f32>
// CHECK: %[[VAL_75:.*]] = mhlo.multiply %[[VAL_8]], %[[VAL_74]] : tensor<f32>
// CHECK: %[[VAL_76:.*]] = mhlo.select %[[VAL_16]], %[[VAL_69]], %[[VAL_75]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_77:.*]] = mhlo.select %[[VAL_14]], %[[VAL_8]], %[[VAL_76]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_78:.*]] = mhlo.atan2 %[[VAL_5]], %[[VAL_77]] : tensor<f32>
// CHECK: %[[VAL_79:.*]] = mhlo.compare LT, %[[VAL_7]], %[[VAL_31]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_80:.*]] = mhlo.constant dense<9.99999995E+11> : tensor<f32>
// CHECK: %[[VAL_81:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_80]] : tensor<f32>
// CHECK: %[[VAL_82:.*]] = mhlo.compare LT, %[[VAL_6]], %[[VAL_81]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_83:.*]] = mhlo.constant dense<9.99999997E-7> : tensor<f32>
// CHECK: %[[VAL_84:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_83]] : tensor<f32>
// CHECK: %[[VAL_85:.*]] = mhlo.constant dense<1.000000e+02> : tensor<f32>
// CHECK: %[[VAL_86:.*]] = mhlo.multiply %[[VAL_13]], %[[VAL_85]] : tensor<f32>
// CHECK: %[[VAL_87:.*]] = mhlo.select %[[VAL_82]], %[[VAL_84]], %[[VAL_86]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_88:.*]] = mhlo.compare GE, %[[VAL_8]], %[[VAL_87]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_89:.*]] = mhlo.select %[[VAL_88]], %[[VAL_8]], %[[VAL_6]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_90:.*]] = mhlo.select %[[VAL_88]], %[[VAL_87]], %[[VAL_13]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_91:.*]] = mhlo.compare GE, %[[VAL_89]], %[[VAL_90]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_92:.*]] = mhlo.log %[[VAL_23]] : tensor<f32>
// CHECK: %[[VAL_93:.*]] = mhlo.log %[[VAL_89]] : tensor<f32>
// CHECK: %[[VAL_94:.*]] = mhlo.add %[[VAL_92]], %[[VAL_93]] : tensor<f32>
// CHECK: %[[VAL_95:.*]] = mhlo.constant dense<0x7F800000> : tensor<f32>
// CHECK: %[[VAL_96:.*]] = mhlo.compare EQ, %[[VAL_8]], %[[VAL_95]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_97:.*]] = mhlo.not %[[VAL_96]] : tensor<i1>
// CHECK: %[[VAL_98:.*]] = mhlo.and %[[VAL_88]], %[[VAL_97]] : tensor<i1>
// CHECK: %[[VAL_99:.*]] = mhlo.divide %[[VAL_6]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_100:.*]] = mhlo.select %[[VAL_98]], %[[VAL_99]], %[[VAL_31]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_101:.*]] = mhlo.multiply %[[VAL_100]], %[[VAL_100]] : tensor<f32>
// CHECK: %[[VAL_102:.*]] = mhlo.log_plus_one %[[VAL_101]] : tensor<f32>
// CHECK: %[[VAL_103:.*]] = mhlo.multiply %[[VAL_17]], %[[VAL_102]] : tensor<f32>
// CHECK: %[[VAL_104:.*]] = mhlo.add %[[VAL_94]], %[[VAL_103]] : tensor<f32>
// CHECK: %[[VAL_105:.*]] = mhlo.constant dense<1.17549435E-38> : tensor<f32>
// CHECK: %[[VAL_106:.*]] = mhlo.sqrt %[[VAL_105]] : tensor<f32>
// CHECK: %[[VAL_107:.*]] = mhlo.constant dense<4.000000e+00> : tensor<f32>
// CHECK: %[[VAL_108:.*]] = mhlo.multiply %[[VAL_106]], %[[VAL_107]] : tensor<f32>
// CHECK: %[[VAL_109:.*]] = mhlo.compare LT, %[[VAL_8]], %[[VAL_108]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_110:.*]] = mhlo.compare LT, %[[VAL_6]], %[[VAL_15]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_111:.*]] = mhlo.and %[[VAL_109]], %[[VAL_110]] : tensor<i1>
// CHECK: %[[VAL_112:.*]] = mhlo.multiply %[[VAL_18]], %[[VAL_40]] : tensor<f32>
// CHECK: %[[VAL_113:.*]] = mhlo.add %[[VAL_60]], %[[VAL_15]] : tensor<f32>
// CHECK: %[[VAL_114:.*]] = mhlo.divide %[[VAL_112]], %[[VAL_113]] : tensor<f32>
// CHECK: %[[VAL_115:.*]] = mhlo.negate %[[VAL_114]] : tensor<f32>
// CHECK: %[[VAL_116:.*]] = mhlo.compare GE, %[[VAL_6]], %[[VAL_15]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_117:.*]] = mhlo.multiply %[[VAL_17]], %[[VAL_63]] : tensor<f32>
// CHECK: %[[VAL_118:.*]] = mhlo.divide %[[VAL_117]], %[[VAL_64]] : tensor<f32>
// CHECK: %[[VAL_119:.*]] = mhlo.multiply %[[VAL_17]], %[[VAL_71]] : tensor<f32>
// CHECK: %[[VAL_120:.*]] = mhlo.add %[[VAL_118]], %[[VAL_119]] : tensor<f32>
// CHECK: %[[VAL_121:.*]] = mhlo.constant dense<1.500000e+00> : tensor<f32>
// CHECK: %[[VAL_122:.*]] = mhlo.compare LE, %[[VAL_60]], %[[VAL_121]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_123:.*]] = mhlo.divide %[[VAL_117]], %[[VAL_66]] : tensor<f32>
// CHECK: %[[VAL_124:.*]] = mhlo.add %[[VAL_118]], %[[VAL_123]] : tensor<f32>
// CHECK: %[[VAL_125:.*]] = mhlo.subtract %[[VAL_60]], %[[VAL_15]] : tensor<f32>
// CHECK: %[[VAL_126:.*]] = mhlo.select %[[VAL_122]], %[[VAL_124]], %[[VAL_125]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_127:.*]] = mhlo.select %[[VAL_116]], %[[VAL_120]], %[[VAL_126]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_128:.*]] = mhlo.select %[[VAL_111]], %[[VAL_115]], %[[VAL_127]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_129:.*]] = mhlo.multiply %[[VAL_128]], %[[VAL_113]] : tensor<f32>
// CHECK: %[[VAL_130:.*]] = mhlo.sqrt %[[VAL_129]] : tensor<f32>
// CHECK: %[[VAL_131:.*]] = mhlo.divide %[[VAL_8]], %[[VAL_130]] : tensor<f32>
// CHECK: %[[VAL_132:.*]] = mhlo.add %[[VAL_128]], %[[VAL_130]] : tensor<f32>
// CHECK: %[[VAL_133:.*]] = mhlo.log_plus_one %[[VAL_132]] : tensor<f32>
// CHECK: %[[VAL_134:.*]] = mhlo.select %[[VAL_111]], %[[VAL_131]], %[[VAL_133]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_135:.*]] = mhlo.select %[[VAL_91]], %[[VAL_104]], %[[VAL_134]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_136:.*]] = mhlo.negate %[[VAL_135]] : tensor<f32>
// CHECK: %[[VAL_137:.*]] = mhlo.select %[[VAL_79]], %[[VAL_136]], %[[VAL_135]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_138:.*]] = mhlo.complex %[[VAL_78]], %[[VAL_137]] : tensor<complex<f32>>
// CHECK: %[[VAL_139:.*]] = mhlo.imag %[[VAL_138]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_140:.*]] = mhlo.real %[[VAL_138]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_141:.*]] = mhlo.negate %[[VAL_140]] : tensor<f32>
// CHECK: %[[VAL_142:.*]] = mhlo.complex %[[VAL_139]], %[[VAL_141]] : tensor<complex<f32>>
// CHECK: return %[[VAL_142]] : tensor<complex<f32>>
// CHECK: }
%result = "chlo.asinh"(%arg) : (tensor<complex<f32>>) -> tensor<complex<f32>>
func.return %result : tensor<complex<f32>>
}
@ -579,47 +696,176 @@ func.func @erf_bf16(%arg : tensor<bf16>) -> tensor<bf16> {
// -----
// CHECK-LABEL: @acosh
// CHECK-SAME: %[[ARG:.*]]: tensor<f16>
// CHECK-SAME: %[[VAL_0:.*]]: tensor<f16>) -> tensor<f16> {
func.func @acosh(%arg: tensor<f16>) -> tensor<f16> {
// CHECK: %[[MINUSONE:.*]] = mhlo.constant dense<-1.000000e+00>
// CHECK: %[[CMP:.*]] = mhlo.compare LT, %[[ARG]], %[[MINUSONE]]
// CHECK-DAG: %[[NAN:.*]] = mhlo.constant dense<0x7E00>
// CHECK-DAG: %[[MAX:.*]] = mhlo.constant dense<6.550400e+04>
// CHECK: %[[SQRTMAX:.*]] = mhlo.sqrt %[[MAX]]
// CHECK: %[[OVERFLOW:.*]] = mhlo.compare GE, %[[ARG]], %[[SQRTMAX]]
// CHECK: %[[LOGARG:.*]] = mhlo.log %[[ARG]]
// CHECK: %[[TWO:.*]] = mhlo.constant dense<2.000000e+00>
// CHECK: %[[LOGTWO:.*]] = mhlo.log %[[TWO]]
// CHECK: %[[OFLRES:.*]] = mhlo.add %[[LOGARG]], %[[LOGTWO]]
// CHECK: %[[ONE:.*]] = mhlo.constant dense<1.000000e+00>
// CHECK: %[[ARGPONE:.*]] = mhlo.add %[[ONE]], %[[ARG]]
// CHECK: %[[MINUSONE2:.*]] = mhlo.constant dense<-1.000000e+00>
// CHECK: %[[ARGMONE:.*]] = mhlo.add %[[MINUSONE2]], %[[ARG]]
// CHECK: %[[MUL:.*]] = mhlo.multiply %[[ARGPONE]], %[[ARGMONE]]
// CHECK: %[[SQRT:.*]] = mhlo.sqrt %[[MUL]]
// CHECK: %[[APSQRT:.*]] = mhlo.add %[[ARG]], %[[SQRT]]
// CHECK: %[[LOGAPMUL:.*]] = mhlo.log %[[APSQRT]]
// CHECK: %[[SEL1:.*]] = mhlo.select %[[OVERFLOW]], %[[OFLRES]], %[[LOGAPMUL]]
// CHECK: %[[RESULT:.*]] = mhlo.select %[[CMP]], %[[NAN]], %[[SEL1]]
// CHECK: return %[[RESULT]]
// CHECK: %[[VAL_1:.*]] = mhlo.constant dense<6.550400e+04> : tensor<f16>
// CHECK: %[[VAL_2:.*]] = mhlo.constant dense<2.000000e+00> : tensor<f16>
// CHECK: %[[VAL_3:.*]] = mhlo.divide %[[VAL_1]], %[[VAL_2]] : tensor<f16>
// CHECK: %[[VAL_4:.*]] = mhlo.compare GE, %[[VAL_0]], %[[VAL_3]] : (tensor<f16>, tensor<f16>) -> tensor<i1>
// CHECK: %[[VAL_5:.*]] = mhlo.log %[[VAL_2]] : tensor<f16>
// CHECK: %[[VAL_6:.*]] = mhlo.log %[[VAL_0]] : tensor<f16>
// CHECK: %[[VAL_7:.*]] = mhlo.add %[[VAL_5]], %[[VAL_6]] : tensor<f16>
// CHECK: %[[VAL_8:.*]] = mhlo.constant dense<1.000000e+00> : tensor<f16>
// CHECK: %[[VAL_9:.*]] = mhlo.subtract %[[VAL_0]], %[[VAL_8]] : tensor<f16>
// CHECK: %[[VAL_10:.*]] = mhlo.sqrt %[[VAL_9]] : tensor<f16>
// CHECK: %[[VAL_11:.*]] = mhlo.add %[[VAL_0]], %[[VAL_8]] : tensor<f16>
// CHECK: %[[VAL_12:.*]] = mhlo.sqrt %[[VAL_11]] : tensor<f16>
// CHECK: %[[VAL_13:.*]] = mhlo.add %[[VAL_12]], %[[VAL_10]] : tensor<f16>
// CHECK: %[[VAL_14:.*]] = mhlo.multiply %[[VAL_10]], %[[VAL_13]] : tensor<f16>
// CHECK: %[[VAL_15:.*]] = mhlo.log_plus_one %[[VAL_14]] : tensor<f16>
// CHECK: %[[VAL_16:.*]] = mhlo.select %[[VAL_4]], %[[VAL_7]], %[[VAL_15]] : tensor<i1>, tensor<f16>
// CHECK: return %[[VAL_16]] : tensor<f16>
// CHECK: }
%1 = "chlo.acosh"(%arg) : (tensor<f16>) -> tensor<f16>
func.return %1 : tensor<f16>
}
// -----
// CHECK-LABEL: @acosh_complex_f32
// CHECK-SAME: %[[ARG:.*]]: tensor<complex<f32>>
// CHECK-LABEL: @acosh_complex_f32(
// CHECK-SAME: %[[VAL_0:.*]]: tensor<complex<f32>>) -> tensor<complex<f32>> {
func.func @acosh_complex_f32(%arg : tensor<complex<f32>>) -> tensor<complex<f32>> {
// CHECK-NEXT: %[[TMP_0:.*]] = mhlo.constant dense<(1.000000e+00,0.000000e+00)>
// CHECK-NEXT: %[[TMP_1:.*]] = mhlo.add %[[ARG]], %[[TMP_0]]
// CHECK-NEXT: %[[TMP_2:.*]] = mhlo.constant dense<(1.000000e+00,0.000000e+00)>
// CHECK-NEXT: %[[TMP_3:.*]] = mhlo.subtract %[[ARG]], %[[TMP_2]]
// CHECK-NEXT: %[[TMP_4:.*]] = mhlo.multiply %[[TMP_1]], %[[TMP_3]]
// CHECK-NEXT: %[[TMP_5:.*]] = mhlo.sqrt %[[TMP_4]]
// CHECK-NEXT: %[[TMP_6:.*]] = mhlo.add %[[ARG]], %[[TMP_5]]
// CHECK-NEXT: %[[TMP_7:.*]] = mhlo.log %[[TMP_6]]
// CHECK-NEXT: return %[[TMP_7]]
// CHECK: %[[VAL_1:.*]] = mhlo.imag %[[VAL_0]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_2:.*]] = mhlo.constant dense<0.000000e+00> : tensor<f32>
// CHECK: %[[VAL_3:.*]] = mhlo.compare LT, %[[VAL_1]], %[[VAL_2]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_4:.*]] = mhlo.real %[[VAL_0]] : (tensor<complex<f32>>) -> tensor<f32>
// CHECK: %[[VAL_5:.*]] = mhlo.constant dense<0.000000e+00> : tensor<f32>
// CHECK: %[[VAL_6:.*]] = mhlo.compare LT, %[[VAL_1]], %[[VAL_5]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_7:.*]] = mhlo.abs %[[VAL_1]] : tensor<f32>
// CHECK: %[[VAL_8:.*]] = mhlo.abs %[[VAL_4]] : tensor<f32>
// CHECK: %[[VAL_9:.*]] = mhlo.constant dense<3.40282347E+38> : tensor<f32>
// CHECK: %[[VAL_10:.*]] = mhlo.sqrt %[[VAL_9]] : tensor<f32>
// CHECK: %[[VAL_11:.*]] = mhlo.constant dense<8.000000e+00> : tensor<f32>
// CHECK: %[[VAL_12:.*]] = mhlo.divide %[[VAL_10]], %[[VAL_11]] : tensor<f32>
// CHECK: %[[VAL_13:.*]] = mhlo.constant dense<9.99999995E+11> : tensor<f32>
// CHECK: %[[VAL_14:.*]] = mhlo.multiply %[[VAL_12]], %[[VAL_13]] : tensor<f32>
// CHECK: %[[VAL_15:.*]] = mhlo.compare LT, %[[VAL_8]], %[[VAL_14]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_16:.*]] = mhlo.constant dense<9.99999997E-7> : tensor<f32>
// CHECK: %[[VAL_17:.*]] = mhlo.multiply %[[VAL_12]], %[[VAL_16]] : tensor<f32>
// CHECK: %[[VAL_18:.*]] = mhlo.constant dense<1.000000e+02> : tensor<f32>
// CHECK: %[[VAL_19:.*]] = mhlo.multiply %[[VAL_12]], %[[VAL_18]] : tensor<f32>
// CHECK: %[[VAL_20:.*]] = mhlo.select %[[VAL_15]], %[[VAL_17]], %[[VAL_19]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_21:.*]] = mhlo.compare GE, %[[VAL_7]], %[[VAL_20]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_22:.*]] = mhlo.select %[[VAL_21]], %[[VAL_7]], %[[VAL_8]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_23:.*]] = mhlo.select %[[VAL_21]], %[[VAL_20]], %[[VAL_12]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_24:.*]] = mhlo.compare GE, %[[VAL_22]], %[[VAL_23]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_25:.*]] = mhlo.constant dense<2.000000e+00> : tensor<f32>
// CHECK: %[[VAL_26:.*]] = mhlo.log %[[VAL_25]] : tensor<f32>
// CHECK: %[[VAL_27:.*]] = mhlo.log %[[VAL_22]] : tensor<f32>
// CHECK: %[[VAL_28:.*]] = mhlo.add %[[VAL_26]], %[[VAL_27]] : tensor<f32>
// CHECK: %[[VAL_29:.*]] = mhlo.constant dense<5.000000e-01> : tensor<f32>
// CHECK: %[[VAL_30:.*]] = mhlo.constant dense<0x7F800000> : tensor<f32>
// CHECK: %[[VAL_31:.*]] = mhlo.compare EQ, %[[VAL_7]], %[[VAL_30]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_32:.*]] = mhlo.not %[[VAL_31]] : tensor<i1>
// CHECK: %[[VAL_33:.*]] = mhlo.and %[[VAL_21]], %[[VAL_32]] : tensor<i1>
// CHECK: %[[VAL_34:.*]] = mhlo.divide %[[VAL_8]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_35:.*]] = mhlo.select %[[VAL_33]], %[[VAL_34]], %[[VAL_5]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_36:.*]] = mhlo.multiply %[[VAL_35]], %[[VAL_35]] : tensor<f32>
// CHECK: %[[VAL_37:.*]] = mhlo.log_plus_one %[[VAL_36]] : tensor<f32>
// CHECK: %[[VAL_38:.*]] = mhlo.multiply %[[VAL_29]], %[[VAL_37]] : tensor<f32>
// CHECK: %[[VAL_39:.*]] = mhlo.add %[[VAL_28]], %[[VAL_38]] : tensor<f32>
// CHECK: %[[VAL_40:.*]] = mhlo.constant dense<1.17549435E-38> : tensor<f32>
// CHECK: %[[VAL_41:.*]] = mhlo.sqrt %[[VAL_40]] : tensor<f32>
// CHECK: %[[VAL_42:.*]] = mhlo.constant dense<4.000000e+00> : tensor<f32>
// CHECK: %[[VAL_43:.*]] = mhlo.multiply %[[VAL_41]], %[[VAL_42]] : tensor<f32>
// CHECK: %[[VAL_44:.*]] = mhlo.compare LT, %[[VAL_7]], %[[VAL_43]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_45:.*]] = mhlo.constant dense<1.000000e+00> : tensor<f32>
// CHECK: %[[VAL_46:.*]] = mhlo.compare LT, %[[VAL_8]], %[[VAL_45]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_47:.*]] = mhlo.and %[[VAL_44]], %[[VAL_46]] : tensor<i1>
// CHECK: %[[VAL_48:.*]] = mhlo.add %[[VAL_8]], %[[VAL_45]] : tensor<f32>
// CHECK: %[[VAL_49:.*]] = mhlo.subtract %[[VAL_8]], %[[VAL_45]] : tensor<f32>
// CHECK: %[[VAL_50:.*]] = mhlo.multiply %[[VAL_48]], %[[VAL_49]] : tensor<f32>
// CHECK: %[[VAL_51:.*]] = mhlo.abs %[[VAL_48]] : tensor<f32>
// CHECK: %[[VAL_52:.*]] = mhlo.maximum %[[VAL_51]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_53:.*]] = mhlo.minimum %[[VAL_51]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_54:.*]] = mhlo.compare EQ, %[[VAL_52]], %[[VAL_53]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_55:.*]] = mhlo.sqrt %[[VAL_25]] : tensor<f32>
// CHECK: %[[VAL_56:.*]] = mhlo.multiply %[[VAL_55]], %[[VAL_52]] : tensor<f32>
// CHECK: %[[VAL_57:.*]] = mhlo.divide %[[VAL_53]], %[[VAL_52]] : tensor<f32>
// CHECK: %[[VAL_58:.*]] = mhlo.multiply %[[VAL_57]], %[[VAL_57]] : tensor<f32>
// CHECK: %[[VAL_59:.*]] = mhlo.add %[[VAL_45]], %[[VAL_58]] : tensor<f32>
// CHECK: %[[VAL_60:.*]] = mhlo.sqrt %[[VAL_59]] : tensor<f32>
// CHECK: %[[VAL_61:.*]] = mhlo.compare EQ, %[[VAL_60]], %[[VAL_45]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_62:.*]] = mhlo.compare GT, %[[VAL_58]], %[[VAL_5]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_63:.*]] = mhlo.and %[[VAL_61]], %[[VAL_62]] : tensor<i1>
// CHECK: %[[VAL_64:.*]] = mhlo.multiply %[[VAL_52]], %[[VAL_58]] : tensor<f32>
// CHECK: %[[VAL_65:.*]] = mhlo.divide %[[VAL_64]], %[[VAL_25]] : tensor<f32>
// CHECK: %[[VAL_66:.*]] = mhlo.add %[[VAL_52]], %[[VAL_65]] : tensor<f32>
// CHECK: %[[VAL_67:.*]] = mhlo.multiply %[[VAL_52]], %[[VAL_60]] : tensor<f32>
// CHECK: %[[VAL_68:.*]] = mhlo.select %[[VAL_63]], %[[VAL_66]], %[[VAL_67]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_69:.*]] = mhlo.select %[[VAL_54]], %[[VAL_56]], %[[VAL_68]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_70:.*]] = mhlo.abs %[[VAL_49]] : tensor<f32>
// CHECK: %[[VAL_71:.*]] = mhlo.maximum %[[VAL_70]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_72:.*]] = mhlo.minimum %[[VAL_70]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_73:.*]] = mhlo.compare EQ, %[[VAL_71]], %[[VAL_72]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_74:.*]] = mhlo.multiply %[[VAL_55]], %[[VAL_71]] : tensor<f32>
// CHECK: %[[VAL_75:.*]] = mhlo.divide %[[VAL_72]], %[[VAL_71]] : tensor<f32>
// CHECK: %[[VAL_76:.*]] = mhlo.multiply %[[VAL_75]], %[[VAL_75]] : tensor<f32>
// CHECK: %[[VAL_77:.*]] = mhlo.add %[[VAL_45]], %[[VAL_76]] : tensor<f32>
// CHECK: %[[VAL_78:.*]] = mhlo.sqrt %[[VAL_77]] : tensor<f32>
// CHECK: %[[VAL_79:.*]] = mhlo.compare EQ, %[[VAL_78]], %[[VAL_45]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_80:.*]] = mhlo.compare GT, %[[VAL_76]], %[[VAL_5]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_81:.*]] = mhlo.and %[[VAL_79]], %[[VAL_80]] : tensor<i1>
// CHECK: %[[VAL_82:.*]] = mhlo.multiply %[[VAL_71]], %[[VAL_76]] : tensor<f32>
// CHECK: %[[VAL_83:.*]] = mhlo.divide %[[VAL_82]], %[[VAL_25]] : tensor<f32>
// CHECK: %[[VAL_84:.*]] = mhlo.add %[[VAL_71]], %[[VAL_83]] : tensor<f32>
// CHECK: %[[VAL_85:.*]] = mhlo.multiply %[[VAL_71]], %[[VAL_78]] : tensor<f32>
// CHECK: %[[VAL_86:.*]] = mhlo.select %[[VAL_81]], %[[VAL_84]], %[[VAL_85]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_87:.*]] = mhlo.select %[[VAL_73]], %[[VAL_74]], %[[VAL_86]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_88:.*]] = mhlo.add %[[VAL_69]], %[[VAL_87]] : tensor<f32>
// CHECK: %[[VAL_89:.*]] = mhlo.multiply %[[VAL_29]], %[[VAL_88]] : tensor<f32>
// CHECK: %[[VAL_90:.*]] = mhlo.add %[[VAL_89]], %[[VAL_45]] : tensor<f32>
// CHECK: %[[VAL_91:.*]] = mhlo.divide %[[VAL_50]], %[[VAL_90]] : tensor<f32>
// CHECK: %[[VAL_92:.*]] = mhlo.negate %[[VAL_91]] : tensor<f32>
// CHECK: %[[VAL_93:.*]] = mhlo.compare GE, %[[VAL_8]], %[[VAL_45]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_94:.*]] = mhlo.multiply %[[VAL_7]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_95:.*]] = mhlo.multiply %[[VAL_29]], %[[VAL_94]] : tensor<f32>
// CHECK: %[[VAL_96:.*]] = mhlo.add %[[VAL_69]], %[[VAL_48]] : tensor<f32>
// CHECK: %[[VAL_97:.*]] = mhlo.divide %[[VAL_95]], %[[VAL_96]] : tensor<f32>
// CHECK: %[[VAL_98:.*]] = mhlo.add %[[VAL_87]], %[[VAL_49]] : tensor<f32>
// CHECK: %[[VAL_99:.*]] = mhlo.multiply %[[VAL_29]], %[[VAL_98]] : tensor<f32>
// CHECK: %[[VAL_100:.*]] = mhlo.add %[[VAL_97]], %[[VAL_99]] : tensor<f32>
// CHECK: %[[VAL_101:.*]] = mhlo.constant dense<1.500000e+00> : tensor<f32>
// CHECK: %[[VAL_102:.*]] = mhlo.compare LE, %[[VAL_89]], %[[VAL_101]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_103:.*]] = mhlo.subtract %[[VAL_87]], %[[VAL_49]] : tensor<f32>
// CHECK: %[[VAL_104:.*]] = mhlo.divide %[[VAL_95]], %[[VAL_103]] : tensor<f32>
// CHECK: %[[VAL_105:.*]] = mhlo.add %[[VAL_97]], %[[VAL_104]] : tensor<f32>
// CHECK: %[[VAL_106:.*]] = mhlo.subtract %[[VAL_89]], %[[VAL_45]] : tensor<f32>
// CHECK: %[[VAL_107:.*]] = mhlo.select %[[VAL_102]], %[[VAL_105]], %[[VAL_106]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_108:.*]] = mhlo.select %[[VAL_93]], %[[VAL_100]], %[[VAL_107]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_109:.*]] = mhlo.select %[[VAL_47]], %[[VAL_92]], %[[VAL_108]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_110:.*]] = mhlo.multiply %[[VAL_109]], %[[VAL_90]] : tensor<f32>
// CHECK: %[[VAL_111:.*]] = mhlo.sqrt %[[VAL_110]] : tensor<f32>
// CHECK: %[[VAL_112:.*]] = mhlo.divide %[[VAL_7]], %[[VAL_111]] : tensor<f32>
// CHECK: %[[VAL_113:.*]] = mhlo.add %[[VAL_109]], %[[VAL_111]] : tensor<f32>
// CHECK: %[[VAL_114:.*]] = mhlo.log_plus_one %[[VAL_113]] : tensor<f32>
// CHECK: %[[VAL_115:.*]] = mhlo.select %[[VAL_47]], %[[VAL_112]], %[[VAL_114]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_116:.*]] = mhlo.select %[[VAL_24]], %[[VAL_39]], %[[VAL_115]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_117:.*]] = mhlo.negate %[[VAL_116]] : tensor<f32>
// CHECK: %[[VAL_118:.*]] = mhlo.select %[[VAL_6]], %[[VAL_116]], %[[VAL_117]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_119:.*]] = mhlo.negate %[[VAL_118]] : tensor<f32>
// CHECK: %[[VAL_120:.*]] = mhlo.maximum %[[VAL_8]], %[[VAL_7]] : tensor<f32>
// CHECK: %[[VAL_121:.*]] = mhlo.compare GE, %[[VAL_120]], %[[VAL_12]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_122:.*]] = mhlo.compare LE, %[[VAL_8]], %[[VAL_45]] : (tensor<f32>, tensor<f32>) -> tensor<i1>
// CHECK: %[[VAL_123:.*]] = mhlo.add %[[VAL_89]], %[[VAL_8]] : tensor<f32>
// CHECK: %[[VAL_124:.*]] = mhlo.multiply %[[VAL_29]], %[[VAL_123]] : tensor<f32>
// CHECK: %[[VAL_125:.*]] = mhlo.divide %[[VAL_94]], %[[VAL_96]] : tensor<f32>
// CHECK: %[[VAL_126:.*]] = mhlo.add %[[VAL_125]], %[[VAL_103]] : tensor<f32>
// CHECK: %[[VAL_127:.*]] = mhlo.multiply %[[VAL_124]], %[[VAL_126]] : tensor<f32>
// CHECK: %[[VAL_128:.*]] = mhlo.sqrt %[[VAL_127]] : tensor<f32>
// CHECK: %[[VAL_129:.*]] = mhlo.divide %[[VAL_124]], %[[VAL_96]] : tensor<f32>
// CHECK: %[[VAL_130:.*]] = mhlo.divide %[[VAL_124]], %[[VAL_98]] : tensor<f32>
// CHECK: %[[VAL_131:.*]] = mhlo.add %[[VAL_129]], %[[VAL_130]] : tensor<f32>
// CHECK: %[[VAL_132:.*]] = mhlo.sqrt %[[VAL_131]] : tensor<f32>
// CHECK: %[[VAL_133:.*]] = mhlo.multiply %[[VAL_7]], %[[VAL_132]] : tensor<f32>
// CHECK: %[[VAL_134:.*]] = mhlo.select %[[VAL_122]], %[[VAL_128]], %[[VAL_133]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_135:.*]] = mhlo.select %[[VAL_121]], %[[VAL_7]], %[[VAL_134]] : tensor<i1>, tensor<f32>
// CHECK: %[[VAL_136:.*]] = mhlo.atan2 %[[VAL_135]], %[[VAL_4]] : tensor<f32>
// CHECK: %[[VAL_137:.*]] = mhlo.complex %[[VAL_119]], %[[VAL_136]] : tensor<complex<f32>>
// CHECK: %[[VAL_138:.*]] = mhlo.negate %[[VAL_137]] : tensor<complex<f32>>
// CHECK: %[[VAL_139:.*]] = mhlo.select %[[VAL_3]], %[[VAL_138]], %[[VAL_137]] : tensor<i1>, tensor<complex<f32>>
// CHECK: return %[[VAL_139]] : tensor<complex<f32>>
// CHECK: }
%result = "chlo.acosh"(%arg) : (tensor<complex<f32>>) -> tensor<complex<f32>>
func.return %result : tensor<complex<f32>>
}

6
third_party/xla/xla/mlir_hlo/tests/Dialect/mhlo/hlo-legalize-to-stablehlo-experimental.mlir vendored
View File

@ -7,7 +7,7 @@
// CHECK-LABEL: "op_all_reduce_tuple"
func.func @op_all_reduce_tuple(%arg0: tensor<8xf32>, %arg1: tensor<f32>) -> (tensor<8xf32>, tensor<f32>) {
// CHECK: "stablehlo.custom_call"(%[[ARG0:.*]], %[[ARG1:.*]]) <{
// CHECK-SAME: backend_config = "", call_target_name = "mhlo.all_reduce", called_computations = [@all_reduce]
// CHECK-SAME: call_target_name = "mhlo.all_reduce", called_computations = [@all_reduce]
// CHECK-SAME: }> {
// CHECK-SAME{LITERAL}: mhlo.attributes = {replica_groups = dense<> : tensor<0x0xi64>}
// CHECK-SAME: } : (tensor<8xf32>, tensor<f32>) -> (tensor<8xf32>, tensor<f32>)
@ -29,7 +29,7 @@ func.func @op_all_reduce_tuple(%arg0: tensor<8xf32>, %arg1: tensor<f32>) -> (ten
// CHECK-LABEL: "op_all_to_all_tuple"
func.func @op_all_to_all_tuple(%arg0: tensor<128x4xf32>, %arg1: tensor<128x4xf32>) -> (tensor<128x4xf32>, tensor<128x4xf32>) {
// CHECK: "stablehlo.custom_call"(%arg0, %arg1) <{backend_config = "", call_target_name = "mhlo.all_to_all"}> {
// CHECK: "stablehlo.custom_call"(%arg0, %arg1) <{call_target_name = "mhlo.all_to_all"}> {
// CHECK-SAME{LITERAL}: mhlo.attributes = {replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>}
// CHECK-SAME: } : (tensor<128x4xf32>, tensor<128x4xf32>)
// expected-error@+1 {{failed to legalize operation 'mhlo.all_to_all' that was explicitly marked illegal}}
@ -43,7 +43,7 @@ func.func @op_all_to_all_tuple(%arg0: tensor<128x4xf32>, %arg1: tensor<128x4xf32
// CHECK-LABEL: "attr_precision_packed_nibble"
func.func @attr_precision_packed_nibble(%arg0: tensor<8x16xf32>, %arg1: tensor<16x8xf32>) -> tensor<8x8xf32> {
// CHECK: "stablehlo.custom_call"(%arg0, %arg1) <{backend_config = "", call_target_name = "mhlo.dot"}> {
// CHECK: "stablehlo.custom_call"(%arg0, %arg1) <{call_target_name = "mhlo.dot"}> {
// CHECK-SAME: mhlo.attributes = {precision_config = ["PACKED_NIBBLE"]}
// CHECK-SAME: } : (tensor<8x16xf32>, tensor<16x8xf32>) -> tensor<8x8xf32>
// expected-error@+1 {{failed to legalize operation 'mhlo.dot' that was explicitly marked illegal}}

33
third_party/xla/xla/mlir_hlo/tests/Dialect/mhlo/hlo-legalize-to-stablehlo.mlir vendored
View File

@ -156,10 +156,9 @@ func.func @attr_custom_call_api_version_status_returning_unified(%arg0: tensor<f
// CHECK-LABEL: "attr_custom_call_api_version_typed_ffi"
func.func @attr_custom_call_api_version_typed_ffi(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{backend_config = "", call_target_name = "mhlo.custom_call"}> {
// CHECK-SAME: mhlo.attributes = {api_version = 4 : i32, backend_config = {foo = "bar"}, call_target_name = "foo"},
// CHECK-SAME: mhlo.version = 1 : i64
// CHECK-SAME: } : (tensor<f32>) -> tensor<f32>
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{api_version = 4 : i32, backend_config = {foo = "bar"},
// CHECK-SAME: call_target_name = "foo"}>
// CHECK-SAME: (tensor<f32>) -> tensor<f32>
%0 = "mhlo.custom_call"(%arg0) {
call_target_name = "foo",
backend_config = {foo = "bar"},
@ -746,7 +745,9 @@ func.func @op_custom_call_api_version_original(%arg0: tensor<f32>) -> tensor<f32
// CHECK-LABEL: "op_custom_call_custom_call_schedule_none"
func.func @op_custom_call_custom_call_schedule_none(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{backend_config = "", call_target_name = "foo"}> : (tensor<f32>) -> tensor<f32>
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{
// CHECK-SAME: call_target_name = "foo"}>
// CHECK-SAME: (tensor<f32>) -> tensor<f32>
%0 = "mhlo.custom_call"(%arg0) {
call_target_name = "foo",
custom_call_schedule = #mhlo<custom_call_schedule NONE>
@ -756,10 +757,10 @@ func.func @op_custom_call_custom_call_schedule_none(%arg0: tensor<f32>) -> tenso
// CHECK-LABEL: "op_custom_call_custom_call_schedule_none_ffi"
func.func @op_custom_call_custom_call_schedule_none_ffi(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{backend_config = "", call_target_name = "mhlo.custom_call"}> {
// CHECK-SAME: mhlo.attributes = {api_version = 4 : i32, backend_config = {foo = "bar"}, call_target_name = "foo"},
// CHECK-SAME: mhlo.version = 1 : i64
// CHECK-SAME: } : (tensor<f32>) -> tensor<f32>
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{
// CHECK-SAME: api_version = 4 : i32, backend_config = {foo = "bar"},
// CHECK-SAME: call_target_name = "foo"}>
// CHECK-SAME: (tensor<f32>) -> tensor<f32>
%0 = "mhlo.custom_call"(%arg0) {
call_target_name = "foo",
backend_config = {foo = "bar"},
@ -1582,10 +1583,9 @@ func.func @op_subtract(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
// CHECK-LABEL: "op_tan"
func.func @op_tan(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{backend_config = "", call_target_name = "mhlo.tan"}> {
// CHECK-SAME{LITERAL}: mhlo.attributes = {}
// CHECK-SAME{LITERAL}: mhlo.version = 1 : i64
// CHECK-SAME: } : (tensor<f32>) -> tensor<f32>
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{
// CHECK-SAME: call_target_name = "mhlo.tan"}> {mhlo.attributes = {}, mhlo.version = 1 : i64}
// CHECK-SAME: (tensor<f32>) -> tensor<f32>
%0 = "mhlo.tan"(%arg0) : (tensor<f32>) -> tensor<f32>
func.return %0 : tensor<f32>
}
@ -1599,10 +1599,9 @@ func.func @op_tanh(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK-LABEL: "op_topk"
func.func @op_topk(%arg0: tensor<5x10xf32>) -> (tensor<5x8xf32>, tensor<5x8xi32>) {
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{backend_config = "", call_target_name = "mhlo.topk"}> {
// CHECK-SAME{LITERAL}: mhlo.attributes = {k = 8 : i64, largest = true}
// CHECK-SAME{LITERAL}: mhlo.version = 1 : i64
// CHECK-SAME: } : (tensor<5x10xf32>) -> (tensor<5x8xf32>, tensor<5x8xi32>)
// CHECK: "stablehlo.custom_call"([[ARG0:%arg[0-9]+]]) <{
// CHECK-SAME: call_target_name = "mhlo.topk"}> {mhlo.attributes = {k = 8 : i64, largest = true}, mhlo.version = 1 : i64}
// CHECK-SAME: (tensor<5x10xf32>) -> (tensor<5x8xf32>, tensor<5x8xi32>)
%0:2 = mhlo.topk(%arg0, k=8, largest=true) : tensor<5x10xf32> -> (tensor<5x8xf32>, tensor<5x8xi32>)
func.return %0#0, %0#1 : tensor<5x8xf32>, tensor<5x8xi32>
}

2
third_party/xla/xla/mlir_hlo/tests/Dialect/mhlo/stablehlo-legalize-to-hlo.mlir vendored
View File

@ -1567,7 +1567,6 @@ func.func @op_subtract(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
func.func @op_tan_mhlo_v1(%arg0: tensor<f32>) -> tensor<f32> {
// CHECK: "mhlo.tan"([[ARG0:%arg[0-9]+]]) : (tensor<f32>) -> tensor<f32>
%0 = "stablehlo.custom_call"(%arg0) {
backend_config = "",
call_target_name = "mhlo.tan",
mhlo.attributes = {},
mhlo.version = 1 : i64
@ -1951,7 +1950,6 @@ func.func @op_custom_call_botched_mhlo_backend_config_version(%arg0: tensor<f32>
func.func @op_topk_mhlo_v1(%arg0: tensor<5x10xf32>) -> (tensor<5x8xf32>, tensor<5x8xi32>) {
// CHECK: "mhlo.topk"([[ARG0:%arg[0-9]+]]) <{k = 8 : i64, largest = true}> : (tensor<5x10xf32>) -> (tensor<5x8xf32>, tensor<5x8xi32>)
%0:2 = "stablehlo.custom_call"(%arg0) {
backend_config = "",
call_target_name = "mhlo.topk",
mhlo.attributes = {k = 8 : i64, largest = true},
mhlo.version = 1 : i64

2
third_party/xla/xla/python/refine_polymorphic_shapes.cc vendored
View File

@ -145,7 +145,7 @@ struct CheckShapeAssertionsPass
return op.emitError()
<< attr.getName() << " is not a supported attribute";
}
if (!op.getBackendConfig().empty())
if (!op.hasEmptyBackendConfig())
return op.emitError() << "expects an empty backend_config";
if (op.getCallTargetName() != shapeAssertionName)
return op.emitError() << "expects @shape_assertion";