Metal: expand llvm.powi into multiplies

[timesselens]

With an integer n, @fastmath x^n compiles to llvm.powi (Base.FastMath.pow_fast), which Metal cannot compile:

julia> f(y, x) = (i = thread_position_in_grid().x; @inbounds y[i] = @fastmath(x[i]^3); nothing)
julia> @metal threads=2 f(y, x)
ERROR: Compilation to native code failed; see below for details.
caused by: NSError: Compilation failed due to an interrupted connection: XPC_ERROR_CONNECTION_INTERRUPTED.

AIR has no integer power, and Apple's back-end can't select llvm.powi either, not even with a constant exponent. Apple's own front-end emits the intrinsic for MSL's __builtin_powif, and metal-tt then aborts with unable to legalize instruction: ... G_FPOWI. The intrinsic also reaches the back-end from IR that a Julia override can't intercept, such as Enzyme's derivative of powi. So this PR lowers it in GPUCompiler, next to the other intrinsics AIR lacks, and doesn't add a pow_fast override to Metal.jl.

The expansion uses exponentiation by squaring, as LLVM does for CPUs:

• A constant exponent is unrolled into multiplies, like SelectionDAG's ExpandPowI.
• Any other exponent calls an alwaysinline loop over the exponent's bits, like compiler-rt's __powisf2. lower_air! then inlines it.

A negative exponent takes the reciprocal, and x^0 = 1, also for NaN. Because the multiplies happen in the same order as on the CPU, the GPU results match @fastmath x^n there bit for bit, except that the GPU flushes subnormal results to zero. LLVM's own GPU back-ends that lack powi (DXIL, Vulkan SPIR-V, AMDGPU's GlobalISel) lower it to pow(x, sitofp(n)) instead. That isn't an option here, because MSL leaves pow undefined for negative bases.

Tests: FileCheck tests on code_native cover a constant exponent, a run-time exponent and a vector overload with an i64 exponent. They fail on main. A Metal.jl integration test that compares against the CPU will follow once this is released.

Related: JuliaGPU/CUDA.jl#3065 is the same failure on CUDA. JuliaGPU/CUDA.jl#3098 fixed it with pow_fast overrides that go through the float power (__nv_fast_powf), so IR-level powi isn't covered there either.

[maleadt]

LGTM, thanks! Slightly abbreviated and rebased. I'll open a counterpart PR on Metal.jl with execution tests once this has been merged and tagged.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 86.54%. Comparing base (8496909) to head (7de6cec).
⚠️ Report is 1 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #946      +/-   ##
==========================================
+ Coverage   86.41%   86.54%   +0.13%     
==========================================
  Files          29       29              
  Lines        5674     5729      +55     
==========================================
+ Hits         4903     4958      +55     
  Misses        771      771              

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.