Fix Metal sort of a view with a negative stride

[robertomeroni]

Proposed changes

Fixes #4225. #4226 addresses the same bug by copying negative-stride inputs to a contiguous buffer
before the sort; the review there asks for the root cause instead, which is what this does. On Metal,
mx.sort / mx.argsort / mx.partition / mx.topk return zeros for every segment but the first
when the input is a view with a negative stride on an axis that is not being sorted:

a = mx.array(np.arange(12, dtype=np.int32).reshape(3, 4))[::-1, :]
mx.sort(a, axis=-1)   # [[8, 9, 10, 11], [0, 0, 0, 0], [0, 0, 0, 0]]

block_sort_nc and mb_block_sort get each segment's base offset from elem_to_loc(tid.y, ...).
elem_to_loc deduces IdxT from its first argument, so tid.y makes it uint: a negative stride
wraps to ~4.29e9 instead of stepping backwards, the segment reads out of bounds, and Metal returns
zeros. Asking for int64_t matches what the CUDA kernels already do with int64_t(blockIdx.y).

block_sort has the same signedness issue at sort.h:283-284, left unpatched here: slicing and
as_strided clear flags().contiguous for a negative stride, so the only producer that reaches it is
an independent Split mis-flagging bug whose CPU path is wrong for the same input. Happy to add the
two lines if you would rather have all three sites consistent.

                 before                                after
sort      [[8,9,10,11],[0,0,0,0],[0,0,0,0]]   [[8,9,10,11],[4,5,6,7],[0,1,2,3]]
topk      [[10,11],[0,0],[0,0]]               [[10,11],[6,7],[2,3]]
partition [[8,9,10,11],[0,0,0,0],[0,0,0,0]]   [[8,9,10,11],[4,5,6,7],[0,1,2,3]]

Added to test_sort, covering int32 and float32 at both the single-block and multi-block sizes;
all four subtests fail on the unpatched kernel. On an M3 Pro python/tests is 810 passed / 21 skipped
and the C++ suite 266 cases / 3600 assertions.

Checklist

• I have read the CONTRIBUTING document
• I have run pre-commit run --all-files to format my code / installed pre-commit prior to committing changes
• I have added tests that prove my fix is effective or that my feature works
• I have updated the necessary documentation (if needed)

---

AI assistance was used in developing this change. I reviewed and tested every line.

[zcbenz]

The added tests pass without the change.

[robertomeroni]

@zcbenz they fail on the Metal stream. On the CPU stream they pass either way, because the CPU path was
already correct. This is a Metal kernel bug

[zcbenz]

Ah sorry didn't realize I was running a cpu only build, this looks good to me!

[robertomeroni]

no worries :)