Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
67 commits
Select commit Hold shift + click to select a range
647f80d
spelling: accumulates
jsoref Jan 2, 2024
73272bb
spelling: across
jsoref Jan 2, 2024
b698ce3
spelling: additional
jsoref Jan 2, 2024
33be1a2
spelling: against
jsoref Jan 2, 2024
7057f13
spelling: among
jsoref Jan 2, 2024
27f640e
spelling: array
jsoref Jan 2, 2024
5097cd7
spelling: at least
jsoref Jan 2, 2024
f28bcda
spelling: available
jsoref Jan 2, 2024
9ad62ea
spelling: axes
jsoref Jan 2, 2024
acde6e4
spelling: basically
jsoref Jan 2, 2024
c8c2d6c
spelling: bfloat
jsoref Jan 2, 2024
94f71e5
spelling: bounds
jsoref Jan 2, 2024
1604608
spelling: broadcast
jsoref Jan 2, 2024
edc830d
spelling: buffer
jsoref Jan 2, 2024
d81a841
spelling: class
jsoref Jan 2, 2024
35c5530
spelling: coefficients
jsoref Jan 2, 2024
76ac49c
spelling: collision
jsoref Jan 2, 2024
272f695
spelling: combinations
jsoref Jan 2, 2024
d26e758
spelling: committing
jsoref Jan 2, 2024
2a5a690
spelling: computation
jsoref Jan 2, 2024
956cac5
spelling: consider
jsoref Jan 2, 2024
0f34dc1
spelling: constructing
jsoref Jan 2, 2024
30e9fdc
spelling: conversions
jsoref Jan 2, 2024
e34205b
spelling: correctly
jsoref Jan 2, 2024
b647328
spelling: corresponding
jsoref Jan 2, 2024
8cae2e1
spelling: declaration
jsoref Jan 2, 2024
c249816
spelling: default
jsoref Jan 2, 2024
61aa866
spelling: dependency
jsoref Jan 2, 2024
a14e197
spelling: destination
jsoref Jan 2, 2024
1c9eea2
spelling: destructor
jsoref Jan 2, 2024
f4fcfb4
spelling: dimensions
jsoref Jan 2, 2024
07cb2f7
spelling: divided
jsoref Jan 2, 2024
f242b42
spelling: element-wise
jsoref Jan 2, 2024
3cce74a
spelling: elements
jsoref Jan 2, 2024
586ac7d
spelling: endianness
jsoref Jan 2, 2024
dc0661b
spelling: equivalent
jsoref Jan 2, 2024
d1db947
spelling: explicitly
jsoref Jan 2, 2024
64baf05
spelling: github
jsoref Jan 2, 2024
e54c872
spelling: indices
jsoref Jan 2, 2024
fe3c914
spelling: irregularly
jsoref Jan 2, 2024
6c1cd87
spelling: memory
jsoref Jan 2, 2024
0112fca
spelling: metallib
jsoref Jan 2, 2024
0778209
spelling: negative
jsoref Jan 2, 2024
0219a05
spelling: notable
jsoref Jan 2, 2024
0ddb83c
spelling: optional
jsoref Jan 2, 2024
116a5b3
spelling: otherwise
jsoref Jan 2, 2024
8f39b89
spelling: overridden
jsoref Jan 2, 2024
c18cb06
spelling: partially
jsoref Jan 2, 2024
979ae32
spelling: partition
jsoref Jan 2, 2024
f890871
spelling: perform
jsoref Jan 2, 2024
039ee91
spelling: perturbations
jsoref Jan 2, 2024
94eeae4
spelling: positively
jsoref Jan 2, 2024
684e946
spelling: primitive
jsoref Jan 2, 2024
61823aa
spelling: repeat
jsoref Jan 2, 2024
e254286
spelling: repeats
jsoref Jan 2, 2024
538f523
spelling: respect
jsoref Jan 2, 2024
f48a352
spelling: respectively
jsoref Jan 2, 2024
69142ef
spelling: result
jsoref Jan 2, 2024
6b0d08a
spelling: rounding
jsoref Jan 2, 2024
2fb5fe3
spelling: separate
jsoref Jan 2, 2024
31b24c2
spelling: skipping
jsoref Jan 2, 2024
3a54373
spelling: structure
jsoref Jan 2, 2024
9557a8f
spelling: the
jsoref Jan 2, 2024
072f7e0
spelling: transpose
jsoref Jan 2, 2024
59b779b
spelling: unnecessary
jsoref Jan 2, 2024
381712a
spelling: unneeded
jsoref Jan 2, 2024
c373f86
spelling: unsupported
jsoref Jan 2, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions benchmarks/python/blas/bench_gemv.py
Original file line number Diff line number Diff line change
Expand Up @@ -133,7 +133,7 @@ def get_gbyte_size(in_vec_len, out_vec_len, np_dtype):
return float(N_iter_bench * N_iter_func * n_elem * item_size) / float(1024**3)


