[ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_…#32439
Closed
draganmladjenovic wants to merge 1 commit into
Closed
[ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_…#32439draganmladjenovic wants to merge 1 commit into
draganmladjenovic wants to merge 1 commit into
Conversation
Contributor
Author
|
@xla-rotation PTAL |
Contributor
Author
|
@ddunl Gentle ping. |
Contributor
Author
|
@ddunl Ping! |
Member
|
Apologies, missed this one. I'm not sure I'm the right person to review - reassigning! |
beckerhe
approved these changes
Oct 27, 2025
copybara-service Bot
pushed a commit
to tensorflow/tensorflow
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR openxla/xla#32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a100377d00d30dbc79e34c977b9219c54bda4b by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 FUTURE_COPYBARA_INTEGRATE_REVIEW=openxla/xla#32439 from ROCm:enable_bitcode_lld 46a100377d00d30dbc79e34c977b9219c54bda4b PiperOrigin-RevId: 824375626
copybara-service Bot
pushed a commit
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR #32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a1003 by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 FUTURE_COPYBARA_INTEGRATE_REVIEW=#32439 from ROCm:enable_bitcode_lld 46a1003 PiperOrigin-RevId: 824375626
copybara-service Bot
pushed a commit
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR #32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a1003 by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 FUTURE_COPYBARA_INTEGRATE_REVIEW=#32439 from ROCm:enable_bitcode_lld 46a1003 PiperOrigin-RevId: 824375626
copybara-service Bot
pushed a commit
to tensorflow/tensorflow
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR openxla/xla#32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a100377d00d30dbc79e34c977b9219c54bda4b by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 FUTURE_COPYBARA_INTEGRATE_REVIEW=openxla/xla#32439 from ROCm:enable_bitcode_lld 46a100377d00d30dbc79e34c977b9219c54bda4b PiperOrigin-RevId: 824375626
copybara-service Bot
pushed a commit
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR #32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a1003 by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 FUTURE_COPYBARA_INTEGRATE_REVIEW=#32439 from ROCm:enable_bitcode_lld 46a1003 PiperOrigin-RevId: 824375626
copybara-service Bot
pushed a commit
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR #32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a1003 by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 FUTURE_COPYBARA_INTEGRATE_REVIEW=#32439 from ROCm:enable_bitcode_lld 46a1003 PiperOrigin-RevId: 824375626
copybara-service Bot
pushed a commit
to tensorflow/tensorflow
that referenced
this pull request
Oct 27, 2025
…la_gpu_use_… Imported from GitHub PR openxla/xla#32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a100377d00d30dbc79e34c977b9219c54bda4b by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes #32439 PiperOrigin-RevId: 824476138
hsharsha
pushed a commit
to ROCm/tensorflow-upstream
that referenced
this pull request
Oct 30, 2025
…_lib and xla_gpu_use_… Imported from GitHub PR openxla/xla#32439 …inprocess_lld 📝 Summary of Changes Enable embedded device libs and in-process lld by default. 🎯 Justification Moves amdgpu backend to be more filesystem layout independent. 🚀 Kind of Contribution 🐛 Bug Fix 📊 Benchmark (for Performance Improvements) N\A 🧪 Unit Tests: None 🧪 Execution Tests: None Copybara import of the project: -- 46a100377d00d30dbc79e34c977b9219c54bda4b by Dragan Mladjenovic <[email protected]>: [ROCm] Fix and enable xla_gpu_use_embeded_device_lib and xla_gpu_use_inprocess_lld Merging this change closes tensorflow#32439 PiperOrigin-RevId: 824476138
draganmladjenovic
added a commit
to ROCm/xla
that referenced
this pull request
Dec 22, 2025
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.
i-chaochen
pushed a commit
to ROCm/xla
that referenced
this pull request
Dec 22, 2025
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.
hsharsha
added a commit
to ROCm/xla
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]>
nurmukhametov
pushed a commit
to ROCm/xla
that referenced
this pull request
Jan 23, 2026
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.
nurmukhametov
pushed a commit
to ROCm/xla
that referenced
this pull request
Jan 23, 2026
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.
nurmukhametov
pushed a commit
to ROCm/xla
that referenced
this pull request
Jan 23, 2026
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.
nurmukhametov
pushed a commit
to ROCm/xla
that referenced
this pull request
Jan 23, 2026
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.
nurmukhametov
pushed a commit
to ROCm/xla
that referenced
this pull request
Jan 23, 2026
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.
nurmukhametov
pushed a commit
to ROCm/xla
that referenced
this pull request
Jan 26, 2026
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.
i-chaochen
added a commit
to ROCm/xla
that referenced
this pull request
Feb 9, 2026
* [ROCm] Build infrastructure and CI scripts * 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. * Mark nvshmem tests as cuda-only (#458) * Skipped CanNotEmitTritonCustomCallOnPreAmpereGpu test for ROCM. * Make device_count_ atomic (#343) * Make device_count_ atomic * Use relaxed memory order * Fix build error * [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. * [ROCm] Pass warp size to Triton compilation pipeline * [ROCm] Add FNUZ FP8 type support in Triton * [ROCm] Temporary workaround for column reduction warp size * 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 * Fix build break in tfrt_gpu_buffer_test using absl_testing::StatusIs (#534) * Port transpose changes from v0.8.0 to v0.8.2 (#526) It should be dropped after the rebase on top of 330a305 * [ROCm] Fix failing test TritonEmitterTest/RocmWarpSizeIsSetCorrectly (#545) * [ROCm] Fix failing test TritonEmitterTest/RocmWarpSizeIsSetCorrectly Define valid tile parameters and non-zero shared memory. * Update xla/backends/gpu/codegen/triton/fusion_emitter_device_test.cc Co-authored-by: Maxime France-Pillois <[email protected]> * Update xla/backends/gpu/codegen/triton/fusion_emitter_device_test.cc Co-authored-by: Maxime France-Pillois <[email protected]> --------- Co-authored-by: Maxime France-Pillois <[email protected]> * Fix MIOpen linking for RNN kernels Add explicit linkopts to miopen cc_library target to ensure libMIOpen.so is properly linked at runtime. This fixes AttributeError: module 'jaxlib.gpu_rnn' has no attribute 'compute_rnn_workspace_reserve_space_sizes' in experimental_rnn_test in JAX. Without this change, the _rnn.so shared library fails to load MIOpen symbols properly, causing RNN test failures. * Force rbe incompatible tests to be executed locally (#485) * [ROCm] Add missing cuda-only tag * enable mx datatype for rocm (#462) * enable mx datatype for rocm * add // TF_ROCM_VERSION >= 70000 * fix unit test build * Add rocprofiler-sdk (v3) integration with roctracer fallback Integrate rocprofiler-sdk for ROCm profiling with fallback to roctracer (v1) when rocprofiler-sdk is not available. * [ROCm] Always process convolutions through MIOpen backend for decomposition Override AddConvAndGemmAutotuningPass in AMDGPUCompiler to ensure convolutions are always sent to MIOpen for processing, regardless of xla_gpu_autotune_level. This is required because MIOpen handles decomposition of unsupported fused convolutions back to regular convs, which must happen even when autotuning is disabled. Fixes cudnn_fused_conv_rewriter_autotune_disabled_test failures on ROCm. * Changed error value for SplitK test in fusion_emitter_device_legacy_port_test.cc (#538) * [ROCm] Add PJRT_Triton_Extension support (#548) This change is PJRT_Triton_Extension support for ROCm as counterpart of that for CUDA. Pallas Triton calls are lowered to HSACO directly rather than PTX on ROCm platform. * Fix expected output in fusion_emitter_int4_device_test for ROCm. * skip conditional graph tests * Fixed missing rtne in Triton to pass support_test. * [ROCm] Add rocm-only tag to triton_rocm target Fix dependency validation by tagging triton_rocm as rocm-only since it depends on the rocm-only amdgpu_backend target. * Avoid upcast of lib func operands to F32 for F16 type. * Modify fusion_emitter_large_test to work on ROCm. (#568) * Modify fusion_emitter_large_test to work on ROCm. * Fix fall-through warning in support_legacy.cc * Fixed dot_algorithms_test. Updated support_legacy and test itself. * Modified triton_fusion_numerics_verifier_test to work on ROCm. * [ROCm] Use shared AsBlasLtEpilogue in GemmWorkspaceRewriter Replace the duplicate with the shared function to fix the issue and prevent future divergence. The duplicate AsBlasLtEpilogue in gemm_workspace_rewriter.cc was missing SILU epilogue support, breaking ROCm Swish fusion tests. This duplicate was introduced in PR openxla#35132. * Sync mgpu tests with xla_mgpu config * [ROCm] Fix RocmWarpSizeIsSetCorrectly test to use new dump file naming After commit 4ce9326, Triton pass dumps use the naming pattern {module}.{kernel}.{pass_manager_name}.txt instead of *.triton-passes.log. Update the test to match the new convention. * Enable hlo_runner_main_gpu for rocm * enable hipblaslt as a default choice and disable nccl comm split to avoid hanging * Add flag to control swish activation fusion. (#577) Add flag to control swish activation fusion. * Improve test strategy for swish fusion flag (#585) Move tests to a more suitable file. * Revert "Fix infinite recursion in HloInstruction::Accept/Visit const wrappers (#470)" This reverts commit 21a2d57. * Disable hipblaslt as default choice * Execute test directly if running on system without GPU (#608) * Execute test directly if running on system without GPU * Address review comments * Address review comments * Remove non-existent test targets from ROCm CI exclusion list The following targets no longer exist in their respective BUILD files and were causing Bazel target pattern parsing failures. * Bundle librocm_smi64.so for MI200 lit tests MI200 lit tests use hlo-opt which links against ROCm libraries. When running on remote workers without ROCm installed, hlo-opt fails with: "error while loading shared libraries: librocm_smi64.so.1" The _tools_on_path rule bundles libraries into lit_lib/ by extracting them from CcInfo.linking_context.linker_inputs[].dynamic_library. However, ROCm's cc_library targets with .so files in srcs don't populate dynamic_library (unlike CUDA which uses cc_import). Add a new rocm_smi_import target using cc_import, which properly exposes the shared library via CcInfo. Use this target in lit.bzl so librocm_smi64.so.1 gets bundled into lit_lib/ and is available at runtime via hlo-opt's rpath. --------- Co-authored-by: Pham Binh <[email protected]> Co-authored-by: Alex <[email protected]> Co-authored-by: Zoran Jovanovic <[email protected]> Co-authored-by: Dragan Mladjenovic <[email protected]> Co-authored-by: Harsha H S <[email protected]> Co-authored-by: Maxime France-Pillois <[email protected]> Co-authored-by: magaonka-amd <[email protected]> Co-authored-by: Xuefei Jiang <[email protected]> Co-authored-by: cj401-amd <[email protected]> Co-authored-by: zoranjovanovic-ns <[email protected]> Co-authored-by: Jian Li <[email protected]> Co-authored-by: Chao Chen <[email protected]> Co-authored-by: Alexandros Theodoridis <[email protected]> Co-authored-by: Milica Makevic <[email protected]>
magaonka-amd
pushed a commit
to magaonka-amd/xla
that referenced
this pull request
Feb 16, 2026
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.
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
3 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.
…inprocess_lld
📝 Summary of Changes
Enable embedded device libs and in-process lld by default.
🎯 Justification
Moves amdgpu backend to be more filesystem layout independent.
🚀 Kind of Contribution
🐛 Bug Fix
📊 Benchmark (for Performance Improvements)
N\A
🧪 Unit Tests:
None
🧪 Execution Tests:
None