Add math mode option for custom Metal kernels

[shubhxho]

Summary

Closes #3592.

This PR introduces an explicit math_mode option to mx.fast.metal_kernel:

mx.fast.metal_kernel(
    name="my_kernel",
    input_names=["x"],
    output_names=["y"],
    source=source,
    math_mode="safe",  # "safe", "relaxed", or "fast"
)

Custom Metal kernels now default to math_mode="safe" to preserve IEEE-compliant behavior. This is particularly important for operations involving special floating-point values, such as ensuring:

exp(-inf) == 0

Preserving this behavior is critical for masked softmax implementations used in causal and sliding-window attention, where masked logits are commonly represented as -inf.

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)

[zcbenz]

It seems that the test is not working, if I change {"math_mode": "safe"} to {"math_mode": "fast"} it still passes, I'm not sure if there is a good way to detect fast math in Metal though.