def bench_with_in_len(ax, in_vec_len, out_vector_lens, dtype, tranpose):
def bench_with_in_len(ax, in_vec_len, out_vector_lens, dtype, transpose):

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this is a guess...

np_dtype = getattr(np, dtype)
mlx_gb_s = []
mlx_gflops = []
Expand Down Expand Up @@ -164,7 +164,7 @@ def bench_with_in_len(ax, in_vec_len, out_vector_lens, dtype, tranpose):
ax.legend()


def bench_with_out_len(ax, out_vec_len, in_vector_lens, dtype, tranpose):
def bench_with_out_len(ax, out_vec_len, in_vector_lens, dtype, transpose):
np_dtype = getattr(np, dtype)
mlx_gb_s = []
mlx_gflops = []
Expand Down
2 changes: 1 addition & 1 deletion benchmarks/python/comparative/compare.py
Original file line number Diff line number Diff line change
Expand Up @@ -62,7 +62,7 @@ def predicate(x):


if __name__ == "__main__":
parser = argparse.ArgumentParser(description="Run comparisons agains PyTorch")
parser = argparse.ArgumentParser(description="Run comparisons against PyTorch")
parser.add_argument(
"--filter", "-f", help="Regex filter to select benchmarks", nargs="+"
)
Expand Down
4 changes: 2 additions & 2 deletions cmake/extension.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ include(CMakeParseArguments)
# OUTPUT_DIRECTORY: Where to place ${TITLE}.metallib
# SOURCES: List of source files
# INCLUDE_DIRS: List of include dirs
# DEPS: List of depedency files (like headers)
# DEPS: List of dependency files (like headers)
#
macro(mlx_build_metallib)
# Parse args
Expand All @@ -32,7 +32,7 @@ macro(mlx_build_metallib)
# Collect compile options
set(MTLLIB_COMPILE_OPTIONS -Wall -Wextra -fno-fast-math)

