Skip to content

Commit 51b2768

Browse files
olilarkinzcbenz
andauthored
Add metal::set_metallib_path() to override the metallib search path (#3597)
Co-authored-by: Cheng <git@zcbenz.com>
1 parent 602b535 commit 51b2768

4 files changed

Lines changed: 40 additions & 0 deletions

File tree

mlx/backend/metal/device.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -162,6 +162,20 @@ std::pair<MTL::Library*, NS::Error*> load_swiftpm_library(
162162
}
163163

164164
MTL::Library* load_default_library(MTL::Device* device) {
165+
// Check override path before automatic lookup
166+
if (!get_metallib_path().empty()) {
167+
auto [lib, error] =
168+
load_library_from_path(device, get_metallib_path().c_str());
169+
if (!lib) {
170+
throw std::runtime_error(
171+
fmt::format(
172+
"Can not load metallib from specified location \"{}\": {}.",
173+
get_metallib_path(),
174+
error->localizedDescription()->utf8String()));
175+
}
176+
return lib;
177+
}
178+
165179
NS::Error* error[5];
166180
MTL::Library* lib;
167181
// First try the colocated mlx.metallib

mlx/backend/metal/metal.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,12 @@
77

88
namespace mlx::core::metal {
99

10+
namespace {
11+
12+
std::string g_metallib_path;
13+
14+
} // namespace
15+
1016
bool is_available() {
1117
return true;
1218
}
@@ -46,4 +52,12 @@ void stop_capture() {
4652
manager->stopCapture();
4753
}
4854

55+
void set_metallib_path(const std::string& path) {
56+
g_metallib_path = path;
57+
}
58+
59+
const std::string& get_metallib_path() {
60+
return g_metallib_path;
61+
}
62+
4963
} // namespace mlx::core::metal

mlx/backend/metal/metal.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,4 +22,9 @@ MLX_API const
2222
std::unordered_map<std::string, std::variant<std::string, size_t>>&
2323
device_info();
2424

25+
/* Set a custom path to mlx.metallib. Must be called before any MLX operation.
26+
*/
27+
MLX_API void set_metallib_path(const std::string& path);
28+
MLX_API const std::string& get_metallib_path();
29+
2530
} // namespace mlx::core::metal

mlx/backend/metal/no_metal.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,13 @@ device_info() {
2121
"[metal::device_info] Cannot get device info without metal backend");
2222
};
2323

24+
void set_metallib_path(const std::string& path) {}
25+
26+
const std::string& get_metallib_path() {
27+
throw std::runtime_error(
28+
"[metal::get_metallib_path] Cannot get metallib path without metal backend");
29+
}
30+
2431
} // namespace metal
2532

2633
} // namespace mlx::core

0 commit comments

Comments
 (0)