Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions third_party/llvm/workspace.bzl
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
248 changes: 31 additions & 217 deletions third_party/shardy/temporary.patch
Original file line number Diff line number Diff line change
@@ -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<Attribute> 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<NamedAttribute> dictFields;
- dictFields.push_back(builder.getNamedAttr("axes", axesArrayAttr));
-
- if (!sdyMeshAttr.getDeviceIds().empty()) {
- auto type = RankedTensorType::get(
- {static_cast<int64_t>(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<MeshOp>())) {
- 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<[<mesh<["b"=2]>, [{"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<mesh="m2", origin=[]> (%[[T0]], %arg1)
+// CHECK-NEXT: %[[ADD:.*]] = mpmd.fragment<mesh="m2", origin=[]> (%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<mesh="m2", origin=[]> (%[[T0]], %arg1)
+// CHECK-NEXT: %[[ADD:.*]] = mpmd.fragment<mesh="m2", origin=[]> (%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<double>(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<f32>, %arg2: tensor<f32>):
-+ %1 = "stablehlo.add"(%arg1, %arg2) : (tensor<f32>, tensor<f32>) -> tensor<f32>
-+ "stablehlo.return"(%1) : (tensor<f32>) -> ()
-+ }) {
-+ replica_groups = #stablehlo.replica_group_mesh_axes<mesh = @sdy_mesh, axes = [#stablehlo.axis_ref<name = "x">]>
-+ } : (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<f32>, %arg2: tensor<f32>):
-+ %1 = "stablehlo.add"(%arg1, %arg2) : (tensor<f32>, tensor<f32>) -> tensor<f32>
-+ "stablehlo.return"(%1) : (tensor<f32>) -> ()
-+ }) {
-+ replica_groups = #stablehlo.replica_group_mesh_axes<mesh = @sdy_mesh_dev, axes = [#stablehlo.axis_ref<name = "x">]>
-+ } : (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<f32>, %arg2: tensor<f32>):
-+ %1 = "stablehlo.add"(%arg1, %arg2) : (tensor<f32>, tensor<f32>) -> tensor<f32>
-+ "stablehlo.return"(%1) : (tensor<f32>) -> ()
-+ }) {
-+ replica_groups = #stablehlo.replica_group_mesh_axes<mesh = @sdy_mesh_max, 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<ArrayAttr>("axes");
-+ if (!axesAttr) return {};
-+
-+ SmallVector<stablehlo::MeshAxisAttr> axes;
-+ for (Attribute attr : axesAttr) {
-+ auto axisDict = dyn_cast<DictionaryAttr>(attr);
-+ if (!axisDict) return {};
-+ auto nameAttr = axisDict.getAs<StringAttr>("name");
-+ auto sizeAttr = axisDict.getAs<IntegerAttr>("size");
-+ if (!nameAttr || !sizeAttr) return {};
-+ axes.push_back(stablehlo::MeshAxisAttr::get(ctx, nameAttr.getValue(),
-+ sizeAttr.getInt()));
-+ }
-+
-+ // Extract device_ids
-+ DenseIntElementsAttr deviceIds =
-+ dictAttr.getAs<DenseIntElementsAttr>("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::MeshAttr>("stablehlo.mesh");
- }
- if (!meshAttr) {
-+ if (auto dictAttr =
-+ symbolOp->getAttrOfType<DictionaryAttr>("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
4 changes: 2 additions & 2 deletions third_party/shardy/workspace.bzl
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down
12 changes: 12 additions & 0 deletions third_party/triton/common/llvm_cl913905807.patch
Original file line number Diff line number Diff line change
@@ -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 ";

1 change: 1 addition & 0 deletions third_party/triton/common/series.bzl
Original file line number Diff line number Diff line change
Expand Up @@ -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
]
2 changes: 1 addition & 1 deletion xla/codegen/intrinsic/cpp/eigen_unary_test.cc
Original file line number Diff line number Diff line change
Expand Up @@ -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"));
Expand Down
3 changes: 1 addition & 2 deletions xla/codegen/intrinsic/cpp/tanh_test.cc
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
16 changes: 8 additions & 8 deletions xla/service/cpu/tests/cpu_intrinsic_test.cc
Original file line number Diff line number Diff line change
Expand Up @@ -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
)"},
Expand All @@ -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
)"},
Expand All @@ -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",
Expand Down Expand Up @@ -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,
Expand Down
Loading
Loading