Skip to content

fix: avoid clearing in-flight pipeline states in custom kernel cache #3434

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.

Sign up for GitHub

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

Open
Ziqiao-git wants to merge 1 commit into
base: main
Choose a base branch
from
+9 −15
Open

fix: avoid clearing in-flight pipeline states in custom kernel cache #3434

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
24 changes: 9 additions & 15 deletions mlx/backend/metal/custom_kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -361,22 +361,16 @@ void CustomKernel::eval_gpu(

auto& d = metal::device(s.device);

{
// Clear kernels from the device library cache if needed
auto& kernel_cache = cache();
if (auto it = kernel_cache.libraries.find(name_);
it != kernel_cache.libraries.end()) {
if (it->second != source_) {
auto& d = metal::device(s.device);
d.clear_library(name_);
it->second = source_;
}
} else {
kernel_cache.libraries.emplace(name_, source_);
}
}
// Use a source-dependent library key so different source variants
// (e.g. from write_signature picking different dtype qualifiers across
// calls) coexist without evicting each other. Clearing a library while
// its pipeline states are still referenced by an in-flight command
// buffer causes use-after-free.
auto source_hash = std::hash<std::string>{}(source_);
auto lib_key = name_ + "_" + std::to_string(source_hash);

auto lib = d.get_library(name_, [this] { return metal::utils() + source_; });
auto lib =
d.get_library(lib_key, [this] { return metal::utils() + source_; });
auto kernel = d.get_kernel(name_, lib);
auto& compute_encoder = metal::get_command_encoder(s);
compute_encoder.set_compute_pipeline_state(kernel);
Expand Down