PR openxla#24900: [ROCM] fix invalid memory access in rocm executor

alekstheod · Google-ML-Automation · commit ead4f4740a83 · 2025-04-09T08:09:56.000-07:00
Imported from GitHub PR openxla#24900 Fix asan memory access violation: ``` exec ${PAGER:-/usr/bin/less} "$0" || exit 1 Executing tests from //xla/service:elemental_ir_emitter_test_gpu_amd_any ----------------------------------------------------------------------------- Running test /home/atheodor/projects/tmp/xla_asan/execroot/xla/bazel-out/k8-opt/bin/xla/service/elemental_ir_emitter_test_gpu_amd_any.runfiles/xla/xla/service/elemental_ir_emitter_test_gpu_amd_any --gtest_shuffle --gtest_fail_if_no_test_linked on GPU 3 Note: Randomizing tests' orders with a seed of 19906 . [==========] Running 118 tests from 13 test suites. [----------] Global test environment set-up. [----------] 10 tests from ElementalIrEmitterExecutionTypedTest/7, where TypeParam = ml_dtypes::float8_internal::float8_e5m2 [ RUN ] ElementalIrEmitterExecutionTypedTest/7.ConvertFloatsToFloat ================================================================= ==2457579==ERROR: AddressSanitizer: use-after-poison on address 0x506000843a08 at pc 0x7f401151be6a bp 0x7ffd1e3c3410 sp 0x7ffd1e3c3408 READ of size 8 at 0x506000843a08 thread T0 #0 0x7f401151be69 in stream_executor::gpu::RocmExecutor::UnloadGpuBinary(stream_executor::ModuleHandle) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/stream_executor/rocm/rocm_executor.cc:596:23 #1 0x7f401151b036 in stream_executor::gpu::RocmExecutor::UnloadModule(stream_executor::ModuleHandle) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/stream_executor/rocm/rocm_executor.cc:496:10 #2 0x7f405dee713b in stream_executor::ScopedModuleHandle::~ScopedModuleHandle() /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/stream_executor/scoped_module_handle.h:48:7 #3 0x7f405dee713b in std::pair<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>::~pair() /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/stl_iterator.h:2488:12 #4 0x7f405dee713b in void __gnu_cxx::new_allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>::destroy<std::pair<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>>(std::pair<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>*) /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/ext/new_allocator.h:168:10 openxla#5 0x7f405dee713b in void std::allocator_traits<std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>>::destroy<std::pair<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>>(std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>&, std::pair<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>*) /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/alloc_traits.h:535:8 openxla#6 0x7f405dee713b in void absl::lts_20230802::container_internal::map_slot_policy<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>::destroy<std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>>(std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>*, absl::lts_20230802::container_internal::map_slot_type<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/container_memory.h:419:7 openxla#7 0x7f405dee713b in void absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>::destroy<std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>>(std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>*, absl::lts_20230802::container_internal::map_slot_type<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/flat_hash_map.h:578:5 openxla#8 0x7f405dee713b in void absl::lts_20230802::container_internal::common_policy_traits<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>, void>::destroy<std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>>(std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>*, absl::lts_20230802::container_internal::map_slot_type<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/common_policy_traits.h:50:5 openxla#9 0x7f405dee713b in absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>, absl::lts_20230802::container_internal::HashEq<stream_executor::StreamExecutor*, void>::Hash, absl::lts_20230802::container_internal::HashEq<stream_executor::StreamExecutor*, void>::Eq, std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>>::destroy_slots() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_set.h:1946:9 openxla#10 0x7f405dee713b in absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::StreamExecutor*, stream_executor::ScopedModuleHandle>, absl::lts_20230802::container_internal::HashEq<stream_executor::StreamExecutor*, void>::Hash, absl::lts_20230802::container_internal::HashEq<stream_executor::StreamExecutor*, void>::Eq, std::allocator<std::pair<stream_executor::StreamExecutor* const, stream_executor::ScopedModuleHandle>>>::~raw_hash_set() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_set.h:1885:5 #11 0x7f405dee8580 in xla::gpu::GpuExecutable::~GpuExecutable() /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/gpu/gpu_executable.cc:155:1 #12 0x7f405dee8d4d in xla::gpu::GpuExecutable::~GpuExecutable() /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/gpu/gpu_executable.cc:151:33 openxla#13 0x7f407b818b3f in std::default_delete<xla::Executable>::operator()(xla::Executable*) const /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/unique_ptr.h:85:2 #14 0x7f407b818b3f in std::unique_ptr<xla::Executable, std::default_delete<xla::Executable>>::~unique_ptr() /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/unique_ptr.h:361:4 openxla#15 0x7f407b818b3f in xla::(anonymous namespace)::HloRunnerExecutable::~HloRunnerExecutable() /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:69:7 openxla#16 0x7f407b818b3f in xla::(anonymous namespace)::HloRunnerExecutable::~HloRunnerExecutable() /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:69:7 openxla#17 0x7f407b7e6503 in std::default_delete<xla::OpaqueExecutable>::operator()(xla::OpaqueExecutable*) const /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/unique_ptr.h:85:2 openxla#18 0x7f407b7e6503 in std::unique_ptr<xla::OpaqueExecutable, std::default_delete<xla::OpaqueExecutable>>::~unique_ptr() /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/bits/unique_ptr.h:361:4 openxla#19 0x7f407b7e6503 in xla::HloRunner::ExecuteWithMovedDeviceBuffersAndBufferAssignment(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, xla::BufferAssignmentProto const*, std::vector<xla::ScopedShapedBuffer, std::allocator<xla::ScopedShapedBuffer>>, bool, xla::ExecutionProfile*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:400:1 openxla#20 0x7f407b7e57c3 in xla::HloRunner::Execute(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal const* const>, bool, xla::ExecutionProfile*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:221:3 #21 0x55b8a3cb4622 in xla::HloRunnerInterface::Execute(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal const* const>, bool) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/service/hlo_runner_interface.h:244:12 openxla#22 0x55b8a3cb4622 in xla::HloRunnerAgnosticReferenceMixin<xla::HloRunnerAgnosticTestBase>::RunAndCompareInternal(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal* const>, std::optional<xla::ErrorSpec> const&, bool, std::function<void (xla::HloModule*)> const&, std::function<void (xla::HloModule*)> const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/tests/hlo_runner_agnostic_reference_mixin.h:238:5 openxla#23 0x55b8a3cbf766 in xla::HloRunnerAgnosticReferenceMixin<xla::HloRunnerAgnosticTestBase>::RunAndCompare(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal* const>, std::optional<xla::ErrorSpec> const&, std::function<void (xla::HloModule*)> const&, std::function<void (xla::HloModule*)> const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/tests/hlo_runner_agnostic_reference_mixin.h:94:9 openxla#24 0x55b8a3cbf235 in xla::HloRunnerAgnosticReferenceMixin<xla::HloRunnerAgnosticTestBase>::RunAndCompare(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, std::optional<xla::ErrorSpec> const&, std::function<void (xla::HloModule*)> const&, std::function<void (xla::HloModule*)> const&, std::optional<long>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/tests/hlo_runner_agnostic_reference_mixin.h:140:12 openxla#25 0x55b8a3cceda8 in xla::(anonymous namespace)::ElementalIrEmitterExecutionTest::RunTypeConversionTest(std::basic_string_view<char, std::char_traits<char>>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/elemental_ir_emitter_test.cc:76:5 openxla#26 0x55b8a3cd8cf3 in xla::(anonymous namespace)::ElementalIrEmitterExecutionTypedTest_ConvertFloatsToFloat_Test<ml_dtypes::float8_internal::float8_e5m2>::TestBody() /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/elemental_ir_emitter_test.cc:472:36 openxla#27 0x7f407b2f09dd in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2664:10 openxla#28 0x7f407b2f09dd in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2700:14 openxla#29 0x7f407b2f0708 in testing::Test::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2739:5 openxla#30 0x7f407b2f371b in testing::TestInfo::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2885:11 openxla#31 0x7f407b2f65ab in testing::TestSuite::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:3063:30 openxla#32 0x7f407b322eba in testing::internal::UnitTestImpl::RunAllTests() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:6054:44 openxla#33 0x7f407b32179d in bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2664:10 openxla#34 0x7f407b32179d in bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2700:14 openxla#35 0x7f407b321203 in testing::UnitTest::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:5594:10 openxla#36 0x7f407b3f59b8 in RUN_ALL_TESTS() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/include/gtest/gtest.h:2334:73 openxla#37 0x7f407b3f59b8 in main /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/tests/xla_internal_test_main.cc:65:10 openxla#38 0x7f4004766d8f in __libc_start_call_main csu/../sysdeps/nptl/libc_start_call_main.h:58:16 openxla#39 0x7f4004766e3f in __libc_start_main csu/../csu/libc-start.c:392:3 openxla#40 0x55b8a3b9be44 in _start (/home/atheodor/projects/tmp/xla_asan/execroot/xla/bazel-out/k8-opt/bin/xla/service/elemental_ir_emitter_test_gpu_amd_any+0x10ce44) (BuildId: 1c37d17e488373aad7bf33204cb4234e) 0x506000843a08 is located 40 bytes inside of 56-byte region [0x5060008439e0,0x506000843a18) allocated by thread T0 here: #0 0x55b8a3c3607f in malloc (/home/atheodor/projects/tmp/xla_asan/execroot/xla/bazel-out/k8-opt/bin/xla/service/elemental_ir_emitter_test_gpu_amd_any+0x1a707f) (BuildId: 1c37d17e488373aad7bf33204cb4234e) #1 0x7f4004a4a98b in operator new(unsigned long) (/lib/x86_64-linux-gnu/libstdc++.so.6+0xae98b) (BuildId: e37fe1a879783838de78cbc8c80621fa685d58a2) #2 0x7f40115449aa in absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::initialize_slots() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_set.h:2505:5 #3 0x7f40115449aa in absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::resize(unsigned long) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_set.h:2515:5 #4 0x7f40115443fa in absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::prepare_insert(unsigned long) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_set.h:2672:7 openxla#5 0x7f40115442df in std::pair<unsigned long, bool> absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::find_or_prepare_insert<stream_executor::ModuleHandle>(stream_executor::ModuleHandle const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_set.h:2659:13 openxla#6 0x7f4011524701 in std::pair<absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::iterator, bool> absl::lts_20230802::container_internal::raw_hash_map<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::try_emplace_impl<stream_executor::ModuleHandle const&>(stream_executor::ModuleHandle const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_map.h:202:22 openxla#7 0x7f4011524701 in std::pair<absl::lts_20230802::container_internal::raw_hash_set<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::iterator, bool> absl::lts_20230802::container_internal::raw_hash_map<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::try_emplace<stream_executor::ModuleHandle, 0>(stream_executor::ModuleHandle const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_map.h:139:12 openxla#8 0x7f4011524701 in decltype(absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>::value(std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>* std::addressof<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>(std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>&)(decltype(__declval<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>(0)) std::declval<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>&>()()))) absl::lts_20230802::container_internal::raw_hash_map<absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>, absl::lts_20230802::hash_internal::Hash<stream_executor::ModuleHandle>, std::equal_to<stream_executor::ModuleHandle>, std::allocator<std::pair<stream_executor::ModuleHandle const, std::pair<ihipModule_t*, unsigned long>>>>::operator[]<stream_executor::ModuleHandle, absl::lts_20230802::container_internal::FlatHashMapPolicy<stream_executor::ModuleHandle, std::pair<ihipModule_t*, unsigned long>>>(stream_executor::ModuleHandle const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_absl/absl/container/internal/raw_hash_map.h:184:28 openxla#9 0x7f4011524701 in stream_executor::gpu::RocmExecutor::LoadModuleFromHsaco(char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/stream_executor/rocm/rocm_executor.cc:717:39 openxla#10 0x7f4011524387 in stream_executor::gpu::RocmExecutor::LoadModule(stream_executor::MultiModuleLoaderSpec const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/stream_executor/rocm/rocm_executor.cc:705:12 #11 0x7f405deeae34 in xla::gpu::GpuExecutable::ResolveConstantGlobals(stream_executor::Stream*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/gpu/gpu_executable.cc:499:5 #12 0x7f405def050a in xla::gpu::GpuExecutable::ExecuteAsyncOnStreamImpl(xla::ServiceExecutableRunOptions const*, std::variant<absl::lts_20230802::Span<xla::ShapedBuffer const* const>, absl::lts_20230802::Span<xla::ExecutionInput>>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/gpu/gpu_executable.cc:703:5 openxla#13 0x7f405deefc6f in xla::gpu::GpuExecutable::ExecuteAsyncOnStream(xla::ServiceExecutableRunOptions const*, std::vector<xla::ExecutionInput, std::allocator<xla::ExecutionInput>>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/gpu/gpu_executable.cc:661:10 #14 0x7f401607a78e in xla::Executable::ExecuteAsyncOnStreamWrapper(xla::ServiceExecutableRunOptions const*, std::vector<xla::ExecutionInput, std::allocator<xla::ExecutionInput>>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/executable.cc:229:7 openxla#15 0x7f4016079fd3 in xla::Executable::ExecuteOnStreamWrapper(xla::ServiceExecutableRunOptions const*, std::vector<xla::ExecutionInput, std::allocator<xla::ExecutionInput>>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/executable.cc:153:7 openxla#16 0x7f407b7ea78b in xla::HloRunner::ExecuteWithExecutionInputs(xla::Executable*, std::vector<xla::ExecutionInput, std::allocator<xla::ExecutionInput>>, xla::ExecutionProfile*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:448:3 openxla#17 0x7f407b7ecde2 in xla::HloRunner::ExecuteWithMovedDeviceBuffers(xla::Executable*, std::vector<xla::ScopedShapedBuffer, std::allocator<xla::ScopedShapedBuffer>>, xla::ExecutionProfile*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:415:3 openxla#18 0x7f407b7e642a in xla::HloRunner::ExecuteWithMovedDeviceBuffersAndBufferAssignment(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, xla::BufferAssignmentProto const*, std::vector<xla::ScopedShapedBuffer, std::allocator<xla::ScopedShapedBuffer>>, bool, xla::ExecutionProfile*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:398:10 openxla#19 0x7f407b7e57c3 in xla::HloRunner::Execute(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal const* const>, bool, xla::ExecutionProfile*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/hlo_runner.cc:221:3 openxla#20 0x55b8a3cb4622 in xla::HloRunnerInterface::Execute(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal const* const>, bool) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/service/hlo_runner_interface.h:244:12 #21 0x55b8a3cb4622 in xla::HloRunnerAgnosticReferenceMixin<xla::HloRunnerAgnosticTestBase>::RunAndCompareInternal(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal* const>, std::optional<xla::ErrorSpec> const&, bool, std::function<void (xla::HloModule*)> const&, std::function<void (xla::HloModule*)> const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/tests/hlo_runner_agnostic_reference_mixin.h:238:5 openxla#22 0x55b8a3cbf766 in xla::HloRunnerAgnosticReferenceMixin<xla::HloRunnerAgnosticTestBase>::RunAndCompare(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, absl::lts_20230802::Span<xla::Literal* const>, std::optional<xla::ErrorSpec> const&, std::function<void (xla::HloModule*)> const&, std::function<void (xla::HloModule*)> const&) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/tests/hlo_runner_agnostic_reference_mixin.h:94:9 openxla#23 0x55b8a3cbf235 in xla::HloRunnerAgnosticReferenceMixin<xla::HloRunnerAgnosticTestBase>::RunAndCompare(std::unique_ptr<xla::HloModule, std::default_delete<xla::HloModule>>, std::optional<xla::ErrorSpec> const&, std::function<void (xla::HloModule*)> const&, std::function<void (xla::HloModule*)> const&, std::optional<long>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/./xla/tests/hlo_runner_agnostic_reference_mixin.h:140:12 openxla#24 0x55b8a3cceda8 in xla::(anonymous namespace)::ElementalIrEmitterExecutionTest::RunTypeConversionTest(std::basic_string_view<char, std::char_traits<char>>) /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/elemental_ir_emitter_test.cc:76:5 openxla#25 0x55b8a3cd8cf3 in xla::(anonymous namespace)::ElementalIrEmitterExecutionTypedTest_ConvertFloatsToFloat_Test<ml_dtypes::float8_internal::float8_e5m2>::TestBody() /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/service/elemental_ir_emitter_test.cc:472:36 openxla#26 0x7f407b2f09dd in void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2664:10 openxla#27 0x7f407b2f09dd in void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2700:14 openxla#28 0x7f407b2f0708 in testing::Test::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2739:5 openxla#29 0x7f407b2f371b in testing::TestInfo::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2885:11 openxla#30 0x7f407b2f65ab in testing::TestSuite::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:3063:30 openxla#31 0x7f407b322eba in testing::internal::UnitTestImpl::RunAllTests() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:6054:44 openxla#32 0x7f407b32179d in bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2664:10 openxla#33 0x7f407b32179d in bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:2700:14 openxla#34 0x7f407b321203 in testing::UnitTest::Run() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/src/gtest.cc:5594:10 openxla#35 0x7f407b3f59b8 in RUN_ALL_TESTS() /home/atheodor/projects/tmp/xla_asan/execroot/xla/external/com_google_googletest/googletest/include/gtest/gtest.h:2334:73 openxla#36 0x7f407b3f59b8 in main /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/tests/xla_internal_test_main.cc:65:10 openxla#37 0x7f4004766d8f in __libc_start_call_main csu/../sysdeps/nptl/libc_start_call_main.h:58:16 SUMMARY: AddressSanitizer: use-after-poison /home/atheodor/projects/tmp/xla_asan/execroot/xla/xla/stream_executor/rocm/rocm_executor.cc:596:23 in stream_executor::gpu::RocmExecutor::UnloadGpuBinary(stream_executor::ModuleHandle) Shadow bytes around the buggy address: 0x506000843780: fa fa fa fa fd fd fd fd fd fd fd fa fa fa fa fa 0x506000843800: fd fd fd fd fd fd fd fa fa fa fa fa fd fd fd fd 0x506000843880: fd fd fd fa fa fa fa fa fd fd fd fd fd fd fd fd 0x506000843900: fa fa fa fa 00 00 00 00 00 00 00 fa fa fa fa fa 0x506000843980: fd fd fd fd fd fd fd fa fa fa fa fa 00 00 00 00 =>0x506000843a00: f7[f7]f7 fa fa fa fa fa 00 00 00 00 00 00 00 00 0x506000843a80: fa fa fa fa 00 00 00 00 00 00 00 fa fa fa fa fa 0x506000843b00: 00 00 00 00 00 00 00 fa fa fa fa fa 00 00 00 00 0x506000843b80: 00 00 00 fa fa fa fa fa 00 00 00 00 00 00 00 fa 0x506000843c00: fa fa fa fa 00 00 00 00 00 00 00 fa fa fa fa fa 0x506000843c80: 00 00 00 00 00 00 00 fa fa fa fa fa fd fd fd fd Shadow byte legend (one shadow byte represents 8 application bytes): Addressable: 00 Partially addressable: 01 02 03 04 05 06 07 Heap left redzone: fa Freed heap region: fd Stack left redzone: f1 Stack mid redzone: f2 Stack right redzone: f3 Stack after return: f5 Stack use after scope: f8 Global redzone: f9 Global init order: f6 Poisoned by user: f7 Container overflow: fc Array cookie: ac Intra object redzone: bb ASan internal: fe Left alloca redzone: ca Right alloca redzone: cb ==2457579==ABORTING ``` Copybara import of the project: -- 8f74d4c by alekstheod <atheodor@amd.com>: Fix asan report memory access vialation in rocm_executor Merging this change closes openxla#24900 COPYBARA_INTEGRATE_REVIEW=openxla#24900 from ROCm:ci_fix_invalid_memory_access_in_rocm_executor 8f74d4c PiperOrigin-RevId: 745592235
diff --git a/xla/stream_executor/rocm/rocm_executor.cc b/xla/stream_executor/rocm/rocm_executor.cc
@@ -584,7 +584,7 @@ bool RocmExecutor::UnloadGpuBinary(ModuleHandle module_handle) {
     VLOG(3) << "No loaded  HSACO module for " << module_handle;
     return false;
   }
-  auto& module = module_it->second.first;
+  auto module = module_it->second.first;
   auto& refcount = module_it->second.second;
   VLOG(3) << "Found HSACO module " << module << " with refcount " << refcount;
   if (--refcount == 0) {

Original file line number	Diff line number	Diff line change
`@@ -584,7 +584,7 @@ bool RocmExecutor::UnloadGpuBinary(ModuleHandle module_handle) {`
`584`	`584`	`VLOG(3) << "No loaded HSACO module for " << module_handle;`
`585`	`585`	`return false;`
`586`	`586`	`}`
`587`		`- auto& module = module_it->second.first;`
	`587`	`+ auto module = module_it->second.first;`
`588`	`588`	`auto& refcount = module_it->second.second;`
`589`	`589`	`VLOG(3) << "Found HSACO module " << module << " with refcount " << refcount;`
`590`	`590`	`if (--refcount == 0) {`