Exclude //xla/tests:collective_ops_test_amdgpu_any#468
Merged
alekstheod merged 1 commit intoDec 5, 2025
Conversation
hsharsha
approved these changes
Dec 5, 2025
cj401-amd
pushed a commit
that referenced
this pull request
Dec 5, 2025
hsharsha
added a commit
that referenced
this pull request
Jan 16, 2026
* Fix AMD GPU alloca address space errors (#433) * Fix AMD GPU alloca address space errors AMD GPUs require stack allocations (alloca instructions) to be in address space 5 (private/local memory), not address space 0 (generic memory) * Optimize AMDGPU allocas by keeping AS5 pointers throughout * Fix AMDGPU allocas to use address space 5 in MLIR lowering * [AIXLA-171] Jaxlib 0.8.0 requried backports (#437) * Rocm jaxlib v0.5.0 warpsize global (#177) * cherry-picked warp size passing to triton calls, and globally enabled warpsize=64 * Fix. --------- Co-authored-by: Pavel Emeliyanenko <[email protected]> (cherry picked from commit f013645) * Added support for waves_per_eu function attribute. (#181) (cherry picked from commit bc1d816) * Add MIOPEN_FIND_ENFORCE For ROCm 7 for convolution gemms (#312) * Add MIOPEN_FIND_ENFORCE For ROCm 7 for convolution gemms * Exclude failing CollectiveOpsE2E tests * Make device_count_ atomic (#343) * Make device_count_ atomic * Use relaxed memory order * Fix build error * added rocm7 support to EnablePeerAccess (#347) * added rocm7 support to EnablePeerAccess * use wrap namespace, clang-format and add comments * [ROCm] Disable Cudnn fusions (#358) * Get CI green * Triton fixes porting from v0.6.0 (#389) * rocprof-sdk addition, upstream PR: openxla/pull/29769 Squash following commits.. Update rocprofiler-sdk (v3) along with roctracer (v1) for rocm-jaxlib-v0.6.0 (#302) * update for integration of rocprofiler-sdk (along with roctracer as a backup based on bazel_options from CLI) (cherry picked from commit 7775dd0) use VLOG(2) to replace LOG(INFO), so PGLE has no verbose info (#357) (cherry picked from commit 5950125) update with kernel details for rocm-7.x (#364) * update with kernel details for rocm-7.x (cherry picked from commit 5597c0d) update to remove previously hard-coded rocprofiler-sdk path (#369) * update to remove previously hard-coded rocprofiler-sdk path and add skip_rocprofiler_sdk to avoid loading `rocprofiler-sdk` (cherry picked from commit ff74b5f) * fixed buffer comparator test * misc fixes ported from rocm-jaxlib-v0.6.0 --------- Co-authored-by: Pavel Emeliyanenko <[email protected]> (cherry picked from commit f013645) (cherry picked from commit b03cd94) Added support for waves_per_eu function attribute. (#181) (cherry picked from commit bc1d816) (cherry picked from commit d3f94e9) removed two line change (revert of half of the openxla#25959 commit (cherry picked from commit 109e138) Fixes for jax 0.6.0 (#207) * Add fixes for jax plugin 0.6.0 Drop NEEDED linking to unnecessary libs. These are loaded by amdhipruntime and not us. Fix missing NEEDED on MIOpen shared object. * Minor rocblas related changes for rocm 70 (cherry picked from commit 0de7d49) --------- Co-authored-by: Zoran Jovanovic <[email protected]> (cherry picked from commit 28f10a0) Add hipBLASLt support for gfx11. (#301) (cherry picked from commit f814bff) Add bf16 starting from gfx11, bugfix & optimize RocmComputeCapability (#303) * Bugfix and improve device_description.h::RocmComputeCompatibility * Enable ALG_DOT_BF16* on rocm with HW support (cherry picked from commit 510ea06) [ROCm] Use bundled bitcode files (#196) Also trim bitcode file list to ockl.bc and ocml.bc only. (cherry picked from commit fc9e3c3) Add MIOPEN_FIND_ENFORCE For ROCm 7 for convolution gemms (#312) * Add MIOPEN_FIND_ENFORCE For ROCm 7 for convolution gemms * Exclude failing CollectiveOpsE2E tests (cherry picked from commit fb6ddfb) Restore RocmComputeCapability:: gfx11_rx7900() and gfx12_rx8900() methods (#333) At least gfx11_rx7900() is still needed for TF build. (cherry picked from commit 13c3de1) Make device_count_ atomic (#343) * Make device_count_ atomic * Use relaxed memory order * Fix build error (cherry picked from commit 8513f2d) fix hardcoded max registers (#345) (cherry picked from commit f3e170a) fix hardcoded ecc enabled (#348) (cherry picked from commit 9cfa74a) remove reserved memory (#349) (cherry picked from commit 0015d0e) Add rocm_dev config for remote caching (#353) (cherry picked from commit c815420) added rocm7 support to EnablePeerAccess (#347) * added rocm7 support to EnablePeerAccess * use wrap namespace, clang-format and add comments (cherry picked from commit 85548a7) [ROCm] Disable Cudnn fusions (#358) (cherry picked from commit edab8b2) * Ported all triton related changes from v0.6.0 to v0.7.1 (cherry picked from commit 1851bcc) Disable softmax triton fusion if triton gemm is off (#281) * Disable softmax rewriter triton if triton gemm is disabled * Add specific flag to enable triton softmax fusion * Address review comments (cherry picked from commit 51a7f4b) [ROCm][Triton] Disable transposed load in certain conditions (cherry picked from commit 50860e9) Enable unit tests that pass after fixing some Triton related issues. (#285) * Enable unit tests that pass after fixing some Triton related issues. * fusion_emitter_device_legacy_test still fails on MI200 (cherry picked from commit 97dd565) Rocm jaxlib v0.6.0 triton support ut (#279) * Fixed triton/support_test - no fmfa. * Fix issue with rounding mode in accelerate amd matmul. * Fixed issues with usage of mfma in support_test. (cherry picked from commit 44f7d87) Restore gpu_triton_custom_call_test (#262) (cherry picked from commit 32eafa4) Skipped CanNotEmitTritonCustomCallOnPreAmpereGpu test for ROCM. (cherry picked from commit 56ec7ec) (cherry picked from commit b1f3e9f) fixed createTritonAMDGPULowerInstructionSchedHintsPass (#179) (cherry picked from commit 8517a3a) (cherry picked from commit c62e47d) fixed bazel build issue --------- Co-authored-by: Chunyu Jin <[email protected]> Co-authored-by: zoranjovanovic-ns <[email protected]> Co-authored-by: Alex <[email protected]> * Fix asan error reported when running set_dimension_size_test under rocm * PR openxla#33909: [ROCm] Make multi gpu tests exclusive if executed locally Imported from GitHub PR openxla#33909 📝 Summary of Changes Make multigpu tests being executed exclusively. 🎯 Justification Multigpu tests have to run exclusively to not clash the utilization of gpus in parallel. This PR marks multigpu tests as such by adding 'multi_gpu' tag, and in case of rocm make them run exclusively if executed locally. 🚀 Kind of Contribution Please remove what does not apply: ♻️ Cleanup 📊 Benchmark (for Performance Improvements) Not relevant, no logic change 🧪 Unit Tests: Not relevant, no logic change 🧪 Execution Tests: Not relevant, no logic change Copybara import of the project: -- 7e84a78 by Alexandros Theodoridis <[email protected]>: Mark multigpu tests as such, make them exclusive if executed locally -- 23fc24c by Alexandros Theodoridis <[email protected]>: Remove invalid multi_gpu tag from cpu tests -- a2a35de by Alexandros Theodoridis <[email protected]>: Restore build_config not exposing has_tag -- babe5c7 by Alexandros Theodoridis <[email protected]>: Unmark test which is not multigpu -- 3e30cf6 by Alexandros Theodoridis <[email protected]>: Add missing tag documentation Merging this change closes openxla#33909 COPYBARA_INTEGRATE_REVIEW=openxla#33909 from ROCm:ci_make_multi_gpu_tests_execute_sequentially_when_locally_but_not_in_rbe 3e30cf6 PiperOrigin-RevId: 836288457 * Cleanup nccl communicator shared memory files * Add missing parallel_gpu_execute target * Fix ci build command, push args to the end * Use upstream version of ci run script * Add cuda-only tag to cuda_test target * Add noasan notsan filters * Fix buildifier error --------- Co-authored-by: zoranjovanovic-ns <[email protected]> Co-authored-by: Pavel Emeliyanenko <[email protected]> Co-authored-by: Alex <[email protected]> Co-authored-by: spiao <[email protected]> Co-authored-by: Dragan Mladjenovic <[email protected]> Co-authored-by: Harsha HS <[email protected]> Co-authored-by: Zahid Iqbal <[email protected]> Co-authored-by: Chunyu Jin <[email protected]> Co-authored-by: Alexandros Theodoridis <[email protected]> * Backport fixes from upstream (#443) * Consider runfiles in lit tests * Address review comments * Introduce pool name for rbe * Introduce rocm rbe pools * First check for multigpu tag * Address review comments * Fix buildifier issue * Assigne specific pool only to test targets * Make rbe amdgpu pools configurable * Trigger CI/CD pipeline * Fix build error * Add excluded tests * Trigger CI/CD pipeline * Switch back to rocminfo to detect number of gpus * Backport rocm_configure * Trigger CI/CD pipeline * Add ignore tests list * Trigger CI/CD pipeline * Add disk cache * Enable rbe for this branch only * Add one more failing test * Update tsan ignore list * Remove printing out the script commands * Update local test jobs * Trigger CI/CD pipeline * Trigger CI/CD pipeline * Disable rbe * Add failing test * Enable rbe * Ignore TransferLiteralFromDevice * Ignore transfer from outfeed * Fix hermetic builds * Print bazel command * Add back communicator test * Remove require amd gpu filter * Remove filter nccl_communicator_test * Remove no_oss filter * Ignore //xla/pjrt/gpu/tfrt:tfrt_gpu_buffer_test * Exclude failing tsts * Mark multigpu tests as such, make them exclusive if executed locally * Remove invalid multi_gpu tag from cpu tests * Restore build_config not exposing has_tag * Unmark test which is not multigpu * Add missing tag documentation * Run multigpu tests locally * Revert tests to fix, use taget * Remove disk_cache * Return disk cache and ignore tsan warning * Fix tsan issues (#450) * Fix tsan issues * Fix hermetic build llvm clash * Set timeouts, ignore tsan warning * Ignore execute on stream * Port PackedTranspose performance regression fix (#441) * [XLA:GPU] Rename warp to shmem_group in PackedTranspose Also calculate their count as kNumThreadsPerBlock / kNumShmemBanks to avoid inconsistency when manually specified. This change is NFC for any GPU in upstream. However, it fixes a performance regression in downstream for AMD GPUs caused by inconsistency between shmem_group size, kNumThreadsPerBlock and kNumShmemBanks. It ended up in a situation downstream where half of the launched threads per block were not utilized at all. Update packed transpose tests to verify correct thread utilization. * [XLA:GPU] Fuse shmem write loops for transposes in PackedTranspose Replace per-transpose loops with a single unified loop that processes all transposes simultaneously, computing indices once and reusing them across all operations. Update packed_transpose_multiple_heroes.hlo test to verify the single-loop structure with multiple iter_args. * Support split for multigpu and single gpu pipelines (#453) * Support split for multigpu and single gpu pipelines * Use placeholder config * Include all multigpu tests * Add back multigpu tests * Ignore ExecuteReplicatedImpl * Migrate settings from run_multi_gpu tests to bazelrc * Include sanitizer wrapper * Exclude failing test * Migrate more settings and enable collective tests * add missing cuda-only tag * Add back failing test * Mark nvshmem tests as cuda-only (#458) * bugfix - consider the situation where the best time is infinite (#440) * bugfix - consider the situation where the best time is infinite * added an unit test where best_tiled_run_time_data contains Infinite. * Enable rbe for testing (#463) * Enable rbe for testing * Force multigpu tests to run locally * Reduce disk_cache size (#467) * Exclude //xla/tests:collective_ops_test_amdgpu_any (#468) * Add register spilling detection AMD v0.8.0 (#464) * register spilling by disassembling object file * added time measurement to the spilling check * adapt the num_warps so that the hlo could be compiled on both amd and nvidia * pass though is_autotuning_compilation flag to the function CompileToHsaco * implementation of register spilling by reading meta data of hasco file using llvm-readobj * utilize amd code object manager library for parsing HSACO metadata * Fix infinite recursion in HloInstruction::Accept/Visit const wrappers (#470) The const wrapper methods for Accept() and Visit() were calling themselves instead of the template versions, causing infinite recursion and stack overflow. * enable mx datatype for rocm (#462) * enable mx datatype for rocm * add // TF_ROCM_VERSION >= 70000 * fix unit test build * Rocm jaxlib v0.8.0 fix fp support ut 2 (#466) * Fix expected output in fusion_emitter_int4_device_test for ROCm. * Enable triton/fusion_emitter_int4_device_test on ci. * Rocm jaxlib v0.8.0 fix triton support (#465) * Fixed triton support_test on 0.7.1. (cherry picked from commit be9da6d) * Fixed triton/support_test on ROCm 0.8.0. Updated number of IsTritonSupport functions. * Remove invalid ignore targets (#475) * revert xla_gpu_enable_triton_softmax_fusion related changes (#474) CI status is irrelevant to this PR * update rocprofiler-sdk (v3) and roctracer (v1) (#473) * update rocprofiler-sdk (v3) and roctracer (v1) for 0.8.0 * update for skip_rocprofiler_sdk * update for the fixing asan check * fix to show kernel details in trace file * Refine rocprofiler-sdk (v3) integration based on Arech8's comments along with rocprofiler_flush_buffer explicitly * update rocprofiler-sdk based on comments * Force rbe incompatible tests to be executed locally (#485) * [ROCm] Restore CudnnFusedConvRewriter (#372) (#392) (#481) * [ROCm] Restore CudnnFusedConvRewriter (#372) (#392) * [ROCm] Allways run GpuConvAlgorithmPicker * [ROCm] Restore CudnnFusedConvRewriter Introduce CudnnFusedConvDecomposer to revert back fused convs if no fused algorithm could be found with ConvAlgorithmPicker. Remove unfused fallback paths from RocmFusedConvRunner. Co-authored-by: Dragan Mladjenovic <[email protected]> * Enable u-tests * Fix * Remove gemm algorithm picker * Backport gemmalgorithm picker * Fix --------- Co-authored-by: Harsha H S <[email protected]> Co-authored-by: Dragan Mladjenovic <[email protected]> * [XLA:GPU] Consistently check which bitcasts we can fuse. (#486) github.com/openxla/pull/30864 added a check that we shouldn't fuse bitcasts that change bit width, but only in one place. We need to do that consistently, because there are multiple places in PriorityFusion that have special handling for bitcasts. PiperOrigin-RevId: 833281452 (cherry picked from commit 8c64c14) Co-authored-by: Oleg Shyshkov <[email protected]> * take silu epilogues into account in autotuning (#488) * Rocm jaxlib v0.8.0 fix fusion emitter device test 2 (#482) * Clean up disk cache using bash trap on exit (#480) * Increase the step for iota test (#487) * Make cublaslt_test platform independent (#489) * Exclude iota_test for tsan builds (#490) * Fix hipSolver FFI errors (#483) * Fix hipSolver FFI errors * Remove srcs and add linkopts instead * Add -L for different ROCm configs * Add similar fixes for rocprofiler-sdk * Revert moving rocprofiler sdk to data --------- Co-authored-by: Alexandros Theodoridis <[email protected]> * Introduce jax utest script (#498) * Make it compatible with rocm plugin * Introduce jax utest script * Sync filter for jax tests * [ROCm] Initialze collectives to nullptr to force its allocation later (#502) * Add SiLU epilogue unit tests (#491) * [ROCm] Enable embeded bitcode libs and inprocess lld (#507) Added TF_ROCM_INPROCESS_LLD and TF_ROCM_EMBEDDED_DEVICE_LIB form 0.6.0 otherwise identical to openxla#32439. Env vars only needed for 0.8.0. * Fixing self_adj_test and fixing sort_rewriter on ROCM (#493) * fixing self_adj_test and fixing sort_rewriter on ROCM * fixing sortrewriter pass * fixing gpu_compiler_test after adding 2nd SortRewriter pass * Multioutput fusion test fix (#496) * fixing warp size for multioutput fusion test * added build fix * Fix TopK algorithm for RDNA architectures by properly handling WAVEFRONT_SIZE and removing __AMDGCN_WAVEFRONT_SIZE that got deprecated in newer versions of ROCm. Fixes sort_ops_test_gpu and sort_ops_test_gpu_mlir_bridge_test UT failures. (#509) * [rocm-jaxlib-v0.8.0] Addressing CI warnings (#508) * Remove unused variable * Add fallthrough attribute in the IsSupportedDotAlgorithm * Specify all cases in the CudnnFusedDecomposer * Specify mx_mode in GemmConfig * Remove unused cleanup_on_error * Specify missing mx_mode in GemmConfig::For (#516) * PR openxla#36046: [ROCm] Fix failing unit tests on ROCm platform Imported from GitHub PR openxla#36046 📝 Summary of Changes - layout_assignment tests are marked cuda-only. - sample_file_test needs higher autotuner level for MIOpen to return conv algorithm. Earlier this was coming from GetDebugOptionsForTest. - buffer_debug_log test is made gpu agnostic by using cannonical gpu name. - cublas_gemm_rewriter_test_amdgpu_any fix unit test to remove padding for ROCm as introduced in openxla#33854 - gpu_kernel_tiling_test_amdgpu_any is updated to respect higher launch dimensions now supported by hipruntime - Mark dynamic_shared_memory_test as cuda-only - Add arch specific checks for barriers to sorting.hlo 🎯 Justification Fixes failing unit tests on ROCm platform 🚀 Kind of Contribution 🐛 Bug Fix, 🧪 Tests Copybara import of the project: -- 472cd54 by Harsha HS <[email protected]>: [ROCm] Fix failing unit tests on ROCm platform - layout_assignment tests are marked cuda-only. - sample_file_test needs higher autotuner level for MIOpen to return conv algorithm. Earlier this was coming from GetDebugOptionsForTest. - buffer_debug_log test is made gpu agnostic by using cannonical gpu name. -- 3bb9422 by Harsha HS <[email protected]>: Fix tests which started to fail due to openxla#33854 -- 850d955 by Harsha HS <[email protected]>: HIP now respects highter launch dimension similar to CUDA -- b504a7e by Harsha HS <[email protected]>: Make dynamic_shared_memory_test cuda only -- 1e4e57a by Harsha HS <[email protected]>: Add arch specific checks to sorting.hlo -- ce1241c by Harsha HS <[email protected]>: Address review comments Merging this change closes openxla#36046 FUTURE_COPYBARA_INTEGRATE_REVIEW=openxla#36046 from ROCm:ci_fix_upstream_ut_20260107 ce1241c PiperOrigin-RevId: 856068530 * PR openxla#35744: [ROCm] Fix TopK algorithm for RDNA Imported from GitHub PR openxla#35744 📝 Summary of Changes xla/stream_executor/rocm/topk_kernel_rocm_common.cu.h now implements proper handling of WAVEFRONT_SIZE based on the architecture. __AMDGCN_WAVEFRONT_SIZE is removed as it is deprecated in newer versions of ROCm. 🎯 Justification Fixes sort_ops_test_gpu and sort_ops_test_gpu_mlir_bridge_test UT failures. 🚀 Kind of Contribution 🐛 Bug Fix 🧪 Unit Tests: Fixes sort_ops_test_gpu and sort_ops_test_gpu_mlir_bridge_test UT failures. **NOTE:** The changes are verified on gfx1100, gfx1201, and gfx90a. Copybara import of the project: -- 7097b9c by Aleksa Arsic <[email protected]>: Fix TopK algorithm for RDNA architectures by properly handling WAVEFRONT_SIZE and removing __AMDGCN_WAVEFRONT_SIZE that got deprecated in newer versions of ROCm. Fixes sort_ops_test_gpu and sort_ops_test_gpu_mlir_bridge_test UT failures. -- 6e1aaab by Harsha HS <[email protected]>: [ROCm] Fix TopK kernel Serialization checks -- 9fa6021 by Harsha HS <[email protected]>: [ROCm] Get Warpsize from device descriptor -- deb4932 by Harsha H S <[email protected]>: Include link to depricated macro Merging this change closes openxla#35744 COPYBARA_INTEGRATE_REVIEW=openxla#35744 from ROCm:ci_fix_topk_algorithm_for_rdna deb4932 PiperOrigin-RevId: 855689898 * PR openxla#36283: [ROCm] Skip conditional tests on ROCm as they are not supported by HIP Graphs Imported from GitHub PR openxla#36283 📝 Summary of Changes Conditionals are not supported by HIP graphs and those tests are skipped. 🚀 Kind of Contribution 🐛 Bug Fix Copybara import of the project: -- 1c607b6 by Harsha HS <[email protected]>: [ROCm] Skip conditional tests on ROCm as they are not supported by HIP Graphs Merging this change closes openxla#36283 COPYBARA_INTEGRATE_REVIEW=openxla#36283 from ROCm:ci_skip_conditionals_on_rocm_20260112 1c607b6 PiperOrigin-RevId: 855601660 * Include more passing tests --------- Co-authored-by: Pham Binh <[email protected]> Co-authored-by: zoranjovanovic-ns <[email protected]> Co-authored-by: Pavel Emeliyanenko <[email protected]> Co-authored-by: Alex <[email protected]> Co-authored-by: spiao <[email protected]> Co-authored-by: Dragan Mladjenovic <[email protected]> Co-authored-by: Harsha HS <[email protected]> Co-authored-by: Zahid Iqbal <[email protected]> Co-authored-by: Chunyu Jin <[email protected]> Co-authored-by: Alexandros Theodoridis <[email protected]> Co-authored-by: Aleksei Nurmukhametov <[email protected]> Co-authored-by: Xuefei Jiang <[email protected]> Co-authored-by: Eetu Sjöblom <[email protected]> Co-authored-by: Oleg Shyshkov <[email protected]> Co-authored-by: charleshofer <[email protected]> Co-authored-by: pemeliya <[email protected]> Co-authored-by: Aleksa Arsic <[email protected]> Co-authored-by: mmakevic-amd <[email protected]>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Exclude //xla/tests:collective_ops_test_amdgpu_any