diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl index de600e09b4d74..dc71501f7e2e7 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" - LLVM_COMMIT = "87d42c13cd6b119240781f31e5869981d500a186" - LLVM_SHA256 = "66b7f645a9a67d6692c5bea033a536972fcd2e80aa31408d579aae5811521869" + LLVM_COMMIT = "7ddee0b619f658cef905a69427ef9531fd1d229d" + LLVM_SHA256 = "0ebcba094207b7104182c5112e1c3b0f5623e7779b901c1d2007e4983c3582b1" tf_http_archive( name = name, diff --git a/third_party/shardy/temporary.patch b/third_party/shardy/temporary.patch index 44811b6e68285..03c176c31080b 100644 --- a/third_party/shardy/temporary.patch +++ b/third_party/shardy/temporary.patch @@ -1,227 +1,41 @@ -diff --git a/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc b/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc -index 7c46ca9d..2722b45d 100644 ---- a/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc -+++ b/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc -@@ -26,7 +26,6 @@ limitations under the License. - #include "mlir/IR/Builders.h" - #include "mlir/IR/BuiltinAttributes.h" - #include "mlir/IR/BuiltinOps.h" --#include "mlir/IR/MLIRContext.h" - #include "mlir/IR/SymbolTable.h" - #include "mlir/Support/LLVM.h" - #include "shardy/dialect/sdy/ir/constants.h" -@@ -73,35 +72,6 @@ MeshOp createNewMeshOp(Location loc, mlir::stablehlo::MeshAttr mesh, - return createMesh("mesh", sdyMeshAttr); - } - --DictionaryAttr getStablehloMeshAttrAsDict(MeshAttr sdyMeshAttr) { -- MLIRContext* ctx = sdyMeshAttr.getContext(); -- Builder builder(ctx); -- -- SmallVector axesAttrs; -- for (MeshAxisAttr axisAttr : sdyMeshAttr.getAxes()) { -- NamedAttribute nameAttr = -- builder.getNamedAttr("name", builder.getStringAttr(axisAttr.getName())); -- NamedAttribute sizeAttr = builder.getNamedAttr( -- "size", builder.getI64IntegerAttr(axisAttr.getSize())); -- axesAttrs.push_back(builder.getDictionaryAttr({nameAttr, sizeAttr})); -- } -- ArrayAttr axesArrayAttr = builder.getArrayAttr(axesAttrs); -- -- SmallVector dictFields; -- dictFields.push_back(builder.getNamedAttr("axes", axesArrayAttr)); -- -- if (!sdyMeshAttr.getDeviceIds().empty()) { -- auto type = RankedTensorType::get( -- {static_cast(sdyMeshAttr.getDeviceIds().size())}, -- builder.getI64Type()); -- auto deviceIds = -- DenseIntElementsAttr::get(type, sdyMeshAttr.getDeviceIds()); -- dictFields.push_back(builder.getNamedAttr("device_ids", deviceIds)); -- } -- -- return builder.getDictionaryAttr(dictFields); --} -- - TensorShardingAttr replaceMesh(TensorShardingAttr sharding, - StringAttr meshName) { - return TensorShardingAttr::get( -@@ -226,14 +196,6 @@ struct LiftInlinedMeshesPass - moduleOp.walk([&](stablehlo::CollectiveBroadcastOp op) { - processMeshInReplicaGroups(op); - }); -- -- // Attach discardable `stablehlo.mesh` attributes to all named meshes. -- for (auto meshOp : llvm::make_early_inc_range(moduleOp.getOps())) { -- if (!meshOp->hasAttr("stablehlo.mesh")) { -- meshOp->setAttr("stablehlo.mesh", -- getStablehloMeshAttrAsDict(meshOp.getMesh())); -- } -- } - } - }; - -diff --git a/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir b/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir -index f45829cd..de18acb6 100644 ---- a/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir -+++ b/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir -@@ -256,17 +256,3 @@ func.func private @foo(%arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mes - %0 = stablehlo.negate %arg0 {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"x"}, {}]>]>} : tensor<8x8xf32> - return %0 : tensor<8x8xf32> - } -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["a"=4]> {stablehlo.mesh = {axes = [{name = "a", size = 4 : i64}]}} --sdy.mesh @mesh = <["a"=4]> -- --// CHECK: sdy.mesh @mesh_0 = <["b"=2]> {stablehlo.mesh = {axes = [{name = "b", size = 2 : i64}]}} -- --// CHECK-LABEL: func @tagged_stablehlo_mesh_attribute --func.func @tagged_stablehlo_mesh_attribute(%arg0: tensor<4x4xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"a"}, {}]>}, %arg1: tensor<4x4xf32>) -> tensor<4x4xf32> { -- // CHECK-NEXT: stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[<@mesh_0, [{"b"}, {}]>]>} -- %0 = stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[, [{"b"}, {}]>]>} : tensor<4x4xf32> -- return %0 : tensor<4x4xf32> --} -diff --git a/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir b/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir -index f021e075..a80a25a3 100644 ---- a/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir -+++ b/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir -@@ -1,4 +1,4 @@ --// RUN: sdy_opt %s -split-input-file -sdy-propagation-pipeline | FileCheck %s -+// RUN: sdy_opt %s -split-input-file -sdy-propagation-pipeline 2>&1 | FileCheck %s - - sdy.mesh @mesh = <["a"=2, "b"=2, "c"=2]> - +diff --git a/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_reductions.mlir b/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_reductions.mlir +index 3412522c..65f3334f 100644 +--- a/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_reductions.mlir ++++ b/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_reductions.mlir +@@ -15,7 +15,7 @@ + func.func @assign_outputs_with_empty_src_set(%arg0: !m3_4x8, %arg1: !m2_4x8) + -> (tensor<4x8xf32>) attributes {topology=#topology} { + // CHECK: %[[T0:.*]] = mpmd.transfer %arg0 {{.*}}m3{{.*}} -> {{.*}}m2 +-// CHECK-NEXT: %[[ADD:.*]] = mpmd.fragment (%[[T0]], %arg1) ++// CHECK-NEXT: %[[ADD:.*]] = mpmd.fragment (%arg1, %[[T0]]) + // CHECK: return %[[ADD]] + %4 = mpmd.unassign %arg0 : (!m3_4x8) -> tensor<4x8xf32> + %5 = mpmd.unassign %arg1 : (!m2_4x8) -> tensor<4x8xf32> +diff --git a/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_transfers.mlir b/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_transfers.mlir +index 33953a32..50877eca 100644 +--- a/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_transfers.mlir ++++ b/shardy/dialect/mpmd/transforms/import/test/infer_mesh_pipeline_create_transfers.mlir +@@ -13,7 +13,7 @@ func.func @assign_outputs_with_empty_src_set(%arg0: !m3_8x16, %arg1: !m2_8x16) + <"m3": <["z"=2]>> + >} { + // CHECK: %[[T0:.*]] = mpmd.transfer %arg0 {{.*}}m3{{.*}} -> {{.*}}m2 +-// CHECK-NEXT: %[[ADD:.*]] = mpmd.fragment (%[[T0]], %arg1) ++// CHECK-NEXT: %[[ADD:.*]] = mpmd.fragment (%arg1, %[[T0]]) + // CHECK: return %[[ADD]] + %4 = mpmd.unassign %arg0 : (!m3_8x16) -> tensor<8x16xf32> + %5 = mpmd.unassign %arg1 : (!m2_8x16) -> tensor<8x16xf32> diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl -index bf16ff8f..de600e09 100644 +index de600e09..dc71501f 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" -- LLVM_COMMIT = "0e6d57648584c3895abd3021fb8fa2a767146d5a" -- LLVM_SHA256 = "bbd2acb1760bb6a3566a0ffa272f0828c4f2d8233be1759741023abb73910c44" -+ LLVM_COMMIT = "87d42c13cd6b119240781f31e5869981d500a186" -+ LLVM_SHA256 = "66b7f645a9a67d6692c5bea033a536972fcd2e80aa31408d579aae5811521869" +- LLVM_COMMIT = "87d42c13cd6b119240781f31e5869981d500a186" +- LLVM_SHA256 = "66b7f645a9a67d6692c5bea033a536972fcd2e80aa31408d579aae5811521869" ++ LLVM_COMMIT = "7ddee0b619f658cef905a69427ef9531fd1d229d" ++ LLVM_SHA256 = "0ebcba094207b7104182c5112e1c3b0f5623e7779b901c1d2007e4983c3582b1" tf_http_archive( name = name, -diff --git a/third_party/stablehlo/temporary.patch b/third_party/stablehlo/temporary.patch -index 6c4788de..98e44999 100644 ---- a/third_party/stablehlo/temporary.patch -+++ b/third_party/stablehlo/temporary.patch -@@ -565,112 +565,6 @@ diff --ruN a/stablehlo/stablehlo/reference/Tensor.cpp b/stablehlo/stablehlo/refe - return std::complex(value.real().convertToDouble(), - value.imag().convertToDouble()); - }); --diff --ruN a/stablehlo/stablehlo/tests/transforms/mesh_axes_replica_group_compatibility.mlir b/stablehlo/stablehlo/tests/transforms/mesh_axes_replica_group_compatibility.mlir ----- stablehlo/stablehlo/tests/transforms/mesh_axes_replica_group_compatibility.mlir --+++ stablehlo/stablehlo/tests/transforms/mesh_axes_replica_group_compatibility.mlir --@@ -1,4 +1,4 @@ ---// RUN: stablehlo-opt %s -stablehlo-compatibility-expander="target=1.14.0" | FileCheck %s --+// RUN: stablehlo-opt %s -stablehlo-compatibility-expander="target=1.14.0" -allow-unregistered-dialect | FileCheck %s -- -- module { -- --@@ -40,4 +40,48 @@ -- -- return %0 : tensor<4xf32> -- } --+ "sdy.mesh"() {sym_name = "sdy_mesh", stablehlo.mesh = {axes = [{name = "x", size = 2 : i64}, {name = "y", size = 2 : i64}]}} : () -> () --+ --+ // CHECK-LABEL: @all_reduce_sdy_mesh --+ func.func @all_reduce_sdy_mesh(%arg0: tensor<4xf32>) -> tensor<4xf32> { --+ // CHECK: replica_groups = dense<{{\[\[}}0, 2], [1, 3]]> : tensor<2x2xi64> --+ %0 = "stablehlo.all_reduce"(%arg0) ({ --+ ^bb0(%arg1: tensor, %arg2: tensor): --+ %1 = "stablehlo.add"(%arg1, %arg2) : (tensor, tensor) -> tensor --+ "stablehlo.return"(%1) : (tensor) -> () --+ }) { --+ replica_groups = #stablehlo.replica_group_mesh_axes]> --+ } : (tensor<4xf32>) -> tensor<4xf32> --+ return %0 : tensor<4xf32> --+ } --+ --+ "sdy.mesh"() {sym_name = "sdy_mesh_dev", stablehlo.mesh = {axes = [{name = "x", size = 2 : i64}, {name = "y", size = 2 : i64}], device_ids = dense<[0, 2, 1, 3]> : tensor<4xi64>}} : () -> () --+ --+ // CHECK-LABEL: @all_reduce_sdy_mesh_dev --+ func.func @all_reduce_sdy_mesh_dev(%arg0: tensor<4xf32>) -> tensor<4xf32> { --+ // CHECK: replica_groups = dense<{{\[\[}}0, 1], [2, 3]]> : tensor<2x2xi64> --+ %0 = "stablehlo.all_reduce"(%arg0) ({ --+ ^bb0(%arg1: tensor, %arg2: tensor): --+ %1 = "stablehlo.add"(%arg1, %arg2) : (tensor, tensor) -> tensor --+ "stablehlo.return"(%1) : (tensor) -> () --+ }) { --+ replica_groups = #stablehlo.replica_group_mesh_axes]> --+ } : (tensor<4xf32>) -> tensor<4xf32> --+ return %0 : tensor<4xf32> --+ } --+ --+ "sdy.mesh"() {sym_name = "sdy_mesh_max", stablehlo.mesh = {axes = [], device_ids = dense<0> : tensor<1xi64>}} : () -> () --+ --+ // CHECK-LABEL: @all_reduce_sdy_mesh_max --+ func.func @all_reduce_sdy_mesh_max(%arg0: tensor<4xf32>) -> tensor<4xf32> { --+ // CHECK: replica_groups = dense<0> : tensor<1x1xi64> --+ %0 = "stablehlo.all_reduce"(%arg0) ({ --+ ^bb0(%arg1: tensor, %arg2: tensor): --+ %1 = "stablehlo.add"(%arg1, %arg2) : (tensor, tensor) -> tensor --+ "stablehlo.return"(%1) : (tensor) -> () --+ }) { --+ replica_groups = #stablehlo.replica_group_mesh_axes --+ } : (tensor<4xf32>) -> tensor<4xf32> --+ return %0 : tensor<4xf32> --+ } -- } --diff --ruN a/stablehlo/stablehlo/transforms/StablehloCompatibilityExpander.cpp b/stablehlo/stablehlo/transforms/StablehloCompatibilityExpander.cpp ----- stablehlo/stablehlo/transforms/StablehloCompatibilityExpander.cpp --+++ stablehlo/stablehlo/transforms/StablehloCompatibilityExpander.cpp --@@ -311,6 +311,31 @@ -- } -- }; -- --+stablehlo::MeshAttr resolveMeshAttrFromDict(DictionaryAttr dictAttr) { --+ MLIRContext* ctx = dictAttr.getContext(); --+ --+ // Extract axes --+ auto axesAttr = dictAttr.getAs("axes"); --+ if (!axesAttr) return {}; --+ --+ SmallVector axes; --+ for (Attribute attr : axesAttr) { --+ auto axisDict = dyn_cast(attr); --+ if (!axisDict) return {}; --+ auto nameAttr = axisDict.getAs("name"); --+ auto sizeAttr = axisDict.getAs("size"); --+ if (!nameAttr || !sizeAttr) return {}; --+ axes.push_back(stablehlo::MeshAxisAttr::get(ctx, nameAttr.getValue(), --+ sizeAttr.getInt())); --+ } --+ --+ // Extract device_ids --+ DenseIntElementsAttr deviceIds = --+ dictAttr.getAs("device_ids"); --+ --+ return stablehlo::MeshAttr::get(ctx, axes, deviceIds); --+} --+ -- // ReplicaGroupMeshAxesExpander handles the conversion of -- // #stablehlo.replica_group_mesh_axes to a dense I64ElementsAttr for target -- // versions that do not support it. --@@ -341,6 +366,12 @@ -- symbolOp->getAttrOfType("stablehlo.mesh"); -- } -- if (!meshAttr) { --+ if (auto dictAttr = --+ symbolOp->getAttrOfType("stablehlo.mesh")) { --+ meshAttr = resolveMeshAttrFromDict(dictAttr); --+ } --+ } --+ if (!meshAttr) { -- return op->emitError("Failed to find mesh attribute on symbol"); -- } -- } - diff --ruN a/stablehlo/stablehlo/transforms/StablehloLegalizeDeprecatedOps.cpp b/stablehlo/stablehlo/transforms/StablehloLegalizeDeprecatedOps.cpp - --- stablehlo/stablehlo/transforms/StablehloLegalizeDeprecatedOps.cpp - +++ stablehlo/stablehlo/transforms/StablehloLegalizeDeprecatedOps.cpp diff --git a/third_party/shardy/workspace.bzl b/third_party/shardy/workspace.bzl index bd7e8c235d674..0e1409b71b19a 100644 --- a/third_party/shardy/workspace.bzl +++ b/third_party/shardy/workspace.bzl @@ -3,8 +3,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls") def repo(): - SHARDY_COMMIT = "7213adc9d3d62cddb37db0368d3530f7eb8275a6" - SHARDY_SHA256 = "337ee4c35cd288c7dbf2717aaaacb5c57e6d4ebd15e2633b41b4c5f18dd2bc1c" + SHARDY_COMMIT = "e124ab482131bf11f805d73b566e09b4a8c3dda3" + SHARDY_SHA256 = "f1542df0bf8d4647612854ce832c6bb11ba0b879fbe38edd8ff18bfd86339abe" tf_http_archive( name = "shardy", diff --git a/third_party/triton/common/llvm_cl913905807.patch b/third_party/triton/common/llvm_cl913905807.patch new file mode 100644 index 0000000000000..41a0fd00e6ac7 --- /dev/null +++ b/third_party/triton/common/llvm_cl913905807.patch @@ -0,0 +1,12 @@ +--- a/third_party/amd/lib/Analysis/RangeAnalysis.cpp ++++ b/third_party/amd/lib/Analysis/RangeAnalysis.cpp +@@ -772,7 +772,7 @@ + } else { + // Else, propagate pred operands. + auto operLat = *getLatticeElementFor(point, oper); +- changed = argLat->join(operLat); ++ changed = argLat->join(operLat.getValue()); + LLVM_DEBUG({ + if (changed == ChangeResult::Change) { + DBGS() << "Operand lattice "; + diff --git a/third_party/triton/common/series.bzl b/third_party/triton/common/series.bzl index b991cc7c7acae..a2b3e5d1ac814 100644 --- a/third_party/triton/common/series.bzl +++ b/third_party/triton/common/series.bzl @@ -35,5 +35,6 @@ common_patch_list = [ "//third_party/triton:common/llvm_cl904544082.patch", "//third_party/triton:common/check_null_encoding.patch", "//third_party/triton:common/llvm_cl910324404.patch", + "//third_party/triton:common/llvm_cl913905807.patch", # Add new patches just above this line ] diff --git a/xla/codegen/intrinsic/cpp/eigen_unary_test.cc b/xla/codegen/intrinsic/cpp/eigen_unary_test.cc index 9b106967d6604..46889af809f33 100644 --- a/xla/codegen/intrinsic/cpp/eigen_unary_test.cc +++ b/xla/codegen/intrinsic/cpp/eigen_unary_test.cc @@ -109,7 +109,7 @@ TEST(EigenUnaryTest, FastTanhfIsVectorized) { EXPECT_THAT(ir, ContainsRegex("fmul <16 x float>")); EXPECT_THAT(ir, ContainsRegex("fmul <8 x double>")); - EXPECT_THAT(ir, ContainsRegex("<16 x float>.*0x3E4DF2A3C0000000")); + EXPECT_THAT(ir, ContainsRegex("<16 x float>.*f0x326F951E")); EXPECT_THAT(ir, Not(ContainsRegex("llvm.x86"))); EXPECT_THAT(ir, Not(ContainsRegex("llvm.aarch64"))); EXPECT_THAT(ir, ContainsRegex("xla.unused.tanh.v16f32")); diff --git a/xla/codegen/intrinsic/cpp/tanh_test.cc b/xla/codegen/intrinsic/cpp/tanh_test.cc index 8f8a2d61a38b5..256c8c449502a 100644 --- a/xla/codegen/intrinsic/cpp/tanh_test.cc +++ b/xla/codegen/intrinsic/cpp/tanh_test.cc @@ -41,8 +41,7 @@ TEST(TanhTest, FloatTanhVectorized) { module->print(stream, nullptr); EXPECT_THAT(ir, ContainsRegex("fmul <4 x float>")); - EXPECT_THAT( - ir, ContainsRegex("fcmp olt <4 x float>.*float 0x3F3A36E2E0000000.*")); + EXPECT_THAT(ir, ContainsRegex("fcmp olt <4 x float>.*float 4.000000e-04.*")); } } // namespace } // namespace codegen diff --git a/xla/service/cpu/tests/cpu_intrinsic_test.cc b/xla/service/cpu/tests/cpu_intrinsic_test.cc index 09e6ec53270d6..e6c60f469e0f3 100644 --- a/xla/service/cpu/tests/cpu_intrinsic_test.cc +++ b/xla/service/cpu/tests/cpu_intrinsic_test.cc @@ -145,14 +145,14 @@ IntrinsicTestSpec CpuUnaryIntrinsicTestCases[] = { IntrinsicTestSpec{ HloOpcode::kExp, F32, true, kTriple_x86_64, "", - R"(CHECK: fmul fast <4 x float> splat (float 0xBF2BD01060000000)"}, + R"(CHECK: fmul fast <4 x float> {{.*}}splat (float f0x3FB8AA3B)"}, // Check that we see inlined vectorized exp.f64 code IntrinsicTestSpec{HloOpcode::kExp, F64, true, kTriple_x86_64, "", R"( CHECK-NOT: define {{[a-z]* ?}}<4 x double> @xla.exp.v4f32 CHECK-NOT: define {{[a-z]* ?}}<4 x double> @xla.exp.v4f64 - CHECK: fmul <2 x double> {{.*}}splat (double 0x3FF71547652B82FE) + CHECK: fmul <2 x double> {{.*}}splat (double f0x3FF71547652B82FE) CHECK-NOT: define {{[a-z]* ?}}<2 x double> @xla.exp.v2f32 CHECK-NOT: define {{[a-z]* ?}}<4 x double> @xla.exp.v4f64 )"}, @@ -161,7 +161,7 @@ IntrinsicTestSpec CpuUnaryIntrinsicTestCases[] = { R"( CHECK-NOT: define {{[a-z]* ?}}<2 x double> @xla.exp.v2f64 CHECK-NOT: define {{[a-z]* ?}}<4 x float> @xla.exp.v4f32 - CHECK: fmul <4 x double> {{.*}}splat (double 0x3FF71547652B82FE) + CHECK: fmul <4 x double> {{.*}}splat (double f0x3FF71547652B82FE) CHECK-NOT: define {{[a-z]* ?}}<4 x float> @xla.exp.v4f32 CHECK-NOT: define {{[a-z]* ?}}<2 x double> @xla.exp.v2f64 )"}, @@ -171,11 +171,11 @@ IntrinsicTestSpec CpuUnaryIntrinsicTestCases[] = { IntrinsicTestSpec{ HloOpcode::kExp, F32, true, kTriple_x86_64, "+avx", - R"(CHECK: fmul fast <8 x float> splat (float 0xBF2BD01060000000)"}, + R"(CHECK: fmul fast <8 x float> {{.*}}splat (float f0x3FB8AA3B)"}, IntrinsicTestSpec{ HloOpcode::kExp, F32, true, kTriple_android_arm, "+neon", - R"(CHECK: fmul fast <4 x float> splat (float 0xBF2BD01060000000)"}, + R"(CHECK: fmul fast <4 x float> {{.*}}splat (float f0x3FB8AA3B)"}, IntrinsicTestSpec{ HloOpcode::kRsqrt, F32, true, kTriple_x86_64, "+avx", @@ -209,15 +209,15 @@ IntrinsicTestSpec CpuUnaryIntrinsicTestCases[] = { IntrinsicTestSpec{ HloOpcode::kLog, F32, true, kTriple_x86_64, "", - R"(CHECK: fadd fast <4 x float> splat (float 0x3FBDE4A340000000)"}, + R"(CHECK: fadd fast <4 x float> splat (float f0x3DEF251A)"}, IntrinsicTestSpec{ HloOpcode::kLog, F32, true, kTriple_x86_64, "+avx", - R"(CHECK: fadd fast <8 x float> splat (float 0x3FBDE4A340000000)"}, + R"(CHECK: fadd fast <8 x float> splat (float f0x3DEF251A)"}, IntrinsicTestSpec{ HloOpcode::kLog, F32, true, kTriple_android_arm, "", - R"(CHECK: fadd fast <4 x float> splat (float 0x3FBDE4A340000000)"}}; + R"(CHECK: fadd fast <4 x float> splat (float f0x3DEF251A)"}}; INSTANTIATE_TEST_SUITE_P(CpuUnaryIntrinsicTestInstantiation, CpuUnaryIntrinsicTest, diff --git a/xla/service/llvm_ir/llvm_util_test.cc b/xla/service/llvm_ir/llvm_util_test.cc index 3f9719b52b6f1..bdcc3569c54dc 100644 --- a/xla/service/llvm_ir/llvm_util_test.cc +++ b/xla/service/llvm_ir/llvm_util_test.cc @@ -83,22 +83,22 @@ TEST_F(EmitReducePrecisionIrExecutionTest, EmitReducePrecisionIR_F16ToF8e5m2) { EmitReducePrecisionIrTestCase test_cases[] = { // clang-format off - {0.0, "half 0xH0000"}, - {0x1.0p-14, "half 0xH0400"}, - {0.250, "half 0xH3400"}, - {1.0, "half 0xH3C00"}, - {0x1.2p0, "half 0xH3C00"}, - {0x1.Cp15, "half 0xH7B00"}, - {-0x1.Cp15, "half 0xHFB00"}, - {0x1.Dp15, "half 0xH7B00"}, - {0x1.Ep15, "half 0xH7C00"}, - {0x1.0p16, "half 0xH7C00"}, - {inf, "half 0xH7C00"}, - {-inf, "half 0xHFC00"}, - {qnan, "half 0xH7E00"}, - {-qnan, "half 0xHFE00"}, - {snan, "half 0xH7F00"}, - {-snan, "half 0xHFF00"}, + {0.0, "half 0.000000e+00"}, + {0x1.0p-14, "half 6.103520e-05"}, + {0.250, "half 2.500000e-01"}, + {1.0, "half 1.000000e+00"}, + {0x1.2p0, "half 1.000000e+00"}, + {0x1.Cp15, "half 5.734400e+04"}, + {-0x1.Cp15, "half -5.734400e+04"}, + {0x1.Dp15, "half 5.734400e+04"}, + {0x1.Ep15, "half +inf"}, + {0x1.0p16, "half +inf"}, + {inf, "half +inf"}, + {-inf, "half -inf"}, + {qnan, "half +qnan"}, + {-qnan, "half -qnan"}, + {snan, "half +nan(0x100)"}, + {-snan, "half -nan(0x100)"}, // clang-format on }; @@ -130,22 +130,22 @@ TEST_F(EmitReducePrecisionIrExecutionTest, EmitReducePrecisionIR_F16ToF8e4m3) { EmitReducePrecisionIrTestCase test_cases[] = { // clang-format off - {0.0, "half 0xH0000"}, - {0x1.0p-6, "half 0xH2400"}, - {0.125, "half 0xH3000"}, - {1.0, "half 0xH3C00"}, - {0x1.1p0, "half 0xH3C00"}, - {0x1.Ep7, "half 0xH5B80"}, - {-0x1.Ep7, "half 0xHDB80"}, - {0x1.E8p7, "half 0xH5B80"}, - {0x1.Fp7, "half 0xH7C00"}, - {0x1.0p8, "half 0xH7C00"}, - {inf, "half 0xH7C00"}, - {-inf, "half 0xHFC00"}, - {qnan, "half 0xH7E00"}, - {-qnan, "half 0xHFE00"}, - {snan, "half 0xH7E00"}, - {-snan, "half 0xHFE00"}, + {0.0, "half 0.000000e+00"}, + {0x1.0p-6, "half 1.562500e-02"}, + {0.125, "half 1.250000e-01"}, + {1.0, "half 1.000000e+00"}, + {0x1.1p0, "half 1.000000e+00"}, + {0x1.Ep7, "half 2.400000e+02"}, + {-0x1.Ep7, "half -2.400000e+02"}, + {0x1.E8p7, "half 2.400000e+02"}, + {0x1.Fp7, "half +inf"}, + {0x1.0p8, "half +inf"}, + {inf, "half +inf"}, + {-inf, "half -inf"}, + {qnan, "half +qnan"}, + {-qnan, "half -qnan"}, + {snan, "half +qnan"}, + {-snan, "half -qnan"}, // clang-format on }; @@ -177,22 +177,22 @@ TEST_F(EmitReducePrecisionIrExecutionTest, EmitReducePrecisionIR_F16ToF8e3m4) { EmitReducePrecisionIrTestCase test_cases[] = { // clang-format off - {0.0, "half 0xH0000"}, - {0x1.0p-2, "half 0xH3400"}, - {0.5, "half 0xH3800"}, - {1.0, "half 0xH3C00"}, - {0x1.08p0, "half 0xH3C00"}, - {0x1.Fp3, "half 0xH4BC0"}, - {-0x1.Fp3, "half 0xHCBC0"}, - {0x1.F4p3, "half 0xH4BC0"}, - {0x1.F8p3, "half 0xH7C00"}, - {0x1.0p4, "half 0xH7C00"}, - {inf, "half 0xH7C00"}, - {-inf, "half 0xHFC00"}, - {qnan, "half 0xH7E00"}, - {-qnan, "half 0xHFE00"}, - {snan, "half 0xH7E00"}, - {-snan, "half 0xHFE00"}, + {0.0, "half 0.000000e+00"}, + {0x1.0p-2, "half 2.500000e-01"}, + {0.5, "half 5.000000e-01"}, + {1.0, "half 1.000000e+00"}, + {0x1.08p0, "half 1.000000e+00"}, + {0x1.Fp3, "half 1.550000e+01"}, + {-0x1.Fp3, "half -1.550000e+01"}, + {0x1.F4p3, "half 1.550000e+01"}, + {0x1.F8p3, "half +inf"}, + {0x1.0p4, "half +inf"}, + {inf, "half +inf"}, + {-inf, "half -inf"}, + {qnan, "half +qnan"}, + {-qnan, "half -qnan"}, + {snan, "half +qnan"}, + {-snan, "half -qnan"}, // clang-format on }; @@ -223,18 +223,18 @@ TEST_F(EmitReducePrecisionIrExecutionTest, EmitReducePrecisionIrTestCase test_cases[] = { // clang-format off - {0.0, "half 0xH0000"}, - {0x1.0p-6, "half 0xH2400"}, - {0.125, "half 0xH3000"}, - {1.0, "half 0xH3C00"}, - {0x1.1p0, "half 0xH3C00"}, - {0x1.Cp8, "half 0xH5F00"}, - {-0x1.Cp8, "half 0xHDF00"}, - {0x1.Dp8, "half 0xH5F00"}, - {0x1.Ep8, "half 0xH5F80"}, - {0x1.0p9, "half 0xH6000"}, - {inf, "half 0xH7C00"}, - {-inf, "half 0xHFC00"}, + {0.0, "half 0.000000e+00"}, + {0x1.0p-6, "half 1.562500e-02"}, + {0.125, "half 1.250000e-01"}, + {1.0, "half 1.000000e+00"}, + {0x1.1p0, "half 1.000000e+00"}, + {0x1.Cp8, "half 4.480000e+02"}, + {-0x1.Cp8, "half -4.480000e+02"}, + {0x1.Dp8, "half 4.480000e+02"}, + {0x1.Ep8, "half 4.800000e+02"}, + {0x1.0p9, "half 5.120000e+02"}, + {inf, "half +inf"}, + {-inf, "half -inf"}, // clang-format on };