From 647f80d09cde3001445b9bab66120c5987c49e03 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:27:30 -0500 Subject: [PATCH 01/67] spelling: accumulates Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/reduce.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/reduce.cpp b/mlx/backend/metal/reduce.cpp index 6a2ce084b0..379e54154d 100644 --- a/mlx/backend/metal/reduce.cpp +++ b/mlx/backend/metal/reduce.cpp @@ -176,7 +176,7 @@ void strided_reduce_general_dispatch( // We spread outputs over the x dimension and inputs over the y dimension // Threads with the same lid.x in a given threadgroup work on the same - // output and each thread in the y dimension accumlates for that output + // output and each thread in the y dimension accumulates for that output uint threadgroup_dim_x = std::min(out_size, 128ul); uint threadgroup_dim_y = kernel->maxTotalThreadsPerThreadgroup() / threadgroup_dim_x; From 73272bb87d7c4c484cf451cd77a86b7d542b6418 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:27:22 -0500 Subject: [PATCH 02/67] spelling: across Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/arg_reduce.metal | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/arg_reduce.metal b/mlx/backend/metal/kernels/arg_reduce.metal index 31bcbfa05c..467e768d61 100644 --- a/mlx/backend/metal/kernels/arg_reduce.metal +++ b/mlx/backend/metal/kernels/arg_reduce.metal @@ -114,7 +114,7 @@ template // 4. Reduce among them and go to 3 // 4. Reduce in each simd_group // 6. Write in the thread local memory - // 6. Reduce them accross thread group + // 6. Reduce them across thread group // 7. Write the output without need for atomic Op op; From b698ce39f741a9365fe03d65809923c8ff1ea2e1 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:27:41 -0500 Subject: [PATCH 03/67] spelling: additional Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- examples/extensions/axpby/axpby.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 9aae931a30..35f493fac3 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -291,7 +291,7 @@ pointwise. This is captured in the templated function :meth:`axpby_impl`. 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]; } } diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index 56a09f34e7..5b60cb81d0 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -98,7 +98,7 @@ void axpby_impl( 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]; } } From 33be1a27afcc436119fa9f5018d461b89e3b8cfa Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:27:50 -0500 Subject: [PATCH 04/67] spelling: against Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- benchmarks/python/comparative/compare.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/benchmarks/python/comparative/compare.py b/benchmarks/python/comparative/compare.py index 4adde50bc6..a9d3df22d7 100644 --- a/benchmarks/python/comparative/compare.py +++ b/benchmarks/python/comparative/compare.py @@ -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="+" ) From 7057f13b5f53e6fc5b97f7bddeae43e092551672 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:28:02 -0500 Subject: [PATCH 05/67] spelling: among Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/gemv.metal | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/backend/metal/kernels/gemv.metal b/mlx/backend/metal/kernels/gemv.metal index 3b4c0a30af..165868401f 100644 --- a/mlx/backend/metal/kernels/gemv.metal +++ b/mlx/backend/metal/kernels/gemv.metal @@ -28,7 +28,7 @@ struct GEMVKernel { static_assert(BN == SIMD_SIZE, "gemv block must have a width of SIMD_SIZE"); // - The matrix of size (M = out_vec_size, N = in_vec_size) is divided up - // into blocks of (BM * TM, BN * TN) divided amoung threadgroups + // into blocks of (BM * TM, BN * TN) divided among threadgroups // - Every thread works on a block of (TM, TN) // - We assume each thead group is launched with (BN, BM, 1) threads // @@ -166,7 +166,7 @@ template < struct GEMVTKernel { // - The matrix of size (M = in_vec_size, N = out_vec_size) is divided up - // into blocks of (BM * TM, BN * TN) divided amoung threadgroups + // into blocks of (BM * TM, BN * TN) divided among threadgroups // - Every thread works on a block of (TM, TN) // - We assume each thead group is launched with (BN, BM, 1) threads // From 27f640e7c9c3942e801202aea015dd91f59e7395 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 23:03:10 -0500 Subject: [PATCH 06/67] spelling: array Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- tests/ops_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/ops_tests.cpp b/tests/ops_tests.cpp index 2c6c8e8efb..1fa3035a26 100644 --- a/tests/ops_tests.cpp +++ b/tests/ops_tests.cpp @@ -1411,7 +1411,7 @@ TEST_CASE("test broadcast") { x.eval(); CHECK_EQ(x.strides(), std::vector{0, 0, 1}); - // Broadcast on transposed arrray works + // Broadcast on transposed array works x = array({0, 1, 2, 3, 4, 5}, {2, 3}); x = broadcast_to(transpose(x), {2, 3, 2}); CHECK_EQ(x.shape(), std::vector{2, 3, 2}); From 5097cd7bca601c949fa1af8c8cc2e75db70aa86d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 23:03:31 -0500 Subject: [PATCH 07/67] spelling: at least Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/reduce.cpp | 2 +- mlx/ops.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/mlx/backend/metal/reduce.cpp b/mlx/backend/metal/reduce.cpp index 379e54154d..9da5c79bf3 100644 --- a/mlx/backend/metal/reduce.cpp +++ b/mlx/backend/metal/reduce.cpp @@ -40,7 +40,7 @@ void all_reduce_dispatch( // Set grid dimensions // We make sure each thread has enough to do by making it read in - // atleast n_reads inputs + // at least n_reads inputs int n_reads = REDUCE_N_READS; // mod_in_size gives us the groups of n_reads needed to go over the entire diff --git a/mlx/ops.cpp b/mlx/ops.cpp index f4f6b922d5..8a898f16dd 100644 --- a/mlx/ops.cpp +++ b/mlx/ops.cpp @@ -247,7 +247,7 @@ array tri(int n, int m, int k, Dtype type, StreamOrDevice s /* = {} */) { array tril(array x, int k, StreamOrDevice s /* = {} */) { if (x.ndim() < 2) { - throw std::invalid_argument("[tril] array must be atleast 2-D"); + throw std::invalid_argument("[tril] array must be at least 2-D"); } auto mask = tri(x.shape(-2), x.shape(-1), k, x.dtype(), s); return where(mask, x, zeros_like(x, s), s); @@ -255,7 +255,7 @@ array tril(array x, int k, StreamOrDevice s /* = {} */) { array triu(array x, int k, StreamOrDevice s /* = {} */) { if (x.ndim() < 2) { - throw std::invalid_argument("[triu] array must be atleast 2-D"); + throw std::invalid_argument("[triu] array must be at least 2-D"); } auto mask = tri(x.shape(-2), x.shape(-1), k - 1, x.dtype(), s); return where(mask, zeros_like(x, s), x, s); From f28bcda866b5ed609f5b36818713cbe69d9c4f02 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:48:56 -0500 Subject: [PATCH 08/67] spelling: available Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- examples/extensions/axpby/axpby.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index 5b60cb81d0..c0834d4412 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -192,7 +192,7 @@ void Axpby::eval_cpu(const std::vector& 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& inputs, array& out) { From 9ad62ea6a008304f8df299e4ec28aa2be38a686d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:49:20 -0500 Subject: [PATCH 09/67] spelling: axes Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/ops.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/ops.cpp b/mlx/ops.cpp index 8a898f16dd..ffed3942e3 100644 --- a/mlx/ops.cpp +++ b/mlx/ops.cpp @@ -350,7 +350,7 @@ array squeeze( ax = ax < 0 ? ax + a.ndim() : ax; if (ax < 0 || ax >= a.ndim()) { std::ostringstream msg; - msg << "[squeeze] Invalid axies " << ax << " for array with " << a.ndim() + msg << "[squeeze] Invalid axes " << ax << " for array with " << a.ndim() << " dimensions."; throw std::invalid_argument(msg.str()); } @@ -405,7 +405,7 @@ array expand_dims( ax = ax < 0 ? ax + out_ndim : ax; if (ax < 0 || ax >= out_ndim) { std::ostringstream msg; - msg << "[squeeze] Invalid axies " << ax << " for output array with " + msg << "[squeeze] Invalid axes " << ax << " for output array with " << a.ndim() << " dimensions."; throw std::invalid_argument(msg.str()); } From acde6e476aaf948f737f6266cc2d5c0f42b2f8b0 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:49:37 -0500 Subject: [PATCH 10/67] spelling: basically Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/ops.cpp b/python/src/ops.cpp index 8da45e0eba..ea44868d67 100644 --- a/python/src/ops.cpp +++ b/python/src/ops.cpp @@ -3050,7 +3050,7 @@ void init_ops(py::module_& m) { Round to the given number of decimals. - Bascially performs: + Basically performs: .. code-block:: python From c8c2d6c4de6c3813ceb00b2401ebe51e46fa9b74 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 23:04:10 -0500 Subject: [PATCH 11/67] spelling: bfloat Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- examples/extensions/axpby/axpby.metal | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 35f493fac3..35c655f6c6 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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); instantiate_axpby(complex64, complex64_t); This kernel will be compiled into a metal library ``mlx_ext.metallib`` as we diff --git a/examples/extensions/axpby/axpby.metal b/examples/extensions/axpby/axpby.metal index 64980578f5..03b373c999 100644 --- a/examples/extensions/axpby/axpby.metal +++ b/examples/extensions/axpby/axpby.metal @@ -59,5 +59,5 @@ template instantiate_axpby(float32, float); instantiate_axpby(float16, half); -instantiate_axpby(bflot16, bfloat16_t); +instantiate_axpby(bfloat16, bfloat16_t); instantiate_axpby(complex64, complex64_t); \ No newline at end of file From 94f71e58329c3f939d9b6e409ca3ffaac8c768e9 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:50:30 -0500 Subject: [PATCH 12/67] spelling: bounds Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/random.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/random.cpp b/mlx/random.cpp index ef11f8c654..63e39cdcc0 100644 --- a/mlx/random.cpp +++ b/mlx/random.cpp @@ -232,7 +232,7 @@ array truncated_normal( auto u = uniform(a, b, shape, dtype, key, s); auto out = multiply(sqrt2, erfinv(u, s), s); - // Clip in bouds + // Clip in bounds return maximum(minimum(upper_t, out, s), lower_t, s); } From 16046082f6e675f31eae72e65666cea03966d53c Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:50:36 -0500 Subject: [PATCH 13/67] spelling: broadcast Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/tests/test_blas.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/tests/test_blas.py b/python/tests/test_blas.py index b2a7626818..ffbbbb0a72 100644 --- a/python/tests/test_blas.py +++ b/python/tests/test_blas.py @@ -191,7 +191,7 @@ def test_matmul_batched(self): self.assertListEqual(list(c_npy.shape), list(c_mlx.shape)) self.assertTrue(np.allclose(c_mlx, c_npy, atol=1e-6)) - # Batched matmul with simple broadast + # Batched matmul with simple broadcast a_npy = np.random.normal(0.0, 1.0 / 128, (32, 128, 16)).astype(np.float32) b_npy = np.random.normal(0.0, 1.0 / 128, (16, 16)).astype(np.float32) c_npy = a_npy @ b_npy @@ -213,7 +213,7 @@ def test_matmul_batched(self): self.assertListEqual(list(e_npy.shape), list(e_mlx.shape)) self.assertTrue(np.allclose(e_mlx, e_npy, atol=1e-6)) - # Batched and transposed matmul with simple broadast + # Batched and transposed matmul with simple broadcast a_npy = np.random.normal(0.0, 1.0 / 128, (32, 128, 16)).astype(np.float32) b_npy = np.random.normal(0.0, 1.0 / 128, (128, 16)).astype(np.float32) a_mlx = mx.array(a_npy) From edc830dcd3b8f985b52836d6bc6ead6e2e772fc1 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:50:55 -0500 Subject: [PATCH 14/67] spelling: buffer Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/indexing.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/backend/metal/indexing.cpp b/mlx/backend/metal/indexing.cpp index eb9a8efb65..1f905db1a4 100644 --- a/mlx/backend/metal/indexing.cpp +++ b/mlx/backend/metal/indexing.cpp @@ -102,7 +102,7 @@ void Gather::eval_gpu(const std::vector& inputs, array& out) { static_cast(idx_strides_buf.raw_ptr()) + i * idx_ndim); } - // Allocate the argument bufer + // Allocate the argument buffer auto arg_buf = allocator::malloc_or_wait(arg_enc->encodedLength()); // Register data with the encoder @@ -246,7 +246,7 @@ void Scatter::eval_gpu(const std::vector& inputs, array& out) { static_cast(idx_strides_buf.raw_ptr()) + i * idx_ndim); } - // Allocate the argument bufer + // Allocate the argument buffer auto arg_buf = allocator::malloc_or_wait(arg_enc->encodedLength()); // Register data with the encoder From d81a841ae891af6ab193d77b1c17d8e1ad5c40ea Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 23:04:38 -0500 Subject: [PATCH 15/67] spelling: class Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/allocator.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/allocator.h b/mlx/allocator.h index 2c3adadf4c..ce0c1cd008 100644 --- a/mlx/allocator.h +++ b/mlx/allocator.h @@ -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; From 35c5530ea1d790d2ca9ebe393495fe20a834abba Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:51:14 -0500 Subject: [PATCH 16/67] spelling: coefficients Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 35c655f6c6..1387fec32b 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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: From 76ac49cac3c3755ea84dd3f0536fa271a66a5e7e Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 23:04:49 -0500 Subject: [PATCH 17/67] spelling: collision Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- tests/ops_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/ops_tests.cpp b/tests/ops_tests.cpp index 1fa3035a26..044edf7fd3 100644 --- a/tests/ops_tests.cpp +++ b/tests/ops_tests.cpp @@ -1733,7 +1733,7 @@ TEST_CASE("test scatter") { out = scatter(in, inds, updates, 0); CHECK(array_equal(out, reshape(arange(16, float32), {4, 4})).item()); - // Irregular strided index and reduce collison test + // Irregular strided index and reduce collision test in = zeros({10}, float32); inds = broadcast_to(array(3), {10}); updates = ones({10, 1}, float32); From 272f6952ba187aa0b11fea63cd32f4c90603c1f1 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:29:05 -0500 Subject: [PATCH 18/67] spelling: combinations Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/tests/test_ops.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tests/test_ops.py b/python/tests/test_ops.py index 89edfdefa1..782249b56d 100644 --- a/python/tests/test_ops.py +++ b/python/tests/test_ops.py @@ -88,7 +88,7 @@ def test_scalar_inputs(self): self.assertEqual(a.dtype, mx.float32) self.assertEqual(a.item(), 3.0) - # Check comibinations with mlx arrays + # Check combinations with mlx arrays a = mx.add(mx.array(True), False) self.assertEqual(a.dtype, mx.bool_) self.assertEqual(a.item(), True) From d26e758aa0b142ee3db6f2269fcf8ce7459988aa Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:29:13 -0500 Subject: [PATCH 19/67] spelling: committing Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 1387fec32b..4157a34f74 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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. From 2a5a690b92a64d1811c12abeeadbbcbebd0693ce Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:29:23 -0500 Subject: [PATCH 20/67] spelling: computation Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- examples/cpp/tutorial.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/cpp/tutorial.cpp b/examples/cpp/tutorial.cpp index 5dc0e0472f..091dfab2d7 100644 --- a/examples/cpp/tutorial.cpp +++ b/examples/cpp/tutorial.cpp @@ -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 From 956cac58fcc3b027b90694f1a724f92f84ac3e76 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:29:32 -0500 Subject: [PATCH 21/67] spelling: consider Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/fft.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/fft.cpp b/mlx/fft.cpp index 6cb33048d0..96d0424ab1 100644 --- a/mlx/fft.cpp +++ b/mlx/fft.cpp @@ -45,7 +45,7 @@ array fft_impl( throw std::invalid_argument(msg.str()); } - // In the following shape manipulations there are three cases to consdier: + // In the following shape manipulations there are three cases to consider: // 1. In a complex to complex transform (fftn / ifftn) the output // and input shapes are the same. // 2. In a real to complex transform (rfftn) n specifies the input dims From 0f34dc11afb3ff6823090921e1eca1d8e318c1a7 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:29:42 -0500 Subject: [PATCH 22/67] spelling: constructing Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- examples/extensions/axpby/axpby.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 4157a34f74..f660761fe4 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -305,7 +305,7 @@ if we encounter an unexpected type. /** Fall back implementation for evaluation on CPU */ void Axpby::eval(const std::vector& 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]; diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index c0834d4412..f13b8352bf 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -105,7 +105,7 @@ void axpby_impl( /** Fall back implementation for evaluation on CPU */ void Axpby::eval(const std::vector& 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]; From 30e9fdcf74b1bf58d932b0aa8020474027ded8fd Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:29:57 -0500 Subject: [PATCH 23/67] spelling: conversions Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/complex.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/complex.h b/mlx/backend/metal/kernels/complex.h index c9fedb7977..ac966a2937 100644 --- a/mlx/backend/metal/kernels/complex.h +++ b/mlx/backend/metal/kernels/complex.h @@ -45,7 +45,7 @@ struct complex64_t { typename = typename enable_if>::type> constexpr complex64_t(T x) constant : real(x), imag(0) {} - // Converstions from complex64_t + // Conversions from complex64_t template < typename T, typename = typename enable_if>::type> From e34205b9cc937cb92cd3c78aaf7d2e65f3606425 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:51:27 -0500 Subject: [PATCH 24/67] spelling: correctly Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- tests/creations_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/creations_tests.cpp b/tests/creations_tests.cpp index edb40a9fe1..ea28638afc 100644 --- a/tests/creations_tests.cpp +++ b/tests/creations_tests.cpp @@ -7,7 +7,7 @@ using namespace mlx::core; TEST_CASE("test arange") { - // Check type is inferred correclty + // Check type is inferred correctly { auto x = arange(10); CHECK_EQ(x.dtype(), int32); From b647328f9e80c6f2b4e833aea3fddec005f2c436 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:30:06 -0500 Subject: [PATCH 25/67] spelling: corresponding Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/array.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/array.cpp b/python/src/array.cpp index f8a1a27cd6..e2ae2fa47c 100644 --- a/python/src/array.cpp +++ b/python/src/array.cpp @@ -559,7 +559,7 @@ void init_array(py::module_& m) { If the array has more than one dimension then the result is a nested list of lists. - The value type of the list correpsonding to the last dimension is either + The value type of the list corresponding to the last dimension is either ``bool``, ``int`` or ``float`` depending on the ``dtype`` of the array. )pbdoc") .def("__array__", &mlx_array_to_np) From 8cae2e18f43189c61a1e0945341ba392e16199a1 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:30:49 -0500 Subject: [PATCH 26/67] spelling: declaration Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- examples/extensions/axpby/axpby.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index f660761fe4..114cafcabe 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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(); diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index f13b8352bf..3914ebeccd 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -254,7 +254,7 @@ void Axpby::eval_gpu(const std::vector& 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(); From c2498164889cb5eec311993fcc1415dfb6a3a7e3 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:30:59 -0500 Subject: [PATCH 27/67] spelling: default Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/random.h | 2 +- python/mlx/optimizers.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/random.h b/mlx/random.h index e684464bcb..360bdbdb11 100644 --- a/mlx/random.h +++ b/mlx/random.h @@ -16,7 +16,7 @@ class KeySequence { void seed(uint64_t seed); array next(); - // static defualt + // static default static KeySequence& default_() { static KeySequence ks(0); return ks; diff --git a/python/mlx/optimizers.py b/python/mlx/optimizers.py index 17a16c459c..601d87b03e 100644 --- a/python/mlx/optimizers.py +++ b/python/mlx/optimizers.py @@ -253,7 +253,7 @@ class AdaDelta(Optimizer): rho (float, optional): The coefficient :math:`\rho` used for computing a running average of squared gradients. Default: ``0.9`` eps (float, optional): The term :math:`\epsilon` added to the denominator to improve - numerical stability. Ddefault: `1e-8` + numerical stability. Default: `1e-8` """ def __init__(self, learning_rate: float, rho: float = 0.9, eps: float = 1e-6): From 61aa86623c6e16088eb3815cc378037c5be3db9d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:31:06 -0500 Subject: [PATCH 28/67] spelling: dependency Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- cmake/extension.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake/extension.cmake b/cmake/extension.cmake index 383656d375..aa2243304f 100644 --- a/cmake/extension.cmake +++ b/cmake/extension.cmake @@ -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 From a14e19704f483b74c42a19eeefdcf16ed7fbcb04 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:32:07 -0500 Subject: [PATCH 29/67] spelling: destination Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/ops.cpp b/python/src/ops.cpp index ea44868d67..3ea35d6daa 100644 --- a/python/src/ops.cpp +++ b/python/src/ops.cpp @@ -1742,7 +1742,7 @@ void init_ops(py::module_& m) { "a"_a, py::pos_only(), "source"_a, - "destiantion"_a, + "destination"_a, py::kw_only(), "stream"_a = none, R"pbdoc( From 1c9eea2272e1df91dc533d980d6e8a1184392daa Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:31:57 -0500 Subject: [PATCH 30/67] spelling: destructor Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/array.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/array.cpp b/mlx/array.cpp index 0a7b52a946..a70cb43a0f 100644 --- a/mlx/array.cpp +++ b/mlx/array.cpp @@ -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 { From f4fcfb4e7ce36fa1c23e5735e916d1ab3b882262 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:32:18 -0500 Subject: [PATCH 31/67] spelling: dimensions Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/matmul.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/matmul.cpp b/mlx/backend/metal/matmul.cpp index 864181da90..0bce599d3a 100644 --- a/mlx/backend/metal/matmul.cpp +++ b/mlx/backend/metal/matmul.cpp @@ -61,7 +61,7 @@ inline void mps_matmul( // 2. Only one of a or b has batch_size_out matrices worth of data and // the other has matrix worth of data - // The matrix dimsenisons of a and b are sure to be regularly strided + // The matrix dimensions of a and b are sure to be regularly strided if (batch_size_out > 1) { // No broadcasting defaults auto batch_size_a = a.data_size() / (M * K); From 07cb2f70c797a5330fe4c60d63d62cacc39abcfb Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:52:01 -0500 Subject: [PATCH 32/67] spelling: divided Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 2 +- examples/extensions/axpby/axpby.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 114cafcabe..683014a3f6 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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); } diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index 3914ebeccd..07720c3710 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -287,7 +287,7 @@ void Axpby::eval_gpu(const std::vector& 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); } From f242b4249334716f0ff09bdfed2c3673f9c29a09 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:34:38 -0500 Subject: [PATCH 33/67] spelling: element-wise Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 6 +++--- examples/extensions/axpby/axpby.cpp | 4 ++-- examples/extensions/axpby/axpby.h | 2 +- examples/extensions/bindings.cpp | 2 +- tests/autograd_tests.cpp | 4 ++-- 5 files changed, 9 insertions(+), 9 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 683014a3f6..90aca6b07e 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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 @@ -284,7 +284,7 @@ pointwise. This is captured in the templated function :meth:`axpby_impl`. T alpha = static_cast(alpha_); T beta = static_cast(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()); @@ -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`` diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index 07720c3710..5d033c70c5 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -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 @@ -91,7 +91,7 @@ void axpby_impl( T alpha = static_cast(alpha_); T beta = static_cast(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()); diff --git a/examples/extensions/axpby/axpby.h b/examples/extensions/axpby/axpby.h index 9ff6af0b14..3f470fd506 100644 --- a/examples/extensions/axpby/axpby.h +++ b/examples/extensions/axpby/axpby.h @@ -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 diff --git a/examples/extensions/bindings.cpp b/examples/extensions/bindings.cpp index 661ddcbaf2..d05e6b636c 100644 --- a/examples/extensions/bindings.cpp +++ b/examples/extensions/bindings.cpp @@ -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`` diff --git a/tests/autograd_tests.cpp b/tests/autograd_tests.cpp index 85dad9160c..a7b7e7fca6 100644 --- a/tests/autograd_tests.cpp +++ b/tests/autograd_tests.cpp @@ -989,7 +989,7 @@ TEST_CASE("test as_strided grads") { } TEST_CASE("test jvp from vjp") { - // Unary elementwise ops + // Unary element-wise ops { auto x = random::uniform({5, 10}); eval(x); @@ -1022,7 +1022,7 @@ TEST_CASE("test jvp from vjp") { CHECK(compute_derivs(mlx::core::rsqrt)); } - // Binary elementwise ops + // Binary element-wise ops { auto x = random::uniform({5, 10}); auto y = random::uniform({5, 10}); From 3cce74a8119121490cb90119e3e51a483adc07b5 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:34:44 -0500 Subject: [PATCH 34/67] spelling: elements Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/array.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/array.cpp b/python/src/array.cpp index e2ae2fa47c..1c6f724f44 100644 --- a/python/src/array.cpp +++ b/python/src/array.cpp @@ -507,7 +507,7 @@ void init_array(py::module_& m) { array_class .def_property_readonly( - "size", &array::size, R"pbdoc(Number of elments in the array.)pbdoc") + "size", &array::size, R"pbdoc(Number of elements in the array.)pbdoc") .def_property_readonly( "ndim", &array::ndim, R"pbdoc(The array's dimension.)pbdoc") .def_property_readonly( From 586ac7d6563b63f4bccf9c08460afd779217b3e2 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:52:24 -0500 Subject: [PATCH 35/67] spelling: endianness Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/common/load.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/mlx/backend/common/load.cpp b/mlx/backend/common/load.cpp index 6cf8ffe534..91f4cee62d 100644 --- a/mlx/backend/common/load.cpp +++ b/mlx/backend/common/load.cpp @@ -13,7 +13,7 @@ namespace mlx::core { namespace { template -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]; }; @@ -39,13 +39,13 @@ void Load::eval(const std::vector& inputs, array& out) { if (swap_endianness_) { switch (out.itemsize()) { case 2: - swap_endianess<2>(out.data(), out.data_size()); + swap_endianness<2>(out.data(), out.data_size()); break; case 4: - swap_endianess<4>(out.data(), out.data_size()); + swap_endianness<4>(out.data(), out.data_size()); break; case 8: - swap_endianess<8>(out.data(), out.data_size()); + swap_endianness<8>(out.data(), out.data_size()); break; } } From dc0661be5a4a3df4c5557abdf3692f024d28b3f9 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:35:08 -0500 Subject: [PATCH 36/67] spelling: equivalent Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/ops.cpp b/python/src/ops.cpp index 3ea35d6daa..db27a14a1a 100644 --- a/python/src/ops.cpp +++ b/python/src/ops.cpp @@ -1263,7 +1263,7 @@ void init_ops(py::module_& m) { If the axis is not specified the array is treated as a flattened 1-D array prior to performing the take. - As an example, if the ``axis=1`` this is equialent to ``a[:, indices, ...]``. + As an example, if the ``axis=1`` this is equivalent to ``a[:, indices, ...]``. Args: a (array): Input array. From d1db9479aab919f6516d4aa87e6d66f1194109e0 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:35:22 -0500 Subject: [PATCH 37/67] spelling: explicitly Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/utils.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/utils.h b/mlx/backend/metal/utils.h index 6fa08e42a1..3788508022 100644 --- a/mlx/backend/metal/utils.h +++ b/mlx/backend/metal/utils.h @@ -18,7 +18,7 @@ void set_array_buffer( auto offset = a.data() - static_cast(const_cast(a_buf)->contents()); enc->setBuffer(a_buf, offset, idx); - // MTL::Resource usage through argument buffer needs to be explicity + // MTL::Resource usage through argument buffer needs to be explicitly // flagged to enable hazard tracking compute_encoder->useResource(a_buf, MTL::ResourceUsageRead); } From 64baf0516cfe4abc0e19754f24275e43676a4a2a Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:23:29 -0500 Subject: [PATCH 38/67] spelling: github Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/README.md b/docs/README.md index f197ecf432..01d41d697a 100644 --- a/docs/README.md +++ b/docs/README.md @@ -26,7 +26,7 @@ python -m http.server and point your browser to `http://localhost:`. -### Push to Github Pages +### Push to GitHub Pages Check-out the `gh-pages` branch (`git switch gh-pages`) and build the docs. Then force add the `build/html` directory: From e54c872bc3d4db24d7281bf0f0ba986f291e02bf Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:36:19 -0500 Subject: [PATCH 39/67] spelling: indices Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/gemm/gemm.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/gemm/gemm.h b/mlx/backend/metal/kernels/gemm/gemm.h index f551947dd0..729c7b49e2 100644 --- a/mlx/backend/metal/kernels/gemm/gemm.h +++ b/mlx/backend/metal/kernels/gemm/gemm.h @@ -93,7 +93,7 @@ struct BlockLoader { tmp_idx[j] = bj + j < src_tile_dim.x ? j : 0; } - // Read all valid indcies into tmp_val + // Read all valid indices into tmp_val #pragma clang loop unroll(full) for (short j = 0; j < vec_size; j++) { tmp_val[j] = src[i * src_ld + tmp_idx[j]]; From fe3c914d419172654becf2214be69f7c3c19590d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:36:38 -0500 Subject: [PATCH 40/67] spelling: irregularly Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- tests/ops_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/ops_tests.cpp b/tests/ops_tests.cpp index 044edf7fd3..f6443bc7e6 100644 --- a/tests/ops_tests.cpp +++ b/tests/ops_tests.cpp @@ -1750,7 +1750,7 @@ TEST_CASE("test scatter") { out = scatter_max(array(1), {}, array(2), std::vector{}); CHECK_EQ(out.item(), 2); - // Irregularaly strided updates test + // Irregularly strided updates test in = ones({3, 3}); updates = broadcast_to(array({0, 0, 0}), {1, 3, 3}); inds = array({0}); From 6c1cd874351438c3e18d6f3adfab821443e06667 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:37:20 -0500 Subject: [PATCH 41/67] spelling: memory Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/allocator.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/allocator.cpp b/mlx/backend/metal/allocator.cpp index a55690947d..af4dd2e36b 100644 --- a/mlx/backend/metal/allocator.cpp +++ b/mlx/backend/metal/allocator.cpp @@ -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}; } From 0112fca3714f98d1bb5f2e3cf97482b4a51c374c Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:55:12 -0500 Subject: [PATCH 42/67] spelling: metallib Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- cmake/extension.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake/extension.cmake b/cmake/extension.cmake index aa2243304f..ffb02ee418 100644 --- a/cmake/extension.cmake +++ b/cmake/extension.cmake @@ -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 From 07782096d82076796811fc5d9322aa40d0ddd9df Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:38:00 -0500 Subject: [PATCH 43/67] spelling: negative Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/tests/test_array.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tests/test_array.py b/python/tests/test_array.py index b6471cdbdb..9016f3ae49 100644 --- a/python/tests/test_array.py +++ b/python/tests/test_array.py @@ -952,7 +952,7 @@ def test_slice_negative_step(self): b_mx = a_mx[25:-50:-3] self.assertTrue(np.array_equal(b_np, b_mx)) - # Negatie slice and ascending bounds + # Negative slice and ascending bounds b_np = a_np[0:20:-3] b_mx = a_mx[0:20:-3] self.assertTrue(np.array_equal(b_np, b_mx)) From 0219a05d1e95777159dbb0ccf042bb6844cfc2f6 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:38:20 -0500 Subject: [PATCH 44/67] spelling: notable Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/index.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/src/index.rst b/docs/src/index.rst index 207238f374..9f0445a188 100644 --- a/docs/src/index.rst +++ b/docs/src/index.rst @@ -19,7 +19,7 @@ The main differences between MLX and NumPy are: The design of MLX is inspired by frameworks like `PyTorch `_, `Jax `_, and -`ArrayFire `_. A noteable difference from these +`ArrayFire `_. 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 From 0ddb83cac58052fd573d70aab8498ee7ef560908 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:38:38 -0500 Subject: [PATCH 45/67] spelling: optional Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/random.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/random.cpp b/python/src/random.cpp index 943370db56..f648a27142 100644 --- a/python/src/random.cpp +++ b/python/src/random.cpp @@ -212,7 +212,7 @@ void init_random(py::module_& parent_module) { upper (scalar or array): Upper bound of the domain. shape (list(int), optional): The shape of the output. Default is ``()``. - dtype (Dtype, optinoal): The data type of the output. + dtype (Dtype, optional): The data type of the output. Default is ``float32``. key (array, optional): A PRNG key. Default: None. From 116a5b3d89b5ab63e9a1236e995847f1fbc74b53 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:42:30 -0500 Subject: [PATCH 46/67] spelling: otherwise Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/gemm/conv.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/gemm/conv.h b/mlx/backend/metal/kernels/gemm/conv.h index 2c4a7074a6..abf6f39d08 100644 --- a/mlx/backend/metal/kernels/gemm/conv.h +++ b/mlx/backend/metal/kernels/gemm/conv.h @@ -105,7 +105,7 @@ struct Conv2DInputBlockLoader { } } - // Zero pad otherwize + // Zero pad otherwise else { #pragma clang loop unroll(full) for (short j = 0; j < vec_size; ++j) { From 8f39b8975980d597ab0a97a4615cc5563fc7244d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:42:44 -0500 Subject: [PATCH 47/67] spelling: overridden Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/primitives.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/primitives.h b/mlx/primitives.h index 747b26c109..f9508096fd 100644 --- a/mlx/primitives.h +++ b/mlx/primitives.h @@ -84,7 +84,7 @@ class Primitive { /** Print the primitive. */ virtual void print(std::ostream& os) = 0; - /** Equivalence check defaults to false unless overriden by the primitive */ + /** Equivalence check defaults to false unless overridden by the primitive */ virtual bool is_equivalent(const Primitive& other) const { return false; } From c18cb065b43d8499d8e44cbbaf73914563f868bf Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:42:55 -0500 Subject: [PATCH 48/67] spelling: partially Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/gemv.metal | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/backend/metal/kernels/gemv.metal b/mlx/backend/metal/kernels/gemv.metal index 165868401f..d85d72d9e5 100644 --- a/mlx/backend/metal/kernels/gemv.metal +++ b/mlx/backend/metal/kernels/gemv.metal @@ -42,7 +42,7 @@ struct GEMVKernel { // Edge case handling: // - The threadgroup with the largest tid will have blocks that exceed the matrix // * The blocks that start outside the matrix are never read (thread results remain zero) - // * The last thread that partialy overlaps with the matrix is shifted inwards + // * The last thread that partially overlaps with the matrix is shifted inwards // such that the thread block fits exactly in the matrix MLX_MTL_CONST short tgp_mem_size = BN * TN * 2; @@ -180,7 +180,7 @@ struct GEMVTKernel { // Edge case handling: // - The threadgroup with the largest tid will have blocks that exceed the matrix // * The blocks that start outside the matrix are never read (thread results remain zero) - // * The last thread that partialy overlaps with the matrix is shifted inwards + // * The last thread that partially overlaps with the matrix is shifted inwards // such that the thread block fits exactly in the matrix From 979ae32b983d9810627f0649011672753701952e Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:43:09 -0500 Subject: [PATCH 49/67] spelling: partition Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/sort.metal | 6 +++--- mlx/backend/metal/sort.cpp | 4 ++-- python/src/ops.cpp | 2 +- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/mlx/backend/metal/kernels/sort.metal b/mlx/backend/metal/kernels/sort.metal index 3aa54de3e9..50b1cfbb67 100644 --- a/mlx/backend/metal/kernels/sort.metal +++ b/mlx/backend/metal/kernels/sort.metal @@ -592,7 +592,7 @@ template < bool ARG_SORT, short BLOCK_THREADS, short N_PER_THREAD> -[[kernel, max_total_threads_per_threadgroup(BLOCK_THREADS)]] void mb_block_partiton( +[[kernel, max_total_threads_per_threadgroup(BLOCK_THREADS)]] void mb_block_partition( device idx_t* block_partitions [[buffer(0)]], const device val_t* dev_vals [[buffer(1)]], const device idx_t* dev_idxs [[buffer(2)]], @@ -777,8 +777,8 @@ template < const device size_t* nc_strides [[buffer(7)]], \ uint3 tid [[threadgroup_position_in_grid]], \ uint3 lid [[thread_position_in_threadgroup]]); \ - template [[host_name("mb_block_partiton_" #vtname "_" #itname "_bn" #bn "_tn" #tn)]] \ - [[kernel]] void mb_block_partiton( \ + template [[host_name("mb_block_partition_" #vtname "_" #itname "_bn" #bn "_tn" #tn)]] \ + [[kernel]] void mb_block_partition( \ device itype* block_partitions [[buffer(0)]], \ const device vtype* dev_vals [[buffer(1)]], \ const device itype* dev_idxs [[buffer(2)]], \ diff --git a/mlx/backend/metal/sort.cpp b/mlx/backend/metal/sort.cpp index befbf2d813..9eb9960e02 100644 --- a/mlx/backend/metal/sort.cpp +++ b/mlx/backend/metal/sort.cpp @@ -165,10 +165,10 @@ void multi_block_sort( dev_idxs_out = ping ? dev_idxs_0 : dev_idxs_1; ping = !ping; - // Do partiton + // Do partition { std::ostringstream kname; - kname << "mb_block_partiton_" << type_to_name(dev_vals_in) << "_" + kname << "mb_block_partition_" << type_to_name(dev_vals_in) << "_" << type_to_name(dev_idxs_in) << "_bn" << bn << "_tn" << tn; auto kernel = d.get_kernel(kname.str()); diff --git a/python/src/ops.cpp b/python/src/ops.cpp index db27a14a1a..cab5f28c6a 100644 --- a/python/src/ops.cpp +++ b/python/src/ops.cpp @@ -2253,7 +2253,7 @@ void init_ops(py::module_& m) { will be of elements less or equal to the element at the ``kth`` index and all indices after will be of elements greater or equal to the element at the ``kth`` index. - axis (int or None, optional): Optional axis to partiton over. + axis (int or None, optional): Optional axis to partition over. If ``None``, this partitions over the flattened array. If unspecified, it defaults to ``-1``. From f890871ed777625b474e4caae6acbbadfa28b0f0 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:55:57 -0500 Subject: [PATCH 50/67] spelling: perform Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/conv.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/backend/metal/conv.cpp b/mlx/backend/metal/conv.cpp index e25599cafd..3377939ba0 100644 --- a/mlx/backend/metal/conv.cpp +++ b/mlx/backend/metal/conv.cpp @@ -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 copies = {in_padded, in_strided}; mlx_matmul( s, @@ -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 copies = {in_padded, in_strided}; mlx_matmul( s, From 039ee912381ac03bd300ad0c779d0de90e5e7ec0 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:43:22 -0500 Subject: [PATCH 51/67] spelling: perturbations Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/tests/test_blas.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/tests/test_blas.py b/python/tests/test_blas.py index ffbbbb0a72..a59052936e 100644 --- a/python/tests/test_blas.py +++ b/python/tests/test_blas.py @@ -53,10 +53,10 @@ def test_matmul_unaligned(self): for dtype in self.dtypes: np_dtype = getattr(np, dtype) base_shapes = [4, 8, 16, 32, 64, 128] - pertubations = [-2, -1, 0, 1, 2] + perturbations = [-2, -1, 0, 1, 2] for dim in base_shapes: - for p in pertubations: + for p in perturbations: shape_a = (dim + p, dim + p) shape_b = (dim + p, dim + p) self.__gemm_test(shape_a, shape_b, np_dtype) From 94eeae4b28278383f62fb0eadc29ff590fe47796 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:43:30 -0500 Subject: [PATCH 52/67] spelling: positively Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/ops.cpp b/mlx/ops.cpp index ffed3942e3..1787cbd955 100644 --- a/mlx/ops.cpp +++ b/mlx/ops.cpp @@ -478,7 +478,7 @@ array slice( // If strides are negative, slice and then make a copy with axes flipped if (negatively_strided_axes.size() > 0) { - // First, take the slice of the positvely strided axes + // First, take the slice of the positively strided axes auto out = array( out_shape, a.dtype(), From 684e946b443b4ac747a27b585605122012bafde5 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:43:44 -0500 Subject: [PATCH 53/67] spelling: primitive Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 4 ++-- examples/extensions/axpby/axpby.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 90aca6b07e..4c31b020cb 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -602,7 +602,7 @@ us the following :meth:`Axpby::jvp` and :meth:`Axpby::vjp` implementations. const std::vector& 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 + // 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 @@ -642,7 +642,7 @@ own :class:`Primitive`. .. code-block:: C++ - /** Vectorize primitve along given axis */ + /** Vectorize primitive along given axis */ std::pair Axpby::vmap( const std::vector& inputs, const std::vector& axes) { diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index 5d033c70c5..dacdbb9ab6 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -312,7 +312,7 @@ array Axpby::jvp( const std::vector& 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 + // 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 @@ -345,7 +345,7 @@ std::vector Axpby::vjp( return vjps; } -/** Vectorize primitve along given axis */ +/** Vectorize primitive along given axis */ std::pair Axpby::vmap( const std::vector& inputs, const std::vector& axes) { From 61823aa54f39db12d4e14c0df2faae2c8ae5b793 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:44:11 -0500 Subject: [PATCH 54/67] spelling: repeat Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/ops.h | 2 +- python/src/ops.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/ops.h b/mlx/ops.h index a99465f3e6..c888c80cde 100644 --- a/mlx/ops.h +++ b/mlx/ops.h @@ -214,7 +214,7 @@ array concatenate(const std::vector& arrays, StreamOrDevice s = {}); array stack(const std::vector& arrays, int axis, StreamOrDevice s = {}); array stack(const std::vector& arrays, StreamOrDevice s = {}); -/** Repeate an array along an axis. */ +/** Repeat an array along an axis. */ array repeat(const array& arr, int repeats, int axis, StreamOrDevice s = {}); array repeat(const array& arr, int repeats, StreamOrDevice s = {}); diff --git a/python/src/ops.cpp b/python/src/ops.cpp index cab5f28c6a..5e3d0534ce 100644 --- a/python/src/ops.cpp +++ b/python/src/ops.cpp @@ -2426,7 +2426,7 @@ void init_ops(py::module_& m) { R"pbdoc( repeat(array: array, repeats: int, axis: Optional[int] = None, *, stream: Union[None, Stream, Device] = None) -> array - Repeate an array along a specified axis. + Repeat an array along a specified axis. Args: array (array): Input array. From e254286d95f0d3a9a11a43b892127d55917904f6 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:44:17 -0500 Subject: [PATCH 55/67] spelling: repeats Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/src/ops.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/src/ops.cpp b/python/src/ops.cpp index 5e3d0534ce..1f60c64448 100644 --- a/python/src/ops.cpp +++ b/python/src/ops.cpp @@ -2432,7 +2432,7 @@ void init_ops(py::module_& m) { array (array): Input array. repeats (int): The number of repetitions for each element. axis (int, optional): The axis in which to repeat the array along. If - unspecified it uses the flattened array of the input and repeates + unspecified it uses the flattened array of the input and repeats along axis 0. stream (Stream, optional): Stream or device. Defaults to ``None``. From 538f52343ac62d84d60e4579a64e7c5dd12a271d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:44:25 -0500 Subject: [PATCH 56/67] spelling: respect Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/transforms.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/transforms.h b/mlx/transforms.h index ff47638bfa..caf648163a 100644 --- a/mlx/transforms.h +++ b/mlx/transforms.h @@ -80,7 +80,7 @@ ValueAndGradFn value_and_grad( /** * Returns a function which computes the value and gradient of the input - * function with repsect to a single input array. + * function with respect to a single input array. **/ ValueAndGradFn inline value_and_grad( const std::function(const std::vector&)>& fun, @@ -132,7 +132,7 @@ std::function(const std::vector&)> inline grad( /** * Returns a function which computes the gradient of the input function with - * repsect to a single input array. + * respect to a single input array. * * The function being differentiated takes a vector of arrays and returns an * array. The optional `argnum` index specifies which the argument to compute From f48a35246f0d2e4e2e0760511a6cdec628e8dad3 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:44:33 -0500 Subject: [PATCH 57/67] spelling: respectively Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- python/mlx/nn/layers/normalization.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/mlx/nn/layers/normalization.py b/python/mlx/nn/layers/normalization.py index d5e1a1c6e7..9c77667e7a 100644 --- a/python/mlx/nn/layers/normalization.py +++ b/python/mlx/nn/layers/normalization.py @@ -198,7 +198,7 @@ class BatchNorm(Module): batch, ``C`` is the number of features or channels, and ``L`` is the sequence length. The output has the same shape as the input. For four-dimensional arrays, the shape is ``NHWC``, where ``H`` and ``W`` are - the height and width respecitvely. + the height and width respectively. For more information on Batch Normalization, see the original paper `Batch Normalization: Accelerating Deep Network Training by Reducing Internal From 69142efdefd4cf75728e320552d097d656ce39b4 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:44:43 -0500 Subject: [PATCH 58/67] spelling: result Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/gemm/conv.h | 2 +- mlx/backend/metal/kernels/gemm/gemm.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/backend/metal/kernels/gemm/conv.h b/mlx/backend/metal/kernels/gemm/conv.h index abf6f39d08..1db3ebac89 100644 --- a/mlx/backend/metal/kernels/gemm/conv.h +++ b/mlx/backend/metal/kernels/gemm/conv.h @@ -334,7 +334,7 @@ struct Conv2DBlockMMA { } simdgroup_barrier(mem_flags::mem_none); -// Multiply and accumulate into resulr simdgroup matrices +// Multiply and accumulate into result simdgroup matrices #pragma clang loop unroll(full) for (short i = 0; i < TM; i++) { #pragma clang loop unroll(full) diff --git a/mlx/backend/metal/kernels/gemm/gemm.h b/mlx/backend/metal/kernels/gemm/gemm.h index 729c7b49e2..0d87df8f86 100644 --- a/mlx/backend/metal/kernels/gemm/gemm.h +++ b/mlx/backend/metal/kernels/gemm/gemm.h @@ -241,7 +241,7 @@ struct BlockMMA { } simdgroup_barrier(mem_flags::mem_none); -// Multiply and accumulate into resulr simdgroup matrices +// Multiply and accumulate into result simdgroup matrices #pragma clang loop unroll(full) for (short i = 0; i < TM; i++) { #pragma clang loop unroll(full) From 6b0d08af507bca2016b6dac94cdf89868fac70d2 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:44:59 -0500 Subject: [PATCH 59/67] spelling: rounding Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/types/fp16.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/types/fp16.h b/mlx/types/fp16.h index 58e1bc088e..c174afebcb 100644 --- a/mlx/types/fp16.h +++ b/mlx/types/fp16.h @@ -68,7 +68,7 @@ struct _MLX_Float16 { inf_scale.u = uint32_t(0x77800000); zero_scale.u = uint32_t(0x08800000); - // Combine with magic and let addition do rouding + // Combine with magic and let addition do rounding magic_bits.u = x_expo_32; magic_bits.f += (std::abs(x) * inf_scale.f) * zero_scale.f; From 2fb5fe3a526e3c457f0631ce8772a8e1394ac41c Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:45:18 -0500 Subject: [PATCH 60/67] spelling: separate Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/reduce.metal | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/reduce.metal b/mlx/backend/metal/kernels/reduce.metal index 85ff41f44d..4182184c2a 100644 --- a/mlx/backend/metal/kernels/reduce.metal +++ b/mlx/backend/metal/kernels/reduce.metal @@ -65,7 +65,7 @@ template in += grid_size * N_READS; } - // Sepate case for the last set as we close the reduction size + // Separate case for the last set as we close the reduction size size_t curr_idx = (gid + r * (size_t)grid_size) * N_READS; if (curr_idx < in_size) { int max_reads = in_size - curr_idx; From 31b24c2289a230f6574e36783710ac8b58ed3226 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:45:29 -0500 Subject: [PATCH 61/67] spelling: skipping Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- tests/arg_reduce_tests.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tests/arg_reduce_tests.cpp b/tests/arg_reduce_tests.cpp index 7e3750e7b2..b571c8f61b 100644 --- a/tests/arg_reduce_tests.cpp +++ b/tests/arg_reduce_tests.cpp @@ -76,7 +76,7 @@ TEST_CASE("test arg reduce small") { {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0}); if (!metal::is_available()) { - INFO("Skiping arg reduction gpu tests"); + INFO("Skipping arg reduction gpu tests"); return; } @@ -106,7 +106,7 @@ TEST_CASE("test arg reduce small") { TEST_CASE("test arg reduce against cpu") { if (!metal::is_available()) { - INFO("Skiping arg reduction gpu tests"); + INFO("Skipping arg reduction gpu tests"); return; } @@ -148,7 +148,7 @@ void test_arg_reduce_small_bool( TEST_CASE("test arg reduce bool") { if (!metal::is_available()) { - INFO("Skiping arg reduction gpu tests"); + INFO("Skipping arg reduction gpu tests"); return; } auto x = array( @@ -201,7 +201,7 @@ TEST_CASE("test arg reduce irregular strides") { Device::cpu, x, ArgReduce::ArgMin, {4, 2}, 2, {0, 0, 1, 1, 1, 1, 2, 2}); if (!metal::is_available()) { - INFO("Skiping arg reduction gpu tests"); + INFO("Skipping arg reduction gpu tests"); return; } } From 3a543733ef57ecc306ac952a01fa71c837ba341d Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:45:48 -0500 Subject: [PATCH 62/67] spelling: structure Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 4c31b020cb..83b22d4e25 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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 @@ -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``. From 9557a8fa6f259363fb0fb33811f303088375292e Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:23:39 -0500 Subject: [PATCH 63/67] spelling: the Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- docs/src/dev/extensions.rst | 4 ++-- examples/extensions/axpby/axpby.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/docs/src/dev/extensions.rst b/docs/src/dev/extensions.rst index 83b22d4e25..0a134e7f5c 100644 --- a/docs/src/dev/extensions.rst +++ b/docs/src/dev/extensions.rst @@ -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. @@ -601,7 +601,7 @@ us the following :meth:`Axpby::jvp` and :meth:`Axpby::vjp` implementations. const std::vector& tangents, const std::vector& argnums) { // Forward mode diff that pushes along the tangents - // The jvp transform on the the primitive can built with ops + // 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 diff --git a/examples/extensions/axpby/axpby.cpp b/examples/extensions/axpby/axpby.cpp index dacdbb9ab6..6da2ff5918 100644 --- a/examples/extensions/axpby/axpby.cpp +++ b/examples/extensions/axpby/axpby.cpp @@ -311,7 +311,7 @@ array Axpby::jvp( const std::vector& tangents, const std::vector& argnums) { // Forward mode diff that pushes along the tangents - // The jvp transform on the the primitive can built with ops + // 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 From 072f7e0b8c94772be7809575db8e7735ae11a8a9 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:46:12 -0500 Subject: [PATCH 64/67] spelling: transpose Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- benchmarks/python/blas/bench_gemv.py | 4 ++-- mlx/ops.cpp | 2 +- python/tests/test_blas.py | 8 ++++---- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/benchmarks/python/blas/bench_gemv.py b/benchmarks/python/blas/bench_gemv.py index 5f491ffc8c..2b564a78af 100644 --- a/benchmarks/python/blas/bench_gemv.py +++ b/benchmarks/python/blas/bench_gemv.py @@ -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): np_dtype = getattr(np, dtype) mlx_gb_s = [] mlx_gflops = [] @@ -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 = [] diff --git a/mlx/ops.cpp b/mlx/ops.cpp index 1787cbd955..014707b384 100644 --- a/mlx/ops.cpp +++ b/mlx/ops.cpp @@ -517,7 +517,7 @@ array slice( // Gather moves the axis up, remainder needs to be squeezed out_reshape[i] = indices[i].size(); - // Gather moves the axis up, needs to be tranposed + // Gather moves the axis up, needs to be transposed t_axes[ax] = i; } diff --git a/python/tests/test_blas.py b/python/tests/test_blas.py index a59052936e..8a7d632fa8 100644 --- a/python/tests/test_blas.py +++ b/python/tests/test_blas.py @@ -81,12 +81,12 @@ def test_matmul_shapes(self): for B, M, N, K in shapes: - with self.subTest(tranpose="nn"): + with self.subTest(transpose="nn"): shape_a = (B, M, K) shape_b = (B, K, N) self.__gemm_test(shape_a, shape_b, np_dtype) - with self.subTest(tranpose="nt"): + with self.subTest(transpose="nt"): shape_a = (B, M, K) shape_b = (B, N, K) self.__gemm_test( @@ -97,7 +97,7 @@ def test_matmul_shapes(self): f_mx_b=lambda x: mx.transpose(x, (0, 2, 1)), ) - with self.subTest(tranpose="tn"): + with self.subTest(transpose="tn"): shape_a = (B, K, M) shape_b = (B, K, N) self.__gemm_test( @@ -108,7 +108,7 @@ def test_matmul_shapes(self): f_mx_a=lambda x: mx.transpose(x, (0, 2, 1)), ) - with self.subTest(tranpose="tt"): + with self.subTest(transpose="tt"): shape_a = (B, K, M) shape_b = (B, N, K) self.__gemm_test( From 59b779b745bf439074b1c9080f8d72ccb94c75ce Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:57:13 -0500 Subject: [PATCH 65/67] spelling: unnecessary Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- examples/extensions/axpby/axpby.h | 2 +- mlx/primitives.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/examples/extensions/axpby/axpby.h b/examples/extensions/axpby/axpby.h index 3f470fd506..2b85dadb26 100644 --- a/examples/extensions/axpby/axpby.h +++ b/examples/extensions/axpby/axpby.h @@ -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& inputs, array& out) override; diff --git a/mlx/primitives.h b/mlx/primitives.h index f9508096fd..deeb498fad 100644 --- a/mlx/primitives.h +++ b/mlx/primitives.h @@ -49,7 +49,7 @@ class 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. */ virtual void eval_cpu(const std::vector& inputs, array& out) = 0; From 381712a8af7aae9d6d787c820700dc54dd3de63f Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 22:57:19 -0500 Subject: [PATCH 66/67] spelling: unneeded Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/backend/metal/kernels/gemm/gemm.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/backend/metal/kernels/gemm/gemm.h b/mlx/backend/metal/kernels/gemm/gemm.h index 0d87df8f86..95d2e6497f 100644 --- a/mlx/backend/metal/kernels/gemm/gemm.h +++ b/mlx/backend/metal/kernels/gemm/gemm.h @@ -99,7 +99,7 @@ struct BlockLoader { tmp_val[j] = src[i * src_ld + tmp_idx[j]]; } - // Zero out uneeded values + // Zero out unneeded values #pragma clang loop unroll(full) for (short j = 0; j < vec_size; j++) { tmp_val[j] = bj + j < src_tile_dim.x ? tmp_val[j] : T(0); From c373f86b661ca9aeafe6add865c098e777473243 Mon Sep 17 00:00:00 2001 From: Josh Soref <2119212+jsoref@users.noreply.github.com> Date: Mon, 1 Jan 2024 23:09:56 -0500 Subject: [PATCH 67/67] spelling: unsupported Signed-off-by: Josh Soref <2119212+jsoref@users.noreply.github.com> --- mlx/io/load.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/mlx/io/load.cpp b/mlx/io/load.cpp index 856cf17a2c..74e0784f8b 100644 --- a/mlx/io/load.cpp +++ b/mlx/io/load.cpp @@ -155,7 +155,7 @@ array load(std::shared_ptr in_stream, StreamOrDevice s) { // Read and check version if (read_magic_and_ver[6] != 1 && read_magic_and_ver[6] != 2) { throw std::runtime_error( - "[load] Unsupport npy format version in " + in_stream->label()); + "[load] Unsupported npy format version in " + in_stream->label()); } // Read header len and header