Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

PIN to 2023_11_29 #4

Merged
merged 110 commits into from
Dec 13, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
110 commits
Select commit Hold shift + click to select a range
aa20614
[XLA:GPU] NFC: use matchers in PriorityFusionTest.
chsigg Nov 26, 2023
048e5e4
[XLA:GPU] NFC: make IsReadCoalesced() slightly easier to read.
chsigg Nov 26, 2023
0154c7b
Make cudnn_fused_conv_rewriter_test work in OSS.
akuegel Nov 27, 2023
f4d1db4
NFC: Reduce number of function arguments in reduction.cc.
jreiffers Nov 27, 2023
3249850
Enable tests that require v100 in OSS.
akuegel Nov 27, 2023
77f55b9
Move ReductionCodegenState to reduction.cc and rename to ReductionGro…
jreiffers Nov 27, 2023
12adbd1
[XLA:GPU] Add can_fuse cache.
olegshyshkov Nov 27, 2023
3ccca51
Integrate LLVM at llvm/llvm-project@5e5a22caf88a
gribozavr Nov 27, 2023
59b1d67
[XLA:GPU] (NFC) Remove producer_user_count_.
olegshyshkov Nov 27, 2023
bb07b21
[XLA:GPU][NFC] Clean up tile analysis tests to use `TF_ASSERT_OK_AND_…
bchetioui Nov 27, 2023
ab750b1
Update TFRT dependency to use revision
tensorflower-gardener Nov 27, 2023
eb3f22f
[XLA:GPU][NFC] Transform dimension propagation to the functional para…
tdanyluk Nov 27, 2023
2e383a1
[XLA:GPU] Fix for "Transform dimension propagation to the functional …
tdanyluk Nov 27, 2023
7506f69
[XLA] Collective pipeliner goes through small reduces as they do not …
Nov 27, 2023
5320016
Adds back the solver parameter string.
tensorflower-gardener Nov 27, 2023
c9e2046
[stream_executor] Add Case conditional command to CommandBuffer #6973
ezhulenev Nov 27, 2023
b71268c
[stream_executor] Add For conditional command to CommandBuffer
ezhulenev Nov 27, 2023
dee190a
[stream_executor] NFC: Do not leak internal stream executor header
ezhulenev Nov 27, 2023
2c5ce09
[XLA:GPU] Triton GEMM: enable fusion of inputs concatenated along non…
Nov 27, 2023
fea3aa5
Deduplicate some code by using a refactored portion of MaybeFollowIns…
tensorflower-gardener Nov 27, 2023
48d695d
Search for possible predecessors instead of only caching one hops.
tensorflower-gardener Nov 27, 2023
ea34cf8
PR #7269: Fix an incorrect static_assert
elfringham Nov 27, 2023
e02661f
PR #7277: [ROCM] rocm 6.0 fixes followup
pemeliya Nov 27, 2023
7c23e8c
Add c-api support for GetCompiledMemoryStats.
pschuh Nov 28, 2023
ecae151
Update visibility for tf.data.
yangustc07 Nov 28, 2023
889fa5c
[xla:gpu] NFC: move thunk->CMD conversion to runtime3 #6528
anlunx Nov 28, 2023
6d09f41
[stream_executor] NFC: Do not leak internal stream executor header
ezhulenev Nov 28, 2023
6643428
[stream_executor] Do not use KernelArgs directly in GpuExecutor::Launch
tyb0807 Nov 28, 2023
e795dc3
[stream_executor] Add While conditional command to CommandBuffer #6973
ezhulenev Nov 28, 2023
537ef5e
[stream_executor] NFC: Do not leak internal stream executor header
ezhulenev Nov 28, 2023
2521abf
Adds tests for AbsOp, DivOp, ExpOp, MaxOp, MulOp, PowOp, SubOp in `xl…
yishuangP Nov 28, 2023
33fc605
[stream_executor] NFC: Restore private visibility of stream executor …
ezhulenev Nov 28, 2023
8e3830c
[XLA/GPU]Add methods to construct GpuCompilationEnvironment.
tensorflower-gardener Nov 28, 2023
9ea69b0
[XLA:GPU] Remove xla_gpu_single_wave_autotuning flag
tdanyluk Nov 28, 2023
370895f
Add a vector->SCF pass to hlo_xla_runtime_pipeline.
pifon2a Nov 28, 2023
80fe7b1
[XLA:GPU] Add cache of runtime data for unfused instruction.
olegshyshkov Nov 28, 2023
d44ee9f
Retire deallocation dialect.
jreiffers Nov 28, 2023
9afe429
PR #7287: Fix TF32 in gemm calls
benbarsdell Nov 28, 2023
3ee4ca8
[XLA:GPU] Add cache for fusion runtime.
olegshyshkov Nov 28, 2023
a51b04c
[XLA:GPU] Fix crash in (Triton) tiling of reductions.
jvstokes Nov 28, 2023
d0aed34
PR #7323: [ROCm] Restore private visibility of stream executor privat…
i-chaochen Nov 28, 2023
95ddb7a
[XLA:CPU] Refactor local collectives into a separate file behind an i…
hawkinsp Nov 28, 2023
bff1238
[XLA:GPU][NFC] Split GemmRewriterTriton into 4 parts
tdanyluk Nov 28, 2023
3032620
Failed compatibility tests were disabled for older version plugins an…
Nov 28, 2023
0ca186f
`third_party/gpus` changes from PR #7277 that were missed
ddunl Nov 28, 2023
7a7eddc
[xla:runtime] NFC: Remove runner library
ezhulenev Nov 29, 2023
1347b18
PR #7136: [XLA:GPU] Add `Allocate` command to command buffer
shawnwang18 Nov 29, 2023
6046ab5
Update XLA GPU config with NVCC compiler.
tensorflower-gardener Nov 29, 2023
6b4d712
For convolutions that can be interpreted as dots, rely on DotHandler …
tensorflower-gardener Nov 29, 2023
85a4f70
[stream_executor] Record cond_builder before evaluating loop conditio…
ezhulenev Nov 29, 2023
c58bc73
[stream_executor] Add Memset command to CommandBuffer
ezhulenev Nov 29, 2023
fede5db
[stream_executor] Add support for updating Memcpy command parameters
ezhulenev Nov 29, 2023
6c4ff70
Converts the AutoShardingSolverRequest object into a proto (which pro…
tensorflower-gardener Nov 29, 2023
8d49e11
Integrate StableHLO at openxla/stablehlo@83f095e7
ghpvnist Nov 29, 2023
12a88f3
[stream_executor] NFC: Guard new features with CUDA_VERSION check
ezhulenev Nov 29, 2023
efcf825
[XLA] Add support for while loop simplifier to remove duplicated "dyn…
Nov 29, 2023
1317a62
[XLA:GPU] Add a pass merging producer fusions in to Triton Softmax fu…
jvstokes Nov 29, 2023
5cd07b6
[XLA] Fix the MHLO TopK `assemblyFormat` to correctly support the opt…
dimitar-asenov Nov 29, 2023
e9500e3
[XLA:GPU][NFC] Beautify code related to Triton fusions
tdanyluk Nov 29, 2023
ec0b228
Fix a compile error caused by in_process_collectives
apaszke Nov 29, 2023
422b8e7
[TileAnalysis] Add indexing computation for reshape.
pifon2a Nov 29, 2023
b98258d
Basic simplifier for indexing maps.
jreiffers Nov 29, 2023
c9c16d6
[TileAnalysis] Add indexing computation for HloBitcast.
pifon2a Nov 29, 2023
757f635
[XLA:GPU] Append ".0" suffix to all instructions names.
olegshyshkov Nov 29, 2023
fecd5e7
[xla:gpu] NFC: Migrate custom CUTLASS kernel to GemmUniversal templates
ezhulenev Nov 29, 2023
d555040
[XLA:GPU][NFC] Refactor tiling propagation.
Nov 29, 2023
4e6e1c5
[XLA] Lower verbosity in MSA.
berkinilbeyi Nov 29, 2023
0e751bf
Record free_gpu_system_memory using addressable_devices() to avoid po…
tensorflower-gardener Nov 29, 2023
7eacd1b
Integrate LLVM at llvm/llvm-project@3287ae8f6520
gribozavr Nov 29, 2023
c4f0a9d
[XLA:CPU] Add a direct implementation of AllGather, rather than lower…
hawkinsp Nov 29, 2023
8e8b121
User sharding annotations can sometimes be invalid wrt to the shape o…
tensorflower-gardener Nov 29, 2023
856d51c
Some cleanup: Inline a function used once, remove some dead code, and…
tensorflower-gardener Nov 29, 2023
c28baac
Reenable `BitcastConvert` tests in JAX CI now that the tests pass again
ddunl Nov 29, 2023
222c77a
Update TFRT dependency to use revision
tensorflower-gardener Nov 29, 2023
89f3b13
Adds the ability to dump solver request protos.
tensorflower-gardener Nov 29, 2023
57edcfa
Fastpath for setting disable-jit.
pschuh Nov 29, 2023
c756093
When iterative solving is turned on, shardings annotations from a pre…
tensorflower-gardener Nov 29, 2023
17d9572
Move large functions in auto_sharding_dot_handler.cc out of class def…
tensorflower-gardener Nov 29, 2023
9e172db
[XLA]
tensorflower-gardener Nov 29, 2023
954527c
Changed Version of Bazel to version 6.4.0
tensorflower-gardener Nov 29, 2023
69f26cf
[XLA:CPU] Add a direct implementation of ReduceScatter, instead of lo…
hawkinsp Nov 29, 2023
aa1f21c
PR #7370: [XLA:GPU] fix command_buffer_thunk_test failure
shawnwang18 Nov 29, 2023
ae7fead
[xla:ffi] Added error reporting to existing decoders
superbobry Nov 29, 2023
f69ec88
Reverts 69f26cf6a03e9e7fc4c45c138dc3917cc3aa36f8
hawkinsp Nov 29, 2023
3400945
Remove obselete TODO
ddunl Nov 29, 2023
baf6958
Reverts c4f0a9d1ba9adc9ed720ef5853fcbbdb4d749df7
hawkinsp Nov 29, 2023
a680ac3
Remove tensorflow namespace from `tsl/platform/status_matchers.h`
ddunl Nov 29, 2023
3981d17
[XLA:LatencyHidingScheduler] Prefer picking the instruction with less…
tensorflower-gardener Nov 29, 2023
b328fe8
[xla:ffi] Parameterized ffi::BaseBuffer with its rank
superbobry Nov 29, 2023
4bf4b1e
[XLA] Allow moving for HloSharding in xla::HloInstruction::set_sharding.
tensorflower-gardener Nov 29, 2023
6449702
Add EvalOrPattern to StablehloRefineShapes pass
tensorflower-gardener Nov 29, 2023
aab9273
[XLA] Allow tuple_shardings move in the HloSharding ctor.
tensorflower-gardener Nov 29, 2023
135b0ea
[xla:gpu] Move external allocation implementation details from Stream…
ezhulenev Nov 29, 2023
3521147
[XLA] Do not create temporary vector in HloSharding::GetSubSharding.
tensorflower-gardener Nov 29, 2023
d6a1cf3
[XLA] Remove an unnecessary HloSharding assignment.
tensorflower-gardener Nov 30, 2023
4695163
Use clang+NVCC compilers for all XLA GPU jobs.
tensorflower-gardener Nov 30, 2023
3f03d1c
[XLA] Move HloSharding where possible.
tensorflower-gardener Nov 30, 2023
c5d313d
[XLA] Change the prototype of ReturnImprovedShardingImpl to minmize c…
tensorflower-gardener Nov 30, 2023
ecfe268
[XLA] Optimize HloValue::ComputeUses().
tensorflower-gardener Nov 30, 2023
da00f2b
[xla:gpu] Own command buffer allocations at a Thunk level
ezhulenev Nov 30, 2023
e86ce2e
Integrate LLVM at llvm/llvm-project@f688e0901213
tensorflower-gardener Nov 30, 2023
c6f0a6e
[stream_executor] Use new CUDA runtime API for TopK
tyb0807 Nov 30, 2023
a36626c
[xla:gpu] Add the AOT compilation pipeline for thunk runtime #7360
anlunx Nov 30, 2023
46d4355
Update TFRT dependency to use revision
tensorflower-gardener Nov 30, 2023
6eed44a
[XLA:GPU] Improve errors if callbacks are not provided to the Collect…
hawkinsp Nov 30, 2023
dbe129e
Relocates the evaluation output into the core solver.
tensorflower-gardener Nov 30, 2023
59dea45
Lower tensor.from_elements and shape.broadcast ops in ShapeLegalizeToHLO
tensorflower-gardener Nov 30, 2023
91b8ddb
Internal infrastructure change
GleasonK Nov 30, 2023
58e6b42
[stream_executor][NFC] More doc for While and For command
tyb0807 Nov 30, 2023
9e419d5
Merge commit '58e6b428e22e40c4100a7b66790fbe86dc9d7845' into HEAD
Dec 13, 2023
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
29 changes: 2 additions & 27 deletions .bazelrc
Original file line number Diff line number Diff line change
Expand Up @@ -526,34 +526,9 @@ build:rbe_linux_cuda --repo_env=TF_TENSORRT_CONFIG_REPO="@sigbuild-r2.16-clang_c
build:rbe_linux_cuda --repo_env=TF_NCCL_CONFIG_REPO="@sigbuild-r2.16-clang_config_nccl"
test:rbe_linux_cuda --test_env=LD_LIBRARY_PATH="/usr/local/cuda/lib64:/usr/local/cuda/extras/CUPTI/lib64"

build:rbe_linux_cuda_nvcc --config=cuda
build:rbe_linux_cuda_nvcc --action_env=TF_NVCC_CLANG="1"
build:rbe_linux_cuda_nvcc --config=rbe_linux_cuda
build:rbe_linux_cuda_nvcc --repo_env TF_NCCL_USE_STUB=1
build:rbe_linux_cuda_nvcc --@xla//xla/python:enable_gpu=true
build:rbe_linux_cuda_nvcc --@xla//xla/python:jax_cuda_pip_rpaths=true
build:rbe_linux_cuda_nvcc --define=xla_python_enable_gpu=true
build:rbe_linux_cuda_nvcc --config=tensorrt
build:rbe_linux_cuda_nvcc --repo_env=TF_CUDA_COMPUTE_CAPABILITIES="sm_50,sm_60,sm_70,sm_75,compute_80"
build:rbe_linux_cuda_nvcc --action_env=TF_CUDA_VERSION="12"
build:rbe_linux_cuda_nvcc --action_env=TF_CUDNN_VERSION="8"
build:rbe_linux_cuda_nvcc --action_env=CUDA_TOOLKIT_PATH="/usr/local/cuda-12.2"
build:rbe_linux_cuda_nvcc --action_env=LD_LIBRARY_PATH="/usr/local/cuda/lib64:/usr/local/cuda/extras/CUPTI/lib64"
build:rbe_linux_cuda_nvcc --crosstool_top="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_cuda//crosstool:toolchain"
build:rbe_linux_cuda_nvcc --config=rbe_linux
build:rbe_linux_cuda_nvcc --host_crosstool_top="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_cuda//crosstool:toolchain"
build:rbe_linux_cuda_nvcc --extra_toolchains="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_cuda//crosstool:toolchain-linux-x86_64"
build:rbe_linux_cuda_nvcc --extra_execution_platforms="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_platform//:platform"
build:rbe_linux_cuda_nvcc --host_platform="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_platform//:platform"
build:rbe_linux_cuda_nvcc --platforms="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_platform//:platform"
build:rbe_linux_cuda_nvcc --repo_env=TF_PYTHON_CONFIG_REPO="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_python3.9"
build:rbe_linux_cuda_nvcc --python_path="/usr/bin/python3"
# These you may need to change for your own GCP project.
common:rbe_linux_cuda_nvcc --remote_instance_name=projects/tensorflow-testing/instances/default_instance
build:rbe_linux_cuda_nvcc --repo_env=REMOTE_GPU_TESTING=1
build:rbe_linux_cuda_nvcc --repo_env=TF_CUDA_CONFIG_REPO="@ubuntu20.04-gcc9_manylinux2014-cuda12.2-cudnn8.9_config_cuda"
build:rbe_linux_cuda_nvcc --repo_env=TF_TENSORRT_CONFIG_REPO="@ubuntu20.04-gcc9_manylinux2014-cuda12.2-cudnn8.9_config_tensorrt"
build:rbe_linux_cuda_nvcc --repo_env=TF_NCCL_CONFIG_REPO="@ubuntu20.04-gcc9_manylinux2014-cuda12.2-cudnn8.9_config_nccl"
test:rbe_linux_cuda_nvcc --test_env=LD_LIBRARY_PATH="/usr/local/cuda/lib64:/usr/local/cuda/extras/CUPTI/lib64"
build:rbe_linux_cuda_nvcc --action_env=TF_NVCC_CLANG="1"

# TODO(kanglan): Remove rbe_win and rbe_win_py3* after b/289091160 is fixed
build:rbe_win --config=rbe_base
Expand Down
2 changes: 1 addition & 1 deletion .bazelversion
Original file line number Diff line number Diff line change
@@ -1,2 +1,2 @@
6.1.0
6.4.0
# NOTE: Update Bazel version in tensorflow/tools/ci_build/release/common.sh.oss
3 changes: 1 addition & 2 deletions .kokoro/jax/build.sh
Original file line number Diff line number Diff line change
Expand Up @@ -74,7 +74,6 @@ build_and_test_on_rbe_gpu() {
# Runs non-multiaccelerator tests with one GPU apiece.
# It appears --run_under needs an absolute path.

# TODO(ddunleavy): reenable `LaxTest.testBitcastConvertType`
bazel \
test \
--verbose_failures=true \
Expand All @@ -87,7 +86,7 @@ build_and_test_on_rbe_gpu() {
--test_output=errors \
--test_env=JAX_SKIP_SLOW_TESTS=1 \
--test_env=TF_CPP_MIN_LOG_LEVEL=0 \
--test_env=JAX_EXCLUDE_TEST_TARGETS="PmapTest.testSizeOverflow|LaxTest.testBitcastConvertType" \
--test_env=JAX_EXCLUDE_TEST_TARGETS="PmapTest.testSizeOverflow" \
--test_tag_filters=-multiaccelerator \
-- //tests:gpu_tests //tests:backend_independent_tests
}
Expand Down
10 changes: 1 addition & 9 deletions .kokoro/linux/build.sh
Original file line number Diff line number Diff line change
Expand Up @@ -26,10 +26,6 @@ function is_linux_gpu_job() {
[[ "$KOKORO_JOB_NAME" =~ tensorflow/xla/linux/.*gpu.* ]]
}

function is_use_nvcc() {
[[ "${USE_NVCC:-}" == "true" ]]
}

# Pull the container (in case it was updated since the instance started) and
# store its SHA in the Sponge log.
docker pull "$DOCKER_IMAGE"
Expand All @@ -54,11 +50,7 @@ if is_linux_gpu_job ; then
TAGS_FILTER="$TAGS_FILTER,gpu,requires-gpu-nvidia,-no_gpu"
ADDITIONAL_FLAGS="$ADDITIONAL_FLAGS --run_under=//tools/ci_build/gpu_build:parallel_gpu_execute"
RC_FILE="/usertools/gpu.bazelrc"
if is_use_nvcc ; then
RBE_CONFIG="rbe_linux_cuda_nvcc"
else
RBE_CONFIG="rbe_linux_cuda"
fi
RBE_CONFIG="rbe_linux_cuda_nvcc"
echo "***NOTE: nvidia-smi lists the highest CUDA version the driver supports, which may be different than the version of CUDA actually used!!***"
nvidia-smi
else
Expand Down
11 changes: 11 additions & 0 deletions third_party/llvm/generated.patch
Original file line number Diff line number Diff line change
@@ -1 +1,12 @@
Auto generated patch. Do not edit or delete it, even if empty.
diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel
--- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel
@@ -594,6 +594,7 @@
name = "__support_bit",
hdrs = ["src/__support/bit.h"],
deps = [
+ ":__support_cpp_type_traits",
":__support_macros_attributes",
],
)
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 = "af7a1453526a88a0e242baf156244aa4ae42ae4b"
LLVM_SHA256 = "f9f75e4823c2f09a8141ab4db40ee2c79aef96017782a9338e26621ee547d3d5"
LLVM_COMMIT = "f688e0901213726feb9b26cedc61919413cbf59c"
LLVM_SHA256 = "b8885c22a9b77f9c91a316b21d71414a7b48dae38513f170da1554002e85b030"

tf_http_archive(
name = name,
Expand Down
196 changes: 73 additions & 123 deletions third_party/stablehlo/temporary.patch
Original file line number Diff line number Diff line change
Expand Up @@ -156,127 +156,6 @@ diff --ruN a/stablehlo/stablehlo/CMakeLists.txt b/stablehlo/stablehlo/CMakeLists
add_subdirectory(integrations)
add_subdirectory(reference)
add_subdirectory(tests)
diff --ruN a/stablehlo/stablehlo/experimental/BUILD b/stablehlo/stablehlo/experimental/BUILD
--- stablehlo/stablehlo/experimental/BUILD
+++ stablehlo/stablehlo/experimental/BUILD
@@ -0,0 +1,117 @@
+load("//third_party/llvm/llvm-project/mlir:tblgen.bzl", "gentbl_cc_library")
+load("//tensorflow:tensorflow.google.bzl", "get_compatible_with_portable")
+load("//tensorflow/core/platform:rules_cc.bzl", "cc_library")
+
+package(
+ default_applicable_licenses = ["//third_party/stablehlo:license"], # copybara:comment
+ default_visibility = ["//learning/brain/mlir:stablehlo_friends"],
+ licenses = ["notice"],
+)
+
+filegroup(
+ name = "stablehlo_experimental_filegroup",
+ srcs = glob(["**"]),
+)
+
+cc_library(
+ name = "experimental_base",
+ srcs = [
+ "dialect/Base.cpp",
+ ],
+ hdrs = [
+ "dialect/Base.h",
+ ],
+ compatible_with = get_compatible_with_portable(),
+ includes = ["../.."],
+ deps = [
+ "//third_party/llvm/llvm-project/llvm:Support",
+ "//third_party/llvm/llvm-project/mlir:IR",
+ ],
+)
+
+cc_library(
+ name = "experimental_stablehlo_ops",
+ srcs = [
+ "dialect/StablehloOps.cpp",
+ ],
+ hdrs = [
+ "dialect/StablehloOps.h",
+ ],
+ compatible_with = get_compatible_with_portable(),
+ includes = ["../.."],
+ deps = [
+ ":experimental_base",
+ "//third_party/llvm/llvm-project/llvm:Support",
+ "//third_party/llvm/llvm-project/mlir:FuncDialect",
+ "//third_party/llvm/llvm-project/mlir:IR",
+ "//third_party/llvm/llvm-project/mlir:Support",
+ "//third_party/stablehlo:stablehlo_ops",
+ ],
+)
+
+gentbl_cc_library(
+ name = "experimental_stablehlo_pass_inc_gen",
+ compatible_with = get_compatible_with_portable(),
+ tbl_outs = [
+ (
+ [
+ "-gen-pass-decls",
+ ],
+ "transforms/Passes.h.inc",
+ ),
+ ],
+ tblgen = "//third_party/llvm/llvm-project/mlir:mlir-tblgen",
+ td_file = "transforms/Passes.td",
+ deps = ["//third_party/llvm/llvm-project/mlir:PassBaseTdFiles"],
+)
+
+cc_library(
+ name = "experimental_stablehlo_passes",
+ srcs = [
+ "transforms/StablehloCanonicalizeDynamism.cpp",
+ "transforms/StablehloRefineShapes.cpp",
+ ],
+ hdrs = [
+ "transforms/Passes.h",
+ ],
+ compatible_with = get_compatible_with_portable(),
+ includes = ["../.."],
+ deps = [
+ ":experimental_stablehlo_ops",
+ ":experimental_stablehlo_pass_inc_gen",
+ "//third_party/llvm/llvm-project/llvm:Support",
+ "//third_party/llvm/llvm-project/mlir:FuncDialect",
+ "//third_party/llvm/llvm-project/mlir:IR",
+ "//third_party/llvm/llvm-project/mlir:InferTypeOpInterface",
+ "//third_party/llvm/llvm-project/mlir:Pass",
+ "//third_party/llvm/llvm-project/mlir:Support",
+ "//third_party/llvm/llvm-project/mlir:TransformUtils",
+ "//third_party/llvm/llvm-project/mlir:Transforms",
+ "//third_party/stablehlo:base",
+ "//third_party/stablehlo:chlo_ops",
+ "//third_party/stablehlo:stablehlo_ops",
+ "//third_party/stablehlo:stablehlo_ops_inc_gen",
+ "//third_party/stablehlo:stablehlo_type_inference",
+ ],
+)
+
+cc_binary(
+ name = "experimental-stablehlo-opt",
+ srcs = [
+ "tools/StablehloOptMain.cpp",
+ ],
+ compatible_with = get_compatible_with_portable(),
+ includes = ["../.."],
+ deps = [
+ ":experimental_stablehlo_passes",
+ "//third_party/llvm/llvm-project/mlir:AllExtensions",
+ "//third_party/llvm/llvm-project/mlir:AllPassesAndDialects",
+ "//third_party/llvm/llvm-project/mlir:MlirOptLib",
+ "//third_party/llvm/llvm-project/mlir:TosaDialect",
+ "//third_party/stablehlo:interpreter_ops",
+ "//third_party/stablehlo:register",
+ "//third_party/stablehlo:stablehlo_passes",
+ "//third_party/stablehlo:test_utils",
+ "//third_party/stablehlo:tosa_passes",
+ ],
+)
diff --ruN a/stablehlo/stablehlo/experimental/BUILD.bazel b/stablehlo/stablehlo/experimental/BUILD.bazel
--- stablehlo/stablehlo/experimental/BUILD.bazel
+++ stablehlo/stablehlo/experimental/BUILD.bazel
Expand Down Expand Up @@ -1291,7 +1170,7 @@ diff --ruN a/stablehlo/stablehlo/experimental/dialect/StablehloOps.h b/stablehlo
diff --ruN a/stablehlo/stablehlo/experimental/tests/BUILD.bazel b/stablehlo/stablehlo/experimental/tests/BUILD.bazel
--- stablehlo/stablehlo/experimental/tests/BUILD.bazel
+++ stablehlo/stablehlo/experimental/tests/BUILD.bazel
@@ -0,0 +1,58 @@
@@ -0,0 +1,59 @@
+# Copyright 2023 The StableHLO Authors. All Rights Reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
Expand Down Expand Up @@ -1338,6 +1217,7 @@ diff --ruN a/stablehlo/stablehlo/experimental/tests/BUILD.bazel b/stablehlo/stab
+ "lit.site.cfg.py",
+ "//:stablehlo-opt",
+ "//:stablehlo-translate",
+ "//stablehlo/experimental:experimental-stablehlo-opt",
+ "@llvm-project//llvm:FileCheck",
+ "@llvm-project//llvm:not",
+ ] + glob(["%s.bc" % src]),
Expand Down Expand Up @@ -2509,7 +2389,7 @@ diff --ruN a/stablehlo/stablehlo/experimental/transforms/StablehloCanonicalizeDy
diff --ruN a/stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.cpp b/stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.cpp
--- stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.cpp
+++ stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.cpp
@@ -0,0 +1,1293 @@
@@ -0,0 +1,1308 @@
+/* Copyright 2022 The StableHLO Authors.
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -2818,6 +2698,20 @@ diff --ruN a/stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.c
+ }
+};
+
+struct EvalOrOpPattern : public OpRewritePattern<OrOp> {
+ using OpRewritePattern::OpRewritePattern;
+ LogicalResult matchAndRewrite(OrOp op,
+ PatternRewriter& rewriter) const override {
+ auto resultType = op.getType();
+ if (!resultType.getElementType().isInteger(1))
+ return rewriter.notifyMatchFailure(op, "expected boolean element type");
+
+ return evalElementwise(rewriter, op, [&](APSInt lhsInt, APSInt rhsInt) {
+ return getAPSInt(resultType.getElementType(), lhsInt != 0 || rhsInt != 0);
+ });
+ }
+};
+
+struct EvalRemOpPattern : public OpRewritePattern<RemOp> {
+ using OpRewritePattern::OpRewritePattern;
+ LogicalResult matchAndRewrite(RemOp op,
Expand Down Expand Up @@ -3764,6 +3658,7 @@ diff --ruN a/stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.c
+ patterns.add<EvalMaxOpPattern>(&getContext());
+ patterns.add<EvalMinOpPattern>(&getContext());
+ patterns.add<EvalMulOpPattern>(&getContext());
+ patterns.add<EvalOrOpPattern>(&getContext());
+ patterns.add<EvalRemOpPattern>(&getContext());
+ patterns.add<EvalReshapeOpPattern>(&getContext());
+ patterns.add<EvalSelectOpPattern>(&getContext());
Expand Down Expand Up @@ -3803,4 +3698,59 @@ diff --ruN a/stablehlo/stablehlo/experimental/transforms/StablehloRefineShapes.c
+} // namespace experimental
+} // namespace stablehlo
+} // namespace mlir
diff --ruN a/stablehlo/stablehlo/tests/stablehlo_refine_shapes.mlir b/stablehlo/stablehlo/tests/stablehlo_refine_shapes.mlir
--- stablehlo/stablehlo/tests/stablehlo_refine_shapes.mlir
+++ stablehlo/stablehlo/tests/stablehlo_refine_shapes.mlir
@@ -340,6 +340,19 @@
%1 = stablehlo.constant dense<2> : tensor<i64>
%2 = stablehlo.multiply %0, %1 : tensor<i64>
func.return %2 : tensor<i64>
+}
+
+// -----
+
+// CHECK-LABEL: func @eval_or
+func.func @eval_or() -> tensor<i1> {
+ // CHECK-NOT: stablehlo.or
+ // CHECK: [[RESULT:%.*]] = stablehlo.constant dense<true> : tensor<i1>
+ // CHECK: return [[RESULT]]
+ %0 = stablehlo.constant dense<true> : tensor<i1>
+ %1 = stablehlo.constant dense<false> : tensor<i1>
+ %2 = stablehlo.or %0, %1 : tensor<i1>
+ func.return %2 : tensor<i1>
}

// -----
diff --ruN a/stablehlo/stablehlo/transforms/StablehloRefineShapes.cpp b/stablehlo/stablehlo/transforms/StablehloRefineShapes.cpp
--- stablehlo/stablehlo/transforms/StablehloRefineShapes.cpp
+++ stablehlo/stablehlo/transforms/StablehloRefineShapes.cpp
@@ -304,6 +304,20 @@
}
};

+struct EvalOrOpPattern : public OpRewritePattern<OrOp> {
+ using OpRewritePattern::OpRewritePattern;
+ LogicalResult matchAndRewrite(OrOp op,
+ PatternRewriter& rewriter) const override {
+ auto resultType = op.getType();
+ if (!resultType.getElementType().isInteger(1))
+ return rewriter.notifyMatchFailure(op, "expected boolean element type");
+
+ return evalElementwise(rewriter, op, [&](APSInt lhsInt, APSInt rhsInt) {
+ return getAPSInt(resultType.getElementType(), lhsInt != 0 || rhsInt != 0);
+ });
+ }
+};
+
struct EvalRemOpPattern : public OpRewritePattern<RemOp> {
using OpRewritePattern::OpRewritePattern;
LogicalResult matchAndRewrite(RemOp op,
@@ -1165,6 +1179,7 @@
patterns.add<EvalMaxOpPattern>(&getContext());
patterns.add<EvalMinOpPattern>(&getContext());
patterns.add<EvalMulOpPattern>(&getContext());
+ patterns.add<EvalOrOpPattern>(&getContext());
patterns.add<EvalRemOpPattern>(&getContext());
patterns.add<EvalReshapeOpPattern>(&getContext());
patterns.add<EvalSelectOpPattern>(&getContext());

4 changes: 2 additions & 2 deletions third_party/stablehlo/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", "tf_mirror_urls")

def repo():
# LINT.IfChange
STABLEHLO_COMMIT = "e3276cd896751bfbebd7b18112f81547fbc2bc9c"
STABLEHLO_SHA256 = "948d0265a9ea4214ecfa854564793b08161d693e9cfbd73686f2df9e38034ada"
STABLEHLO_COMMIT = "83f095e7217c897f1eccac5652600ceb944cb0e0"
STABLEHLO_SHA256 = "00e442f7e9c8a52a1ac774ce997f8b5a99d12450c4dfe1594df816dcbad5126f"
# LINT.ThenChange(Google-internal path)

tf_http_archive(
Expand Down
29 changes: 2 additions & 27 deletions third_party/tsl/.bazelrc
Original file line number Diff line number Diff line change
Expand Up @@ -526,34 +526,9 @@ build:rbe_linux_cuda --repo_env=TF_TENSORRT_CONFIG_REPO="@sigbuild-r2.16-clang_c
build:rbe_linux_cuda --repo_env=TF_NCCL_CONFIG_REPO="@sigbuild-r2.16-clang_config_nccl"
test:rbe_linux_cuda --test_env=LD_LIBRARY_PATH="/usr/local/cuda/lib64:/usr/local/cuda/extras/CUPTI/lib64"

build:rbe_linux_cuda_nvcc --config=cuda
build:rbe_linux_cuda_nvcc --action_env=TF_NVCC_CLANG="1"
build:rbe_linux_cuda_nvcc --config=rbe_linux_cuda
build:rbe_linux_cuda_nvcc --repo_env TF_NCCL_USE_STUB=1
build:rbe_linux_cuda_nvcc --@xla//xla/python:enable_gpu=true
build:rbe_linux_cuda_nvcc --@xla//xla/python:jax_cuda_pip_rpaths=true
build:rbe_linux_cuda_nvcc --define=xla_python_enable_gpu=true
build:rbe_linux_cuda_nvcc --config=tensorrt
build:rbe_linux_cuda_nvcc --repo_env=TF_CUDA_COMPUTE_CAPABILITIES="sm_50,sm_60,sm_70,sm_75,compute_80"
build:rbe_linux_cuda_nvcc --action_env=TF_CUDA_VERSION="12"
build:rbe_linux_cuda_nvcc --action_env=TF_CUDNN_VERSION="8"
build:rbe_linux_cuda_nvcc --action_env=CUDA_TOOLKIT_PATH="/usr/local/cuda-12.2"
build:rbe_linux_cuda_nvcc --action_env=LD_LIBRARY_PATH="/usr/local/cuda/lib64:/usr/local/cuda/extras/CUPTI/lib64"
build:rbe_linux_cuda_nvcc --crosstool_top="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_cuda//crosstool:toolchain"
build:rbe_linux_cuda_nvcc --config=rbe_linux
build:rbe_linux_cuda_nvcc --host_crosstool_top="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_cuda//crosstool:toolchain"
build:rbe_linux_cuda_nvcc --extra_toolchains="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_cuda//crosstool:toolchain-linux-x86_64"
build:rbe_linux_cuda_nvcc --extra_execution_platforms="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_platform//:platform"
build:rbe_linux_cuda_nvcc --host_platform="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_platform//:platform"
build:rbe_linux_cuda_nvcc --platforms="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_platform//:platform"
build:rbe_linux_cuda_nvcc --repo_env=TF_PYTHON_CONFIG_REPO="@ubuntu20.04-clang_manylinux2014-cuda12.2-cudnn8.9_config_python3.9"
build:rbe_linux_cuda_nvcc --python_path="/usr/bin/python3"
# These you may need to change for your own GCP project.
common:rbe_linux_cuda_nvcc --remote_instance_name=projects/tensorflow-testing/instances/default_instance
build:rbe_linux_cuda_nvcc --repo_env=REMOTE_GPU_TESTING=1
build:rbe_linux_cuda_nvcc --repo_env=TF_CUDA_CONFIG_REPO="@ubuntu20.04-gcc9_manylinux2014-cuda12.2-cudnn8.9_config_cuda"
build:rbe_linux_cuda_nvcc --repo_env=TF_TENSORRT_CONFIG_REPO="@ubuntu20.04-gcc9_manylinux2014-cuda12.2-cudnn8.9_config_tensorrt"
build:rbe_linux_cuda_nvcc --repo_env=TF_NCCL_CONFIG_REPO="@ubuntu20.04-gcc9_manylinux2014-cuda12.2-cudnn8.9_config_nccl"
test:rbe_linux_cuda_nvcc --test_env=LD_LIBRARY_PATH="/usr/local/cuda/lib64:/usr/local/cuda/extras/CUPTI/lib64"
build:rbe_linux_cuda_nvcc --action_env=TF_NVCC_CLANG="1"

# TODO(kanglan): Remove rbe_win and rbe_win_py3* after b/289091160 is fixed
build:rbe_win --config=rbe_base
Expand Down
Loading
Loading