Compare commits
30 Commits
main
...
sign-warns
| Author | SHA1 | Date | |
|---|---|---|---|
| 24828b1b2f | |||
| 9f649b5658 | |||
| 18aa921388 | |||
| 8d13a0bc6b | |||
| ac75c87fd7 | |||
| 7107802e09 | |||
| c5913131cf | |||
| 19ab7911f6 | |||
| 4a1b1796b7 | |||
| b48d298205 | |||
| 8277e71ea9 | |||
| b0d985416a | |||
| 8d10f3ec75 | |||
| 6343622c67 | |||
| 979abf462b | |||
| 981d2fdaf0 | |||
| 5a306d3495 | |||
| 5baa361779 | |||
| 1bac0db7e3 | |||
| a1212b4e44 | |||
| 45a8b226af | |||
| 76ef1e98f3 | |||
| 63d91557e0 | |||
| 310e501e6a | |||
| cacc3ab7fd | |||
| 53525cba23 | |||
| 3d67b717a0 | |||
| 953b2f5be2 | |||
| 26f7155537 | |||
| 66fcb9fe94 |
+5
-1
@@ -20,9 +20,13 @@ project(
|
||||
LANGUAGES C CXX
|
||||
VERSION ${MLX_PROJECT_VERSION})
|
||||
|
||||
if(CMAKE_CXX_COMPILER_ID STREQUAL "AppleClang")
|
||||
add_compile_options(-Wall -Wextra)
|
||||
endif()
|
||||
|
||||
# ----------------------------- Setup -----------------------------
|
||||
set(CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake")
|
||||
set(CMAKE_CXX_STANDARD 17)
|
||||
set(CMAKE_CXX_STANDARD 20)
|
||||
set(CMAKE_CXX_STANDARD_REQUIRED ON)
|
||||
set(CMAKE_POSITION_INDEPENDENT_CODE ON)
|
||||
set(CMAKE_INSTALL_MESSAGE NEVER)
|
||||
|
||||
@@ -14,14 +14,17 @@ void array_basics() {
|
||||
// Get the value out of it:
|
||||
auto s = x.item<float>();
|
||||
assert(s == 1.0);
|
||||
(void)s;
|
||||
|
||||
// Scalars have a size of 1:
|
||||
size_t size = x.size();
|
||||
int64_t size = x.size();
|
||||
assert(size == 1);
|
||||
(void)size;
|
||||
|
||||
// Scalars have 0 dimensions:
|
||||
int ndim = x.ndim();
|
||||
assert(ndim == 0);
|
||||
(void)ndim;
|
||||
|
||||
// The shape should be an empty vector:
|
||||
auto shape = x.shape();
|
||||
@@ -30,6 +33,7 @@ void array_basics() {
|
||||
// The datatype should be float32:
|
||||
auto dtype = x.dtype();
|
||||
assert(dtype == mx::float32);
|
||||
(void)dtype;
|
||||
|
||||
// Specify the dtype when constructing the array:
|
||||
x = mx::array(1, mx::int32);
|
||||
|
||||
+8
-7
@@ -44,11 +44,11 @@ std::vector<array> array::make_arrays(
|
||||
const std::shared_ptr<Primitive>& primitive,
|
||||
const std::vector<array>& inputs) {
|
||||
std::vector<array> outputs;
|
||||
for (size_t i = 0; i < shapes.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(shapes); ++i) {
|
||||
outputs.emplace_back(std::move(shapes[i]), dtypes[i], primitive, inputs);
|
||||
}
|
||||
// For each node in |outputs|, its siblings are the other nodes.
|
||||
for (size_t i = 0; i < outputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(outputs); ++i) {
|
||||
auto siblings = outputs;
|
||||
siblings.erase(siblings.begin() + i);
|
||||
outputs[i].set_siblings(std::move(siblings), i);
|
||||
@@ -145,8 +145,9 @@ void array::set_data(allocator::Buffer buffer, Deleter d) {
|
||||
array_desc_->data_size = size();
|
||||
array_desc_->flags.contiguous = true;
|
||||
array_desc_->flags.row_contiguous = true;
|
||||
auto max_dim = std::max_element(shape().begin(), shape().end());
|
||||
array_desc_->flags.col_contiguous = size() <= 1 || size() == *max_dim;
|
||||
auto max_dim =
|
||||
static_cast<int64_t>(*std::max_element(shape().begin(), shape().end()));
|
||||
array_desc_->flags.col_contiguous = size() <= 1 || size() == max_dim;
|
||||
}
|
||||
|
||||
void array::set_data(
|
||||
@@ -192,7 +193,7 @@ array::~array() {
|
||||
}
|
||||
|
||||
// Break circular reference for non-detached arrays with siblings
|
||||
if (auto n = siblings().size(); n > 0) {
|
||||
if (auto n = std::ssize(siblings()); n > 0) {
|
||||
bool do_detach = true;
|
||||
// If all siblings have siblings.size() references except
|
||||
// the one we are currently destroying (which has siblings.size() + 1)
|
||||
@@ -274,7 +275,7 @@ array::ArrayDesc::~ArrayDesc() {
|
||||
ad.inputs.clear();
|
||||
for (auto& [_, a] : input_map) {
|
||||
bool is_deletable =
|
||||
(a.array_desc_.use_count() <= a.siblings().size() + 1);
|
||||
(a.array_desc_.use_count() <= std::ssize(a.siblings()) + 1);
|
||||
// An array with siblings is deletable only if all of its siblings
|
||||
// are deletable
|
||||
for (auto& s : a.siblings()) {
|
||||
@@ -283,7 +284,7 @@ array::ArrayDesc::~ArrayDesc() {
|
||||
}
|
||||
int is_input = (input_map.find(s.id()) != input_map.end());
|
||||
is_deletable &=
|
||||
s.array_desc_.use_count() <= a.siblings().size() + is_input;
|
||||
s.array_desc_.use_count() <= std::ssize(a.siblings()) + is_input;
|
||||
}
|
||||
if (is_deletable) {
|
||||
for_deletion.push_back(std::move(a.array_desc_));
|
||||
|
||||
+7
-7
@@ -81,22 +81,22 @@ class array {
|
||||
}
|
||||
|
||||
/** The size of the array's datatype in bytes. */
|
||||
size_t itemsize() const {
|
||||
int itemsize() const {
|
||||
return size_of(dtype());
|
||||
}
|
||||
|
||||
/** The number of elements in the array. */
|
||||
size_t size() const {
|
||||
int64_t size() const {
|
||||
return array_desc_->size;
|
||||
}
|
||||
|
||||
/** The number of bytes in the array. */
|
||||
size_t nbytes() const {
|
||||
int64_t nbytes() const {
|
||||
return size() * itemsize();
|
||||
}
|
||||
|
||||
/** The number of dimensions of the array. */
|
||||
size_t ndim() const {
|
||||
int ndim() const {
|
||||
return array_desc_->shape.size();
|
||||
}
|
||||
|
||||
@@ -329,7 +329,7 @@ class array {
|
||||
* corresponding to ``arr[-1, -1, ...]``) then ``data_size = last - first``.
|
||||
* Note, ``data_size`` is in units of ``item_size`` (not bytes).
|
||||
**/
|
||||
size_t data_size() const {
|
||||
int64_t data_size() const {
|
||||
return array_desc_->data_size;
|
||||
}
|
||||
|
||||
@@ -340,7 +340,7 @@ class array {
|
||||
return array_desc_->data->buffer;
|
||||
}
|
||||
|
||||
size_t buffer_size() const {
|
||||
int64_t buffer_size() const {
|
||||
return allocator::allocator().size(buffer());
|
||||
}
|
||||
|
||||
@@ -530,7 +530,7 @@ array::array(
|
||||
Shape shape,
|
||||
Dtype dtype /* = TypeToDtype<T>() */)
|
||||
: array_desc_(std::make_shared<ArrayDesc>(std::move(shape), dtype)) {
|
||||
if (data.size() != size()) {
|
||||
if (std::ssize(data) != size()) {
|
||||
throw std::invalid_argument(
|
||||
"Data size and provided shape mismatch in array construction.");
|
||||
}
|
||||
|
||||
@@ -21,8 +21,8 @@ void AsStrided::eval(const std::vector<array>& inputs, array& out) {
|
||||
|
||||
// Compute the flags given the shape and strides
|
||||
bool row_contiguous = true, col_contiguous = true;
|
||||
size_t r = 1, c = 1;
|
||||
for (int i = strides_.size() - 1, j = 0; i >= 0; i--, j++) {
|
||||
int64_t r = 1, c = 1;
|
||||
for (int i = std::ssize(strides_) - 1, j = 0; i >= 0; i--, j++) {
|
||||
row_contiguous &= (r == strides_[i]) || (shape_[i] == 1);
|
||||
col_contiguous &= (c == strides_[j]) || (shape_[j] == 1);
|
||||
r *= shape_[i];
|
||||
@@ -60,7 +60,8 @@ void CustomTransforms::eval(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
assert(inputs.size() > outputs.size());
|
||||
for (int i = 0, j = inputs.size() - outputs.size(); i < outputs.size();
|
||||
for (int i = 0, j = std::ssize(inputs) - std::ssize(outputs);
|
||||
i < std::ssize(outputs);
|
||||
i++, j++) {
|
||||
outputs[i].copy_shared_buffer(inputs[j]);
|
||||
}
|
||||
@@ -70,7 +71,7 @@ void Depends::eval(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
assert(inputs.size() > outputs.size());
|
||||
for (int i = 0; i < outputs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(outputs); i++) {
|
||||
outputs[i].copy_shared_buffer(inputs[i]);
|
||||
}
|
||||
}
|
||||
@@ -206,11 +207,11 @@ void Split::eval(
|
||||
|
||||
auto compute_new_flags = [](const auto& shape,
|
||||
const auto& strides,
|
||||
size_t in_data_size,
|
||||
int64_t in_data_size,
|
||||
auto flags) {
|
||||
size_t data_size = 1;
|
||||
size_t f_stride = 1;
|
||||
size_t b_stride = 1;
|
||||
int64_t data_size = 1;
|
||||
int64_t f_stride = 1;
|
||||
int64_t b_stride = 1;
|
||||
flags.row_contiguous = true;
|
||||
flags.col_contiguous = true;
|
||||
for (int i = 0, ri = shape.size() - 1; ri >= 0; i++, ri--) {
|
||||
@@ -240,7 +241,7 @@ void Split::eval(
|
||||
|
||||
std::vector<int> indices(1, 0);
|
||||
indices.insert(indices.end(), indices_.begin(), indices_.end());
|
||||
for (int i = 0; i < indices.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(indices); i++) {
|
||||
size_t offset = indices[i] * in.strides()[axis_];
|
||||
auto [new_flags, data_size] = compute_new_flags(
|
||||
outputs[i].shape(), in.strides(), in.data_size(), in.flags());
|
||||
@@ -254,7 +255,7 @@ void Squeeze::eval(const std::vector<array>& inputs, array& out) {
|
||||
const auto& in = inputs[0];
|
||||
Strides strides;
|
||||
for (int i = 0, j = 0; i < in.ndim(); ++i) {
|
||||
if (j < axes_.size() && i == axes_[j]) {
|
||||
if (j < std::ssize(axes_) && i == axes_[j]) {
|
||||
j++;
|
||||
} else {
|
||||
strides.push_back(in.strides(i));
|
||||
@@ -272,7 +273,7 @@ void Transpose::eval(const std::vector<array>& inputs, array& out) {
|
||||
assert(inputs.size() == 1);
|
||||
Strides out_strides(out.ndim());
|
||||
auto& in = inputs[0];
|
||||
for (int ax = 0; ax < axes_.size(); ++ax) {
|
||||
for (int ax = 0; ax < std::ssize(axes_); ++ax) {
|
||||
out_strides[ax] = in.strides()[axes_[ax]];
|
||||
}
|
||||
|
||||
|
||||
@@ -120,7 +120,7 @@ void compiled_allocate_outputs(
|
||||
Strides strides;
|
||||
size_t data_size;
|
||||
array::Flags flags;
|
||||
for (int i = 0; i < inputs.size() && o < outputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs) && o < std::ssize(outputs); ++i) {
|
||||
auto& in = inputs[i];
|
||||
// Conditions for donation
|
||||
// - Correct size
|
||||
@@ -138,7 +138,7 @@ void compiled_allocate_outputs(
|
||||
data_size = in.data_size();
|
||||
}
|
||||
}
|
||||
for (; o < outputs.size(); ++o) {
|
||||
for (; o < std::ssize(outputs); ++o) {
|
||||
outputs[o].set_data(
|
||||
allocator::malloc(data_size * outputs[o].itemsize()),
|
||||
data_size,
|
||||
@@ -147,7 +147,7 @@ void compiled_allocate_outputs(
|
||||
}
|
||||
} else {
|
||||
int o = 0;
|
||||
for (int i = 0; i < inputs.size() && o < outputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs) && o < std::ssize(outputs); ++i) {
|
||||
auto& in = inputs[i];
|
||||
// Conditions for donation
|
||||
// - Row contiguous
|
||||
@@ -162,7 +162,7 @@ void compiled_allocate_outputs(
|
||||
o++;
|
||||
}
|
||||
}
|
||||
for (; o < outputs.size(); ++o) {
|
||||
for (; o < std::ssize(outputs); ++o) {
|
||||
outputs[o].set_data(allocator::malloc(outputs[o].nbytes()));
|
||||
}
|
||||
}
|
||||
@@ -193,7 +193,7 @@ std::tuple<bool, Shape, std::vector<Strides>> compiled_collapse_contiguous_dims(
|
||||
|
||||
// Broadcast the inputs to the output shape.
|
||||
Strides xstrides;
|
||||
size_t j = 0;
|
||||
int j = 0;
|
||||
for (; j < shape.size() - x.ndim(); ++j) {
|
||||
if (shape[j] == 1) {
|
||||
xstrides.push_back(out.strides()[j]);
|
||||
@@ -201,7 +201,7 @@ std::tuple<bool, Shape, std::vector<Strides>> compiled_collapse_contiguous_dims(
|
||||
xstrides.push_back(0);
|
||||
}
|
||||
}
|
||||
for (size_t i = 0; i < x.ndim(); ++i, ++j) {
|
||||
for (int i = 0; i < x.ndim(); ++i, ++j) {
|
||||
if (x.shape(i) == 1) {
|
||||
if (shape[j] == 1) {
|
||||
xstrides.push_back(out.strides()[j]);
|
||||
@@ -224,13 +224,13 @@ bool compiled_use_large_index(
|
||||
const std::vector<array>& outputs,
|
||||
bool contiguous) {
|
||||
if (contiguous) {
|
||||
size_t max_size = 0;
|
||||
int64_t max_size = 0;
|
||||
for (const auto& in : inputs) {
|
||||
max_size = std::max(max_size, in.data_size());
|
||||
}
|
||||
return max_size > UINT32_MAX;
|
||||
} else {
|
||||
size_t max_size = 0;
|
||||
int64_t max_size = 0;
|
||||
for (const auto& o : outputs) {
|
||||
max_size = std::max(max_size, o.size());
|
||||
}
|
||||
|
||||
@@ -27,7 +27,7 @@ void swap_endianness(uint8_t* data_bytes, size_t N) {
|
||||
|
||||
namespace mlx::core {
|
||||
|
||||
void Load::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
void Load::eval_cpu(const std::vector<array>& /* inputs */, array& out) {
|
||||
out.set_data(allocator::malloc(out.nbytes()));
|
||||
auto read_task = [out_ptr = out.data<char>(),
|
||||
size = out.size(),
|
||||
|
||||
@@ -28,7 +28,7 @@ std::pair<Shape, Strides> shapes_without_reduction_axes(
|
||||
|
||||
ReductionPlan get_reduction_plan(const array& x, const std::vector<int>& axes) {
|
||||
// The data is all there and we are reducing over everything
|
||||
if (x.size() == x.data_size() && axes.size() == x.ndim() &&
|
||||
if (x.size() == x.data_size() && std::ssize(axes) == x.ndim() &&
|
||||
x.flags().contiguous) {
|
||||
return ContiguousAllReduce;
|
||||
}
|
||||
@@ -38,7 +38,7 @@ ReductionPlan get_reduction_plan(const array& x, const std::vector<int>& axes) {
|
||||
// Merge consecutive axes
|
||||
Shape shape = {x.shape(axes[0])};
|
||||
Strides strides = {x.strides()[axes[0]]};
|
||||
for (int i = 1; i < axes.size(); i++) {
|
||||
for (int i = 1; i < std::ssize(axes); i++) {
|
||||
if (axes[i] - 1 == axes[i - 1] && x.shape(axes[i]) > 1) {
|
||||
shape.back() *= x.shape(axes[i]);
|
||||
strides.back() = x.strides()[axes[i]];
|
||||
|
||||
@@ -24,8 +24,8 @@ std::tuple<int64_t, Strides> prepare_slice(
|
||||
void shared_buffer_slice(
|
||||
const array& in,
|
||||
const Strides& out_strides,
|
||||
size_t data_offset,
|
||||
size_t data_size,
|
||||
int64_t data_offset,
|
||||
int64_t data_size,
|
||||
array& out) {
|
||||
// Compute row/col contiguity
|
||||
auto [no_bsx_size, is_row_contiguous, is_col_contiguous] =
|
||||
@@ -61,7 +61,7 @@ void slice(
|
||||
if (data_end < 0) {
|
||||
data_end += in.data_size();
|
||||
}
|
||||
size_t data_size = (data_end - data_offset);
|
||||
int64_t data_size = (data_end - data_offset);
|
||||
shared_buffer_slice(in, inp_strides, data_offset, data_size, out);
|
||||
}
|
||||
|
||||
|
||||
@@ -28,7 +28,7 @@ std::tuple<Shape, std::vector<Strides>> collapse_contiguous_dims(
|
||||
if (shape[0] != 1) {
|
||||
to_collapse.push_back(0);
|
||||
}
|
||||
size_t size = shape[0];
|
||||
int64_t size = shape[0];
|
||||
for (int i = 1; i < shape.size(); i++) {
|
||||
bool contiguous = true;
|
||||
size *= shape[i];
|
||||
@@ -64,7 +64,7 @@ std::tuple<Shape, std::vector<Strides>> collapse_contiguous_dims(
|
||||
current_shape *= shape[to_collapse[k]];
|
||||
}
|
||||
out_shape.push_back(current_shape);
|
||||
for (int j = 0; j < strides.size(); j++) {
|
||||
for (int j = 0; j < std::ssize(strides); j++) {
|
||||
const auto& st = strides[j];
|
||||
out_strides[j].push_back(st[to_collapse[k - 1]]);
|
||||
}
|
||||
|
||||
@@ -162,7 +162,7 @@ struct ContiguousIterator {
|
||||
};
|
||||
|
||||
inline auto check_contiguity(const Shape& shape, const Strides& strides) {
|
||||
size_t no_broadcast_data_size = 1;
|
||||
int64_t no_broadcast_data_size = 1;
|
||||
int64_t f_stride = 1;
|
||||
int64_t b_stride = 1;
|
||||
bool is_row_contiguous = true;
|
||||
@@ -183,7 +183,7 @@ inline auto check_contiguity(const Shape& shape, const Strides& strides) {
|
||||
}
|
||||
|
||||
inline bool is_donatable(const array& in, const array& out) {
|
||||
constexpr size_t donation_extra = 16384;
|
||||
constexpr int64_t donation_extra = 16384;
|
||||
|
||||
return in.is_donatable() && in.itemsize() == out.itemsize() &&
|
||||
in.buffer_size() <= out.nbytes() + donation_extra;
|
||||
|
||||
@@ -10,7 +10,7 @@ namespace mlx::core {
|
||||
namespace {
|
||||
|
||||
template <typename T>
|
||||
void arange(T start, T next, array& out, size_t size, Stream stream) {
|
||||
void arange(T start, T next, array& out, int64_t size, Stream stream) {
|
||||
auto ptr = out.data<T>();
|
||||
auto step_size = next - start;
|
||||
auto& encoder = cpu::get_command_encoder(stream);
|
||||
|
||||
@@ -19,12 +19,12 @@ void arg_reduce(const array& in, array& out, const OpT& op, int axis) {
|
||||
auto in_ptr = in.data<InT>();
|
||||
auto out_ptr = out.data<uint32_t>();
|
||||
|
||||
for (uint32_t i = 0; i < out.size(); ++i) {
|
||||
for (int64_t i = 0; i < out.size(); ++i) {
|
||||
auto loc = elem_to_loc(i, shape, strides);
|
||||
auto local_in_ptr = in_ptr + loc;
|
||||
uint32_t ind_v = 0;
|
||||
InT v = (*local_in_ptr);
|
||||
for (uint32_t j = 0; j < axis_size; ++j, local_in_ptr += axis_stride) {
|
||||
for (int64_t j = 0; j < axis_size; ++j, local_in_ptr += axis_stride) {
|
||||
op(j, (*local_in_ptr), &ind_v, &v);
|
||||
}
|
||||
out_ptr[i] = ind_v;
|
||||
|
||||
@@ -17,7 +17,12 @@ namespace mlx::core {
|
||||
namespace {
|
||||
|
||||
template <typename Op>
|
||||
void binary(const array& a, const array& b, array& out, Op op, Stream stream) {
|
||||
void binary(
|
||||
const array& a,
|
||||
const array& b,
|
||||
array& out,
|
||||
Op /* op */,
|
||||
Stream stream) {
|
||||
auto bopt = get_binary_op_type(a, b);
|
||||
set_binary_op_output_data(a, b, out, bopt);
|
||||
|
||||
@@ -81,7 +86,7 @@ void comparison_op(
|
||||
const array& a,
|
||||
const array& b,
|
||||
array& out,
|
||||
Op op,
|
||||
Op /* op */,
|
||||
Stream stream) {
|
||||
auto bopt = get_binary_op_type(a, b);
|
||||
set_binary_op_output_data(a, b, out, bopt);
|
||||
@@ -146,7 +151,7 @@ void binary_float(
|
||||
const array& a,
|
||||
const array& b,
|
||||
array& out,
|
||||
Op op,
|
||||
Op /* op */,
|
||||
Stream stream) {
|
||||
auto bopt = get_binary_op_type(a, b);
|
||||
set_binary_op_output_data(a, b, out, bopt);
|
||||
@@ -187,7 +192,7 @@ void binary_int(
|
||||
const array& a,
|
||||
const array& b,
|
||||
array& out,
|
||||
Op op,
|
||||
Op /* op */,
|
||||
Stream stream) {
|
||||
auto bopt = get_binary_op_type(a, b);
|
||||
set_binary_op_output_data(a, b, out, bopt);
|
||||
|
||||
@@ -99,7 +99,7 @@ void binary_op_dispatch_dims(
|
||||
ContiguousIterator a_it(shape, a_strides, ndim - 2);
|
||||
ContiguousIterator b_it(shape, b_strides, ndim - 2);
|
||||
auto stride = out_strides[ndim - 3];
|
||||
for (size_t elem = 0; elem < a.size(); elem += stride) {
|
||||
for (int64_t elem = 0; elem < std::ssize(a); elem += stride) {
|
||||
binary_op_dims<T, U, Op, 2>(
|
||||
a_ptr + a_it.loc,
|
||||
b_ptr + b_it.loc,
|
||||
@@ -137,21 +137,21 @@ void binary_op(
|
||||
if (bopt == BinaryOpType::ScalarScalar) {
|
||||
std::tie(*out_a_ptr, *out_b_ptr) = op(*a_ptr, *b_ptr);
|
||||
} else if (bopt == BinaryOpType::ScalarVector) {
|
||||
for (size_t i = 0; i < b.data_size(); ++i) {
|
||||
for (int64_t i = 0; i < b.data_size(); ++i) {
|
||||
std::tie(*out_a_ptr, *out_b_ptr) = op(*a_ptr, *b_ptr);
|
||||
out_a_ptr++;
|
||||
out_b_ptr++;
|
||||
b_ptr++;
|
||||
}
|
||||
} else if (bopt == BinaryOpType::VectorScalar) {
|
||||
for (size_t i = 0; i < a.data_size(); ++i) {
|
||||
for (int64_t i = 0; i < a.data_size(); ++i) {
|
||||
std::tie(*out_a_ptr, *out_b_ptr) = op(*a_ptr, *b_ptr);
|
||||
out_a_ptr++;
|
||||
out_b_ptr++;
|
||||
a_ptr++;
|
||||
}
|
||||
} else { // VectorVector
|
||||
for (size_t i = 0; i < a.size(); ++i) {
|
||||
for (int64_t i = 0; i < a.size(); ++i) {
|
||||
std::tie(*out_a_ptr, *out_b_ptr) = op(*a_ptr, *b_ptr);
|
||||
out_a_ptr++;
|
||||
out_b_ptr++;
|
||||
|
||||
@@ -33,8 +33,8 @@ void cholesky_impl(const array& a, array& factor, bool upper, Stream stream) {
|
||||
N = a.shape(-1),
|
||||
size = a.size()]() mutable {
|
||||
char uplo = (upper) ? 'L' : 'U';
|
||||
size_t num_matrices = size / (N * N);
|
||||
for (int i = 0; i < num_matrices; i++) {
|
||||
int64_t num_matrices = size / (N * N);
|
||||
for (int64_t i = 0; i < num_matrices; i++) {
|
||||
// Compute Cholesky factorization.
|
||||
int info;
|
||||
potrf<T>(
|
||||
|
||||
@@ -49,7 +49,7 @@ static CompilerCache& cache() {
|
||||
// GPU compile is always available if the GPU is available and since we are in
|
||||
// this file CPU compile is also available.
|
||||
namespace detail {
|
||||
bool compile_available_for_device(const Device& device) {
|
||||
bool compile_available_for_device(const Device& /* device */) {
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -168,7 +168,7 @@ inline void build_kernel(
|
||||
// Add the input arguments
|
||||
int cnt = 0;
|
||||
int strides_index = 1;
|
||||
for (size_t i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
// Skip constants from the input list
|
||||
if (is_constant(i)) {
|
||||
continue;
|
||||
@@ -238,7 +238,7 @@ inline void build_kernel(
|
||||
} else {
|
||||
os << x.primitive().name();
|
||||
os << "()(";
|
||||
for (int i = 0; i < x.inputs().size() - 1; i++) {
|
||||
for (int i = 0; i < std::ssize(x.inputs()) - 1; i++) {
|
||||
os << "tmp_" << namer.get_name(x.inputs()[i]) << ", ";
|
||||
}
|
||||
os << "tmp_" << namer.get_name(x.inputs().back()) << ");" << std::endl;
|
||||
|
||||
@@ -860,7 +860,7 @@ void explicit_gemm_conv_1D_cpu(
|
||||
const std::vector<int>& padding_lo,
|
||||
const std::vector<int>& padding_hi,
|
||||
const std::vector<int>& wt_strides,
|
||||
const std::vector<int>& wt_dilation,
|
||||
const std::vector<int>& /* wt_dilation */,
|
||||
Stream stream) {
|
||||
const int N = in.shape(0); // Batch size, should be the same as out.shape(0)
|
||||
const int iH = in.shape(1); // Input spatial dim
|
||||
@@ -1003,7 +1003,7 @@ void explicit_gemm_conv_ND_cpu(
|
||||
const std::vector<int>& padding_lo,
|
||||
const std::vector<int>& padding_hi,
|
||||
const std::vector<int>& wt_strides,
|
||||
const std::vector<int>& wt_dilation,
|
||||
const std::vector<int>& /* wt_dilation */,
|
||||
const bool flip,
|
||||
Stream stream) {
|
||||
const int N = in.shape(0); // Batch size, should be the same as out.shape(0)
|
||||
@@ -1023,7 +1023,7 @@ void explicit_gemm_conv_ND_cpu(
|
||||
// Pad input
|
||||
Shape padded_shape(in.shape().size());
|
||||
padded_shape.front() = N;
|
||||
for (size_t i = 0; i < iDim.size(); i++) {
|
||||
for (int i = 0; i < iDim.size(); i++) {
|
||||
padded_shape[i + 1] = iDim[i] + padding_lo[i] + padding_hi[i];
|
||||
}
|
||||
padded_shape.back() = C;
|
||||
@@ -1054,20 +1054,20 @@ void explicit_gemm_conv_ND_cpu(
|
||||
// Make strided view
|
||||
Shape strided_shape(oDim.size() + wDim.size() + 2);
|
||||
strided_shape.front() = N;
|
||||
for (size_t i = 0; i < oDim.size(); i++) {
|
||||
for (int i = 0; i < oDim.size(); i++) {
|
||||
strided_shape[i + 1] = oDim[i];
|
||||
}
|
||||
for (size_t i = 0; i < wDim.size(); i++) {
|
||||
for (int i = 0; i < wDim.size(); i++) {
|
||||
strided_shape[i + 1 + oDim.size()] = wDim[i];
|
||||
}
|
||||
strided_shape.back() = C;
|
||||
|
||||
Strides strided_strides(in.shape().size() * 2 - 2);
|
||||
strided_strides[0] = in_padded.strides()[0];
|
||||
for (size_t i = 0; i < wt_strides.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(wt_strides); i++) {
|
||||
strided_strides[i + 1] = in_padded.strides()[i + 1] * wt_strides[i];
|
||||
}
|
||||
for (size_t i = 1; i < in_padded.strides().size(); i++) {
|
||||
for (int i = 1; i < std::ssize(in_padded.strides()); i++) {
|
||||
strided_strides[i + wt_strides.size()] = in_padded.strides()[i];
|
||||
}
|
||||
|
||||
|
||||
@@ -90,6 +90,7 @@ void Recv::eval_cpu(
|
||||
std::vector<array>& outputs) {
|
||||
assert(inputs.size() == 0);
|
||||
assert(outputs.size() == 1);
|
||||
(void)inputs;
|
||||
|
||||
outputs[0].set_data(allocator::malloc(outputs[0].nbytes()));
|
||||
distributed::detail::recv(group(), outputs[0], src_, stream());
|
||||
|
||||
@@ -70,7 +70,7 @@ void eig_impl(
|
||||
auto eig_tmp = static_cast<T*>(eig_tmp_data.buffer.raw_ptr());
|
||||
auto vec_tmp = static_cast<T*>(vec_tmp_data.buffer.raw_ptr());
|
||||
auto work_buf = array::Data{allocator::malloc(sizeof(T) * lwork)};
|
||||
for (size_t i = 0; i < size / (N * N); ++i) {
|
||||
for (int64_t i = 0; i < size / (N * N); ++i) {
|
||||
geev<T>(
|
||||
&jobl,
|
||||
&jobr,
|
||||
|
||||
@@ -165,7 +165,7 @@ void eigh_impl(
|
||||
EighWork<T> work(jobz, uplo, N);
|
||||
|
||||
// Work loop
|
||||
for (size_t i = 0; i < size / (N * N); ++i) {
|
||||
for (int64_t i = 0; i < size / (N * N); ++i) {
|
||||
work.run(vec_ptr, eig_ptr);
|
||||
vec_ptr += N * N;
|
||||
eig_ptr += N;
|
||||
|
||||
@@ -20,8 +20,8 @@ struct CommandEncoder {
|
||||
CommandEncoder(CommandEncoder&&) = delete;
|
||||
CommandEncoder& operator=(CommandEncoder&&) = delete;
|
||||
|
||||
void set_input_array(const array& a) {}
|
||||
void set_output_array(array& a) {}
|
||||
void set_input_array(const array& /* a */) {}
|
||||
void set_output_array(array& /* a */) {}
|
||||
|
||||
// Hold onto a temporary until any already scheduled tasks which use it as
|
||||
// an input are complete.
|
||||
|
||||
@@ -12,12 +12,12 @@ void matmul(
|
||||
T* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
int64_t ldc,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
int64_t batch_size,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
|
||||
@@ -34,7 +34,7 @@ void matmul_bnns(
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
size_t /* ldc */,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
@@ -52,7 +52,7 @@ void matmul_bnns(
|
||||
#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
|
||||
if (beta != 1.0 && beta != 0.0) {
|
||||
// scale the output
|
||||
for (auto i = 0; i < batch_size * M * N; ++i) {
|
||||
for (size_t i = 0; i < batch_size * M * N; ++i) {
|
||||
out[i] *= beta;
|
||||
}
|
||||
beta = 1.0;
|
||||
@@ -127,7 +127,7 @@ void matmul_bnns(
|
||||
auto bnns_filter =
|
||||
BNNSFilterCreateLayerBroadcastMatMul(&gemm_params, nullptr);
|
||||
|
||||
for (int i = 0; i < batch_size; ++i) {
|
||||
for (size_t i = 0; i < batch_size; ++i) {
|
||||
BNNSFilterApplyTwoInput(
|
||||
bnns_filter,
|
||||
reinterpret_cast<const uint8_t*>(
|
||||
@@ -148,12 +148,12 @@ void matmul<float16_t>(
|
||||
float16_t* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
int64_t ldc,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
int64_t batch_size,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
@@ -183,12 +183,12 @@ void matmul<bfloat16_t>(
|
||||
bfloat16_t* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
int64_t ldc,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
int64_t batch_size,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
|
||||
@@ -13,20 +13,20 @@ void matmul<float>(
|
||||
float* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
int64_t ldc,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
int64_t batch_size,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
const Strides& b_strides) {
|
||||
auto ndim = a_shape.size();
|
||||
size_t M = a_shape[ndim - 2];
|
||||
size_t N = b_shape[ndim - 1];
|
||||
size_t K = a_shape[ndim - 1];
|
||||
int64_t M = a_shape[ndim - 2];
|
||||
int64_t N = b_shape[ndim - 1];
|
||||
int64_t K = a_shape[ndim - 1];
|
||||
|
||||
for (int i = 0; i < batch_size; ++i) {
|
||||
cblas_sgemm(
|
||||
@@ -54,20 +54,20 @@ void matmul<double>(
|
||||
double* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
int64_t ldc,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
int64_t batch_size,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
const Strides& b_strides) {
|
||||
auto ndim = a_shape.size();
|
||||
size_t M = a_shape[ndim - 2];
|
||||
size_t N = b_shape[ndim - 1];
|
||||
size_t K = a_shape[ndim - 1];
|
||||
int64_t M = a_shape[ndim - 2];
|
||||
int64_t N = b_shape[ndim - 1];
|
||||
int64_t K = a_shape[ndim - 1];
|
||||
|
||||
for (int i = 0; i < batch_size; ++i) {
|
||||
cblas_dgemm(
|
||||
@@ -95,20 +95,20 @@ void matmul<complex64_t>(
|
||||
complex64_t* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
size_t ldc,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
int64_t ldc,
|
||||
float alpha,
|
||||
float beta,
|
||||
size_t batch_size,
|
||||
int64_t batch_size,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
const Strides& b_strides) {
|
||||
auto ndim = a_shape.size();
|
||||
size_t M = a_shape[ndim - 2];
|
||||
size_t N = b_shape[ndim - 1];
|
||||
size_t K = a_shape[ndim - 1];
|
||||
int64_t M = a_shape[ndim - 2];
|
||||
int64_t N = b_shape[ndim - 1];
|
||||
int64_t K = a_shape[ndim - 1];
|
||||
auto calpha = static_cast<complex64_t>(alpha);
|
||||
auto cbeta = static_cast<complex64_t>(beta);
|
||||
|
||||
|
||||
@@ -11,9 +11,9 @@ namespace mlx::core {
|
||||
|
||||
// n = 2^k component
|
||||
template <typename T>
|
||||
void hadamard_n(T* out, int n, int m, float scale, size_t size) {
|
||||
void hadamard_n(T* out, int n, int /* m */, float scale, int64_t size) {
|
||||
for (int b = 0; b < size / n; b++) {
|
||||
size_t loc = b * n;
|
||||
int64_t loc = b * n;
|
||||
T* data_ptr = out + loc;
|
||||
int h = 1;
|
||||
int n_over_2 = n / 2;
|
||||
@@ -37,7 +37,7 @@ void hadamard_n(T* out, int n, int m, float scale, size_t size) {
|
||||
|
||||
// m component
|
||||
template <typename T>
|
||||
void hadamard_m(T* out, int n, int m, float scale, size_t size) {
|
||||
void hadamard_m(T* out, int n, int m, float scale, int64_t size) {
|
||||
auto h_matrices = hadamard_matrices();
|
||||
auto& matrix = h_matrices[m];
|
||||
auto start = 1;
|
||||
@@ -45,7 +45,7 @@ void hadamard_m(T* out, int n, int m, float scale, size_t size) {
|
||||
std::vector<bool> hmat_vec;
|
||||
while (end != std::string_view::npos) {
|
||||
auto row = matrix.substr(start, end - start);
|
||||
for (int i = 0; i < row.length(); i++) {
|
||||
for (int i = 0; i < std::ssize(row); i++) {
|
||||
hmat_vec.push_back(row[i] == '+');
|
||||
}
|
||||
start = end + 1;
|
||||
@@ -53,7 +53,7 @@ void hadamard_m(T* out, int n, int m, float scale, size_t size) {
|
||||
}
|
||||
|
||||
for (int b = 0; b < size / m / n; b++) {
|
||||
size_t loc = b * n * m;
|
||||
int64_t loc = b * n * m;
|
||||
T* data_ptr = out + loc;
|
||||
for (int i = 0; i < n; i++) {
|
||||
std::vector<float> out(m);
|
||||
|
||||
@@ -78,7 +78,7 @@ void gather(
|
||||
can_copy = true;
|
||||
|
||||
// Ignore leading 1s
|
||||
int i = 0;
|
||||
int64_t i = 0;
|
||||
for (; i < slice_sizes.size() && slice_sizes[i] == 1; ++i)
|
||||
;
|
||||
|
||||
@@ -91,7 +91,7 @@ void gather(
|
||||
can_copy = true;
|
||||
|
||||
// Ignore trailing 1s
|
||||
int i = slice_sizes.size() - 1;
|
||||
int64_t i = slice_sizes.size() - 1;
|
||||
for (; i >= 0 && slice_sizes[i] == 1; --i)
|
||||
;
|
||||
|
||||
@@ -101,11 +101,11 @@ void gather(
|
||||
can_copy = (src.shape(i) == slice_sizes[i]);
|
||||
}
|
||||
}
|
||||
size_t slice_size = 1;
|
||||
int64_t slice_size = 1;
|
||||
for (auto s : slice_sizes) {
|
||||
slice_size *= s;
|
||||
}
|
||||
size_t ind_size = slice_size == 0 ? 0 : out.size() / slice_size;
|
||||
int64_t ind_size = slice_size == 0 ? 0 : out.size() / slice_size;
|
||||
const T* src_ptr = src.data<T>();
|
||||
T* dst_ptr = out.data<T>();
|
||||
|
||||
@@ -115,10 +115,10 @@ void gather(
|
||||
src_it = ContiguousIterator(slice_sizes, src.strides(), src.ndim());
|
||||
}
|
||||
|
||||
size_t out_idx = 0;
|
||||
for (int idx = 0; idx < ind_size; idx++) {
|
||||
size_t src_idx = 0;
|
||||
for (int ii = 0; ii < inds.size(); ++ii) {
|
||||
int64_t out_idx = 0;
|
||||
for (int64_t idx = 0; idx < ind_size; idx++) {
|
||||
int64_t src_idx = 0;
|
||||
for (int ii = 0; ii < std::ssize(inds); ++ii) {
|
||||
auto ax = axes[ii];
|
||||
auto idx_loc = its[ii].loc;
|
||||
its[ii].step();
|
||||
@@ -134,7 +134,7 @@ void gather(
|
||||
src_ptr + src_idx, src_ptr + src_idx + slice_size, dst_ptr + out_idx);
|
||||
out_idx += slice_size;
|
||||
} else {
|
||||
for (int jj = 0; jj < slice_size; jj++) {
|
||||
for (int64_t jj = 0; jj < slice_size; jj++) {
|
||||
dst_ptr[out_idx++] = src_ptr[src_idx + src_it.loc];
|
||||
src_it.step();
|
||||
}
|
||||
@@ -403,11 +403,11 @@ void scatter(
|
||||
const std::vector<int>& axes) {
|
||||
int nind = inds.size();
|
||||
auto inds_ndim = updates.ndim() - out.ndim();
|
||||
size_t n_updates = nind ? inds[0].size() : 1;
|
||||
int64_t n_updates = nind ? inds[0].size() : 1;
|
||||
|
||||
Shape update_shape(
|
||||
updates.shape().begin() + inds_ndim, updates.shape().end());
|
||||
size_t update_size = 1;
|
||||
int64_t update_size = 1;
|
||||
for (auto us : update_shape) {
|
||||
update_size *= us;
|
||||
}
|
||||
@@ -418,9 +418,9 @@ void scatter(
|
||||
|
||||
auto out_ptr = out.data<InT>();
|
||||
auto upd_ptr = updates.data<InT>();
|
||||
for (int i = 0; i < n_updates; ++i) {
|
||||
size_t out_offset = 0;
|
||||
for (int j = 0; j < inds.size(); ++j) {
|
||||
for (int64_t i = 0; i < n_updates; ++i) {
|
||||
int64_t out_offset = 0;
|
||||
for (int j = 0; j < std::ssize(inds); ++j) {
|
||||
auto ax = axes[j];
|
||||
auto idx_loc = its[j].loc;
|
||||
its[j].step();
|
||||
@@ -429,7 +429,7 @@ void scatter(
|
||||
out_offset += (idx_val * out.strides()[ax]);
|
||||
}
|
||||
update_it.seek(i * update_size);
|
||||
for (int j = 0; j < update_size; ++j) {
|
||||
for (int64_t j = 0; j < update_size; ++j) {
|
||||
OpT{}(upd_ptr[update_it.loc], out_ptr + out_offset + out_it.loc);
|
||||
update_it.step();
|
||||
out_it.step();
|
||||
|
||||
@@ -122,7 +122,7 @@ void inverse_impl(
|
||||
stream);
|
||||
|
||||
const int N = a.shape(-1);
|
||||
const size_t num_matrices = a.size() / (N * N);
|
||||
const int64_t num_matrices = a.size() / (N * N);
|
||||
|
||||
auto& encoder = cpu::get_command_encoder(stream);
|
||||
encoder.set_output_array(inv);
|
||||
@@ -130,13 +130,13 @@ void inverse_impl(
|
||||
auto inv_ptr = inv.data<T>();
|
||||
if (tri) {
|
||||
encoder.dispatch([inv_ptr, N, num_matrices, upper]() {
|
||||
for (int i = 0; i < num_matrices; i++) {
|
||||
for (int64_t i = 0; i < num_matrices; i++) {
|
||||
tri_inv<T>(inv_ptr + N * N * i, N, upper);
|
||||
}
|
||||
});
|
||||
} else {
|
||||
encoder.dispatch([inv_ptr, N, num_matrices]() {
|
||||
for (int i = 0; i < num_matrices; i++) {
|
||||
for (int64_t i = 0; i < num_matrices; i++) {
|
||||
general_inv<T>(inv_ptr + N * N * i, N);
|
||||
}
|
||||
});
|
||||
|
||||
@@ -25,7 +25,7 @@ inline void mask_matrix(
|
||||
const int64_t Y_data_str,
|
||||
const int64_t X_mask_str,
|
||||
const int64_t Y_mask_str,
|
||||
const size_t mask_offset) {
|
||||
const int64_t mask_offset) {
|
||||
int tX = (X + block_size - 1) / block_size;
|
||||
int tY = (Y + block_size - 1) / block_size;
|
||||
|
||||
@@ -61,13 +61,13 @@ inline void segmented_mm(
|
||||
T* out,
|
||||
bool a_transposed,
|
||||
bool b_transposed,
|
||||
size_t lda,
|
||||
size_t ldb,
|
||||
int64_t lda,
|
||||
int64_t ldb,
|
||||
const Shape& a_shape,
|
||||
const Strides& a_strides,
|
||||
const Shape& b_shape,
|
||||
const Strides& b_strides,
|
||||
size_t num_segments,
|
||||
int64_t num_segments,
|
||||
const Shape& segments_shape,
|
||||
const Strides& segments_strides) {
|
||||
int ndim = a_shape.size();
|
||||
@@ -149,9 +149,9 @@ void BlockMaskedMM::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
auto [b_transposed, ldb, b, b_copied] =
|
||||
check_transpose(b_pre, has_op_mask, inputs.back().dtype() != bool_);
|
||||
|
||||
size_t M = a.shape(-2);
|
||||
size_t N = b.shape(-1);
|
||||
size_t K = a.shape(-1);
|
||||
int64_t M = a.shape(-2);
|
||||
int64_t N = b.shape(-1);
|
||||
int64_t K = a.shape(-1);
|
||||
|
||||
if (M == 0 || N == 0) {
|
||||
return;
|
||||
@@ -172,8 +172,8 @@ void BlockMaskedMM::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
int batch_idx,
|
||||
int X,
|
||||
int Y,
|
||||
size_t X_data_str,
|
||||
size_t Y_data_str,
|
||||
int64_t X_data_str,
|
||||
int64_t Y_data_str,
|
||||
const Shape& mask_shape,
|
||||
const Strides& mask_strides,
|
||||
bool is_bool) {
|
||||
@@ -253,7 +253,7 @@ void BlockMaskedMM::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
auto a_ptr = a.data<float>();
|
||||
auto b_ptr = b.data<float>();
|
||||
auto out_ptr = out.data<float>();
|
||||
size_t num_matrices = out.size() / (M * size_t(N));
|
||||
int64_t num_matrices = out.size() / (M * int64_t(N));
|
||||
auto ldc = out.shape(-1);
|
||||
|
||||
encoder.dispatch([a_ptr,
|
||||
@@ -394,9 +394,9 @@ void GatherMM::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
auto [a_transposed, lda, a] = check_transpose(a_pre);
|
||||
auto [b_transposed, ldb, b] = check_transpose(b_pre);
|
||||
|
||||
size_t M = a.shape(-2);
|
||||
size_t N = b.shape(-1);
|
||||
size_t K = a.shape(-1);
|
||||
int64_t M = a.shape(-2);
|
||||
int64_t N = b.shape(-1);
|
||||
int64_t K = a.shape(-1);
|
||||
|
||||
if (M == 0 || N == 0) {
|
||||
return;
|
||||
@@ -413,7 +413,7 @@ void GatherMM::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
|
||||
// Get batch dims
|
||||
auto batch_size_out = out.size() / (M * N);
|
||||
size_t matrix_stride_out = M * N;
|
||||
int64_t matrix_stride_out = M * N;
|
||||
|
||||
auto get_batch_dims = [](const auto& v) {
|
||||
return decltype(v){v.begin(), v.end() - 2};
|
||||
|
||||
@@ -48,7 +48,7 @@ static std::pair<array, bool> compute_dynamic_offset(
|
||||
auto compute_offset =
|
||||
[strides, axes, offset = offset.data<int64_t>()](const auto* indices) {
|
||||
int64_t offset_ = 0;
|
||||
for (int i = 0; i < axes.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(axes); ++i) {
|
||||
offset_ += indices[i] * strides[axes[i]];
|
||||
}
|
||||
offset[0] = offset_;
|
||||
@@ -124,6 +124,7 @@ void Transpose::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
|
||||
void Arange::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
assert(inputs.size() == 0);
|
||||
(void)inputs;
|
||||
out.set_data(allocator::malloc(out.nbytes()));
|
||||
switch (out.dtype()) {
|
||||
case bool_:
|
||||
@@ -193,9 +194,9 @@ void Concatenate::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
flags.row_contiguous = false;
|
||||
flags.col_contiguous = false;
|
||||
flags.contiguous = false;
|
||||
for (int i = 0; i < inputs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(inputs); i++) {
|
||||
array out_slice(inputs[i].shape(), out.dtype(), nullptr, {});
|
||||
size_t data_offset = strides[axis_] * sizes[i];
|
||||
int64_t data_offset = strides[axis_] * sizes[i];
|
||||
out_slice.copy_shared_buffer(
|
||||
out, strides, flags, out_slice.size(), data_offset);
|
||||
copy_cpu_inplace(inputs[i], out_slice, CopyType::GeneralGeneral, stream());
|
||||
@@ -205,7 +206,7 @@ void Concatenate::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
void Contiguous::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
assert(inputs.size() == 1);
|
||||
auto& in = inputs[0];
|
||||
constexpr size_t extra_bytes = 16384;
|
||||
constexpr int64_t extra_bytes = 16384;
|
||||
if (in.buffer_size() <= out.nbytes() + extra_bytes &&
|
||||
(in.flags().row_contiguous ||
|
||||
(allow_col_major_ && in.flags().col_contiguous))) {
|
||||
@@ -254,8 +255,8 @@ void Pad::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
copy_cpu(val, out, CopyType::Scalar, stream());
|
||||
|
||||
// Find offset for start of input values
|
||||
size_t data_offset = 0;
|
||||
for (int i = 0; i < axes_.size(); i++) {
|
||||
int64_t data_offset = 0;
|
||||
for (int i = 0; i < std::ssize(axes_); i++) {
|
||||
auto ax = axes_[i] < 0 ? out.ndim() + axes_[i] : axes_[i];
|
||||
data_offset += out.strides()[ax] * low_pad_size_[i];
|
||||
}
|
||||
@@ -274,10 +275,10 @@ void RandomBits::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
// keys has shape (N1, ..., NK, 2)
|
||||
// out has shape (N1, ..., NK, M1, M2, ...)
|
||||
auto& keys = inputs[0];
|
||||
size_t num_keys = keys.size() / 2;
|
||||
int64_t num_keys = keys.size() / 2;
|
||||
|
||||
size_t elems_per_key = out.size() / num_keys;
|
||||
size_t bytes_per_key = out.itemsize() * elems_per_key;
|
||||
int64_t elems_per_key = out.size() / num_keys;
|
||||
int64_t bytes_per_key = out.itemsize() * elems_per_key;
|
||||
out.set_data(allocator::malloc(out.nbytes()));
|
||||
|
||||
auto kptr = inputs[0].data<uint32_t>();
|
||||
@@ -291,8 +292,8 @@ void RandomBits::eval_cpu(const std::vector<array>& inputs, array& out) {
|
||||
num_keys,
|
||||
kshape = keys.shape(),
|
||||
kstrides = keys.strides()]() mutable {
|
||||
size_t out_skip = (bytes_per_key + 4 - 1) / 4;
|
||||
auto half_size = out_skip / 2;
|
||||
int64_t out_skip = (bytes_per_key + 4 - 1) / 4;
|
||||
uintptr_t half_size = out_skip / 2;
|
||||
bool even = out_skip % 2 == 0;
|
||||
for (int i = 0; i < num_keys; ++i, cptr += bytes_per_key) {
|
||||
auto ptr = reinterpret_cast<uint32_t*>(cptr);
|
||||
|
||||
@@ -13,7 +13,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
|
||||
const int M = a.shape(-2);
|
||||
const int N = a.shape(-1);
|
||||
const int lda = M;
|
||||
size_t num_matrices = a.size() / (M * N);
|
||||
int64_t num_matrices = a.size() / (M * N);
|
||||
|
||||
// Copy A to inplace input and make it col-contiguous
|
||||
array in(a.shape(), a.dtype(), nullptr, {});
|
||||
@@ -54,7 +54,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
|
||||
auto work = allocator::malloc(sizeof(T) * lwork);
|
||||
|
||||
// Loop over matrices
|
||||
for (int i = 0; i < num_matrices; ++i) {
|
||||
for (int64_t i = 0; i < num_matrices; ++i) {
|
||||
// Solve
|
||||
geqrf<T>(
|
||||
&M,
|
||||
@@ -68,7 +68,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
|
||||
}
|
||||
allocator::free(work);
|
||||
|
||||
for (int i = 0; i < num_matrices; ++i) {
|
||||
for (int64_t i = 0; i < num_matrices; ++i) {
|
||||
/// num_reflectors x N
|
||||
for (int j = 0; j < num_reflectors; ++j) {
|
||||
for (int k = 0; k < j; ++k) {
|
||||
@@ -97,7 +97,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
|
||||
work = allocator::malloc(sizeof(T) * lwork);
|
||||
|
||||
// Loop over matrices
|
||||
for (int i = 0; i < num_matrices; ++i) {
|
||||
for (int64_t i = 0; i < num_matrices; ++i) {
|
||||
// Compute Q
|
||||
orgqr<T>(
|
||||
&M,
|
||||
@@ -111,7 +111,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
|
||||
&info);
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_matrices; ++i) {
|
||||
for (int64_t i = 0; i < num_matrices; ++i) {
|
||||
// M x num_reflectors
|
||||
for (int j = 0; j < M; ++j) {
|
||||
for (int k = 0; k < num_reflectors; ++k) {
|
||||
|
||||
@@ -253,12 +253,12 @@ Simd<T, N> pow(Simd<T, N> base, Simd<T, N> exp) {
|
||||
} else {
|
||||
Simd<T, N> res = 1;
|
||||
// Raising an integer to a negative power is undefined
|
||||
if (any(exp < 0)) {
|
||||
if (any(exp < static_cast<T>(0))) {
|
||||
return 0;
|
||||
}
|
||||
while (any(exp > 0)) {
|
||||
while (any(exp > static_cast<T>(0))) {
|
||||
res = select((exp & 1) != 0, res * base, res);
|
||||
base = select(exp > 0, base * base, base);
|
||||
base = select(exp > static_cast<T>(0), base * base, base);
|
||||
exp = exp >> 1;
|
||||
}
|
||||
return res;
|
||||
|
||||
@@ -79,7 +79,8 @@ Simd<T, N> sincos(Simd<T, N> in) {
|
||||
|
||||
// Get the polynom selection mask. There is one polynom for 0 <= x <= Pi/4
|
||||
// and another one for Pi/4<x<=Pi/2. Both branches will be computed.
|
||||
auto poly_mask = (emm2 & 2) != 0;
|
||||
auto poly_mask =
|
||||
(emm2 & static_cast<uint32_t>(2)) != static_cast<uint32_t>(0);
|
||||
|
||||
// The magic pass: "Extended precision modular arithmetic"
|
||||
// x = ((x - y * DP1) - y * DP2) - y * DP3
|
||||
@@ -87,8 +88,8 @@ Simd<T, N> sincos(Simd<T, N> in) {
|
||||
x = fma(y, Simd<float, N>(-2.4187564849853515625e-4f), x);
|
||||
x = fma(y, Simd<float, N>(-3.77489497744594108e-8f), x);
|
||||
|
||||
sign_mask_sin = sign_mask_sin ^ ((emm2 & 4) != 0);
|
||||
auto sign_mask_cos = ((emm2 - 2) & 4) != 0;
|
||||
sign_mask_sin = sign_mask_sin ^ ((emm2 & 4) != static_cast<uint32_t>(0));
|
||||
auto sign_mask_cos = ((emm2 - 2) & 4) != static_cast<uint32_t>(0);
|
||||
|
||||
// Evaluate the first polynom (0 <= x <= Pi/4) in y1,
|
||||
// and the second polynom (Pi/4 <= x <= 0) in y2
|
||||
|
||||
+10
-10
@@ -120,8 +120,8 @@ template <typename T>
|
||||
void sort(array& out, int axis) {
|
||||
// Get axis, shape and stride info
|
||||
axis = axis < 0 ? axis + out.ndim() : axis;
|
||||
size_t in_size = out.size();
|
||||
size_t n_rows = in_size / out.shape(axis);
|
||||
int64_t in_size = out.size();
|
||||
int64_t n_rows = in_size / out.shape(axis);
|
||||
|
||||
auto remaining_shape = out.shape();
|
||||
remaining_shape.erase(remaining_shape.begin() + axis);
|
||||
@@ -136,7 +136,7 @@ void sort(array& out, int axis) {
|
||||
ContiguousIterator src_it(
|
||||
remaining_shape, remaining_strides, remaining_shape.size());
|
||||
auto out_ptr = out.data<T>();
|
||||
for (int i = 0; i < n_rows; i++) {
|
||||
for (int64_t i = 0; i < n_rows; i++) {
|
||||
T* data_ptr = out_ptr + src_it.loc;
|
||||
|
||||
StridedIterator st(data_ptr, axis_stride, 0);
|
||||
@@ -151,7 +151,7 @@ template <typename T, typename IdxT = uint32_t>
|
||||
void argsort(const array& in, array& out, int axis) {
|
||||
// Get axis, shape and stride info
|
||||
axis = axis < 0 ? axis + in.ndim() : axis;
|
||||
size_t n_rows = in.size() / in.shape(axis);
|
||||
int64_t n_rows = in.size() / in.shape(axis);
|
||||
|
||||
auto in_remaining_shape = in.shape();
|
||||
in_remaining_shape.erase(in_remaining_shape.begin() + axis);
|
||||
@@ -176,7 +176,7 @@ void argsort(const array& in, array& out, int axis) {
|
||||
out_remaining_shape, out_remaining_strides, out_remaining_shape.size());
|
||||
auto in_ptr = in.data<T>();
|
||||
auto out_ptr = out.data<IdxT>();
|
||||
for (int i = 0; i < n_rows; i++) {
|
||||
for (int64_t i = 0; i < n_rows; i++) {
|
||||
const T* data_ptr = in_ptr + in_it.loc;
|
||||
IdxT* idx_ptr = out_ptr + out_it.loc;
|
||||
|
||||
@@ -214,8 +214,8 @@ template <typename T>
|
||||
void partition(array& out, int axis, int kth) {
|
||||
// Get axis, shape and stride info
|
||||
axis = axis < 0 ? axis + out.ndim() : axis;
|
||||
size_t in_size = out.size();
|
||||
size_t n_rows = in_size / out.shape(axis);
|
||||
int64_t in_size = out.size();
|
||||
int64_t n_rows = in_size / out.shape(axis);
|
||||
|
||||
auto remaining_shape = out.shape();
|
||||
remaining_shape.erase(remaining_shape.begin() + axis);
|
||||
@@ -232,7 +232,7 @@ void partition(array& out, int axis, int kth) {
|
||||
ContiguousIterator src_it(
|
||||
remaining_shape, remaining_strides, remaining_shape.size());
|
||||
auto out_ptr = out.data<T>();
|
||||
for (int i = 0; i < n_rows; i++) {
|
||||
for (int64_t i = 0; i < n_rows; i++) {
|
||||
T* data_ptr = out_ptr + src_it.loc;
|
||||
src_it.step();
|
||||
|
||||
@@ -248,7 +248,7 @@ template <typename T, typename IdxT = uint32_t>
|
||||
void argpartition(const array& in, array& out, int axis, int kth) {
|
||||
// Get axis, shape and stride info
|
||||
axis = axis < 0 ? axis + in.ndim() : axis;
|
||||
size_t n_rows = in.size() / in.shape(axis);
|
||||
int64_t n_rows = in.size() / in.shape(axis);
|
||||
|
||||
auto in_remaining_shape = in.shape();
|
||||
in_remaining_shape.erase(in_remaining_shape.begin() + axis);
|
||||
@@ -277,7 +277,7 @@ void argpartition(const array& in, array& out, int axis, int kth) {
|
||||
auto in_ptr = in.data<T>();
|
||||
auto out_ptr = out.data<IdxT>();
|
||||
|
||||
for (int i = 0; i < n_rows; i++) {
|
||||
for (int64_t i = 0; i < n_rows; i++) {
|
||||
const T* data_ptr = in_ptr + in_it.loc;
|
||||
IdxT* idx_ptr = out_ptr + out_it.loc;
|
||||
in_it.step();
|
||||
|
||||
@@ -27,7 +27,7 @@ void svd_impl(
|
||||
const int N = a.shape(-1);
|
||||
const int K = std::min(M, N);
|
||||
|
||||
size_t num_matrices = a.size() / (M * N);
|
||||
int64_t num_matrices = a.size() / (M * N);
|
||||
|
||||
// lapack clobbers the input, so we have to make a copy.
|
||||
array in(a.shape(), a.dtype(), nullptr, {});
|
||||
@@ -121,7 +121,7 @@ void svd_impl(
|
||||
auto scratch = array::Data{allocator::malloc(sizeof(T) * lwork)};
|
||||
|
||||
// Loop over matrices.
|
||||
for (int i = 0; i < num_matrices; i++) {
|
||||
for (int64_t i = 0; i < num_matrices; i++) {
|
||||
gesdd<T>(
|
||||
/* jobz = */ jobz,
|
||||
// M and N are swapped since lapack expects column-major.
|
||||
@@ -153,10 +153,10 @@ void svd_impl(
|
||||
|
||||
template <typename T>
|
||||
void compute_svd(
|
||||
const array& a,
|
||||
bool compute_uv,
|
||||
std::vector<array>& outputs,
|
||||
Stream stream) {}
|
||||
const array& /* a */,
|
||||
bool /* compute_uv */,
|
||||
std::vector<array>& /* outputs */,
|
||||
Stream /* stream */) {}
|
||||
|
||||
void SVD::eval_cpu(
|
||||
const std::vector<array>& inputs,
|
||||
|
||||
@@ -136,7 +136,7 @@ void ternary_op(
|
||||
if (topt == TernaryOpType::ScalarScalarScalar) {
|
||||
*out_ptr = op(*a_ptr, *b_ptr, *c_ptr);
|
||||
} else if (topt == TernaryOpType::VectorVectorVector) {
|
||||
for (size_t i = 0; i < out.size(); ++i) {
|
||||
for (int64_t i = 0; i < out.size(); ++i) {
|
||||
*out_ptr = op(*a_ptr, *b_ptr, *c_ptr);
|
||||
a_ptr++;
|
||||
b_ptr++;
|
||||
|
||||
@@ -10,8 +10,8 @@
|
||||
namespace mlx::core {
|
||||
|
||||
template <typename T, typename U = T, typename Op>
|
||||
void unary_op(const T* a, U* out, size_t shape, size_t stride) {
|
||||
for (size_t i = 0; i < shape; i += 1) {
|
||||
void unary_op(const T* a, U* out, int64_t shape, int64_t stride) {
|
||||
for (int64_t i = 0; i < shape; i += 1) {
|
||||
out[i] = Op{}(*a);
|
||||
a += stride;
|
||||
}
|
||||
@@ -38,14 +38,14 @@ void unary_op(const array& a, array& out, Op) {
|
||||
src++;
|
||||
}
|
||||
} else {
|
||||
size_t shape = ndim > 0 ? a.shape().back() : 1;
|
||||
size_t stride = ndim > 0 ? a.strides().back() : 1;
|
||||
int64_t shape = ndim > 0 ? a.shape().back() : 1;
|
||||
int64_t stride = ndim > 0 ? a.strides().back() : 1;
|
||||
if (ndim <= 1) {
|
||||
unary_op<T, U, Op>(src, dst, shape, stride);
|
||||
return;
|
||||
}
|
||||
auto it = ContiguousIterator(a.shape(), a.strides(), ndim - 1);
|
||||
for (size_t elem = 0; elem < a.size(); elem += shape) {
|
||||
for (int64_t elem = 0; elem < a.size(); elem += shape) {
|
||||
unary_op<T, U, Op>(src + it.loc, dst + elem, shape, stride);
|
||||
it.step();
|
||||
}
|
||||
|
||||
@@ -35,9 +35,9 @@ std::vector<array> precompiled_cuda_kernel(
|
||||
const std::vector<ScalarArg>&,
|
||||
std::tuple<int, int, int>,
|
||||
std::tuple<int, int, int>,
|
||||
int shared_memory,
|
||||
std::optional<float> init_value,
|
||||
bool ensure_row_contiguous,
|
||||
int /* shared_memory */,
|
||||
std::optional<float> /* init_value */,
|
||||
bool /* ensure_row_contiguous */,
|
||||
StreamOrDevice) {
|
||||
throw std::runtime_error("[cuda_kernel] No CUDA back-end.");
|
||||
}
|
||||
|
||||
@@ -51,7 +51,7 @@ void Contiguous::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
MLX_PROFILER_RANGE("Contiguous::eval_gpu");
|
||||
assert(inputs.size() == 1);
|
||||
auto& in = inputs[0];
|
||||
constexpr size_t extra_bytes = 16384;
|
||||
constexpr int64_t extra_bytes = 16384;
|
||||
if (in.buffer_size() <= out.nbytes() + extra_bytes &&
|
||||
(in.flags().row_contiguous ||
|
||||
(allow_col_major_ && in.flags().col_contiguous))) {
|
||||
|
||||
@@ -11,7 +11,7 @@ void slice_gpu(
|
||||
array& out,
|
||||
const Shape& start_indices,
|
||||
const Shape& strides,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
slice(in, out, start_indices, strides);
|
||||
}
|
||||
|
||||
@@ -27,7 +27,7 @@ void pad_gpu(
|
||||
|
||||
// Find offset for start of input values
|
||||
size_t data_offset = 0;
|
||||
for (int i = 0; i < axes.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(axes); i++) {
|
||||
auto ax = axes[i] < 0 ? out.ndim() + axes[i] : axes[i];
|
||||
data_offset += out.strides()[ax] * low_pad_size[i];
|
||||
}
|
||||
|
||||
@@ -109,7 +109,7 @@ inline void build_kernel(
|
||||
|
||||
// Read constant / contiguous inputs in tmps
|
||||
std::vector<array> nc_inputs;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
auto& x = inputs[i];
|
||||
auto& xname = namer.get_name(x);
|
||||
|
||||
@@ -134,7 +134,7 @@ inline void build_kernel(
|
||||
}
|
||||
|
||||
// Initialize the indices for non-contiguous inputs
|
||||
for (int i = 0; i < nc_inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(nc_inputs); ++i) {
|
||||
auto& xname = namer.get_name(nc_inputs[i]);
|
||||
os += fmt::format(" {0} index_{1} = ", idx_type, xname);
|
||||
if (ndim == 1) {
|
||||
@@ -174,7 +174,7 @@ inline void build_kernel(
|
||||
os += fmt::format(" for (int d = {0}; d >= 0; --d) {{\n", ndim - 3);
|
||||
}
|
||||
os += " uint l = zpos % output_shape[d];\n";
|
||||
for (int i = 0; i < nc_inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(nc_inputs); ++i) {
|
||||
auto& xname = namer.get_name(nc_inputs[i]);
|
||||
os += fmt::format(" index_{0} += ", xname);
|
||||
if (dynamic_dims) {
|
||||
@@ -195,7 +195,7 @@ inline void build_kernel(
|
||||
}
|
||||
|
||||
// Read non-contiguous inputs into tmps
|
||||
for (int i = 0; i < nc_inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(nc_inputs); ++i) {
|
||||
auto& x = nc_inputs[i];
|
||||
auto& xname = namer.get_name(x);
|
||||
os += fmt::format(
|
||||
@@ -214,7 +214,7 @@ inline void build_kernel(
|
||||
} else {
|
||||
os += x.primitive().name();
|
||||
os += "()(";
|
||||
for (int i = 0; i < x.inputs().size() - 1; i++) {
|
||||
for (int i = 0; i < std::ssize(x.inputs()) - 1; i++) {
|
||||
os += fmt::format("tmp_{0}, ", namer.get_name(x.inputs()[i]));
|
||||
}
|
||||
os += fmt::format("tmp_{0});\n", namer.get_name(x.inputs().back()));
|
||||
@@ -227,7 +227,7 @@ inline void build_kernel(
|
||||
}
|
||||
// Increment indices and close per thread loop
|
||||
if (work_per_thread > 1) {
|
||||
for (int i = 0; i < nc_inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(nc_inputs); ++i) {
|
||||
auto& x = nc_inputs[i];
|
||||
auto& xname = namer.get_name(x);
|
||||
if (!dynamic_dims) {
|
||||
@@ -396,7 +396,7 @@ void Compiled::eval_gpu(
|
||||
int cnt = 0;
|
||||
int stride_idx = 1; // idx 0 is the output strides
|
||||
Strides in_strides;
|
||||
for (int i = 0; i < inputs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(inputs); i++) {
|
||||
if (is_constant_(i)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
@@ -990,7 +990,7 @@ void conv_3D_gpu(
|
||||
const std::vector<int>& wt_dilation,
|
||||
const std::vector<int>& in_dilation,
|
||||
bool flip,
|
||||
std::vector<array>& copies) {
|
||||
std::vector<array>& /* copies */) {
|
||||
// Make conv params
|
||||
MLXConvParams<3> conv_params{
|
||||
/* const int N = */ static_cast<int>(in.shape(0)),
|
||||
|
||||
@@ -68,7 +68,7 @@ std::string write_signature(
|
||||
int index = 0;
|
||||
constexpr int max_constant_array_size = 8;
|
||||
// Add inputs
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
const auto& name = input_names[i];
|
||||
const auto& arr = inputs[i];
|
||||
auto dtype = get_type_string(arr.dtype());
|
||||
@@ -109,7 +109,7 @@ std::string write_signature(
|
||||
}
|
||||
}
|
||||
// Add outputs
|
||||
for (int i = 0; i < output_names.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(output_names); ++i) {
|
||||
const auto& name = output_names[i];
|
||||
const auto& dtype = output_dtypes[i];
|
||||
kernel_source += " device ";
|
||||
@@ -126,8 +126,8 @@ std::string write_signature(
|
||||
kernel_source += " [[buffer(";
|
||||
kernel_source += std::to_string(index);
|
||||
kernel_source += ")]]";
|
||||
if (index < inputs.size() + output_names.size() - 1 ||
|
||||
attributes.size() > 0) {
|
||||
if (index < std::ssize(inputs) + std::ssize(output_names) - 1 ||
|
||||
std::ssize(attributes) > 0) {
|
||||
kernel_source += ",\n";
|
||||
} else {
|
||||
kernel_source += ") {\n";
|
||||
@@ -138,7 +138,7 @@ std::string write_signature(
|
||||
index = 0;
|
||||
for (const auto& attr : attributes) {
|
||||
kernel_source += attr;
|
||||
if (index < attributes.size() - 1) {
|
||||
if (index < std::ssize(attributes) - 1) {
|
||||
kernel_source += ",\n";
|
||||
} else {
|
||||
kernel_source += ") {\n";
|
||||
@@ -381,7 +381,7 @@ void CustomKernel::eval_gpu(
|
||||
auto& compute_encoder = d.get_command_encoder(s.index);
|
||||
compute_encoder.set_compute_pipeline_state(kernel);
|
||||
int index = 0;
|
||||
for (int i = 0; i < checked_inputs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(checked_inputs); i++) {
|
||||
const array& in = checked_inputs[i];
|
||||
auto& shape_info = shape_infos_[i];
|
||||
compute_encoder.set_input_array(in, index);
|
||||
@@ -408,7 +408,7 @@ void CustomKernel::eval_gpu(
|
||||
}
|
||||
|
||||
const auto [tx, ty, tz] = threadgroup_;
|
||||
auto tg_size = tx * ty * tz;
|
||||
unsigned long tg_size = tx * ty * tz;
|
||||
auto max_tg_size = kernel->maxTotalThreadsPerThreadgroup();
|
||||
if (tg_size > max_tg_size) {
|
||||
std::ostringstream msg;
|
||||
|
||||
@@ -127,6 +127,9 @@ std::pair<MTL::Library*, NS::Error*> load_swiftpm_library(
|
||||
}
|
||||
}
|
||||
}
|
||||
#else
|
||||
(void)device;
|
||||
(void)lib_name;
|
||||
#endif
|
||||
return {nullptr, nullptr};
|
||||
}
|
||||
@@ -713,7 +716,7 @@ MTL::LinkedFunctions* Device::get_linked_functions_(
|
||||
auto lfuncs = MTL::LinkedFunctions::linkedFunctions();
|
||||
|
||||
std::vector<NS::Object*> objs(funcs.size());
|
||||
for (int i = 0; i < funcs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(funcs); i++) {
|
||||
objs[i] = funcs[i];
|
||||
}
|
||||
|
||||
|
||||
@@ -137,7 +137,7 @@ struct DeviceStream {
|
||||
// Data updated between command buffers
|
||||
MTL::CommandBuffer* buffer{nullptr};
|
||||
int buffer_ops{0};
|
||||
size_t buffer_sizes{0};
|
||||
int64_t buffer_sizes{0};
|
||||
|
||||
// The command encoder, fence, and temporaries are updated between command
|
||||
// encoders
|
||||
|
||||
@@ -76,7 +76,7 @@ void Fence::wait(Stream stream, const array& x) {
|
||||
auto command_buffer = d.get_command_buffer(idx);
|
||||
command_buffer->encodeWait(static_cast<MTL::Event*>(f.fence), f.count);
|
||||
command_buffer->addCompletedHandler(
|
||||
[fence_ = fence_](MTL::CommandBuffer* cbuf) {});
|
||||
[fence_ = fence_](MTL::CommandBuffer* /* cbuf */) {});
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -96,7 +96,7 @@ void Fence::wait(Stream stream, const array& x) {
|
||||
compute_encoder.dispatch_threads(kernel_dims, kernel_dims);
|
||||
|
||||
d.get_command_buffer(idx)->addCompletedHandler(
|
||||
[fence_ = fence_](MTL::CommandBuffer* cbuf) {});
|
||||
[fence_ = fence_](MTL::CommandBuffer* /* cbuf */) {});
|
||||
}
|
||||
|
||||
void Fence::update(Stream stream, const array& x) {
|
||||
@@ -124,7 +124,7 @@ void Fence::update(Stream stream, const array& x) {
|
||||
command_buffer->encodeSignalEvent(
|
||||
static_cast<MTL::Event*>(f.fence), f.count);
|
||||
command_buffer->addCompletedHandler(
|
||||
[fence_ = fence_](MTL::CommandBuffer* cbuf) {});
|
||||
[fence_ = fence_](MTL::CommandBuffer* /* cbuf */) {});
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -154,7 +154,7 @@ void Fence::update(Stream stream, const array& x) {
|
||||
compute_encoder.dispatch_threads(kernel_dims, kernel_dims);
|
||||
|
||||
d.get_command_buffer(idx)->addCompletedHandler(
|
||||
[fence_ = fence_](MTL::CommandBuffer* cbuf) {});
|
||||
[fence_ = fence_](MTL::CommandBuffer* /* cbuf */) {});
|
||||
}
|
||||
|
||||
} // namespace mlx::core
|
||||
|
||||
+11
-11
@@ -60,7 +60,7 @@ struct FourStepParams {
|
||||
void fft_op(
|
||||
const array& in,
|
||||
array& out,
|
||||
size_t axis,
|
||||
int64_t axis,
|
||||
bool inverse,
|
||||
bool real,
|
||||
const FourStepParams four_step_params,
|
||||
@@ -93,7 +93,7 @@ std::vector<int> plan_stockham_fft(int n) {
|
||||
if (n == 1) {
|
||||
return plan;
|
||||
}
|
||||
for (int i = 0; i < radices.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(radices); i++) {
|
||||
int radix = radices[i];
|
||||
// Manually tuned radices for powers of 2
|
||||
if (is_power_of_2(orig_n) && orig_n < 512 && radix > 4) {
|
||||
@@ -181,7 +181,7 @@ int compute_elems_per_thread(FFTPlan plan) {
|
||||
steps.insert(steps.end(), plan.stockham.begin(), plan.stockham.end());
|
||||
steps.insert(steps.end(), plan.rader.begin(), plan.rader.end());
|
||||
std::set<int> used_radices;
|
||||
for (int i = 0; i < steps.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(steps); i++) {
|
||||
int radix = radices[i % radices.size()];
|
||||
if (steps[i] > 0) {
|
||||
used_radices.insert(radix);
|
||||
@@ -260,7 +260,7 @@ int primitive_root(int n) {
|
||||
|
||||
std::tuple<array, array, array> compute_raders_constants(
|
||||
int rader_n,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
int proot = primitive_root(rader_n);
|
||||
// Fermat's little theorem
|
||||
int inv = mod_exp(proot, rader_n - 2, rader_n);
|
||||
@@ -508,7 +508,7 @@ void four_step_fft(
|
||||
void fft_op(
|
||||
const array& in,
|
||||
array& out,
|
||||
size_t axis,
|
||||
int64_t axis,
|
||||
bool inverse,
|
||||
bool real,
|
||||
const FourStepParams four_step_params,
|
||||
@@ -612,11 +612,11 @@ void fft_op(
|
||||
|
||||
// Start of radix/rader step constants
|
||||
int index = 4;
|
||||
for (int i = 0; i < plan.stockham.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(plan.stockham); i++) {
|
||||
func_consts.push_back(make_int(&plan.stockham[i], index));
|
||||
index += 1;
|
||||
}
|
||||
for (int i = 0; i < plan.rader.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(plan.rader); i++) {
|
||||
func_consts.push_back(make_int(&plan.rader[i], index));
|
||||
index += 1;
|
||||
}
|
||||
@@ -771,8 +771,8 @@ void nd_fft_op(
|
||||
array temp1(temp_shape, complex64, nullptr, {});
|
||||
array temp2(temp_shape, complex64, nullptr, {});
|
||||
std::vector<array> temp_arrs = {temp1, temp2};
|
||||
for (int i = axes.size() - 1; i >= 0; i--) {
|
||||
int reverse_index = axes.size() - i - 1;
|
||||
for (int i = std::ssize(axes) - 1; i >= 0; i--) {
|
||||
int reverse_index = std::ssize(axes) - i - 1;
|
||||
// For 5D and above, we don't want to reallocate our two temporary arrays
|
||||
bool inplace = reverse_index >= 3 && i != 0;
|
||||
// Opposite order for fft vs ifft
|
||||
@@ -780,8 +780,8 @@ void nd_fft_op(
|
||||
size_t axis = axes[index];
|
||||
// Mirror np.fft.(i)rfftn and perform a real transform
|
||||
// only on the final axis.
|
||||
bool step_real = (real && index == axes.size() - 1);
|
||||
const array& in_arr = i == axes.size() - 1 ? in : temp_arrs[1 - i % 2];
|
||||
bool step_real = (real && index == std::ssize(axes) - 1);
|
||||
const array& in_arr = i == std::ssize(axes) - 1 ? in : temp_arrs[1 - i % 2];
|
||||
array& out_arr = i == 0 ? out : temp_arrs[i % 2];
|
||||
fft_op(in_arr, out_arr, axis, inverse, step_real, inplace, s);
|
||||
}
|
||||
|
||||
@@ -43,7 +43,7 @@ std::string gen_hadamard_codelet(int m) {
|
||||
while (end != std::string_view::npos) {
|
||||
source << " tmp[" << index << "] = ";
|
||||
auto row = matrix.substr(start, end - start);
|
||||
for (int i = 0; i < row.length(); i++) {
|
||||
for (int i = 0; i < std::ssize(row); i++) {
|
||||
source << " " << row[i] << " x[" << i << "]";
|
||||
}
|
||||
source << ";" << std::endl;
|
||||
|
||||
@@ -52,7 +52,7 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
auto& s = stream();
|
||||
auto& d = metal::device(s.device);
|
||||
|
||||
size_t slice_size = 1;
|
||||
int64_t slice_size = 1;
|
||||
for (auto s : slice_sizes_) {
|
||||
slice_size *= s;
|
||||
}
|
||||
@@ -94,8 +94,8 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
auto kernel = d.get_kernel(kernel_name, lib);
|
||||
compute_encoder.set_compute_pipeline_state(kernel);
|
||||
|
||||
size_t dim_x = (slice_size + work_per_thread - 1) / work_per_thread;
|
||||
size_t dim_y = indices.size();
|
||||
int64_t dim_x = (slice_size + work_per_thread - 1) / work_per_thread;
|
||||
int64_t dim_y = indices.size();
|
||||
auto group_dims = get_block_dims(dim_x, dim_y, 1);
|
||||
MTL::Size grid_dims = MTL::Size(dim_x, dim_y, 1);
|
||||
|
||||
@@ -110,7 +110,7 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
}
|
||||
|
||||
int idx_ndim = nidx ? inputs[1].ndim() : 0;
|
||||
size_t ndim = src.ndim();
|
||||
int64_t ndim = src.ndim();
|
||||
|
||||
std::string kernel_name = fmt::format(
|
||||
"gather{0}{1}_{2}_{3}_{4}",
|
||||
@@ -149,8 +149,8 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
|
||||
// Launch 3D grid of threads
|
||||
// First two dimensions for the indices, the last one for the slice
|
||||
size_t dim0 = 1;
|
||||
size_t dim1 = 1;
|
||||
int64_t dim0 = 1;
|
||||
int64_t dim1 = 1;
|
||||
if (nidx) {
|
||||
if (inputs[1].ndim() >= 1) {
|
||||
dim0 = inputs[1].shape(0);
|
||||
@@ -159,13 +159,13 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
dim1 = inputs[1].size() / dim0;
|
||||
}
|
||||
}
|
||||
size_t dim2 = slice_size;
|
||||
int64_t dim2 = slice_size;
|
||||
auto group_dims = get_block_dims(dim0, dim1, dim2);
|
||||
MTL::Size grid_dims = MTL::Size(dim0, dim1, dim2);
|
||||
|
||||
// Collect all idx shapes and strides into one place
|
||||
std::vector<int> idx_shapes;
|
||||
std::vector<size_t> idx_strides;
|
||||
std::vector<int64_t> idx_strides;
|
||||
std::vector<char> idx_contigs;
|
||||
for (int i = 0; i < nidx; ++i) {
|
||||
idx_shapes.insert(
|
||||
@@ -246,7 +246,7 @@ void Scatter::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
auto& d = metal::device(s.device);
|
||||
|
||||
int idx_ndim = nidx ? inputs[1].ndim() : 0;
|
||||
size_t idx_size = nidx ? inputs[1].size() : 1;
|
||||
int64_t idx_size = nidx ? inputs[1].size() : 1;
|
||||
|
||||
auto idx_to_out = idx_size / out.size();
|
||||
int nwork;
|
||||
@@ -345,7 +345,7 @@ void Scatter::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
auto& compute_encoder = d.get_command_encoder(s.index);
|
||||
auto kernel = d.get_kernel(kernel_name, lib);
|
||||
|
||||
size_t nthreads = upd.size();
|
||||
int64_t nthreads = upd.size();
|
||||
|
||||
compute_encoder.set_compute_pipeline_state(kernel);
|
||||
|
||||
@@ -354,8 +354,8 @@ void Scatter::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
compute_encoder.set_output_array(out, 2);
|
||||
|
||||
// Set update info
|
||||
size_t upd_ndim = upd.ndim();
|
||||
size_t upd_size = 1;
|
||||
int64_t upd_ndim = upd.ndim();
|
||||
int64_t upd_size = 1;
|
||||
for (int i = idx_ndim; i < upd.ndim(); ++i) {
|
||||
upd_size *= upd.shape(i);
|
||||
}
|
||||
@@ -391,7 +391,7 @@ void Scatter::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
compute_encoder.set_bytes(upd_size, 6);
|
||||
|
||||
// Set output info
|
||||
size_t out_ndim = out.ndim();
|
||||
int64_t out_ndim = out.ndim();
|
||||
if (out_ndim == 0) {
|
||||
// Need placeholders so Metal doesn't complain
|
||||
int shape_ = 0;
|
||||
@@ -448,7 +448,7 @@ void GatherAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
auto& s = stream();
|
||||
auto& d = metal::device(s.device);
|
||||
|
||||
size_t ndim = src.ndim();
|
||||
int64_t ndim = src.ndim();
|
||||
|
||||
bool large = idx.size() > INT32_MAX || src.size() > INT32_MAX;
|
||||
|
||||
@@ -486,8 +486,8 @@ void GatherAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
compute_encoder.set_compute_pipeline_state(kernel);
|
||||
|
||||
// Grid [size post, index size, size pre]
|
||||
size_t size_pre = 1;
|
||||
size_t size_post = 1;
|
||||
int64_t size_pre = 1;
|
||||
int64_t size_post = 1;
|
||||
for (int i = 0; i < axis_; ++i) {
|
||||
size_pre *= idx.shape(i);
|
||||
}
|
||||
@@ -541,7 +541,7 @@ void ScatterAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
auto& s = stream();
|
||||
auto& d = metal::device(s.device);
|
||||
|
||||
size_t ndim = src.ndim();
|
||||
int64_t ndim = src.ndim();
|
||||
|
||||
bool large = idx.size() > INT32_MAX || src.size() > INT32_MAX;
|
||||
|
||||
@@ -602,8 +602,8 @@ void ScatterAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
compute_encoder.set_compute_pipeline_state(kernel);
|
||||
|
||||
// Grid [size post, index size, size pre]
|
||||
size_t size_pre = 1;
|
||||
size_t size_post = 1;
|
||||
int64_t size_pre = 1;
|
||||
int64_t size_post = 1;
|
||||
for (int i = 0; i < axis_; ++i) {
|
||||
size_pre *= idx.shape(i);
|
||||
}
|
||||
|
||||
@@ -344,7 +344,7 @@ void steel_gemm_splitk_axpby(
|
||||
int M,
|
||||
int N,
|
||||
int K,
|
||||
int batch_size_out,
|
||||
int /* batch_size_out */,
|
||||
int lda,
|
||||
int ldb,
|
||||
bool transpose_a,
|
||||
|
||||
@@ -179,8 +179,8 @@ MTL::ComputePipelineState* get_steel_gemm_masked_kernel(
|
||||
metal::Device& d,
|
||||
const std::string& kernel_name,
|
||||
const array&,
|
||||
const std::optional<array>& mask_out,
|
||||
const std::optional<array>& mask_op,
|
||||
const std::optional<array>& /* mask_out */,
|
||||
const std::optional<array>& /* mask_op */,
|
||||
bool,
|
||||
bool,
|
||||
int,
|
||||
|
||||
@@ -134,7 +134,7 @@ void RMSNormVJP::eval_gpu(
|
||||
d.add_temporary(g, s.index);
|
||||
}
|
||||
|
||||
auto axis_size = static_cast<uint32_t>(x.shape().back());
|
||||
auto axis_size = x.shape().back();
|
||||
int n_rows = x.data_size() / axis_size;
|
||||
|
||||
// Allocate the gradient accumulator gw and a temporary to store the
|
||||
@@ -246,7 +246,7 @@ void LayerNorm::eval_gpu(
|
||||
const array& w = inputs[1];
|
||||
const array& b = inputs[2];
|
||||
|
||||
auto axis_size = static_cast<uint32_t>(x.shape().back());
|
||||
auto axis_size = x.shape().back();
|
||||
int n_rows = x.data_size() / axis_size;
|
||||
|
||||
int simd_size = 32;
|
||||
@@ -344,7 +344,7 @@ void LayerNormVJP::eval_gpu(
|
||||
d.add_temporary(g, s.index);
|
||||
}
|
||||
|
||||
auto axis_size = static_cast<uint32_t>(x.shape().back());
|
||||
auto axis_size = x.shape().back();
|
||||
int n_rows = x.data_size() / axis_size;
|
||||
|
||||
// Allocate a temporary to store the gradients for w and allocate the output
|
||||
|
||||
@@ -26,6 +26,7 @@ void arange_set_scalars(T start, T next, metal::CommandEncoder& enc) {
|
||||
|
||||
void Arange::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
assert(inputs.size() == 0);
|
||||
(void)inputs;
|
||||
out.set_data(allocator::malloc(out.nbytes()));
|
||||
if (out.size() == 0) {
|
||||
return;
|
||||
@@ -152,7 +153,7 @@ void ArgReduce::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
}
|
||||
}
|
||||
|
||||
void Load::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
void Load::eval_gpu(const std::vector<array>& /* inputs */, array& /* out */) {
|
||||
throw std::runtime_error("[Load::eval_gpu] Not implemented.");
|
||||
}
|
||||
|
||||
@@ -201,41 +202,45 @@ void RandomBits::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
}
|
||||
|
||||
void QRF::eval_gpu(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) {
|
||||
throw std::runtime_error("[QRF::eval_gpu] Metal QR factorization NYI.");
|
||||
}
|
||||
|
||||
void SVD::eval_gpu(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) {
|
||||
throw std::runtime_error("[SVD::eval_gpu] Metal SVD NYI.");
|
||||
}
|
||||
|
||||
void Inverse::eval_gpu(const std::vector<array>& inputs, array& output) {
|
||||
void Inverse::eval_gpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
array& /* output */) {
|
||||
throw std::runtime_error("[Inverse::eval_gpu] Metal inversion NYI.");
|
||||
}
|
||||
|
||||
void Cholesky::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
void Cholesky::eval_gpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
array& /* out */) {
|
||||
throw std::runtime_error(
|
||||
"[Cholesky::eval_gpu] Metal Cholesky decomposition NYI.");
|
||||
}
|
||||
|
||||
void Eig::eval_gpu(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) {
|
||||
throw std::runtime_error("[Eig::eval_gpu] Metal Eig NYI.");
|
||||
}
|
||||
|
||||
void Eigh::eval_gpu(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) {
|
||||
throw std::runtime_error("[Eigh::eval_gpu] Metal Eigh NYI.");
|
||||
}
|
||||
|
||||
void LUF::eval_gpu(
|
||||
const std::vector<array>& inputs,
|
||||
std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) {
|
||||
throw std::runtime_error("[LUF::eval_gpu] Metal LU factorization NYI.");
|
||||
}
|
||||
|
||||
|
||||
@@ -291,7 +291,7 @@ void init_reduce(
|
||||
const std::string& op_name,
|
||||
CommandEncoder& compute_encoder,
|
||||
metal::Device& d,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
auto [_, out_type] = remap_reduce_types(out, op_name);
|
||||
const std::string func_name = "init_reduce";
|
||||
std::string kname = func_name;
|
||||
@@ -397,7 +397,7 @@ void row_reduce_small(
|
||||
RowReduceArgs& args,
|
||||
CommandEncoder& compute_encoder,
|
||||
metal::Device& d,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
// Set the kernel
|
||||
int n = get_kernel_reduce_ndim(args.reduce_ndim);
|
||||
auto [in_type, out_type] = remap_reduce_types(in, op_name);
|
||||
@@ -453,7 +453,7 @@ void row_reduce_simple(
|
||||
RowReduceArgs& args,
|
||||
CommandEncoder& compute_encoder,
|
||||
metal::Device& d,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
// Set the kernel
|
||||
auto [in_type, out_type] = remap_reduce_types(in, op_name);
|
||||
const std::string func_name = "row_reduce_simple";
|
||||
@@ -493,7 +493,7 @@ void row_reduce_looped(
|
||||
RowReduceArgs& args,
|
||||
CommandEncoder& compute_encoder,
|
||||
metal::Device& d,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
auto [in_type, out_type] = remap_reduce_types(in, op_name);
|
||||
|
||||
// Set the kernel
|
||||
@@ -570,7 +570,7 @@ void strided_reduce_small(
|
||||
ColReduceArgs& args,
|
||||
CommandEncoder& compute_encoder,
|
||||
metal::Device& d,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
auto [in_type, out_type] = remap_reduce_types(in, op_name);
|
||||
|
||||
// Figure out the grid dims
|
||||
@@ -747,7 +747,7 @@ void strided_reduce_looped(
|
||||
ColReduceArgs& args,
|
||||
CommandEncoder& compute_encoder,
|
||||
metal::Device& d,
|
||||
const Stream& s) {
|
||||
const Stream& /* s */) {
|
||||
auto [in_type, out_type] = remap_reduce_types(in, op_name);
|
||||
|
||||
// Prepare the arguments for the kernel
|
||||
@@ -959,7 +959,7 @@ void Reduce::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
// Continue with reduction operation
|
||||
// Minimum of 4 bytes since we use size 4 structs for all reduce
|
||||
// and metal will complain o/w
|
||||
size_t min_bytes = std::max(out.nbytes(), 4ul);
|
||||
size_t min_bytes = std::max<int64_t>(out.nbytes(), 4);
|
||||
out.set_data(allocator::malloc(min_bytes));
|
||||
std::string op_name;
|
||||
switch (reduce_type_) {
|
||||
|
||||
@@ -80,7 +80,7 @@ void ResidencySet::resize(size_t size) {
|
||||
// Remove wired allocations until under capacity
|
||||
auto allocations = wired_set_->allAllocations();
|
||||
auto num_allocations = wired_set_->allocationCount();
|
||||
for (int i = 0; i < num_allocations && current_size > size; ++i) {
|
||||
for (size_t i = 0; i < num_allocations && current_size > size; ++i) {
|
||||
auto buf = static_cast<const MTL::Allocation*>(allocations->object(i));
|
||||
wired_set_->removeAllocation(buf);
|
||||
current_size -= buf->allocatedSize();
|
||||
|
||||
@@ -76,7 +76,7 @@ void Scan::eval_gpu(const std::vector<array>& inputs, array& out) {
|
||||
compute_encoder.set_compute_pipeline_state(kernel);
|
||||
compute_encoder.set_input_array(in, 0);
|
||||
compute_encoder.set_output_array(out, 1);
|
||||
size_t size = in.shape(axis_);
|
||||
int64_t size = in.shape(axis_);
|
||||
compute_encoder.set_bytes(size, 2);
|
||||
|
||||
// Compute the thread grid
|
||||
|
||||
@@ -33,7 +33,7 @@ void concatenate_gpu(
|
||||
auto& d = metal::device(s.device);
|
||||
auto& compute_encoder = d.get_command_encoder(s.index);
|
||||
auto concurrent_ctx = compute_encoder.start_concurrent();
|
||||
for (int i = 0; i < inputs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(inputs); i++) {
|
||||
array out_slice(inputs[i].shape(), out.dtype(), nullptr, {});
|
||||
size_t data_offset = strides[axis] * sizes[i];
|
||||
out_slice.copy_shared_buffer(
|
||||
|
||||
@@ -29,6 +29,10 @@ inline void debug_set_stream_queue_label(MTL::CommandQueue* queue, int index) {
|
||||
std::ostringstream label;
|
||||
label << "Stream " << index;
|
||||
queue->setLabel(make_string(label));
|
||||
#else
|
||||
// appease warnings
|
||||
(void)queue;
|
||||
(void)index;
|
||||
#endif
|
||||
}
|
||||
|
||||
@@ -42,6 +46,9 @@ inline void debug_set_primitive_buffer_label(
|
||||
}
|
||||
label << primitive.name();
|
||||
command_buffer->setLabel(make_string(label));
|
||||
#else
|
||||
(void)command_buffer;
|
||||
(void)primitive;
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
+15
-15
@@ -194,7 +194,7 @@ const char* Compiled::name() const {
|
||||
}
|
||||
|
||||
std::vector<Shape> Compiled::output_shapes(const std::vector<array>& inputs) {
|
||||
size_t nd = 0;
|
||||
int nd = 0;
|
||||
for (auto& in : inputs) {
|
||||
nd = std::max(nd, in.ndim());
|
||||
}
|
||||
@@ -256,7 +256,7 @@ void merge(array& dst, array& src, ParentsMap& parents_map) {
|
||||
auto sources = src.outputs();
|
||||
auto dests = dst.outputs();
|
||||
// For each src parent, point it to the corresponding dst
|
||||
for (int i = 0; i < sources.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(sources); ++i) {
|
||||
merge_one(dests[i], sources[i], parents_map);
|
||||
}
|
||||
}
|
||||
@@ -327,7 +327,7 @@ class CompilerCache {
|
||||
if (in1.size() != in2.size()) {
|
||||
return false;
|
||||
}
|
||||
for (size_t i = 0; i < in1.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(in1); ++i) {
|
||||
if (in1[i].ndim() != in2[i].ndim()) {
|
||||
return false;
|
||||
}
|
||||
@@ -399,7 +399,7 @@ compile_trace(
|
||||
// Run the function on placeholder inputs
|
||||
// to get compute graph
|
||||
std::vector<array> tracer_inputs;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
array in(inputs[i].shape(), inputs[i].dtype(), nullptr, {});
|
||||
in.set_tracer(true);
|
||||
tracer_inputs.push_back(std::move(in));
|
||||
@@ -420,7 +420,7 @@ std::pair<std::vector<array>, ParentsMap> compile_dfs(
|
||||
std::unordered_set<std::uintptr_t> original_input_set;
|
||||
std::unordered_map<std::uintptr_t, std::vector<std::pair<array, int>>>
|
||||
parents_map;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
input_set.insert(inputs[i].id());
|
||||
original_input_set.insert(original_inputs[i].id());
|
||||
}
|
||||
@@ -436,7 +436,7 @@ std::pair<std::vector<array>, ParentsMap> compile_dfs(
|
||||
if (cache.find(id) != cache.end()) {
|
||||
return;
|
||||
}
|
||||
for (int i = 0; i < a.inputs().size(); i++) {
|
||||
for (int i = 0; i < std::ssize(a.inputs()); i++) {
|
||||
auto& in = a.inputs()[i];
|
||||
parents_map[in.id()].push_back({a, i});
|
||||
for (auto& s : a.siblings()) {
|
||||
@@ -534,7 +534,7 @@ void compile_simplify(
|
||||
return false;
|
||||
}
|
||||
|
||||
for (int i = 0; i < a.inputs().size(); i++) {
|
||||
for (int i = 0; i < std::ssize(a.inputs()); i++) {
|
||||
if (a.inputs()[i].id() != b.inputs()[i].id()) {
|
||||
return false;
|
||||
}
|
||||
@@ -599,7 +599,7 @@ void compile_simplify(
|
||||
auto maybe_merge_parents = [&](auto& a) {
|
||||
auto parents = parents_map.find(a.id());
|
||||
if (parents != parents_map.end()) {
|
||||
auto N = parents->second.size();
|
||||
auto N = std::ssize(parents->second);
|
||||
std::vector<bool> mask(N, false);
|
||||
|
||||
auto try_merge = [&](int dst_idx, int src_idx) {
|
||||
@@ -642,11 +642,11 @@ void compile_simplify(
|
||||
it->second.push_back(i);
|
||||
}
|
||||
for (auto& [_, group] : dst_map) {
|
||||
for (int i = 0; i < group.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(group); ++i) {
|
||||
if (mask[group[i]]) {
|
||||
continue;
|
||||
}
|
||||
for (int j = i + 1; j < group.size(); ++j) {
|
||||
for (int j = i + 1; j < std::ssize(group); ++j) {
|
||||
if (mask[group[j]]) {
|
||||
continue;
|
||||
}
|
||||
@@ -847,7 +847,7 @@ void compile_fuse(
|
||||
std::vector<array> old_outputs;
|
||||
// Add to global cache and add any global outputs to outputs
|
||||
// of new primitive
|
||||
for (int j = 0; j < fused_tape.size() - 1; ++j) {
|
||||
for (int j = 0; j < std::ssize(fused_tape) - 1; ++j) {
|
||||
auto& f = fused_tape[j];
|
||||
if (output_map.find(f.id()) != output_map.end()) {
|
||||
old_outputs.push_back(f);
|
||||
@@ -903,7 +903,7 @@ void compile_fuse(
|
||||
new_tape.push_back(compiled_outputs.back());
|
||||
|
||||
// Replace inputs old parents with compiled_outputs
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
auto& pairs = parents_map[inputs[i].id()];
|
||||
pairs.erase(
|
||||
std::remove_if(
|
||||
@@ -918,7 +918,7 @@ void compile_fuse(
|
||||
|
||||
// - Update outputs parents to point to compiled outputs
|
||||
// - Update any overall graph outputs to be compiled outputs
|
||||
for (int o = 0; o < old_outputs.size(); ++o) {
|
||||
for (int o = 0; o < std::ssize(old_outputs); ++o) {
|
||||
merge_one(compiled_outputs[o], old_outputs[o], parents_map);
|
||||
if (auto it = output_map.find(old_outputs[o].id());
|
||||
it != output_map.end()) {
|
||||
@@ -943,7 +943,7 @@ std::vector<array> compile_replace(
|
||||
const std::vector<array>& inputs,
|
||||
bool shapeless) {
|
||||
std::unordered_map<uintptr_t, array> trace_to_real;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
trace_to_real.insert({trace_inputs[i].id(), inputs[i]});
|
||||
}
|
||||
|
||||
@@ -989,7 +989,7 @@ std::vector<array> compile_replace(
|
||||
}
|
||||
auto real_out = array::make_arrays(
|
||||
std::move(shapes), types, a.primitive_ptr(), real_inputs);
|
||||
for (int i = 0; i < trace_out.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(trace_out); ++i) {
|
||||
trace_to_real.insert({trace_out[i].id(), std::move(real_out[i])});
|
||||
}
|
||||
}
|
||||
|
||||
@@ -55,7 +55,8 @@ class EmptyGroup : public GroupImpl {
|
||||
return 1;
|
||||
}
|
||||
|
||||
std::shared_ptr<GroupImpl> split(int color, int key = -1) override {
|
||||
std::shared_ptr<GroupImpl> split(int /* color */, int /* key */ = -1)
|
||||
override {
|
||||
throw std::runtime_error("Cannot split the distributed group further.");
|
||||
}
|
||||
|
||||
|
||||
@@ -36,7 +36,7 @@ void simple_sum(
|
||||
void* input,
|
||||
void* accumulator,
|
||||
int* len,
|
||||
MPI_Datatype* datatype) {
|
||||
MPI_Datatype* /* datatype */) {
|
||||
T* in = (T*)input;
|
||||
T* acc = (T*)accumulator;
|
||||
int N = *len;
|
||||
@@ -55,7 +55,7 @@ void simple_max(
|
||||
void* input,
|
||||
void* accumulator,
|
||||
int* len,
|
||||
MPI_Datatype* datatype) {
|
||||
MPI_Datatype* /* datatype */) {
|
||||
T* in = (T*)input;
|
||||
T* acc = (T*)accumulator;
|
||||
int N = *len;
|
||||
@@ -75,7 +75,7 @@ void simple_min(
|
||||
void* input,
|
||||
void* accumulator,
|
||||
int* len,
|
||||
MPI_Datatype* datatype) {
|
||||
MPI_Datatype* /* datatype */) {
|
||||
T* in = (T*)input;
|
||||
T* acc = (T*)accumulator;
|
||||
int N = *len;
|
||||
|
||||
@@ -27,7 +27,7 @@ std::pair<std::vector<array>, std::vector<int>> AllReduce::vmap(
|
||||
}
|
||||
|
||||
std::vector<array> AllReduce::jvp(
|
||||
const std::vector<array>& primals,
|
||||
const std::vector<array>& /* primals */,
|
||||
const std::vector<array>& tangents,
|
||||
const std::vector<int>&) {
|
||||
switch (reduce_type_) {
|
||||
@@ -44,10 +44,10 @@ std::vector<array> AllReduce::jvp(
|
||||
}
|
||||
|
||||
std::vector<array> AllReduce::vjp(
|
||||
const std::vector<array>& primals,
|
||||
const std::vector<array>& /* primals */,
|
||||
const std::vector<array>& cotangents,
|
||||
const std::vector<int>&,
|
||||
const std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* outputs */) {
|
||||
return cotangents;
|
||||
}
|
||||
|
||||
@@ -58,7 +58,7 @@ std::pair<std::vector<array>, std::vector<int>> AllGather::vmap(
|
||||
}
|
||||
|
||||
std::vector<array> AllGather::jvp(
|
||||
const std::vector<array>& primals,
|
||||
const std::vector<array>& /* primals */,
|
||||
const std::vector<array>& tangents,
|
||||
const std::vector<int>&) {
|
||||
return {all_gather(tangents[0], group(), stream())};
|
||||
|
||||
@@ -90,8 +90,8 @@
|
||||
|
||||
namespace mlx::core::distributed::ring {
|
||||
|
||||
constexpr const size_t ALL_SUM_SIZE = 8 * 1024 * 1024;
|
||||
constexpr const size_t ALL_SUM_BUFFERS = 2;
|
||||
constexpr const int64_t ALL_SUM_SIZE = 8 * 1024 * 1024;
|
||||
constexpr const int64_t ALL_SUM_BUFFERS = 2;
|
||||
constexpr const int CONN_ATTEMPTS = 5;
|
||||
constexpr const int CONN_WAIT = 1000;
|
||||
|
||||
@@ -141,27 +141,27 @@ class SocketThread {
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
std::future<void> send(const T* buffer, size_t size) {
|
||||
std::future<void> send(const T* buffer, int64_t size) {
|
||||
return send_impl(reinterpret_cast<const char*>(buffer), size * sizeof(T));
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
std::future<void> recv(T* buffer, size_t size) {
|
||||
std::future<void> recv(T* buffer, int64_t size) {
|
||||
return recv_impl(reinterpret_cast<char*>(buffer), size * sizeof(T));
|
||||
}
|
||||
|
||||
private:
|
||||
struct SocketTask {
|
||||
SocketTask(void* b, size_t s, std::promise<void>&& p)
|
||||
SocketTask(void* b, int64_t s, std::promise<void>&& p)
|
||||
: buffer(b), size(s), promise(std::move(p)) {}
|
||||
SocketTask(SocketTask&& t)
|
||||
: buffer(t.buffer), size(t.size), promise(std::move(t.promise)) {}
|
||||
void* buffer;
|
||||
size_t size;
|
||||
int64_t size;
|
||||
std::promise<void> promise;
|
||||
};
|
||||
|
||||
std::future<void> send_impl(const char* buffer, size_t size) {
|
||||
std::future<void> send_impl(const char* buffer, int64_t size) {
|
||||
std::promise<void> send_completed_promise;
|
||||
auto send_completed_future = send_completed_promise.get_future();
|
||||
if (size == 0) {
|
||||
@@ -178,7 +178,7 @@ class SocketThread {
|
||||
return send_completed_future;
|
||||
}
|
||||
|
||||
std::future<void> recv_impl(char* buffer, size_t size) {
|
||||
std::future<void> recv_impl(char* buffer, int64_t size) {
|
||||
std::promise<void> recv_completed_promise;
|
||||
auto recv_completed_future = recv_completed_promise.get_future();
|
||||
if (size == 0) {
|
||||
@@ -232,7 +232,7 @@ class SocketThread {
|
||||
|
||||
if (!recvs_.empty()) {
|
||||
auto& task = recvs_.front();
|
||||
ssize_t r = ::recv(fd_, task.buffer, task.size, 0);
|
||||
int64_t r = ::recv(fd_, task.buffer, task.size, 0);
|
||||
if (r > 0) {
|
||||
task.buffer = static_cast<char*>(task.buffer) + r;
|
||||
task.size -= r;
|
||||
@@ -246,7 +246,7 @@ class SocketThread {
|
||||
}
|
||||
if (!sends_.empty()) {
|
||||
auto& task = sends_.front();
|
||||
ssize_t r = ::send(fd_, task.buffer, task.size, 0);
|
||||
int64_t r = ::send(fd_, task.buffer, task.size, 0);
|
||||
if (r > 0) {
|
||||
task.buffer = static_cast<char*>(task.buffer) + r;
|
||||
task.size -= r;
|
||||
@@ -283,12 +283,12 @@ class CommunicationThreads {
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
std::future<void> send(int socket, T* buffer, size_t size) {
|
||||
std::future<void> send(int socket, T* buffer, int64_t size) {
|
||||
return threads_.at(socket).send<T>(buffer, size);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
std::future<void> recv(int socket, T* buffer, size_t size) {
|
||||
std::future<void> recv(int socket, T* buffer, int64_t size) {
|
||||
return threads_.at(socket).recv<T>(buffer, size);
|
||||
}
|
||||
|
||||
@@ -505,7 +505,7 @@ std::vector<int> make_connections(
|
||||
}
|
||||
template <typename T>
|
||||
struct SumOp {
|
||||
void operator()(const T* input, T* output, size_t N) {
|
||||
void operator()(const T* input, T* output, int64_t N) {
|
||||
while (N-- > 0) {
|
||||
*output += *input;
|
||||
input++;
|
||||
@@ -516,7 +516,7 @@ struct SumOp {
|
||||
|
||||
template <typename T>
|
||||
struct MaxOp {
|
||||
void operator()(const T* input, T* output, size_t N) {
|
||||
void operator()(const T* input, T* output, int64_t N) {
|
||||
while (N-- > 0) {
|
||||
*output = std::max(*output, *input);
|
||||
input++;
|
||||
@@ -527,7 +527,7 @@ struct MaxOp {
|
||||
|
||||
template <typename T>
|
||||
struct MinOp {
|
||||
void operator()(const T* input, T* output, size_t N) {
|
||||
void operator()(const T* input, T* output, int64_t N) {
|
||||
while (N-- > 0) {
|
||||
*output = std::min(*output, *input);
|
||||
input++;
|
||||
@@ -542,7 +542,7 @@ class RingGroup : public GroupImpl {
|
||||
public:
|
||||
RingGroup(int rank, std::vector<std::vector<address_t>> nodes, bool verbose)
|
||||
: rank_(rank), verbose_(verbose), pool_(0) {
|
||||
if (rank_ > 0 && rank_ >= nodes.size()) {
|
||||
if (rank_ > 0 && rank_ >= std::ssize(nodes)) {
|
||||
throw std::runtime_error(
|
||||
"[ring] Rank cannot be larger than the size of the group");
|
||||
}
|
||||
@@ -589,7 +589,7 @@ class RingGroup : public GroupImpl {
|
||||
|
||||
// Configure all sockets to use TCP no delay.
|
||||
int one = 1;
|
||||
for (int i = 0; i < sockets_right_.size(); i++) {
|
||||
for (int64_t i = 0; i < std::ssize(sockets_right_); i++) {
|
||||
setsockopt(sockets_right_[i], SOL_TCP, TCP_NODELAY, &one, sizeof(one));
|
||||
setsockopt(sockets_left_[i], SOL_TCP, TCP_NODELAY, &one, sizeof(one));
|
||||
}
|
||||
@@ -646,7 +646,8 @@ class RingGroup : public GroupImpl {
|
||||
output, all_reduce<T, MinOp<T>>(input, output, stream, MinOp<T>()));
|
||||
}
|
||||
|
||||
std::shared_ptr<GroupImpl> split(int color, int key = -1) override {
|
||||
std::shared_ptr<GroupImpl> split(int /* color */, int /* key */ = -1)
|
||||
override {
|
||||
throw std::runtime_error("[ring] Group split not supported.");
|
||||
}
|
||||
|
||||
@@ -658,15 +659,15 @@ class RingGroup : public GroupImpl {
|
||||
nbytes = input.nbytes(),
|
||||
output_ptr = output.data<char>(),
|
||||
this]() {
|
||||
constexpr size_t min_send_size = 262144;
|
||||
size_t n_gathers = std::max(
|
||||
std::min(
|
||||
constexpr int64_t min_send_size = 262144;
|
||||
int64_t n_gathers = std::max<int64_t>(
|
||||
std::min<int64_t>(
|
||||
sockets_right_.size() + sockets_left_.size(),
|
||||
nbytes / min_send_size),
|
||||
size_t(1));
|
||||
size_t bytes_per_gather = ceildiv(nbytes, n_gathers);
|
||||
1);
|
||||
int64_t bytes_per_gather = ceildiv(nbytes, n_gathers);
|
||||
std::vector<std::future<void>> all_gathers;
|
||||
for (int i = 0; i < n_gathers; i++) {
|
||||
for (int64_t i = 0; i < n_gathers; i++) {
|
||||
auto offset = i * bytes_per_gather;
|
||||
all_gathers.emplace_back(pool_.enqueue(std::bind(
|
||||
&RingGroup::all_gather_impl,
|
||||
@@ -742,10 +743,14 @@ class RingGroup : public GroupImpl {
|
||||
auto out_ptr = output.data<char>();
|
||||
auto& encoder = cpu::get_command_encoder(stream);
|
||||
encoder.set_output_array(output);
|
||||
encoder.dispatch([in_ptr, out_ptr, size = input.size(), this, reduce_op]() {
|
||||
encoder.dispatch([in_ptr,
|
||||
out_ptr,
|
||||
size = static_cast<int64_t>(input.size()),
|
||||
this,
|
||||
reduce_op]() {
|
||||
// If the input data cannot be split into size_ segments then copy it and
|
||||
// all reduce a local buffer prefilled with 0s.
|
||||
size_t nbytes = size * sizeof(T);
|
||||
int64_t nbytes = size * sizeof(T);
|
||||
if (size < size_) {
|
||||
// TODO: Maybe allocate dynamically so we don't have the constraint
|
||||
// below?
|
||||
@@ -778,16 +783,16 @@ class RingGroup : public GroupImpl {
|
||||
|
||||
// Split the all reduces so that each member has at least 1 buffer to
|
||||
// send/recv per segment.
|
||||
constexpr size_t min_send_size = 262144;
|
||||
size_t n_reduces = std::max(
|
||||
std::min(
|
||||
constexpr int64_t min_send_size = 262144;
|
||||
int64_t n_reduces = std::max<int64_t>(
|
||||
std::min<int64_t>(
|
||||
sockets_right_.size() + sockets_left_.size(),
|
||||
nbytes / (size_ * min_send_size)),
|
||||
size_t(1));
|
||||
size_t step = ceildiv(size, n_reduces);
|
||||
1);
|
||||
int64_t step = ceildiv(size, n_reduces);
|
||||
std::vector<std::future<void>> all_sums;
|
||||
|
||||
for (int i = 0; i < n_reduces; i++) {
|
||||
for (int64_t i = 0; i < n_reduces; i++) {
|
||||
all_sums.emplace_back(pool_.enqueue(std::bind(
|
||||
&RingGroup::all_reduce_impl<T, ReduceOp>,
|
||||
this,
|
||||
@@ -810,7 +815,7 @@ class RingGroup : public GroupImpl {
|
||||
void all_reduce_impl(
|
||||
T* buffer,
|
||||
T* data,
|
||||
size_t data_size,
|
||||
int64_t data_size,
|
||||
int socket_right,
|
||||
int socket_left,
|
||||
int direction,
|
||||
@@ -821,10 +826,10 @@ class RingGroup : public GroupImpl {
|
||||
|
||||
// We split the data into `size_` segments of size `segment_size` and each
|
||||
// of these in smaller segments of ALL_SUM_SIZE which we 'll call packets.
|
||||
size_t segment_size = ceildiv(data_size, size_);
|
||||
size_t BUFFER_SIZE = std::max(
|
||||
size_t(32768), std::min(ALL_SUM_SIZE / sizeof(T), segment_size / 2));
|
||||
size_t n_packets = ceildiv(segment_size, BUFFER_SIZE);
|
||||
int64_t segment_size = ceildiv(data_size, size_);
|
||||
int64_t BUFFER_SIZE = std::max<int64_t>(
|
||||
32768, std::min<int64_t>(ALL_SUM_SIZE / sizeof(T), segment_size / 2));
|
||||
int64_t n_packets = ceildiv(segment_size, BUFFER_SIZE);
|
||||
|
||||
// Initial segments
|
||||
int send_segment = rank_;
|
||||
@@ -833,21 +838,21 @@ class RingGroup : public GroupImpl {
|
||||
// Plan the whole reduce in terms of sends and recvs as indices in data.
|
||||
// It makes the actual async send and recv a bit simpler to follow when
|
||||
// there are less offset calculations around.
|
||||
std::vector<std::pair<size_t, size_t>> send_plan;
|
||||
std::vector<std::pair<size_t, size_t>> recv_plan;
|
||||
std::vector<std::pair<int64_t, int64_t>> send_plan;
|
||||
std::vector<std::pair<int64_t, int64_t>> recv_plan;
|
||||
|
||||
// Two times the same send/recv operations, first scatter reduce and then
|
||||
// gather.
|
||||
for (int k = 0; k < 2; k++) {
|
||||
for (int i = 0; i < size_ - 1; i++) {
|
||||
size_t send_start = send_segment * segment_size;
|
||||
size_t send_stop =
|
||||
int64_t send_start = send_segment * segment_size;
|
||||
int64_t send_stop =
|
||||
std::min((send_segment + 1) * segment_size, data_size);
|
||||
size_t recv_start = recv_segment * segment_size;
|
||||
size_t recv_stop =
|
||||
int64_t recv_start = recv_segment * segment_size;
|
||||
int64_t recv_stop =
|
||||
std::min((recv_segment + 1) * segment_size, data_size);
|
||||
|
||||
for (size_t j = 0; j < n_packets; j++) {
|
||||
for (int64_t j = 0; j < n_packets; j++) {
|
||||
send_plan.emplace_back(
|
||||
std::min(send_start + j * BUFFER_SIZE, send_stop),
|
||||
std::min(send_start + (j + 1) * BUFFER_SIZE, send_stop));
|
||||
@@ -864,18 +869,18 @@ class RingGroup : public GroupImpl {
|
||||
// Running the plan is fairly simple, we keep a send and a recv in flight
|
||||
// while doing the summation.
|
||||
T* recv_buffers[ALL_SUM_BUFFERS];
|
||||
for (int i = 0; i < ALL_SUM_BUFFERS; i++) {
|
||||
for (int64_t i = 0; i < ALL_SUM_BUFFERS; i++) {
|
||||
recv_buffers[i] = buffer + i * BUFFER_SIZE;
|
||||
}
|
||||
std::future<void> sends[2], recvs[2];
|
||||
int a = 0;
|
||||
int b = (n_packets > 1) ? 1 : 0;
|
||||
for (int i = 0, j = -b; i < send_plan.size(); j++, i++) {
|
||||
for (int i = 0, j = -b; i < std::ssize(send_plan); j++, i++) {
|
||||
sends[a] = comm_.send(
|
||||
socket_send,
|
||||
data + send_plan[i].first,
|
||||
send_plan[i].second - send_plan[i].first);
|
||||
if (2 * i < send_plan.size()) {
|
||||
if (2 * i < std::ssize(send_plan)) {
|
||||
recvs[a] = comm_.recv(
|
||||
socket_recv,
|
||||
recv_buffers[i % ALL_SUM_BUFFERS],
|
||||
@@ -890,7 +895,7 @@ class RingGroup : public GroupImpl {
|
||||
if (j >= 0) {
|
||||
sends[b].wait();
|
||||
recvs[b].wait();
|
||||
if (2 * j < send_plan.size()) {
|
||||
if (2 * j < std::ssize(send_plan)) {
|
||||
reduce_op(
|
||||
recv_buffers[j % ALL_SUM_BUFFERS],
|
||||
data + recv_plan[j].first,
|
||||
@@ -907,8 +912,8 @@ class RingGroup : public GroupImpl {
|
||||
void all_gather_impl(
|
||||
const char* input,
|
||||
char* output,
|
||||
size_t input_size,
|
||||
size_t data_size,
|
||||
int64_t input_size,
|
||||
int64_t data_size,
|
||||
int socket_right,
|
||||
int socket_left,
|
||||
int direction) {
|
||||
@@ -941,11 +946,11 @@ class RingGroup : public GroupImpl {
|
||||
}
|
||||
|
||||
void
|
||||
send(const std::vector<int>& sockets, const char* data, size_t data_size) {
|
||||
size_t segment_size =
|
||||
std::max(size_t(1024), ceildiv(data_size, sockets.size()));
|
||||
send(const std::vector<int>& sockets, const char* data, int64_t data_size) {
|
||||
int64_t segment_size =
|
||||
std::max<int64_t>(1024, ceildiv(data_size, std::ssize(sockets)));
|
||||
std::vector<std::future<void>> sends;
|
||||
for (int i = 0; i < sockets.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(sockets); i++) {
|
||||
if (i * segment_size >= data_size) {
|
||||
break;
|
||||
}
|
||||
@@ -959,11 +964,11 @@ class RingGroup : public GroupImpl {
|
||||
}
|
||||
}
|
||||
|
||||
void recv(const std::vector<int>& sockets, char* data, size_t data_size) {
|
||||
size_t segment_size =
|
||||
std::max(size_t(1024), ceildiv(data_size, sockets.size()));
|
||||
void recv(const std::vector<int>& sockets, char* data, int64_t data_size) {
|
||||
int64_t segment_size =
|
||||
std::max<int64_t>(1024, ceildiv(data_size, std::ssize(sockets)));
|
||||
std::vector<std::future<void>> recvs;
|
||||
for (int i = 0; i < sockets.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(sockets); i++) {
|
||||
if (i * segment_size >= data_size) {
|
||||
break;
|
||||
}
|
||||
|
||||
+1
-1
@@ -166,7 +166,7 @@ bool issubdtype(const Dtype& a, const Dtype& b) {
|
||||
return a == b;
|
||||
}
|
||||
|
||||
bool issubdtype(const Dtype::Category& cat, const Dtype& type) {
|
||||
bool issubdtype(const Dtype::Category& /* cat */, const Dtype& /* type */) {
|
||||
return false;
|
||||
}
|
||||
|
||||
|
||||
+23
-23
@@ -190,8 +190,8 @@ std::tuple<std::vector<PathNode>, size_t, int> greedy_path(
|
||||
|
||||
// Start by iterating over all possible combinations
|
||||
std::vector<std::pair<int, int>> pos_pairs;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int j = i + 1; j < inputs.size(); ++j) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
for (int j = i + 1; j < std::ssize(inputs); ++j) {
|
||||
pos_pairs.emplace_back(i, j);
|
||||
}
|
||||
}
|
||||
@@ -200,13 +200,13 @@ std::tuple<std::vector<PathNode>, size_t, int> greedy_path(
|
||||
std::vector<Contraction> possible_contractions;
|
||||
size_t path_cost = 0;
|
||||
int path_scaling = 0;
|
||||
auto num_in = inputs.size();
|
||||
auto num_in = std::ssize(inputs);
|
||||
for (int i = 0; i < num_in - 1; ++i) {
|
||||
auto add_contraction = [&](int p1, int p2) {
|
||||
CharSet new_term;
|
||||
CharSet contractions(inputs[p1].set.begin(), inputs[p1].set.end());
|
||||
contractions.insert(inputs[p2].set.begin(), inputs[p2].set.end());
|
||||
for (int i = 0; i < inputs.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(inputs); i++) {
|
||||
if (i == p1 || i == p2) {
|
||||
continue;
|
||||
}
|
||||
@@ -321,7 +321,7 @@ std::tuple<std::vector<PathNode>, size_t, int> greedy_path(
|
||||
}
|
||||
|
||||
pos_pairs.clear();
|
||||
for (int i = 0; i < inputs.size() - 1; ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs) - 1; ++i) {
|
||||
pos_pairs.emplace_back(i, inputs.size() - 1);
|
||||
}
|
||||
path_cost += best.cost;
|
||||
@@ -360,7 +360,7 @@ array batch_tensordot(
|
||||
{
|
||||
auto a_shape = a.shape();
|
||||
auto b_shape = b.shape();
|
||||
for (int i = 0; i < a_contract.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(a_contract); ++i) {
|
||||
auto d = std::max(a.shape(a_contract[i]), b.shape(b_contract[i]));
|
||||
a_shape[a_contract[i]] = d;
|
||||
b_shape[b_contract[i]] = d;
|
||||
@@ -430,7 +430,7 @@ array collapse_repeats(array in, Subscript& subscript, StreamOrDevice s) {
|
||||
std::string repeat_str;
|
||||
std::string no_repeat_str;
|
||||
std::unordered_map<char, int> counts;
|
||||
for (int i = 0; i < str.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(str); ++i) {
|
||||
auto [it, _] = counts.insert({str[i], 0});
|
||||
it->second++;
|
||||
}
|
||||
@@ -455,7 +455,7 @@ array collapse_repeats(array in, Subscript& subscript, StreamOrDevice s) {
|
||||
std::vector<array> indices;
|
||||
int n_expand = repeats.size();
|
||||
for (auto [c, v] : repeats) {
|
||||
for (int i = 0; i < str.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(str); ++i) {
|
||||
if (str[i] == c) {
|
||||
slice_sizes[i] = 1;
|
||||
axes.push_back(i);
|
||||
@@ -494,7 +494,7 @@ void preprocess_einsum_inputs(
|
||||
std::vector<array>& operands,
|
||||
StreamOrDevice s) {
|
||||
// Collapse repeat indices
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
auto& in = inputs[i];
|
||||
if (in.set.size() < in.str.size()) {
|
||||
operands[positions[i]] = collapse_repeats(operands[positions[i]], in, s);
|
||||
@@ -514,10 +514,10 @@ void preprocess_einsum_inputs(
|
||||
auto inserted = counts.insert({c, 0});
|
||||
inserted.first->second++;
|
||||
}
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
auto& in = inputs[i];
|
||||
std::vector<int> sum_axes;
|
||||
for (int ax = 0; ax < in.str.size(); ++ax) {
|
||||
for (int ax = 0; ax < std::ssize(in.str); ++ax) {
|
||||
if (counts[in.str[ax]] == 1) {
|
||||
sum_axes.push_back(ax);
|
||||
}
|
||||
@@ -549,12 +549,12 @@ array einsum_naive(
|
||||
}
|
||||
|
||||
// Expand and transpose inputs as needed
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
int pos = positions[i];
|
||||
auto& op = operands[pos];
|
||||
|
||||
// Add missing dimensions at the end
|
||||
if (op.ndim() != char_to_ax.size()) {
|
||||
if (op.ndim() != std::ssize(char_to_ax)) {
|
||||
auto shape = op.shape();
|
||||
shape.insert(shape.end(), char_to_ax.size() - shape.size(), 1);
|
||||
op = reshape(op, std::move(shape), s);
|
||||
@@ -597,7 +597,7 @@ array einsum_naive(
|
||||
|
||||
// Multiply and sum
|
||||
auto out = operands[positions[0]];
|
||||
for (int i = 1; i < positions.size(); ++i) {
|
||||
for (int i = 1; i < std::ssize(positions); ++i) {
|
||||
out = multiply(out, operands[positions[i]], s);
|
||||
}
|
||||
std::vector<int> sum_axes;
|
||||
@@ -675,9 +675,9 @@ std::pair<std::vector<PathNode>, PathInfo> einsum_path_helper(
|
||||
int operand_idx) {
|
||||
bool have_ellipsis = false;
|
||||
int cnt_before = 0, cnt_after = 0;
|
||||
for (int i = 0; i < subscript.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(subscript); i++) {
|
||||
if (!isalpha(subscript[i])) {
|
||||
if (i + 2 >= subscript.size() || subscript[i] != '.' ||
|
||||
if (i + 2 >= std::ssize(subscript) || subscript[i] != '.' ||
|
||||
subscript[i + 1] != '.' || subscript[i + 2] != '.') {
|
||||
std::ostringstream msg;
|
||||
msg << "[" << fn_name << "] Subscripts must be letters, but got '"
|
||||
@@ -732,7 +732,7 @@ std::pair<std::vector<PathNode>, PathInfo> einsum_path_helper(
|
||||
}
|
||||
};
|
||||
|
||||
for (int i = 0; i < operands.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(operands); i++) {
|
||||
check_letters_and_expand_ellipsis(in_subscripts[i], &operands[i], i);
|
||||
}
|
||||
check_letters_and_expand_ellipsis(out_subscript, nullptr, -1);
|
||||
@@ -747,12 +747,12 @@ std::pair<std::vector<PathNode>, PathInfo> einsum_path_helper(
|
||||
|
||||
std::unordered_map<char, ShapeElem> dim_map;
|
||||
std::vector<Subscript> inputs;
|
||||
for (int i = 0; i < in_subscripts.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(in_subscripts); ++i) {
|
||||
auto& in = in_subscripts[i];
|
||||
CharSet in_set(in.begin(), in.end());
|
||||
inputs.emplace_back(in, in_set);
|
||||
|
||||
if (in.size() != operands[i].ndim()) {
|
||||
if (std::ssize(in) != operands[i].ndim()) {
|
||||
std::ostringstream msg;
|
||||
msg << "[" << fn_name << "] Invalid number of subscripts " << in.size()
|
||||
<< " for input " << i << " with " << operands[i].ndim()
|
||||
@@ -763,7 +763,7 @@ std::pair<std::vector<PathNode>, PathInfo> einsum_path_helper(
|
||||
// Check repeat subscripts are valid
|
||||
if (in_set.size() < in.size()) {
|
||||
std::unordered_map<char, ShapeElem> local_dims;
|
||||
for (int j = 0; j < in.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(in); ++j) {
|
||||
auto dim = operands[i].shape(j);
|
||||
auto inserted = local_dims.insert({in[j], dim});
|
||||
if (!inserted.second) {
|
||||
@@ -778,7 +778,7 @@ std::pair<std::vector<PathNode>, PathInfo> einsum_path_helper(
|
||||
}
|
||||
}
|
||||
|
||||
for (int j = 0; j < in.size(); j++) {
|
||||
for (int j = 0; j < std::ssize(in); j++) {
|
||||
auto c = in[j];
|
||||
auto dim = operands[i].shape(j);
|
||||
auto inserted = dim_map.insert({c, dim});
|
||||
@@ -864,7 +864,7 @@ array einsum(
|
||||
std::vector<int> a_contract;
|
||||
std::vector<int> a_batch;
|
||||
std::vector<int> a_concat;
|
||||
for (int i = 0; i < in_a.str.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(in_a.str); ++i) {
|
||||
auto c = in_a.str[i];
|
||||
if (out.set.find(c) == out.set.end()) {
|
||||
// Not in the output, contraction
|
||||
@@ -887,7 +887,7 @@ array einsum(
|
||||
for (auto a_i : a_batch) {
|
||||
b_batch.push_back(in_b.str.find(in_a.str[a_i]));
|
||||
}
|
||||
for (int i = 0; i < in_b.str.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(in_b.str); ++i) {
|
||||
auto c = in_b.str[i];
|
||||
if (out.set.find(c) != out.set.end() &&
|
||||
in_a.set.find(c) == in_a.set.end()) {
|
||||
|
||||
+10
-9
@@ -138,7 +138,7 @@ T deserialize(Reader& is) {
|
||||
T v;
|
||||
auto size = deserialize<uint64_t>(is);
|
||||
v.reserve(size);
|
||||
for (int i = 0; i < size; ++i) {
|
||||
for (size_t i = 0; i < size; ++i) {
|
||||
v.push_back(deserialize<typename T::value_type>(is));
|
||||
}
|
||||
return v;
|
||||
@@ -487,11 +487,11 @@ struct FunctionTable {
|
||||
int n = 1;
|
||||
for (auto& [_, vec] : table) {
|
||||
for (auto& fun : vec) {
|
||||
auto npos = fun.inputs.size() - fun.kwarg_keys.size();
|
||||
auto npos = std::ssize(fun.inputs) - std::ssize(fun.kwarg_keys);
|
||||
os << " " << n++ << ". Function with " << npos
|
||||
<< " positional inputs and " << fun.kwarg_keys.size()
|
||||
<< " positional inputs and " << std::ssize(fun.kwarg_keys)
|
||||
<< " keyword inputs:\n";
|
||||
for (int j = 0; j < fun.inputs.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(fun.inputs); ++j) {
|
||||
auto& in = fun.inputs[j];
|
||||
if (j < npos) {
|
||||
os << " " << j + 1 << ": ";
|
||||
@@ -536,7 +536,7 @@ bool FunctionTable::match(
|
||||
};
|
||||
|
||||
int i = 0;
|
||||
for (; i < args.size(); ++i) {
|
||||
for (; i < std::ssize(args); ++i) {
|
||||
if (!match_inputs(args[i], fun.inputs[i])) {
|
||||
return false;
|
||||
}
|
||||
@@ -627,7 +627,8 @@ void FunctionExporter::export_with_callback(
|
||||
// Callback on the inputs
|
||||
callback({{"type", "inputs"}, {"inputs", to_vector_data(inputs)}});
|
||||
std::vector<std::pair<std::string, std::string>> keyword_inputs;
|
||||
for (int i = inputs.size() - kwarg_keys.size(), j = 0; i < inputs.size();
|
||||
for (int i = std::ssize(inputs) - std::ssize(kwarg_keys), j = 0;
|
||||
i < std::ssize(inputs);
|
||||
++i, ++j) {
|
||||
keyword_inputs.emplace_back(kwarg_keys[j], namer.get_name(inputs[i]));
|
||||
}
|
||||
@@ -928,7 +929,7 @@ std::vector<array> ImportedFunction::operator()(
|
||||
ftable->print_functions(msg);
|
||||
msg << "\nCalled with " << args.size() << " positional inputs and "
|
||||
<< kwargs.size() << " keyword inputs:\n";
|
||||
for (int i = 0; i < args.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(args); ++i) {
|
||||
auto& in = args[i];
|
||||
msg << " " << i + 1 << ": " << in.shape() << " " << in.dtype() << "\n";
|
||||
}
|
||||
@@ -970,7 +971,7 @@ ImportedFunction::ImportedFunction(const std::string& file)
|
||||
std::unordered_map<uint64_t, array> array_map;
|
||||
auto trace_input_ids = deserialize<std::vector<uint64_t>>(is);
|
||||
auto trace_inputs = deserialize<std::vector<array>>(is);
|
||||
for (int i = 0; i < trace_inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(trace_inputs); ++i) {
|
||||
array_map.emplace(trace_input_ids[i], trace_inputs[i]);
|
||||
}
|
||||
auto trace_output_ids = deserialize<std::vector<uint64_t>>(is);
|
||||
@@ -1006,7 +1007,7 @@ ImportedFunction::ImportedFunction(const std::string& file)
|
||||
std::move(types),
|
||||
std::move(prim),
|
||||
std::move(inputs));
|
||||
for (int i = 0; i < arrays.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(arrays); ++i) {
|
||||
auto sid = ids[i];
|
||||
if (sid == id) {
|
||||
tape.push_back(arrays[i]);
|
||||
|
||||
+9
-7
@@ -13,11 +13,11 @@ std::vector<array> Custom::vjp(
|
||||
const std::vector<array>& primals,
|
||||
const std::vector<array>& cotangents,
|
||||
const std::vector<int>& argnums,
|
||||
const std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* outputs */) {
|
||||
auto [_, vjps] = mlx::core::vjp(fallback_, primals, cotangents);
|
||||
std::vector<array> vjp_outs;
|
||||
for (int i = 0, j = 0; i < vjps.size(); ++i) {
|
||||
if (j < argnums.size() && i == argnums[j]) {
|
||||
for (int i = 0, j = 0; i < std::ssize(vjps); ++i) {
|
||||
if (j < std::ssize(argnums) && i == argnums[j]) {
|
||||
vjp_outs.push_back(vjps[i]);
|
||||
j++;
|
||||
}
|
||||
@@ -30,8 +30,8 @@ std::vector<array> Custom::jvp(
|
||||
const std::vector<array>& tangents,
|
||||
const std::vector<int>& argnums) {
|
||||
std::vector<array> all_tangents;
|
||||
for (int i = 0, j = 0; i < primals.size(); i++) {
|
||||
if (j < argnums.size() && i == argnums[j]) {
|
||||
for (int i = 0, j = 0; i < std::ssize(primals); i++) {
|
||||
if (j < std::ssize(argnums) && i == argnums[j]) {
|
||||
all_tangents.emplace_back(tangents[j++]);
|
||||
} else {
|
||||
all_tangents.emplace_back(zeros_like(primals[i]));
|
||||
@@ -127,6 +127,7 @@ std::vector<array> RMSNorm::vjp(
|
||||
assert(primals.size() == 2);
|
||||
assert(outputs.size() == 1);
|
||||
assert(cotangents.size() == 1);
|
||||
(void)outputs;
|
||||
|
||||
auto s = stream();
|
||||
auto fallback = [eps = eps_, s](const std::vector<array>& inputs) {
|
||||
@@ -269,6 +270,7 @@ std::vector<array> LayerNorm::vjp(
|
||||
assert(primals.size() == 3);
|
||||
assert(outputs.size() == 1);
|
||||
assert(cotangents.size() == 1);
|
||||
(void)outputs;
|
||||
|
||||
auto s = stream();
|
||||
auto fallback = [eps = eps_, s](const std::vector<array>& inputs) {
|
||||
@@ -536,7 +538,7 @@ std::vector<array> RoPE::vjp(
|
||||
const std::vector<array>& primals,
|
||||
const std::vector<array>& cotangents,
|
||||
const std::vector<int>& argnums,
|
||||
const std::vector<array>& outputs) {
|
||||
const std::vector<array>& /* outputs */) {
|
||||
auto s = stream();
|
||||
auto fallback = [dims = dims_,
|
||||
traditional = traditional_,
|
||||
@@ -635,7 +637,7 @@ array scaled_dot_product_attention(
|
||||
throw std::invalid_argument(msg.str());
|
||||
}
|
||||
|
||||
const size_t batch_dim = queries.shape(0);
|
||||
const int batch_dim = queries.shape(0);
|
||||
for (const auto& tensor : {keys, values}) {
|
||||
if (tensor.shape(0) != batch_dim) {
|
||||
std::ostringstream msg;
|
||||
|
||||
+21
-14
@@ -46,8 +46,9 @@ class RMSNorm : public Custom {
|
||||
|
||||
static bool use_fallback(Stream stream);
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("NYI");
|
||||
}
|
||||
void eval_gpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
@@ -79,8 +80,9 @@ class RMSNormVJP : public Custom {
|
||||
float eps)
|
||||
: Custom(stream, fallback), eps_(eps) {}
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("NYI");
|
||||
}
|
||||
void eval_gpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
@@ -106,8 +108,9 @@ class LayerNorm : public Custom {
|
||||
|
||||
static bool use_fallback(Stream s);
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("NYI");
|
||||
}
|
||||
void eval_gpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
@@ -138,8 +141,9 @@ class LayerNormVJP : public Custom {
|
||||
float eps)
|
||||
: Custom(stream, fallback), eps_(eps) {}
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("NYI");
|
||||
}
|
||||
void eval_gpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
@@ -174,8 +178,9 @@ class RoPE : public Custom {
|
||||
|
||||
static bool use_fallback(Stream s);
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("NYI");
|
||||
}
|
||||
void eval_gpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
@@ -225,8 +230,9 @@ class ScaledDotProductAttention : public Custom {
|
||||
bool do_causal,
|
||||
Stream s);
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("NYI");
|
||||
}
|
||||
|
||||
@@ -320,8 +326,9 @@ class CustomKernel : public Primitive {
|
||||
is_precompiled_(is_precompiled),
|
||||
shared_memory_(shared_memory) {}
|
||||
|
||||
void eval_cpu(const std::vector<array>& inputs, std::vector<array>& outputs)
|
||||
override {
|
||||
void eval_cpu(
|
||||
const std::vector<array>& /* inputs */,
|
||||
std::vector<array>& /* outputs */) override {
|
||||
throw std::runtime_error("Custom kernels only run on GPU.");
|
||||
}
|
||||
|
||||
|
||||
+2
-2
@@ -20,7 +20,7 @@ array fft_impl(
|
||||
throw std::invalid_argument(
|
||||
"[fftn] Requires array with at least one dimension.");
|
||||
}
|
||||
if (n.size() != axes.size()) {
|
||||
if (n.size() != std::ssize(axes)) {
|
||||
throw std::invalid_argument("[fftn] Shape and axes have different sizes.");
|
||||
}
|
||||
if (axes.empty()) {
|
||||
@@ -59,7 +59,7 @@ array fft_impl(
|
||||
}
|
||||
|
||||
auto in_shape = a.shape();
|
||||
for (int i = 0; i < valid_axes.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(valid_axes); ++i) {
|
||||
in_shape[valid_axes[i]] = n[i];
|
||||
}
|
||||
if (real && inverse) {
|
||||
|
||||
+2
-2
@@ -238,7 +238,7 @@ std::unordered_map<std::string, array> load_arrays(gguf_ctx* ctx) {
|
||||
return array_map;
|
||||
}
|
||||
|
||||
GGUFLoad load_gguf(const std::string& file, StreamOrDevice s) {
|
||||
GGUFLoad load_gguf(const std::string& file, StreamOrDevice /* s */) {
|
||||
bool exists;
|
||||
{
|
||||
std::ifstream f(file.c_str());
|
||||
@@ -440,7 +440,7 @@ void save_gguf(
|
||||
}
|
||||
const char* tensorname = key.c_str();
|
||||
const uint64_t namelen = key.length();
|
||||
const uint32_t num_dim = arr.ndim();
|
||||
const int num_dim = arr.ndim();
|
||||
std::vector<uint64_t> dim(num_dim);
|
||||
for (int i = 0; i < num_dim; i++) {
|
||||
dim[i] = arr.shape()[num_dim - 1 - i];
|
||||
|
||||
@@ -77,8 +77,8 @@ void extract_q8_0_data(
|
||||
array& weights_arr,
|
||||
array& scales_arr,
|
||||
array& biases_arr) {
|
||||
const uint64_t weights_per_block = 32;
|
||||
const uint64_t bytes_per_block = 34; // 2 bytes scale, 32x1 byte weights
|
||||
const int64_t weights_per_block = 32;
|
||||
const int64_t bytes_per_block = 34; // 2 bytes scale, 32x1 byte weights
|
||||
auto data = static_cast<uint8_t*>(tensor.weights_data);
|
||||
auto weights = weights_arr.data<int8_t>();
|
||||
auto scales = scales_arr.data<float16_t>();
|
||||
|
||||
+26
-26
@@ -390,7 +390,7 @@ array unflatten(
|
||||
throw std::invalid_argument(msg.str());
|
||||
}
|
||||
|
||||
size_t size = 1;
|
||||
int64_t size = 1;
|
||||
int infer_idx = -1;
|
||||
for (int i = 0; i < shape.size(); ++i) {
|
||||
if (shape[i] == -1) {
|
||||
@@ -687,10 +687,10 @@ void normalize_dynamic_slice_inputs(
|
||||
<< ".";
|
||||
throw std::invalid_argument(msg.str());
|
||||
}
|
||||
if (start.size() != axes.size()) {
|
||||
if (start.size() != std::ssize(axes)) {
|
||||
std::ostringstream msg;
|
||||
msg << prefix << " Number of starting indices " << start.size()
|
||||
<< " does not match number of axes " << axes.size() << ".";
|
||||
<< " does not match number of axes " << std::ssize(axes) << ".";
|
||||
throw std::invalid_argument(msg.str());
|
||||
}
|
||||
if (!issubdtype(start.dtype(), integer)) {
|
||||
@@ -847,7 +847,7 @@ array slice_update(
|
||||
|
||||
// Broadcast update with unspecified axes
|
||||
auto up_shape = update.shape();
|
||||
auto dim_diff = std::max(src.ndim() - update.ndim(), size_t(0));
|
||||
auto dim_diff = std::max(src.ndim() - update.ndim(), 0);
|
||||
up_shape.insert(
|
||||
up_shape.begin(), src.shape().begin(), src.shape().begin() + dim_diff);
|
||||
for (int d = dim_diff; d < src.ndim(); ++d) {
|
||||
@@ -957,7 +957,7 @@ std::vector<array> meshgrid(
|
||||
"[meshgrid] Invalid indexing value. Valid values are 'xy' and 'ij'.");
|
||||
}
|
||||
|
||||
auto ndim = arrays.size();
|
||||
auto ndim = std::ssize(arrays);
|
||||
std::vector<array> outputs;
|
||||
for (int i = 0; i < ndim; ++i) {
|
||||
Shape shape(ndim, 1);
|
||||
@@ -1135,10 +1135,10 @@ array tile(
|
||||
std::vector<int> reps,
|
||||
StreamOrDevice s /* = {} */) {
|
||||
auto shape = arr.shape();
|
||||
if (reps.size() < shape.size()) {
|
||||
if (std::ssize(reps) < shape.size()) {
|
||||
reps.insert(reps.begin(), shape.size() - reps.size(), 1);
|
||||
}
|
||||
if (reps.size() > shape.size()) {
|
||||
if (std::ssize(reps) > shape.size()) {
|
||||
shape.insert(shape.begin(), reps.size() - shape.size(), 1);
|
||||
}
|
||||
|
||||
@@ -1162,7 +1162,7 @@ array tile(
|
||||
|
||||
array edge_pad(
|
||||
const array& a,
|
||||
const std::vector<int>& axes,
|
||||
const std::vector<int>& /* axes */,
|
||||
const Shape& low_pad_size,
|
||||
const Shape& high_pad_size,
|
||||
const Shape& out_shape,
|
||||
@@ -1214,17 +1214,17 @@ array pad(
|
||||
const array& pad_value /*= array(0)*/,
|
||||
const std::string& mode /*= "constant"*/,
|
||||
StreamOrDevice s /* = {}*/) {
|
||||
if (axes.size() != low_pad_size.size() ||
|
||||
axes.size() != high_pad_size.size()) {
|
||||
if (std::ssize(axes) != low_pad_size.size() ||
|
||||
std::ssize(axes) != high_pad_size.size()) {
|
||||
std::ostringstream msg;
|
||||
msg << "Invalid number of padding sizes passed to pad "
|
||||
<< "with axes of size " << axes.size();
|
||||
<< "with axes of size " << std::ssize(axes);
|
||||
throw std::invalid_argument(msg.str());
|
||||
}
|
||||
|
||||
auto out_shape = a.shape();
|
||||
|
||||
for (int i = 0; i < axes.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(axes); i++) {
|
||||
if (low_pad_size[i] < 0) {
|
||||
std::ostringstream msg;
|
||||
msg << "Invalid low padding size (" << low_pad_size[i]
|
||||
@@ -1365,7 +1365,7 @@ array transpose(
|
||||
for (auto& ax : axes) {
|
||||
ax = ax < 0 ? ax + a.ndim() : ax;
|
||||
}
|
||||
if (axes.size() != a.ndim()) {
|
||||
if (std::ssize(axes) != a.ndim()) {
|
||||
std::ostringstream msg;
|
||||
msg << "[transpose] Recived " << axes.size() << " axes for array with "
|
||||
<< a.ndim() << " dimensions.";
|
||||
@@ -1387,7 +1387,7 @@ array transpose(
|
||||
shape[ax] = 1;
|
||||
}
|
||||
|
||||
for (int i = 0; i < axes.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(axes); ++i) {
|
||||
shape[i] = a.shape()[axes[i]];
|
||||
}
|
||||
return array(
|
||||
@@ -1444,7 +1444,7 @@ std::vector<array> broadcast_arrays(
|
||||
auto shape = BroadcastAxes::output_shape(inputs, ignore_axes);
|
||||
auto check_and_get_shape = [&shape, &ignore_axes](const array& in) {
|
||||
auto out_shape = shape;
|
||||
for (int i = 0; i < ignore_axes.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(ignore_axes); ++i) {
|
||||
auto ax = ignore_axes[i];
|
||||
auto pos_ax = in.ndim() + ax;
|
||||
if (pos_ax < 0 || pos_ax > in.ndim() ||
|
||||
@@ -1478,7 +1478,7 @@ std::vector<array> broadcast_arrays(
|
||||
stop_grad_inputs.push_back(stop_gradient(in, s));
|
||||
}
|
||||
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
auto& in = inputs[i];
|
||||
auto out_shape = check_and_get_shape(in);
|
||||
if (in.shape() == out_shape) {
|
||||
@@ -1486,7 +1486,7 @@ std::vector<array> broadcast_arrays(
|
||||
} else {
|
||||
// broadcasted array goes first followed by other stopgrad inputs
|
||||
std::vector<array> p_inputs = {in};
|
||||
for (int j = 0; j < inputs.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(inputs); ++j) {
|
||||
if (j == i) {
|
||||
continue;
|
||||
}
|
||||
@@ -1530,14 +1530,14 @@ std::vector<array> broadcast_arrays(
|
||||
for (auto& in : inputs) {
|
||||
stop_grad_inputs.push_back(stop_gradient(in, s));
|
||||
}
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
auto& in = inputs[i];
|
||||
if (in.shape() == shape) {
|
||||
outputs.push_back(in);
|
||||
} else {
|
||||
// broadcasted array goes first followed by other stopgrad inputs
|
||||
std::vector<array> p_inputs = {in};
|
||||
for (int j = 0; j < inputs.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(inputs); ++j) {
|
||||
if (j == i) {
|
||||
continue;
|
||||
}
|
||||
@@ -1961,7 +1961,7 @@ array median(
|
||||
auto dtype = at_least_float(a.dtype());
|
||||
std::vector<int> transpose_axes;
|
||||
for (int i = 0, j = 0; i < a.ndim(); ++i) {
|
||||
if (j < sorted_axes.size() && i == sorted_axes[j]) {
|
||||
if (j < std::ssize(sorted_axes) && i == sorted_axes[j]) {
|
||||
j++;
|
||||
continue;
|
||||
}
|
||||
@@ -3010,7 +3010,7 @@ array gather(
|
||||
const Shape& slice_sizes,
|
||||
StreamOrDevice s /* = {} */) {
|
||||
// Checks that indices, dimensions, and slice_sizes are all valid
|
||||
if (indices.size() > a.ndim()) {
|
||||
if (std::ssize(indices) > a.ndim()) {
|
||||
std::ostringstream msg;
|
||||
msg << "[gather] Too many index arrays. Got " << indices.size()
|
||||
<< " index arrays for input with " << a.ndim() << " dimensions.";
|
||||
@@ -3312,7 +3312,7 @@ array scatter(
|
||||
Scatter::ReduceType mode,
|
||||
StreamOrDevice s) {
|
||||
// Checks that indices, dimensions, and slice_sizes are all valid
|
||||
if (indices.size() > a.ndim()) {
|
||||
if (std::ssize(indices) > a.ndim()) {
|
||||
std::ostringstream msg;
|
||||
msg << "[scatter] Too many index arrays. Got " << indices.size()
|
||||
<< " index arrays for input with " << a.ndim() << " dimensions.";
|
||||
@@ -3820,7 +3820,7 @@ array conv_transpose_general(
|
||||
StreamOrDevice s) {
|
||||
std::vector<int> padding_lo(padding.size());
|
||||
std::vector<int> padding_hi(padding.size());
|
||||
for (int i = 0; i < padding.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(padding); ++i) {
|
||||
int wt_size = 1 + dilation[i] * (weight.shape(1 + i) - 1);
|
||||
padding_lo[i] = wt_size - padding[i] - 1;
|
||||
|
||||
@@ -4632,7 +4632,7 @@ array tensordot(
|
||||
int csize = 1;
|
||||
auto x = a;
|
||||
auto y = b;
|
||||
for (int i = 0; i < axes_a.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(axes_a); i++) {
|
||||
if (x.shape(axes_a.at(i)) == y.shape(axes_b.at(i))) {
|
||||
csize *= x.shape(axes_a.at(i));
|
||||
} else {
|
||||
@@ -5560,7 +5560,7 @@ array roll(
|
||||
return a;
|
||||
}
|
||||
|
||||
if (shift.size() < axes.size()) {
|
||||
if (shift.size() < std::ssize(axes)) {
|
||||
std::ostringstream msg;
|
||||
msg << "[roll] At least one shift value per axis is required, "
|
||||
<< shift.size() << " provided for " << axes.size() << " axes.";
|
||||
@@ -5568,7 +5568,7 @@ array roll(
|
||||
}
|
||||
|
||||
array result = a;
|
||||
for (int i = 0; i < axes.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(axes); i++) {
|
||||
int ax = axes[i];
|
||||
if (ax < 0) {
|
||||
ax += a.ndim();
|
||||
|
||||
+179
-111
File diff suppressed because it is too large
Load Diff
+7
-5
@@ -31,9 +31,9 @@
|
||||
return #PRIMITIVE; \
|
||||
}
|
||||
|
||||
#define DEFINE_DEFAULT_IS_EQUIVALENT() \
|
||||
bool is_equivalent(const Primitive& other) const override { \
|
||||
return true; \
|
||||
#define DEFINE_DEFAULT_IS_EQUIVALENT() \
|
||||
bool is_equivalent(const Primitive& /* other */) const override { \
|
||||
return true; \
|
||||
}
|
||||
|
||||
#define DEFINE_INPUT_OUTPUT_SHAPE() \
|
||||
@@ -104,7 +104,7 @@ class Primitive {
|
||||
virtual const char* name() const = 0;
|
||||
|
||||
/** Equivalence check defaults to false unless overridden by the primitive */
|
||||
virtual bool is_equivalent(const Primitive& other) const {
|
||||
virtual bool is_equivalent(const Primitive& /* other */) const {
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -1071,6 +1071,7 @@ class FFT : public UnaryPrimitive {
|
||||
public:
|
||||
explicit FFT(
|
||||
Stream stream,
|
||||
// Note: PocketFFT requires size_t
|
||||
const std::vector<size_t>& axes,
|
||||
bool inverse,
|
||||
bool real)
|
||||
@@ -1526,7 +1527,8 @@ class NumberOfElements : public UnaryPrimitive {
|
||||
DEFINE_VMAP()
|
||||
DEFINE_NAME(NumberOfElements)
|
||||
bool is_equivalent(const Primitive& other) const override;
|
||||
std::vector<Shape> output_shapes(const std::vector<array>& inputs) override {
|
||||
std::vector<Shape> output_shapes(
|
||||
const std::vector<array>& /* inputs */) override {
|
||||
return {{}};
|
||||
}
|
||||
std::tuple<std::vector<int>, bool, Dtype> state() const {
|
||||
|
||||
@@ -89,6 +89,7 @@ inline array uniform(
|
||||
const Shape& shape,
|
||||
const std::optional<array>& key = std::nullopt,
|
||||
StreamOrDevice s = {}) {
|
||||
(void)s;
|
||||
return uniform(shape, float32, key);
|
||||
}
|
||||
|
||||
|
||||
+2
-2
@@ -103,7 +103,7 @@ class Scheduler {
|
||||
default_streams_.at(s.device.type) = s;
|
||||
}
|
||||
|
||||
void notify_new_task(const Stream& stream) {
|
||||
void notify_new_task(const Stream& /* stream */) {
|
||||
{
|
||||
std::lock_guard<std::mutex> lk(mtx);
|
||||
n_active_tasks_++;
|
||||
@@ -111,7 +111,7 @@ class Scheduler {
|
||||
completion_cv.notify_all();
|
||||
}
|
||||
|
||||
void notify_task_completion(const Stream& stream) {
|
||||
void notify_task_completion(const Stream& /* stream */) {
|
||||
{
|
||||
std::lock_guard<std::mutex> lk(mtx);
|
||||
n_active_tasks_--;
|
||||
|
||||
+18
-21
@@ -121,10 +121,10 @@ class SmallVector {
|
||||
std::initializer_list<T> init,
|
||||
const Allocator& allocator = Allocator())
|
||||
: allocator_(allocator) {
|
||||
if (init.size() > capacity()) {
|
||||
if (static_cast<int>(init.size()) > capacity()) {
|
||||
grow(init.size());
|
||||
}
|
||||
assert(capacity() >= init.size()); // sanity check
|
||||
assert(capacity() >= static_cast<int>(init.size())); // sanity check
|
||||
std::uninitialized_move(init.begin(), init.end(), begin_);
|
||||
end_ = begin_ + init.size();
|
||||
}
|
||||
@@ -132,7 +132,7 @@ class SmallVector {
|
||||
template <typename Iter, typename = std::enable_if_t<is_iterator_v<Iter>>>
|
||||
SmallVector(Iter begin, Iter end, const Allocator& allocator = Allocator())
|
||||
: allocator_(allocator) {
|
||||
size_t size = std::distance(begin, end);
|
||||
int size = std::distance(begin, end);
|
||||
if (size > capacity()) {
|
||||
grow(size);
|
||||
}
|
||||
@@ -164,7 +164,7 @@ class SmallVector {
|
||||
if (this == &other) {
|
||||
return *this;
|
||||
}
|
||||
size_t other_size = other.size();
|
||||
int other_size = other.size();
|
||||
if (capacity() < other_size) {
|
||||
// Create large-enough heap-allocated storage.
|
||||
free_storage();
|
||||
@@ -273,13 +273,13 @@ class SmallVector {
|
||||
return std::make_reverse_iterator(begin_);
|
||||
}
|
||||
|
||||
size_t size() const {
|
||||
int size() const {
|
||||
return end_ - begin_;
|
||||
}
|
||||
bool empty() const {
|
||||
return end_ == begin_;
|
||||
}
|
||||
size_t capacity() const {
|
||||
int capacity() const {
|
||||
return end_of_storage_ - begin_;
|
||||
}
|
||||
|
||||
@@ -301,21 +301,21 @@ class SmallVector {
|
||||
return end_[-1];
|
||||
}
|
||||
|
||||
T& at(size_t index) {
|
||||
if (index >= size()) {
|
||||
T& at(int index) {
|
||||
if (index < 0 || index >= size()) {
|
||||
throw std::out_of_range("SmallVector out of range.");
|
||||
}
|
||||
return begin_[index];
|
||||
}
|
||||
const T& at(size_t index) const {
|
||||
const T& at(int index) const {
|
||||
return const_cast<SmallVector*>(this)->at(index);
|
||||
}
|
||||
|
||||
T& operator[](size_t index) {
|
||||
assert(size() > index);
|
||||
T& operator[](int index) {
|
||||
assert(index >= 0 && size() > index);
|
||||
return begin_[index];
|
||||
}
|
||||
const T& operator[](size_t index) const {
|
||||
const T& operator[](int index) const {
|
||||
return const_cast<SmallVector*>(this)->operator[](index);
|
||||
}
|
||||
|
||||
@@ -333,7 +333,7 @@ class SmallVector {
|
||||
emplace_back(std::move(x));
|
||||
}
|
||||
|
||||
void pop_back(size_t count = 1) {
|
||||
void pop_back(int count = 1) {
|
||||
assert(size() >= count);
|
||||
end_ -= count;
|
||||
std::destroy_n(end_, count);
|
||||
@@ -400,7 +400,7 @@ class SmallVector {
|
||||
return erase(pos, pos + 1);
|
||||
}
|
||||
|
||||
void resize(size_t new_size) {
|
||||
void resize(int new_size) {
|
||||
if (new_size > capacity()) {
|
||||
grow(new_size);
|
||||
}
|
||||
@@ -415,7 +415,7 @@ class SmallVector {
|
||||
end_ = new_end;
|
||||
}
|
||||
|
||||
void resize(size_t new_size, const T& initial_value) {
|
||||
void resize(int new_size, const T& initial_value) {
|
||||
if (new_size > capacity()) {
|
||||
grow(new_size);
|
||||
}
|
||||
@@ -428,7 +428,7 @@ class SmallVector {
|
||||
end_ = new_end;
|
||||
}
|
||||
|
||||
void reserve(size_t new_capacity) {
|
||||
void reserve(int new_capacity) {
|
||||
if (new_capacity > capacity()) {
|
||||
grow(new_capacity);
|
||||
}
|
||||
@@ -443,8 +443,8 @@ class SmallVector {
|
||||
private:
|
||||
// Grows the backing store by a factor of two, and at least to {min_capacity}.
|
||||
// TODO: Move to private after removing external code using this method.
|
||||
MLX_NOINLINE void grow(size_t min_capacity = 0) {
|
||||
size_t new_capacity = std::max(min_capacity, 2 * capacity());
|
||||
MLX_NOINLINE void grow(int min_capacity = 0) {
|
||||
int new_capacity = std::max(min_capacity, 2 * capacity());
|
||||
// Round up to power of 2.
|
||||
new_capacity--;
|
||||
new_capacity |= new_capacity >> 1;
|
||||
@@ -452,9 +452,6 @@ class SmallVector {
|
||||
new_capacity |= new_capacity >> 4;
|
||||
new_capacity |= new_capacity >> 8;
|
||||
new_capacity |= new_capacity >> 16;
|
||||
if constexpr (sizeof(size_t) == sizeof(uint64_t)) {
|
||||
new_capacity |= new_capacity >> 32;
|
||||
}
|
||||
new_capacity++;
|
||||
|
||||
T* new_storage = allocator_.allocate(new_capacity);
|
||||
|
||||
+28
-28
@@ -89,7 +89,7 @@ array eval_impl(std::vector<array> outputs, bool async) {
|
||||
auto& [a_ref, idx] = dfs.top();
|
||||
auto& a = a_ref.get();
|
||||
|
||||
if (idx < a.inputs().size()) {
|
||||
if (idx < std::ssize(a.inputs())) {
|
||||
// Add an input, and continue
|
||||
auto& in = a.inputs()[idx++];
|
||||
|
||||
@@ -146,16 +146,16 @@ array eval_impl(std::vector<array> outputs, bool async) {
|
||||
int max_width = env::bfs_max_width();
|
||||
dfs = std::stack<std::pair<std::reference_wrapper<array>, int>>();
|
||||
tape.push_back(synchronizer);
|
||||
for (int i = 0; !cache.empty() && (i < tape.size() || !dfs.empty());) {
|
||||
auto& a = (i >= tape.size()) ? dfs.top().first.get() : tape[i];
|
||||
for (int i = 0; !cache.empty() && (i < std::ssize(tape) || !dfs.empty());) {
|
||||
auto& a = (i >= std::ssize(tape)) ? dfs.top().first.get() : tape[i];
|
||||
int j = 0;
|
||||
if (i >= tape.size()) {
|
||||
if (i >= std::ssize(tape)) {
|
||||
j = dfs.top().second;
|
||||
dfs.pop();
|
||||
} else {
|
||||
i++;
|
||||
}
|
||||
for (; j < a.inputs().size(); ++j) {
|
||||
for (; j < std::ssize(a.inputs()); ++j) {
|
||||
auto& in = a.inputs()[j];
|
||||
if (in.status() != array::Status::unscheduled) {
|
||||
continue;
|
||||
@@ -163,7 +163,7 @@ array eval_impl(std::vector<array> outputs, bool async) {
|
||||
|
||||
// If the width limit is exceeded, push the array on the stack
|
||||
// and go down a level
|
||||
if ((tape.size() - i) >= max_width) {
|
||||
if ((std::ssize(tape) - i) >= max_width) {
|
||||
dfs.emplace(a, j);
|
||||
break;
|
||||
}
|
||||
@@ -343,14 +343,14 @@ std::pair<std::vector<array>, std::vector<array>> vjp(
|
||||
// that have stop_gradient called on them
|
||||
int cotan_index = 0;
|
||||
std::vector<std::pair<int, int>> output_cotan_pairs;
|
||||
for (int i = 0; i < outputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(outputs); ++i) {
|
||||
auto& out = outputs[i];
|
||||
if (out.has_primitive()) {
|
||||
if (auto& p = out.primitive(); typeid(p) == typeid(StopGradient)) {
|
||||
continue;
|
||||
}
|
||||
}
|
||||
if (cotan_index >= cotans.size()) {
|
||||
if (cotan_index >= std::ssize(cotans)) {
|
||||
std::ostringstream msg;
|
||||
msg << "[vjp] Number of outputs to compute gradients for ("
|
||||
<< outputs.size() << ") does not match number of cotangents ("
|
||||
@@ -374,11 +374,11 @@ std::pair<std::vector<array>, std::vector<array>> vjp(
|
||||
// to the tape which need a gradient.
|
||||
std::unordered_set<std::uintptr_t> cache;
|
||||
std::unordered_set<std::uintptr_t> calc_grad;
|
||||
for (int i = 0, j = 0; i < primals_.size(); ++i) {
|
||||
for (int i = 0, j = 0; i < std::ssize(primals_); ++i) {
|
||||
auto& primal = primals_[i];
|
||||
primal.set_tracer(false);
|
||||
cache.insert(primal.id());
|
||||
if (j < argnums.size() && argnums[j] == i) {
|
||||
if (j < std::ssize(argnums) && argnums[j] == i) {
|
||||
j++;
|
||||
calc_grad.insert(primal.id());
|
||||
}
|
||||
@@ -440,7 +440,7 @@ std::pair<std::vector<array>, std::vector<array>> vjp(
|
||||
|
||||
// Get the arguments whose gradients are needed
|
||||
std::vector<int> argnums;
|
||||
for (int i = 0; i < a.inputs().size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(a.inputs()); ++i) {
|
||||
if (calc_grad.find(a.inputs()[i].id()) != calc_grad.end()) {
|
||||
argnums.push_back(i);
|
||||
}
|
||||
@@ -473,7 +473,7 @@ std::pair<std::vector<array>, std::vector<array>> vjp(
|
||||
vjps = a.primitive().vjp(a.inputs(), cotangents, argnums, outputs);
|
||||
}
|
||||
// Accumulate the vector-jacobian products for each input
|
||||
for (int i = 0; i < argnums.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(argnums); ++i) {
|
||||
auto in_id = a.inputs()[argnums[i]].id();
|
||||
if (auto cotan_it = cotan_map.find(in_id); cotan_it != cotan_map.end()) {
|
||||
cotan_it->second = add(cotan_it->second, vjps[i], s);
|
||||
@@ -528,7 +528,7 @@ std::pair<std::vector<array>, std::vector<array>> jvp(
|
||||
throw std::invalid_argument(
|
||||
"[jvp] Number of inputs does not match number of tangents.");
|
||||
}
|
||||
for (int i = 0; i < primals.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(primals); ++i) {
|
||||
if (primals[i].shape() != tangents[i].shape()) {
|
||||
throw std::invalid_argument(
|
||||
"[jvp] Input shape does not match shape of tangent.");
|
||||
@@ -597,7 +597,7 @@ std::pair<std::vector<array>, std::vector<array>> jvp(
|
||||
}
|
||||
|
||||
std::unordered_map<std::uintptr_t, array> tan_map;
|
||||
for (int i = 0; i < primals_.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(primals_); ++i) {
|
||||
tan_map.insert({primals_[i].id(), tangents[i]});
|
||||
}
|
||||
|
||||
@@ -605,7 +605,7 @@ std::pair<std::vector<array>, std::vector<array>> jvp(
|
||||
// Get the arguments used in the jvp
|
||||
std::vector<int> argnums;
|
||||
std::vector<array> tangents;
|
||||
for (int i = 0; i < a.inputs().size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(a.inputs()); ++i) {
|
||||
if (auto it = tan_map.find(a.inputs()[i].id()); it != tan_map.end()) {
|
||||
argnums.push_back(i);
|
||||
tangents.push_back(it->second);
|
||||
@@ -614,7 +614,7 @@ std::pair<std::vector<array>, std::vector<array>> jvp(
|
||||
|
||||
auto jvps = a.primitive().jvp(a.inputs(), tangents, argnums);
|
||||
auto outputs = a.outputs();
|
||||
for (int i = 0; i < jvps.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(jvps); ++i) {
|
||||
tan_map.insert({outputs[i].id(), jvps[i]});
|
||||
}
|
||||
}
|
||||
@@ -658,7 +658,7 @@ ValueAndGradFn value_and_grad(
|
||||
throw std::invalid_argument(
|
||||
"[grad] Repeat argument number not allowed in grad.");
|
||||
}
|
||||
if (*args.begin() < 0 || *args.rbegin() >= inputs.size()) {
|
||||
if (*args.begin() < 0 || *args.rbegin() >= std::ssize(inputs)) {
|
||||
std::ostringstream msg;
|
||||
msg << "[grad] Invalid argument number for function with "
|
||||
<< inputs.size() << " inputs.";
|
||||
@@ -668,7 +668,7 @@ ValueAndGradFn value_and_grad(
|
||||
|
||||
auto gfun = [&fun](const std::vector<array>& inputs) {
|
||||
auto outputs = fun(inputs);
|
||||
for (int i = 1; i < outputs.size(); i++) {
|
||||
for (int i = 1; i < std::ssize(outputs); i++) {
|
||||
auto& out = outputs[i];
|
||||
auto s = out.has_primitive() ? out.primitive().stream()
|
||||
: default_stream(default_device());
|
||||
@@ -701,7 +701,7 @@ std::pair<std::vector<array>, std::vector<array>> vmap_trace(
|
||||
|
||||
// Some error checking and get the vmap axis size
|
||||
size_t vmap_ax_size;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
if (in_axes[i] != -1) {
|
||||
if (inputs[i].ndim() == 0) {
|
||||
throw std::invalid_argument(
|
||||
@@ -717,7 +717,7 @@ std::pair<std::vector<array>, std::vector<array>> vmap_trace(
|
||||
}
|
||||
}
|
||||
// Check that all vmapped axes have the same size
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
if (in_axes[i] != -1) {
|
||||
if (size_t in_ax = inputs[i].shape(in_axes[i]); vmap_ax_size != in_ax) {
|
||||
std::ostringstream msg;
|
||||
@@ -731,7 +731,7 @@ std::pair<std::vector<array>, std::vector<array>> vmap_trace(
|
||||
// Run the function on placeholder inputs
|
||||
// to get the original graph
|
||||
std::vector<array> s_inputs;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
if (in_axes[i] != -1) {
|
||||
auto shape = inputs[i].shape();
|
||||
shape.erase(shape.begin() + in_axes[i]);
|
||||
@@ -759,7 +759,7 @@ std::vector<array> vmap_replace(
|
||||
}
|
||||
|
||||
int vmap_size = -1;
|
||||
for (int i = 0; i < inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(inputs); ++i) {
|
||||
if (in_axes[i] >= 0) {
|
||||
vmap_size = inputs[i].shape(in_axes[i]);
|
||||
break;
|
||||
@@ -772,7 +772,7 @@ std::vector<array> vmap_replace(
|
||||
std::unordered_map<std::uintptr_t, std::pair<array, int>> tmap;
|
||||
std::unordered_set<std::uintptr_t> needs_vmap;
|
||||
std::unordered_set<std::uintptr_t> cache;
|
||||
for (int i = 0; i < s_inputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(s_inputs); ++i) {
|
||||
auto in = s_inputs[i];
|
||||
if (in_axes[i] != -1) {
|
||||
tmap.insert({in.id(), {inputs[i], in_axes[i]}});
|
||||
@@ -843,7 +843,7 @@ std::vector<array> vmap_replace(
|
||||
|
||||
// For each primitive's outputs add its id, the vout id and the vax
|
||||
auto outputs = a.outputs();
|
||||
for (int i = 0; i < v_outputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(v_outputs); ++i) {
|
||||
tmap.insert({outputs[i].id(), {v_outputs[i], v_out_axes[i]}});
|
||||
}
|
||||
}
|
||||
@@ -851,7 +851,7 @@ std::vector<array> vmap_replace(
|
||||
// Populate the outputs and make sure all the output axes are
|
||||
// in the right place
|
||||
std::vector<array> outputs;
|
||||
for (int i = 0; i < s_outputs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(s_outputs); ++i) {
|
||||
if (auto map_it = tmap.find(s_outputs[i].id()); map_it != tmap.end()) {
|
||||
auto& [out, vdim] = map_it->second;
|
||||
if (vdim != out_axes[i]) {
|
||||
@@ -995,7 +995,7 @@ std::function<std::vector<array>(const std::vector<array>&)> custom_function(
|
||||
// using `fun` directly because we may not be able to fully reuse
|
||||
// the outputs of the forward pass.
|
||||
fun_vjp.value_or(
|
||||
[fun](auto primals, auto cotangents, auto outputs) {
|
||||
[fun](auto primals, auto cotangents, auto /* outputs */) {
|
||||
auto [__, vjps] = vjp(fun, primals, cotangents);
|
||||
return vjps;
|
||||
}),
|
||||
@@ -1009,8 +1009,8 @@ std::function<std::vector<array>(const std::vector<array>&)> custom_function(
|
||||
// waste computation.
|
||||
fun_jvp.value_or([fun](auto primals, auto tangents, auto argnums) {
|
||||
std::vector<array> all_tangents;
|
||||
for (int i = 0, j = 0; i < primals.size(); i++) {
|
||||
if (j < argnums.size() && i == argnums[j]) {
|
||||
for (int i = 0, j = 0; i < std::ssize(primals); i++) {
|
||||
if (j < std::ssize(argnums) && i == argnums[j]) {
|
||||
all_tangents.emplace_back(tangents[j++]);
|
||||
} else {
|
||||
all_tangents.emplace_back(zeros_like(primals[i]));
|
||||
|
||||
@@ -24,9 +24,6 @@ struct _MLX_BFloat16 {
|
||||
// Default constructor
|
||||
_MLX_BFloat16() = default;
|
||||
|
||||
// Default copy constructor
|
||||
_MLX_BFloat16(_MLX_BFloat16 const&) = default;
|
||||
|
||||
// Appease std::vector<bool> for being special
|
||||
_MLX_BFloat16& operator=(std::vector<bool>::reference x) {
|
||||
bits_ = x;
|
||||
|
||||
+2
-2
@@ -138,8 +138,8 @@ namespace env {
|
||||
|
||||
int get_var(const char* name, int default_value);
|
||||
|
||||
inline int bfs_max_width() {
|
||||
static int bfs_max_width_ = get_var("MLX_BFS_MAX_WIDTH", 20);
|
||||
inline unsigned int bfs_max_width() {
|
||||
static unsigned int bfs_max_width_ = get_var("MLX_BFS_MAX_WIDTH", 20);
|
||||
return bfs_max_width_;
|
||||
}
|
||||
|
||||
|
||||
@@ -83,7 +83,7 @@ class ArrayPythonIterator {
|
||||
throw nb::stop_iteration();
|
||||
}
|
||||
|
||||
if (idx_ >= 0 && idx_ < splits_.size()) {
|
||||
if (idx_ >= 0 && idx_ < std::ssize(splits_)) {
|
||||
return mx::squeeze(splits_[idx_++], 0);
|
||||
}
|
||||
|
||||
@@ -390,7 +390,7 @@ void init_array(nb::module_& m) {
|
||||
)pbdoc")
|
||||
.def(
|
||||
"__array_namespace__",
|
||||
[](const mx::array& a,
|
||||
[](const mx::array& /* a */,
|
||||
const std::optional<std::string>& api_version) {
|
||||
if (api_version) {
|
||||
throw std::invalid_argument(
|
||||
@@ -501,7 +501,7 @@ void init_array(nb::module_& m) {
|
||||
.def("__dlpack__", [](const mx::array& a) { return mlx_to_dlpack(a); })
|
||||
.def(
|
||||
"__dlpack_device__",
|
||||
[](const mx::array& a) {
|
||||
[](const mx::array& /* a */) {
|
||||
// See
|
||||
// https://github.com/dmlc/dlpack/blob/5c210da409e7f1e51ddf445134a4376fdbd70d7d/include/dlpack/dlpack.h#L74
|
||||
if (mx::metal::is_available()) {
|
||||
|
||||
@@ -50,7 +50,7 @@ mx::array nd_array_to_mlx(
|
||||
// Compute the shape and size
|
||||
mx::Shape shape;
|
||||
shape.reserve(nd_array.ndim());
|
||||
for (int i = 0; i < nd_array.ndim(); i++) {
|
||||
for (int i = 0; i < static_cast<int>(nd_array.ndim()); i++) {
|
||||
shape.push_back(check_shape_dim(nd_array.shape(i)));
|
||||
}
|
||||
auto type = nd_array.dtype();
|
||||
@@ -289,7 +289,7 @@ PyScalarT validate_shape(
|
||||
throw std::invalid_argument("Initialization encountered extra dimension.");
|
||||
}
|
||||
auto s = shape[idx];
|
||||
if (nb::len(list) != s) {
|
||||
if (nb::len(list) != static_cast<size_t>(s)) {
|
||||
throw std::invalid_argument(
|
||||
"Initialization encountered non-uniform length.");
|
||||
}
|
||||
|
||||
@@ -201,7 +201,6 @@ void init_fast(nb::module_& parent_module) {
|
||||
bool has_mask = !std::holds_alternative<std::monostate>(mask);
|
||||
bool has_str_mask =
|
||||
has_mask && std::holds_alternative<std::string>(mask);
|
||||
bool has_arr_mask = has_mask && std::holds_alternative<mx::array>(mask);
|
||||
|
||||
if (has_mask) {
|
||||
if (has_str_mask) {
|
||||
|
||||
+13
-12
@@ -115,7 +115,7 @@ mx::array mlx_gather_nd(
|
||||
std::vector<bool> is_slice(indices.size(), false);
|
||||
int num_slices = 0;
|
||||
// gather all the arrays
|
||||
for (int i = 0; i < indices.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(indices); i++) {
|
||||
auto& idx = indices[i];
|
||||
|
||||
if (nb::isinstance<nb::slice>(idx)) {
|
||||
@@ -142,7 +142,7 @@ mx::array mlx_gather_nd(
|
||||
// reshape them so that the int/array indices are first
|
||||
if (gather_first) {
|
||||
int slice_index = 0;
|
||||
for (int i = 0; i < gather_indices.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(gather_indices); i++) {
|
||||
if (is_slice[i]) {
|
||||
mx::Shape index_shape(max_dims + num_slices, 1);
|
||||
index_shape[max_dims + slice_index] = gather_indices[i].shape(0);
|
||||
@@ -156,7 +156,7 @@ mx::array mlx_gather_nd(
|
||||
}
|
||||
} else {
|
||||
// reshape them so that the int/array indices are last
|
||||
for (int i = 0; i < gather_indices.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(gather_indices); i++) {
|
||||
if (i < num_slices) {
|
||||
mx::Shape index_shape(max_dims + num_slices, 1);
|
||||
index_shape[i] = gather_indices[i].shape(0);
|
||||
@@ -190,7 +190,7 @@ auto mlx_expand_ellipsis(const mx::Shape& shape, const nb::tuple& entries) {
|
||||
bool has_ellipsis = false;
|
||||
|
||||
// Start from dimension 0 till we hit an ellipsis
|
||||
for (; i < entries.size(); i++) {
|
||||
for (; i < std::ssize(entries); i++) {
|
||||
auto idx = entries[i];
|
||||
if (!is_valid_index_type(idx)) {
|
||||
throw std::invalid_argument(
|
||||
@@ -301,7 +301,8 @@ mx::array mlx_get_item_nd(mx::array src, const nb::tuple& entries) {
|
||||
if (have_array) {
|
||||
int last_array;
|
||||
// Then find the last array
|
||||
for (last_array = indices.size() - 1; last_array >= 0; last_array--) {
|
||||
for (last_array = std::ssize(indices) - 1; last_array >= 0;
|
||||
last_array--) {
|
||||
auto& idx = indices[last_array];
|
||||
if (nb::isinstance<mx::array>(idx) || nb::isinstance<nb::int_>(idx)) {
|
||||
break;
|
||||
@@ -333,11 +334,11 @@ mx::array mlx_get_item_nd(mx::array src, const nb::tuple& entries) {
|
||||
nb::slice(nb::none(), nb::none(), nb::none()));
|
||||
}
|
||||
}
|
||||
for (int i = last_array + 1; i < indices.size(); i++) {
|
||||
for (int i = last_array + 1; i < std::ssize(indices); i++) {
|
||||
remaining_indices.push_back(indices[i]);
|
||||
}
|
||||
} else {
|
||||
for (int i = 0; i < indices.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(indices); i++) {
|
||||
auto& idx = indices[i];
|
||||
if (nb::isinstance<mx::array>(idx) || nb::isinstance<nb::int_>(idx)) {
|
||||
break;
|
||||
@@ -352,7 +353,7 @@ mx::array mlx_get_item_nd(mx::array src, const nb::tuple& entries) {
|
||||
remaining_indices.push_back(
|
||||
nb::slice(nb::none(), nb::none(), nb::none()));
|
||||
}
|
||||
for (int i = last_array + 1; i < indices.size(); i++) {
|
||||
for (int i = last_array + 1; i < std::ssize(indices); i++) {
|
||||
remaining_indices.push_back(indices[i]);
|
||||
}
|
||||
}
|
||||
@@ -406,7 +407,7 @@ mx::array mlx_get_item_nd(mx::array src, const nb::tuple& entries) {
|
||||
if (unsqueeze_needed || squeeze_needed) {
|
||||
std::vector<int> squeeze_axes;
|
||||
std::vector<int> unsqueeze_axes;
|
||||
for (int axis = 0; axis < remaining_indices.size(); ++axis) {
|
||||
for (int axis = 0; axis < std::ssize(remaining_indices); ++axis) {
|
||||
auto& idx = remaining_indices[axis];
|
||||
if (unsqueeze_needed && idx.is_none()) {
|
||||
unsqueeze_axes.push_back(axis - squeeze_axes.size());
|
||||
@@ -583,7 +584,7 @@ mlx_scatter_args_nd(
|
||||
}
|
||||
|
||||
// Analyse the types of the indices
|
||||
size_t max_dim = 0;
|
||||
int max_dim = 0;
|
||||
bool arrays_first = false;
|
||||
int num_none = 0;
|
||||
int num_slices = 0;
|
||||
@@ -640,7 +641,7 @@ mlx_scatter_args_nd(
|
||||
std::vector<int> update_shape(non_none_indices, 1);
|
||||
std::vector<int> slice_shapes;
|
||||
|
||||
for (int i = 0; i < indices.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(indices); ++i) {
|
||||
auto& pyidx = indices[i];
|
||||
if (nb::isinstance<nb::slice>(pyidx)) {
|
||||
mx::ShapeElem start, end, stride;
|
||||
@@ -848,7 +849,7 @@ auto mlx_slice_update(
|
||||
int unspecified = src.ndim() - non_none_indices;
|
||||
std::vector<int> squeeze_dims;
|
||||
std::vector<int> expand_dims;
|
||||
for (int i = indices.size() - 1,
|
||||
for (int i = std::ssize(indices) - 1,
|
||||
ax = non_none_indices - 1,
|
||||
upd_ax = upd.ndim() - unspecified - 1;
|
||||
i >= 0;
|
||||
|
||||
+1
-1
@@ -436,7 +436,7 @@ void mlx_savez_helper(
|
||||
nb::cast<std::unordered_map<std::string, mx::array>>(kwargs);
|
||||
auto arrays_list = nb::cast<std::vector<mx::array>>(args);
|
||||
|
||||
for (int i = 0; i < arrays_list.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(arrays_list); i++) {
|
||||
std::string arr_name = "arr_" + std::to_string(i);
|
||||
|
||||
if (arrays_dict.count(arr_name) > 0) {
|
||||
|
||||
@@ -22,7 +22,9 @@ bool DEPRECATE(const char* old_fn, const char* new_fn) {
|
||||
return true;
|
||||
}
|
||||
|
||||
#define DEPRECATE(oldfn, newfn) static bool dep = DEPRECATE(oldfn, newfn)
|
||||
#define DEPRECATE(oldfn, newfn) \
|
||||
static bool dep = DEPRECATE(oldfn, newfn); \
|
||||
(void)dep;
|
||||
|
||||
void init_metal(nb::module_& m) {
|
||||
nb::module_ metal = m.def_submodule("metal", "mlx.metal");
|
||||
|
||||
@@ -107,7 +107,7 @@ nb::callable mlx_func(
|
||||
return nb::steal<nb::callable>((PyObject*)r);
|
||||
}
|
||||
|
||||
void init_mlx_func(nb::module_& m) {
|
||||
void init_mlx_func(nb::module_& /* m */) {
|
||||
gc_func_tp = (PyTypeObject*)PyType_FromSpec(&gc_func_spec);
|
||||
if (!gc_func_tp) {
|
||||
nb::raise("Could not register MLX function type.");
|
||||
|
||||
@@ -100,9 +100,9 @@ void init_stream(nb::module_& m) {
|
||||
.def(
|
||||
"__exit__",
|
||||
[](PyStreamContext& scm,
|
||||
const std::optional<nb::type_object>& exc_type,
|
||||
const std::optional<nb::object>& exc_value,
|
||||
const std::optional<nb::object>& traceback) { scm.exit(); },
|
||||
const std::optional<nb::type_object>& /* exc_type */,
|
||||
const std::optional<nb::object>& /* exc_value */,
|
||||
const std::optional<nb::object>& /* traceback */) { scm.exit(); },
|
||||
"exc_type"_a = nb::none(),
|
||||
"exc_value"_a = nb::none(),
|
||||
"traceback"_a = nb::none());
|
||||
|
||||
+12
-13
@@ -86,7 +86,7 @@ auto py_value_and_grad(
|
||||
<< argnums[0];
|
||||
throw std::invalid_argument(msg.str());
|
||||
}
|
||||
for (int i = 1; i < argnums.size(); ++i) {
|
||||
for (int i = 1; i < std::ssize(argnums); ++i) {
|
||||
if (argnums[i] == argnums[i - 1]) {
|
||||
std::ostringstream msg;
|
||||
msg << error_msg_tag << " Duplicate argument index " << argnums[0]
|
||||
@@ -99,7 +99,7 @@ auto py_value_and_grad(
|
||||
return [fun, argnums, argnames, error_msg_tag, scalar_func_only](
|
||||
nb::args& args, nb::kwargs& kwargs) {
|
||||
// Sanitize the input
|
||||
if (argnums.size() > 0 && argnums.back() >= args.size()) {
|
||||
if (argnums.size() > 0 && argnums.back() >= std::ssize(args)) {
|
||||
std::ostringstream msg;
|
||||
msg << error_msg_tag << " Can't compute the gradient of argument index "
|
||||
<< argnums.back() << " because the function is called with only "
|
||||
@@ -126,8 +126,8 @@ auto py_value_and_grad(
|
||||
std::vector<mx::array> arrays;
|
||||
std::vector<int> counts(1, 0);
|
||||
std::vector<int> gradient_indices;
|
||||
for (int i = 0, j = 0; i < args.size(); ++i) {
|
||||
bool needs_grad = (j < argnums.size() && argnums[j] == i);
|
||||
for (int i = 0, j = 0; i < std::ssize(args); ++i) {
|
||||
bool needs_grad = (j < std::ssize(argnums) && argnums[j] == i);
|
||||
auto argsi = tree_flatten(args[i], /* strict = */ needs_grad);
|
||||
if (needs_grad) {
|
||||
auto old_size = gradient_indices.size();
|
||||
@@ -257,7 +257,7 @@ auto py_value_and_grad(
|
||||
positional_grads = tree_unflatten(args[argnums[0]], gradients, counts[0]);
|
||||
} else if (argnums.size() > 1) {
|
||||
nb::list grads_;
|
||||
for (int i = 0; i < argnums.size(); i++) {
|
||||
for (int i = 0; i < std::ssize(argnums); i++) {
|
||||
grads_.append(tree_unflatten(args[argnums[i]], gradients, counts[i]));
|
||||
}
|
||||
positional_grads = nb::tuple(grads_);
|
||||
@@ -366,14 +366,13 @@ auto py_vmap(
|
||||
// able to reconstruct the python tree of extra return values
|
||||
nb::object py_outputs;
|
||||
|
||||
auto vmap_fn =
|
||||
[&fun, &args, &inputs, &py_outputs](const std::vector<mx::array>& a) {
|
||||
// Call the python function
|
||||
py_outputs = fun(*tree_unflatten(args, a));
|
||||
auto vmap_fn = [&fun, &args, &py_outputs](const std::vector<mx::array>& a) {
|
||||
// Call the python function
|
||||
py_outputs = fun(*tree_unflatten(args, a));
|
||||
|
||||
// Flatten the outputs
|
||||
return tree_flatten(py_outputs, true);
|
||||
};
|
||||
// Flatten the outputs
|
||||
return tree_flatten(py_outputs, true);
|
||||
};
|
||||
|
||||
auto [trace_inputs, trace_outputs] =
|
||||
mx::detail::vmap_trace(vmap_fn, inputs, flat_in_axes);
|
||||
@@ -451,7 +450,7 @@ struct PyCompiledFun {
|
||||
if (nb::isinstance<nb::list>(obj)) {
|
||||
auto l = nb::cast<nb::list>(obj);
|
||||
constants.push_back(list_identifier);
|
||||
for (int i = 0; i < l.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(l); ++i) {
|
||||
recurse(l[i]);
|
||||
}
|
||||
} else if (nb::isinstance<nb::tuple>(obj)) {
|
||||
|
||||
+14
-13
@@ -6,7 +6,8 @@ template <typename T, typename U, typename V>
|
||||
void validate_subtrees(const std::vector<nb::object>& subtrees) {
|
||||
int len = nb::cast<T>(subtrees[0]).size();
|
||||
for (auto& subtree : subtrees) {
|
||||
if ((nb::isinstance<T>(subtree) && nb::cast<T>(subtree).size() != len) ||
|
||||
if ((nb::isinstance<T>(subtree) &&
|
||||
std::ssize(nb::cast<T>(subtree)) != len) ||
|
||||
nb::isinstance<U>(subtree) || nb::isinstance<V>(subtree)) {
|
||||
throw std::invalid_argument(
|
||||
"[tree_map] Additional input tree is not a valid prefix of the first tree.");
|
||||
@@ -24,8 +25,8 @@ nb::object tree_map(
|
||||
nb::list l;
|
||||
std::vector<nb::object> items(subtrees.size());
|
||||
validate_subtrees<nb::list, nb::tuple, nb::dict>(subtrees);
|
||||
for (int i = 0; i < nb::cast<nb::list>(subtrees[0]).size(); ++i) {
|
||||
for (int j = 0; j < subtrees.size(); ++j) {
|
||||
for (int i = 0; i < std::ssize(nb::cast<nb::list>(subtrees[0])); ++i) {
|
||||
for (int j = 0; j < std::ssize(subtrees); ++j) {
|
||||
if (nb::isinstance<nb::list>(subtrees[j])) {
|
||||
items[j] = nb::cast<nb::list>(subtrees[j])[i];
|
||||
} else {
|
||||
@@ -42,7 +43,7 @@ nb::object tree_map(
|
||||
nb::list l;
|
||||
validate_subtrees<nb::tuple, nb::list, nb::dict>(subtrees);
|
||||
for (int i = 0; i < len; ++i) {
|
||||
for (int j = 0; j < subtrees.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(subtrees); ++j) {
|
||||
if (nb::isinstance<nb::tuple>(subtrees[j])) {
|
||||
items[j] = nb::cast<nb::tuple>(subtrees[j])[i];
|
||||
} else {
|
||||
@@ -57,7 +58,7 @@ nb::object tree_map(
|
||||
validate_subtrees<nb::dict, nb::list, nb::tuple>(subtrees);
|
||||
nb::dict d;
|
||||
for (auto item : nb::cast<nb::dict>(subtrees[0])) {
|
||||
for (int j = 0; j < subtrees.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(subtrees); ++j) {
|
||||
if (nb::isinstance<nb::dict>(subtrees[j])) {
|
||||
auto subdict = nb::cast<nb::dict>(subtrees[j]);
|
||||
if (!subdict.contains(item.first)) {
|
||||
@@ -96,8 +97,8 @@ void tree_visit(
|
||||
if (nb::isinstance<nb::list>(subtrees[0])) {
|
||||
std::vector<nb::object> items(subtrees.size());
|
||||
validate_subtrees<nb::list, nb::tuple, nb::dict>(subtrees);
|
||||
for (int i = 0; i < nb::cast<nb::list>(subtrees[0]).size(); ++i) {
|
||||
for (int j = 0; j < subtrees.size(); ++j) {
|
||||
for (int i = 0; i < std::ssize(nb::cast<nb::list>(subtrees[0])); ++i) {
|
||||
for (int j = 0; j < std::ssize(subtrees); ++j) {
|
||||
if (nb::isinstance<nb::list>(subtrees[j])) {
|
||||
items[j] = nb::cast<nb::list>(subtrees[j])[i];
|
||||
} else {
|
||||
@@ -112,7 +113,7 @@ void tree_visit(
|
||||
int len = nb::cast<nb::tuple>(subtrees[0]).size();
|
||||
validate_subtrees<nb::tuple, nb::list, nb::dict>(subtrees);
|
||||
for (int i = 0; i < len; ++i) {
|
||||
for (int j = 0; j < subtrees.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(subtrees); ++j) {
|
||||
if (nb::isinstance<nb::tuple>(subtrees[j])) {
|
||||
items[j] = nb::cast<nb::tuple>(subtrees[j])[i];
|
||||
} else {
|
||||
@@ -125,7 +126,7 @@ void tree_visit(
|
||||
std::vector<nb::object> items(subtrees.size());
|
||||
validate_subtrees<nb::dict, nb::list, nb::tuple>(subtrees);
|
||||
for (auto item : nb::cast<nb::dict>(subtrees[0])) {
|
||||
for (int j = 0; j < subtrees.size(); ++j) {
|
||||
for (int j = 0; j < std::ssize(subtrees); ++j) {
|
||||
if (nb::isinstance<nb::dict>(subtrees[j])) {
|
||||
auto subdict = nb::cast<nb::dict>(subtrees[j]);
|
||||
if (!subdict.contains(item.first)) {
|
||||
@@ -173,13 +174,13 @@ void tree_visit_update(
|
||||
recurse = [&](nb::handle subtree) {
|
||||
if (nb::isinstance<nb::list>(subtree)) {
|
||||
auto l = nb::cast<nb::list>(subtree);
|
||||
for (int i = 0; i < l.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(l); ++i) {
|
||||
l[i] = recurse(l[i]);
|
||||
}
|
||||
return nb::cast<nb::object>(l);
|
||||
} else if (nb::isinstance<nb::tuple>(subtree)) {
|
||||
nb::list l(subtree);
|
||||
for (int i = 0; i < l.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(l); ++i) {
|
||||
l[i] = recurse(l[i]);
|
||||
}
|
||||
return nb::cast<nb::object>(nb::tuple(l));
|
||||
@@ -204,7 +205,7 @@ void tree_visit_update(
|
||||
void tree_fill(nb::object& tree, const std::vector<mx::array>& values) {
|
||||
size_t index = 0;
|
||||
tree_visit_update(
|
||||
tree, [&](nb::handle node) { return nb::cast(values[index++]); });
|
||||
tree, [&](nb::handle /* node */) { return nb::cast(values[index++]); });
|
||||
}
|
||||
|
||||
// Replace all the arrays from the src values with the dst values in the tree
|
||||
@@ -213,7 +214,7 @@ void tree_replace(
|
||||
const std::vector<mx::array>& src,
|
||||
const std::vector<mx::array>& dst) {
|
||||
std::unordered_map<uintptr_t, mx::array> src_to_dst;
|
||||
for (int i = 0; i < src.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(src); ++i) {
|
||||
src_to_dst.insert({src[i].id(), dst[i]});
|
||||
}
|
||||
tree_visit_update(tree, [&](nb::handle node) {
|
||||
|
||||
@@ -57,8 +57,8 @@ std::pair<mx::array, mx::array> to_arrays(
|
||||
// - If neither is an array convert to arrays but leave their types alone
|
||||
auto is_mlx_array = [](const ScalarOrArray& x) {
|
||||
return std::holds_alternative<mx::array>(x) ||
|
||||
std::holds_alternative<ArrayLike>(x) &&
|
||||
nb::hasattr(std::get<ArrayLike>(x).obj, "__mlx_array__");
|
||||
(std::holds_alternative<ArrayLike>(x) &&
|
||||
nb::hasattr(std::get<ArrayLike>(x).obj, "__mlx_array__"));
|
||||
};
|
||||
auto get_mlx_array = [](const ScalarOrArray& x) {
|
||||
if (auto px = std::get_if<mx::array>(&x); px) {
|
||||
|
||||
@@ -145,7 +145,7 @@ TEST_CASE("test jvp") {
|
||||
|
||||
// No dependence between input and output
|
||||
{
|
||||
auto fun = [](array in) { return array({1.0, 1.0}); };
|
||||
auto fun = [](array /* in */) { return array({1.0, 1.0}); };
|
||||
auto out = jvp(fun, array(1.0f), array(1.0f)).second;
|
||||
CHECK(array_equal(out, zeros({2})).item<bool>());
|
||||
}
|
||||
@@ -195,7 +195,7 @@ TEST_CASE("test vjp") {
|
||||
|
||||
// No dependence between input and output
|
||||
{
|
||||
auto fun = [](array in) { return array(1.); };
|
||||
auto fun = [](array /* in */) { return array(1.); };
|
||||
auto out = vjp(fun, zeros({2}), array(1.)).second;
|
||||
CHECK(array_equal(out, zeros({2})).item<bool>());
|
||||
}
|
||||
|
||||
@@ -44,7 +44,7 @@ TEST_CASE("test export basic functions") {
|
||||
}
|
||||
|
||||
TEST_CASE("test export function with no inputs") {
|
||||
auto fun = [](std::vector<array> x) -> std::vector<array> {
|
||||
auto fun = [](std::vector<array> /* x */) -> std::vector<array> {
|
||||
return {zeros({2, 2})};
|
||||
};
|
||||
|
||||
|
||||
@@ -168,7 +168,7 @@ TEST_CASE("test gguf metadata") {
|
||||
CHECK_EQ(loaded_metadata.count("meta"), 1);
|
||||
auto& strs = std::get<std::vector<std::string>>(loaded_metadata["meta"]);
|
||||
CHECK_EQ(strs.size(), 3);
|
||||
for (int i = 0; i < strs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(strs); ++i) {
|
||||
CHECK_EQ(strs[i], data[i]);
|
||||
}
|
||||
}
|
||||
@@ -187,7 +187,7 @@ TEST_CASE("test gguf metadata") {
|
||||
CHECK_EQ(loaded_metadata.size(), 4);
|
||||
auto& strs = std::get<std::vector<std::string>>(loaded_metadata["meta1"]);
|
||||
CHECK_EQ(strs.size(), 3);
|
||||
for (int i = 0; i < strs.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(strs); ++i) {
|
||||
CHECK_EQ(strs[i], data[i]);
|
||||
}
|
||||
auto& arr = std::get<array>(loaded_metadata["meta2"]);
|
||||
|
||||
+2
-2
@@ -1668,7 +1668,7 @@ TEST_CASE("test error functions") {
|
||||
-0.1124629160182849,
|
||||
-0.5204998778130465,
|
||||
-0.7969082124228322};
|
||||
for (int i = 0; i < vals.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(vals); ++i) {
|
||||
x = array(vals.begin()[i]);
|
||||
CHECK_EQ(erf(x).item<float>(), doctest::Approx(expected.begin()[i]));
|
||||
}
|
||||
@@ -1686,7 +1686,7 @@ TEST_CASE("test error functions") {
|
||||
-0.08885599049425769,
|
||||
-0.4769362762044699,
|
||||
-1.1630871536766743};
|
||||
for (int i = 0; i < vals.size(); ++i) {
|
||||
for (int i = 0; i < std::ssize(vals); ++i) {
|
||||
x = array(vals.begin()[i]);
|
||||
CHECK_EQ(erfinv(x).item<float>(), doctest::Approx(expected.begin()[i]));
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user