# Prepare metllib build command
# Prepare metallib build command
add_custom_command(
OUTPUT ${MTLLIB_BUILD_TARGET}
COMMAND xcrun -sdk macosx metal
Expand Down
2 changes: 1 addition & 1 deletion docs/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ python -m http.server <port>

and point your browser to `http://localhost:<port>`.

### Push to Github Pages
### Push to GitHub Pages

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

brand


Check-out the `gh-pages` branch (`git switch gh-pages`) and build
the docs. Then force add the `build/html` directory:
Expand Down
32 changes: 16 additions & 16 deletions docs/src/dev/extensions.rst
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@ Introducing the Example
-----------------------

Let's say that you would like an operation that takes in two arrays,
``x`` and ``y``, scales them both by some coefficents ``alpha`` and ``beta``
``x`` and ``y``, scales them both by some coefficients ``alpha`` and ``beta``
respectively, and then adds them together to get the result
``z = alpha * x + beta * y``. Well, you can very easily do that by just
writing out a function as follows:
Expand Down Expand Up @@ -69,7 +69,7 @@ C++ API:
.. code-block:: C++

/**
* Scale and sum two vectors elementwise
* Scale and sum two vectors element-wise
* z = alpha * x + beta * y
*
* Follow numpy style broadcasting between x and y
Expand Down Expand Up @@ -230,7 +230,7 @@ Let's re-implement our operation now in terms of our :class:`Axpby` primitive.

This operation now handles the following:

#. Upcast inputs and resolve the the output data type.
#. Upcast inputs and resolve the output data type.
#. Broadcast the inputs and resolve the output shape.
#. Construct the primitive :class:`Axpby` using the given stream, ``alpha``, and ``beta``.
#. Construct the output :class:`array` using the primitive and the inputs.
Expand Down Expand Up @@ -284,14 +284,14 @@ pointwise. This is captured in the templated function :meth:`axpby_impl`.
T alpha = static_cast<T>(alpha_);
T beta = static_cast<T>(beta_);

// Do the elementwise operation for each output
// Do the element-wise operation for each output
for (size_t out_idx = 0; out_idx < out.size(); out_idx++) {
// Map linear indices to offsets in x and y
auto x_offset = elem_to_loc(out_idx, x.shape(), x.strides());
auto y_offset = elem_to_loc(out_idx, y.shape(), y.strides());

// We allocate the output to be contiguous and regularly strided
// (defaults to row major) and hence it doesn't need additonal mapping
// (defaults to row major) and hence it doesn't need additional mapping
out_ptr[out_idx] = alpha * x_ptr[x_offset] + beta * y_ptr[y_offset];
}
}
Expand All @@ -305,7 +305,7 @@ if we encounter an unexpected type.

/** Fall back implementation for evaluation on CPU */
void Axpby::eval(const std::vector<array>& inputs, array& out) {
// Check the inputs (registered in the op while contructing the out array)
// Check the inputs (registered in the op while constructing the out array)
assert(inputs.size() == 2);
auto& x = inputs[0];
auto& y = inputs[1];
Expand Down Expand Up @@ -485,7 +485,7 @@ each data type.

instantiate_axpby(float32, float);
instantiate_axpby(float16, half);
instantiate_axpby(bflot16, bfloat16_t);
instantiate_axpby(bfloat16, bfloat16_t);

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

??

instantiate_axpby(complex64, complex64_t);

This kernel will be compiled into a metal library ``mlx_ext.metallib`` as we
Expand Down Expand Up @@ -537,7 +537,7 @@ below.
compute_encoder->setComputePipelineState(kernel);

// Kernel parameters are registered with buffer indices corresponding to
// those in the kernel decelaration at axpby.metal
// those in the kernel declaration at axpby.metal
int ndim = out.ndim();
size_t nelem = out.size();

Expand Down Expand Up @@ -568,7 +568,7 @@ below.
// Fix the 3D size of the launch grid (in terms of threads)
MTL::Size grid_dims = MTL::Size(nelem, 1, 1);

// Launch the grid with the given number of threads divded among
// Launch the grid with the given number of threads divided among
// the given threadgroups
compute_encoder->dispatchThreads(grid_dims, group_dims);
}
Expand All @@ -581,7 +581,7 @@ to give us the active metal compute command encoder instead of building a
new one and calling :meth:`compute_encoder->end_encoding` at the end.
MLX keeps adding kernels (compute pipelines) to the active command encoder
until some specified limit is hit or the compute encoder needs to be flushed
for synchronization. MLX also handles enqueuing and commiting the associated
for synchronization. MLX also handles enqueuing and committing the associated
command buffers as needed. We suggest taking a deeper dive into
:class:`metal::Device` if you would like to study this routine further.

Expand All @@ -601,8 +601,8 @@ us the following :meth:`Axpby::jvp` and :meth:`Axpby::vjp` implementations.
const std::vector<array>& tangents,
const std::vector<int>& argnums) {
// Forward mode diff that pushes along the tangents
// The jvp transform on the the primitive can built with ops
// that are scheduled on the same stream as the primtive
// The jvp transform on the primitive can built with ops
// that are scheduled on the same stream as the primitive

// If argnums = {0}, we only push along x in which case the
// jvp is just the tangent scaled by alpha
Expand Down Expand Up @@ -642,7 +642,7 @@ own :class:`Primitive`.

.. code-block:: C++

/** Vectorize primitve along given axis */
/** Vectorize primitive along given axis */
std::pair<array, int> Axpby::vmap(
const std::vector<array>& inputs,
const std::vector<int>& axes) {
Expand All @@ -666,7 +666,7 @@ Let's look at the overall directory structure first.
| └── setup.py

* ``extensions/axpby/`` defines the C++ extension library
* ``extensions/mlx_sample_extensions`` sets out the strucutre for the
* ``extensions/mlx_sample_extensions`` sets out the structure for the
associated python package
* ``extensions/bindings.cpp`` provides python bindings for our operation
* ``extensions/CMakeLists.txt`` holds CMake rules to build the library and
Expand Down Expand Up @@ -697,7 +697,7 @@ are already provided, adding our :meth:`axpby` becomes very simple!
py::kw_only(),
"stream"_a = py::none(),
R"pbdoc(
Scale and sum two vectors elementwise
Scale and sum two vectors element-wise
``z = alpha * x + beta * y``

Follows numpy style broadcasting between ``x`` and ``y``
Expand Down Expand Up @@ -840,7 +840,7 @@ This will result in a directory structure as follows:
| ...

When you try to install using the command ``python -m pip install .``
(in ``extensions/``), the package will be installed with the same strucutre as
(in ``extensions/``), the package will be installed with the same structure as
``extensions/mlx_sample_extensions`` and the C++ and metal library will be
copied along with the python binding since they are specified as ``package_data``.

Expand Down
2 changes: 1 addition & 1 deletion docs/src/index.rst
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ The main differences between MLX and NumPy are:

The design of MLX is inspired by frameworks like `PyTorch
<https://pytorch.org/>`_, `Jax <https://github.com/google/jax>`_, and
`ArrayFire <https://arrayfire.org/>`_. A noteable difference from these
`ArrayFire <https://arrayfire.org/>`_. A notable difference from these
frameworks and MLX is the *unified memory model*. Arrays in MLX live in shared
memory. Operations on MLX arrays can be performed on any of the supported
device types without performing data copies. Currently supported device types
Expand Down
2 changes: 1 addition & 1 deletion examples/cpp/tutorial.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -57,7 +57,7 @@ void array_basics() {
assert(z.shape(0) == 2);
assert(z.shape(1) == 2);

// To actually run the compuation you must evaluate `z`.
// To actually run the computation you must evaluate `z`.
// Under the hood, mlx records operations in a graph.
// The variable `z` is a node in the graph which points to its operation
// and inputs. When `eval` is called on an array (or arrays), the array and
Expand Down
20 changes: 10 additions & 10 deletions examples/extensions/axpby/axpby.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ namespace mlx::core {
///////////////////////////////////////////////////////////////////////////////

/**
* Scale and sum two vectors elementwise
* Scale and sum two vectors element-wise
* z = alpha * x + beta * y
*
* Follow numpy style broadcasting between x and y
Expand Down Expand Up @@ -91,21 +91,21 @@ void axpby_impl(
T alpha = static_cast<T>(alpha_);
T beta = static_cast<T>(beta_);

// Do the elementwise operation for each output
// Do the element-wise operation for each output
for (size_t out_idx = 0; out_idx < out.size(); out_idx++) {
// Map linear indices to offsets in x and y
auto x_offset = elem_to_loc(out_idx, x.shape(), x.strides());
auto y_offset = elem_to_loc(out_idx, y.shape(), y.strides());

// We allocate the output to be contiguous and regularly strided
// (defaults to row major) and hence it doesn't need additonal mapping
// (defaults to row major) and hence it doesn't need additional mapping
out_ptr[out_idx] = alpha * x_ptr[x_offset] + beta * y_ptr[y_offset];
}
}

/** Fall back implementation for evaluation on CPU */
void Axpby::eval(const std::vector<array>& inputs, array& out) {
// Check the inputs (registered in the op while contructing the out array)
// Check the inputs (registered in the op while constructing the out array)
assert(inputs.size() == 2);
auto& x = inputs[0];
auto& y = inputs[1];
Expand Down Expand Up @@ -192,7 +192,7 @@ void Axpby::eval_cpu(const std::vector<array>& inputs, array& out) {
eval(inputs, out);
}

#else // Accelerate not avaliable
#else // Accelerate not available

/** Evaluate primitive on CPU falling back to common backend */
void Axpby::eval_cpu(const std::vector<array>& inputs, array& out) {
Expand Down Expand Up @@ -254,7 +254,7 @@ void Axpby::eval_gpu(const std::vector<array>& inputs, array& out) {
compute_encoder->setComputePipelineState(kernel);

// Kernel parameters are registered with buffer indices corresponding to
// those in the kernel decelaration at axpby.metal
// those in the kernel declaration at axpby.metal
int ndim = out.ndim();
size_t nelem = out.size();

Expand Down Expand Up @@ -287,7 +287,7 @@ void Axpby::eval_gpu(const std::vector<array>& inputs, array& out) {
// Fix the 3D size of the launch grid (in terms of threads)
MTL::Size grid_dims = MTL::Size(nelem, 1, 1);

// Launch the grid with the given number of threads divded among
// Launch the grid with the given number of threads divided among
// the given threadgroups
compute_encoder->dispatchThreads(grid_dims, group_dims);
}
Expand All @@ -311,8 +311,8 @@ array Axpby::jvp(
const std::vector<array>& tangents,
const std::vector<int>& argnums) {
// Forward mode diff that pushes along the tangents
// The jvp transform on the the primitive can built with ops
// that are scheduled on the same stream as the primtive
// The jvp transform on the primitive can built with ops
// that are scheduled on the same stream as the primitive

// If argnums = {0}, we only push along x in which case the
// jvp is just the tangent scaled by alpha
Expand Down Expand Up @@ -345,7 +345,7 @@ std::vector<array> Axpby::vjp(
return vjps;
}

/** Vectorize primitve along given axis */
/** Vectorize primitive along given axis */
std::pair<array, int> Axpby::vmap(
const std::vector<array>& inputs,
const std::vector<int>& axes) {
Expand Down
4 changes: 2 additions & 2 deletions examples/extensions/axpby/axpby.h
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ namespace mlx::core {
///////////////////////////////////////////////////////////////////////////////

/**
* Scale and sum two vectors elementwise
* Scale and sum two vectors element-wise
* z = alpha * x + beta * y
*
* Follow numpy style broadcasting between x and y
Expand All @@ -39,7 +39,7 @@ class Axpby : public Primitive {
* A primitive must know how to evaluate itself on the CPU/GPU
* for the given inputs and populate the output array.
*
* To avoid unecessary allocations, the evaluation function
* To avoid unnecessary allocations, the evaluation function
* is responsible for allocating space for the array.
*/
void eval_cpu(const std::vector<array>& inputs, array& out) override;
Expand Down
2 changes: 1 addition & 1 deletion examples/extensions/axpby/axpby.metal
Original file line number Diff line number Diff line change
Expand Up @@ -59,5 +59,5 @@ template <typename T>

instantiate_axpby(float32, float);
instantiate_axpby(float16, half);
instantiate_axpby(bflot16, bfloat16_t);
instantiate_axpby(bfloat16, bfloat16_t);
instantiate_axpby(complex64, complex64_t);
2 changes: 1 addition & 1 deletion examples/extensions/bindings.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ PYBIND11_MODULE(mlx_sample_extensions, m) {
py::kw_only(),
"stream"_a = py::none(),
R"pbdoc(
Scale and sum two vectors elementwise
Scale and sum two vectors element-wise
``z = alpha * x + beta * y``

Follows numpy style broadcasting between ``x`` and ``y``
Expand Down
2 changes: 1 addition & 1 deletion mlx/allocator.h
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@ void free(Buffer buffer);
Buffer malloc_or_wait(size_t size);

class Allocator {
/** Abstract base clase for a memory allocator. */
/** Abstract base class for a memory allocator. */
public:
virtual Buffer malloc(size_t size) = 0;
virtual void free(Buffer buffer) = 0;
Expand Down
2 changes: 1 addition & 1 deletion mlx/array.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -129,7 +129,7 @@ array::ArrayDesc::ArrayDesc(
}

// Needed because the Primitive type used in array.h is incomplete and the
// compiler needs to see the call to the desctructor after the type is complete.
// compiler needs to see the call to the destructor after the type is complete.
array::ArrayDesc::~ArrayDesc() = default;

array::ArrayIterator::reference array::ArrayIterator::operator*() const {
Expand Down
8 changes: 4 additions & 4 deletions mlx/backend/common/load.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ namespace mlx::core {
namespace {

template <const uint8_t scalar_size>
void swap_endianess(uint8_t* data_bytes, size_t N) {
void swap_endianness(uint8_t* data_bytes, size_t N) {
struct Elem {
uint8_t bytes[scalar_size];
};
Expand All @@ -39,13 +39,13 @@ void Load::eval(const std::vector<array>& inputs, array& out) {
if (swap_endianness_) {
switch (out.itemsize()) {
case 2:
swap_endianess<2>(out.data<uint8_t>(), out.data_size());
swap_endianness<2>(out.data<uint8_t>(), out.data_size());
break;
case 4:
swap_endianess<4>(out.data<uint8_t>(), out.data_size());
swap_endianness<4>(out.data<uint8_t>(), out.data_size());
break;
case 8:
swap_endianess<8>(out.data<uint8_t>(), out.data_size());
swap_endianness<8>(out.data<uint8_t>(), out.data_size());
break;
}
}
Expand Down
2 changes: 1 addition & 1 deletion mlx/backend/metal/allocator.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -165,7 +165,7 @@ Buffer MetalAllocator::malloc(size_t size) {

// Prepare to allocate new memory as needed
if (!buf) {
// If we are under very high memoory pressure, we don't allocate further
// If we are under very high memory pressure, we don't allocate further
if (device_->currentAllocatedSize() >= block_limit_) {
return Buffer{nullptr};
}
Expand Down
4 changes: 2 additions & 2 deletions mlx/backend/metal/conv.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ void explicit_gemm_conv_1D_gpu(
array in_strided(strided_reshape, in_strided_view.dtype(), nullptr, {});
copy_gpu(in_strided_view, in_strided, CopyType::General, s);

// Peform gemm
// Perform gemm
std::vector<array> copies = {in_padded, in_strided};
mlx_matmul(
s,
Expand Down Expand Up @@ -260,7 +260,7 @@ void explicit_gemm_conv_2D_gpu(
array in_strided(strided_reshape, in_strided_view.dtype(), nullptr, {});
copy_gpu(in_strided_view, in_strided, CopyType::General, s);

// Peform gemm
// Perform gemm
std::vector<array> copies = {in_padded, in_strided};
mlx_matmul(
s,
Expand Down
Loading