### Example Crash Log Header

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

This is an example of the header found in Metal compiler crash log files. It provides process name, version, OS build, and incident ID.

```text
~/Library/Logs/DiagnosticReports/MTLCompilerService-<timestamp>.ips
```

```text
~/Library/Logs/DiagnosticReports/air-nt-<timestamp>.ips
```

--------------------------------

### Install and Smoke Test Metal.jl

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/index.md

Installs the Metal.jl package and performs a basic check to ensure it is working correctly. No additional installs are required as it uses the system's Metal Shading Language compiler.

```julia
using Pkg
Pkg.add("Metal")

using Metal
Metal.versioninfo()
```

--------------------------------

### Full Metal Library Workflow: Disassemble, Edit, Assemble, Load

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

This example demonstrates a complete workflow: disassembling a library, editing the resulting LLVM IR, and then reassembling and loading it. It uses standard input/output for piping between commands.

```bash
./metallib-dis -S input.metallib -o kernel.ll
$EDITOR kernel.ll
./metallib-as kernel.ll -o - | ./metallib-load -
```

--------------------------------

### Run Isolated Test File

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/contributing.md

Isolate and run a specific test file after executing the setup script. This method allows for focused testing of individual files.

```bash
julia --project=test -L test/setup.jl test/metal.jl
```

--------------------------------

### Compiler Service Log Message

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Example log message from the Metal compiler service indicating the start of a compilation request.

```text
MTLCompilerService ... (MTLCompiler) Compilation BEGIN (ParentProcessName=julia_NNN) Build request: MTLBuildFunctions - pipeline
```

--------------------------------

### Add Metal.jl Package via Pkg API

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Install the Metal.jl package using the Pkg API.

```julia
julia> import Pkg; Pkg.add("Metal")
```

--------------------------------

### Add Metal.jl Package

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Install the Metal.jl package using the Julia package manager.

```julia
pkg> add Metal
```

--------------------------------

### Get macOS Version

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves the macOS version supported by Metal.jl.

```julia
macos_version
```

--------------------------------

### Set and Get GPU Device

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Allows setting and retrieving the current GPU device used by Metal.jl.

```julia
Metal.device!
```

```julia
Metal.device
```

--------------------------------

### Get Default Metal Device

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Retrieves the default Metal device available on the system.

```julia
julia> using Metal

julia> dev = device()
<AGXG13XDevice: 0x14c17f200>
    name = Apple M1 Pro
```

--------------------------------

### Get Darwin Version

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves the Darwin kernel version supported by Metal.jl.

```julia
darwin_version
```

--------------------------------

### Get Global Command Queue

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves the global command queue for Metal.jl operations.

```julia
Metal.global_queue
```

--------------------------------

### Metal Framework Failure Log Message

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Example log message from the client-side Metal framework indicating a compilation failure.

```text
julia ... (Metal) MTLCompiler: Compilation failed with XPC_ERROR_CONNECTION_INTERRUPTED on 1 try.
```

--------------------------------

### Get AIR Target Version

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves the target AIR version supported by Metal.jl.

```julia
Metal.air_target
```

--------------------------------

### Metal-TT Compilation Error

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Example of an error message from `metal-tt` indicating a compilation failure.

```text
air-tt: applegpu-nt command failed
```

--------------------------------

### Example Faulting Thread Backtrace

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

This snippet shows a typical backtrace from a Metal compiler crash log, highlighting the faulting thread's frames. The top frame indicates the specific compiler function where the crash occurred.

```text
AGCLLVMAirBuiltins::buildAtomic(llvm::Value**, llvm::StringRef)   [AGXCompilerCore]
AGCLLVMAirBuiltinReplacement<...>::doReplacement(...)             [AGXCompilerCore]
AGCLLVMAirBuiltins::replaceBuiltins()                             [AGXCompilerCore]
...
AIRNTEmitPipelineImageInternal(...)                              [AGXCompilerCore]
MTLCompilerObject::backendCompileModule(...)                     [MTLCompiler]
```

--------------------------------

### Get Metal Target Version

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves the target Metal version supported by Metal.jl.

```julia
Metal.metal_target
```

--------------------------------

### Get MetalLib Target Version

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves the target MetalLib version supported by Metal.jl.

```julia
Metal.metallib_target
```

--------------------------------

### Example Kernel Function

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/kernel.md

A simple kernel function that performs element-wise addition of two input vectors and stores the result in a third output vector. Each thread computes one element of the sum.

```julia
function vadd(a, b, c)
    i = thread_position_in_grid().x
    c[i] = a[i] + b[i]
    return
end
```

--------------------------------

### Main Compiler Entry Point

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/compiler.md

Use the `@metal` macro as the primary way to invoke the compiler.

```julia
using Metal

@metal
```

--------------------------------

### Assemble LLVM Module into Metal Library

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Use `metallib-as` to package an LLVM module (textual IR or bitcode) into a `.metallib` container. The input module must contain exactly one external function to serve as the entry point. This script also handles downgrading IR to Metal-compatible versions.

```bash
./metallib-as reduced.ll -o reduced.metallib
```

--------------------------------

### Load Metal Library onto GPU

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Use `metallib-load` to instantiate a Metal library on the GPU. This process builds compute pipeline states for each function, triggering the backend compilation to machine code. It's crucial for verifying library validity and identifying backend compilation issues.

```bash
./metallib-load -v reduced.metallib
```

--------------------------------

### Versions and Support

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Provides information about the supported versions of macOS, Metal, and related libraries, as well as target specifications.

```APIDOC
## Versions and Support

### Functions

- `macos_version`
- `darwin_version`
- `Metal.metal_support`
- `Metal.metallib_support`
- `Metal.air_support`
- `Metal.metal_target`
- `Metal.metallib_target`
- `Metal.air_target`

### Description

These functions provide details on the compatibility and targeting capabilities of the Metal framework within the Julia environment. They allow users to query the supported versions of macOS and Metal, as well as the specific Metal, Metal Library, and Air (Intermediate Representation) targets that can be utilized.
```

--------------------------------

### Create and Manage MtlArray

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/overview.md

Demonstrates the creation of an MtlArray for GPU memory management and essential operations like copying, filling, and reshaping. Automatic memory management is also shown.

```julia
a = MtlArray{Int}(undef, 1024)

# essential memory operations, like copying, filling, reshaping, ...
b = copy(a)
fill!(b, 0)
@test b == Metal.zeros(Int, 1024)

# automatic memory management
a = nothing
```

--------------------------------

### Launching a Kernel with Dynamic Thread Count

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/kernel.md

Launches the 'vadd' kernel, ensuring the number of threads launched matches the length of the input vectors. This is crucial for correct kernel execution.

```julia
len = prod(size(d_a))
@metal threads=len vadd(d_a, d_b, d_c)
```

--------------------------------

### Function-Style Entry Points for Code Inspection

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/compiler.md

Use function-style entry points such as `Metal.code_llvm` and `Metal.code_air` for inspecting generated code.

```julia
Metal.code_llvm
Metal.code_air
```

--------------------------------

### Create and Manipulate MtlArray

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Demonstrates creating a simple MtlArray and performing an element-wise addition operation on it.

```julia
julia> a = MtlArray([1])
1-element MtlVector{Int64, Metal.PrivateStorage}:
 1

julia> a .+ 1
1-element MtlVector{Int64, Metal.PrivateStorage}:
 2
```

--------------------------------

### Compile MSL to Human-Readable LLVM IR

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Use the 'xcrun metal' command-line tool to compile an MSL kernel into a human-readable LLVM IR (.ll) file. This is essential for inspecting intrinsic names, signatures, and metadata.

```bash
xcrun metal -S -emit-llvm dummy_kernel.metal
```

--------------------------------

### Display Metal.jl Version Info

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Display detailed version information for Metal.jl and its dependencies, including the detected GPU and its status.

```julia
julia> using Metal

julia> Metal.versioninfo()
macOS 26.5.0, Darwin 25.5.0

Toolchain:
- Julia: 1.12.6
- LLVM: 18.1.7

Julia packages:
- Metal.jl: 1.10.0
- GPUArrays: 11.5.4
- GPUCompiler: 1.13.2
- KernelAbstractions: 0.9.41
- ObjectiveC: 5.0.0
- LLVM: 9.8.2
- LLVMDowngrader_jll: 0.6.0+2

1 device:
- Apple M2 Max 30 GPU cores (64.000 KiB allocated)
```

--------------------------------

### Reflection Macros for Device Code

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/compiler.md

Inspect generated device code using macros like `@device_code_lowered`, `@device_code_typed`, `@device_code_warntype`, `@device_code_llvm`, `@device_code_air`, and `@device_code`.

```julia
@device_code_lowered
@device_code_typed
@device_code_warntype
@device_code_llvm
@device_code_air
@device_code
```

--------------------------------

### Augment GPU Traces with Signposts

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/profiling.md

Add custom intervals and events to GPU traces using signposts from ObjectiveC.jl for more detailed profiling. This allows marking specific code sections and points of interest within the trace.

```julia
using ObjectiveC, .OS

@signpost_interval "My Interval" begin
    # code to profile
    @signpost_event "My Event"
end
```

--------------------------------

### Capture GPU Frames with Metal.@capture

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/profiling.md

Enable Metal's frame capture feature by setting the METAL_CAPTURE_ENABLED environment variable to 1. Wrap code of interest in Metal.@capture to generate a detailed, replayable GPU operations trace that can be opened with Xcode.

```shell
$ METAL_CAPTURE_ENABLED=1 julia
```

```julia
using Metal

function vadd(a, b, c)
           i = thread_position_in_grid().x
           c[i] = a[i] + b[i]
           return
       end

a = MtlArray([1]); b = MtlArray([2]); c = similar(a);
... Metal GPU Frame Capture Enabled

Metal.@capture @metal threads=length(c) vadd(a, b, c);
... 
[ Info: GPU frame capture saved to julia_1.gputrace; open the resulting trace in Xcode
```

--------------------------------

### Check MetalLib Support

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Verifies if Metal.jl supports MetalLib files.

```julia
Metal.metallib_support
```

--------------------------------

### Run Metal.jl Test Suite

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/index.md

Executes the test suite for the Metal.jl package to verify its functionality. This is a good way to ensure everything is set up as expected.

```julia
using Pkg
Pkg.test("Metal")
```

--------------------------------

### Copy Data to/from GPU

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/array.md

Transfer data between host (CPU) and device (GPU) using MtlArray constructors or the `copyto!` function.

```julia
a = MtlArray([1,2])
```

```julia
b = Array(a)
```

```julia
copyto!(b, a)
```

--------------------------------

### List Available GPU Devices

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Retrieves a list of all available GPU devices that Metal.jl can use.

```julia
Metal.devices
```

--------------------------------

### Check Metal Library Support

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Verifies if Metal.jl supports Metal libraries.

```julia
Metal.metal_support
```

--------------------------------

### Define and Execute a Metal Kernel for Vector Addition

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Defines a Julia kernel for element-wise vector addition and executes it on the GPU using Metal.jl.

```julia
julia> function vadd(a, b, c)
           i = thread_position_in_grid().x
           c[i] = a[i] + b[i]
           return
       end
vadd (generic function with 1 method)

julia> a = MtlArray([1,1,1,1]); b = MtlArray([2,2,2,2]); c = similar(a);

julia> @metal threads=2 groups=2 vadd(a, b, c);

julia> Array(c)
4-element Vector{Int64}:
 3
 3
 3
 3
```

--------------------------------

### Minimal MSL Kernel for Feature Testing

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Write a minimal Metal Shading Language (MSL) kernel to isolate and test specific features or intrinsics. Ensure the kernel is not so minimal that the optimizer removes the relevant code.

```objective-c
#include <metal_stdlib>

using namespace metal;

kernel void dummy_kernel(device volatile atomic_float* out,
                         uint i [[thread_position_in_grid]])
{
    atomic_store_explicit(&out[i], 0.0f, memory_order_relaxed);
}
```

--------------------------------

### Printing Thread Index with @mtlprintf

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/kernel.md

Use @mtlprintf to print values from within a kernel. The output is logged and available after command buffer completion. Requires macOS 15+.

```julia
function gpu_add2_print!(y, x) 
    index = thread_position_in_grid_1d()
    @mtlprintf("thread %d", index)
    @inbounds y[i] += x[i]
    return nothing
end

A = Metal.ones(Float32, 8);
B = Metal.rand(Float32, 8);

@metal threads=length(A) gpu_add2_print!(A, B)
```

--------------------------------

### Printing

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Macros for printing and showing data from within Metal kernels.

```APIDOC
## Printing

- `@mtlprintf`
- `@mtlprint`
- `@mtlprintln`
- `@mtlshow`
```

--------------------------------

### Check AIR Support

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Verifies if Metal.jl supports AIR (Apple Intermediate Representation).

```julia
Metal.air_support
```

--------------------------------

### Profile GPU Kernel Execution with Metal.@profile

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/profiling.md

Use Metal.@profile to track GPU calls and usage statistics for profiling. The macro saves trace information to a temporary '.trace' folder, which can be opened with Xcode's Instruments app.

```julia
using Metal

function vadd(a, b, c)
           i = thread_position_in_grid().x
           c[i] = a[i] + b[i]
           return
       end
julia> a = MtlArray([1]); b = MtlArray([2]); c = similar(a);

Metal.@profile @metal threads=length(c) vadd(a, b, c);
... 
[ Info: System trace saved to julia_3.trace; open the resulting trace in Instruments
```

--------------------------------

### Synchronize GPU Operations

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Ensures all previously issued Metal commands have been completed.

```julia
synchronize
```

--------------------------------

### Compile Metal Library Offline

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Use `metal-tt` to compile a Metal library (`.metallib`) with a pipeline descriptor (`.mtlp-json`) into a GPU binary (`.gpubin`). This provides a scriptable compilation process with a clean exit code.

```bash
xcrun metal-tt -arch applegpu_g15s -o main.gpubin descriptor.mtlp-json input.metallib
```

--------------------------------

### GPU Printing Macros

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Provides macros for printing and debugging directly from Metal kernels. These include `@mtlprintf`, `@mtlprint`, `@mtlprintln`, and `@mtlshow` for formatted output and value display on the GPU.

```julia
@mtlprintf
@mtlprint
@mtlprintln
@mtlshow
```

--------------------------------

### Lower-Level Compiler API

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/compiler.md

Access lower-level APIs like `Metal.mtlconvert`, `Metal.mtlfunction` for detailed compiler kernel inspection.

```julia
using Metal

Metal.mtlconvert
Metal.mtlfunction
```

--------------------------------

### Handling Out-of-Bounds Kernel Exception

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/kernel.md

Demonstrates how an out-of-bounds access in a kernel triggers a KernelException on the host after synchronization. The detail of the exception message depends on the Julia debug level.

```julia
function kernel(a) 
    a[2] = 1f0   # out-of-bounds: `a` has length 1
    return
end

julia> @metal threads=1 kernel(Metal.zeros(Float32, 1));

julia> synchronize()
ERROR: KernelException: A BoundsError was thrown on device Apple M1: Out-of-bounds array access
For more details, run Julia with "-g2", or launch the kernel with "@metal debug_level=2".
```

--------------------------------

### Disassemble Metal Library to LLVM IR

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Use `metallib-dis` to unpack a `.metallib` container into textual LLVM IR. The `-S` flag ensures textual output, while omitting it yields bitcode.

```bash
./metallib-dis -S broken.metallib -o reduced.ll
```

--------------------------------

### Synchronization

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Provides functions for memory and thread synchronization within kernels.

```APIDOC
## Synchronization

- `MemoryFlags`
- `threadgroup_barrier()`
- `simdgroup_barrier()`
```

--------------------------------

### Run Native Translator Directly

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Execute the native translator (`applegpu-nt`) directly for debugging. This allows for obtaining honest exit statuses and crash logs, though symbol names may be stripped.

```bash
applegpu-nt -arch applegpu_g15s -platform_version macos <ver> <sdk> \
  -sysroot <sysroot> input.metallib -N descriptor.mtlp-json -o out
```

--------------------------------

### Synchronize Device Operations

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Ensures all previously issued commands on the current device have been completed.

```julia
device_synchronize
```

--------------------------------

### Device Array Primitive

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Introduces `MtlDeviceArray`, a lightweight array type for managing GPU data in a dense fashion. It serves as the device-counterpart to `MtlArray` and supports array interface functionalities for on-GPU use.

```julia
MtlDeviceArray
Metal.Const
```

--------------------------------

### Dump Compiled Kernel Artifacts with JULIA_METAL_DUMP_DIR

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/profiling.md

Set the JULIA_METAL_DUMP_DIR environment variable to capture LLVM IR (.ll), AIR (.air), and Metal library (.metallib) files for all compiled kernels. This is useful for debugging compilation failures on different machines.

```shell
$ JULIA_METAL_DUMP_DIR=/tmp/metal-dumps julia
```

```julia
using Metal

@metal threads=length(c) vadd(a, b, c);

readdir("/tmp/metal-dumps")
3-element Vector{String}:
 "jl_8kQ2Xv.air"
 "jl_8kQ2Xv.ll"
 "jl_8kQ2Xv.metallib"
```

--------------------------------

### Synchronization Primitives

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Details synchronization mechanisms for Metal kernels. `threadgroup_barrier` and `simdgroup_barrier` ensure correct execution order and data visibility within thread groups and SIMD groups, respectively.

```julia
MemoryFlags
threadgroup_barrier
simdgroup_barrier
```

--------------------------------

### Error: Invalid or Corrupt Metal Library during Loading

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

This error message signifies that `metallib-load` encountered an issue with the Metal library file itself, suggesting it's either improperly formatted, corrupted, or not a valid Metal library container.

```text
ERROR: Failed to load Metal library; the library is likely invalid or corrupt.
```

--------------------------------

### Custom RNG with Random Standard Library

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/array.md

Integrate with Julia's `Random` standard library for GPU random number generation. Use `Random.rand!` to fill existing arrays.

```julia
using Random
```

```julia
a = Random.rand(Metal.default_rng(), Float32, 1)
```

```julia
Random.rand!(Metal.default_rng(), a)
```

--------------------------------

### Indexing and Dimensions

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Provides functions for querying thread and threadgroup indices and positions within the grid and threadgroup structures.

```APIDOC
## Indexing and Dimensions

This section lists the package's public functionality that corresponds to special Metal functions for use in device code.

### Functions

- `thread_index_in_quadgroup()`
- `thread_index_in_simdgroup()`
- `thread_index_in_threadgroup()`
- `thread_position_in_grid()`
- `thread_position_in_threadgroup()`
- `threadgroup_position_in_grid()`
- `threads_per_grid()`
- `threads_per_simdgroup()`
- `threads_per_threadgroup()`
- `simdgroups_per_threadgroup()`
- `simdgroup_index_in_threadgroup()`
- `quadgroup_index_in_threadgroup()`
- `quadgroups_per_threadgroup()`
- `grid_size()`
- `grid_origin()`
- `thread_execution_width()`
```

--------------------------------

### Scripting Compiler Crash Reduction Loop

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

This script automates the process of testing reduced IR files. It removes old crash logs, assembles and loads the candidate IR, and then checks the new crash log for a specific function to confirm the crash is the same.

```shell
rm -f ~/Library/Logs/DiagnosticReports/MTLCompilerService*.ips
./metallib-as candidate.ll -o - | ./metallib-load -
grep -q buildAtomic ~/Library/Logs/DiagnosticReports/MTLCompilerService*.ips && echo "same crash"
```

--------------------------------

### Run Specific Testset

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/contributing.md

Execute a particular testset by passing its name to the runtests.jl script. This is useful for testing new functionality without running the entire test suite.

```bash
julia --project=test test/runtests.jl metal
```

--------------------------------

### Device Arrays

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Defines `MtlDeviceArray` for managing GPU data and `Metal.Const` for constant data access on the GPU.

```APIDOC
## Device Arrays

Metal.jl provides a primitive, lightweight array type to manage GPU data organized in an plain, dense fashion. This is the device-counterpart to the `MtlArray`, and implements (part of) the array interface as well as other functionality for use _on_ the GPU:

### Types

- `MtlDeviceArray`
- `Metal.Const`
```

--------------------------------

### Global State

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/essentials.md

Manages the global state of Metal.jl, including device selection, listing available devices, and synchronizing operations.

```APIDOC
## Global State

### Functions

- `Metal.device!`
- `Metal.devices`
- `Metal.device`
- `Metal.global_queue`
- `synchronize`
- `device_synchronize`

### Description

This section details the functions for managing the global state of Metal.jl. Users can set the active GPU device using `Metal.device!`, retrieve a list of all available devices with `Metal.devices`, or get the current default device with `Metal.device`. `Metal.global_queue` provides access to the default command queue. The `synchronize` and `device_synchronize` functions are crucial for ensuring that all GPU operations are completed before proceeding.
```

--------------------------------

### Perform Array Operations with MtlArray

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/overview.md

Illustrates using higher-order array operations like map, reduce, and broadcast with MtlArray for kernel-like computations without explicit kernel writing.

```julia
a = Metal.zeros(1024)
b = Metal.ones(1024)
a.^2 .+ sin.(b)
```

--------------------------------

### Show Log Messages for Metal.jl

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

Command to display info and debug log messages for the Metal compiler service and julia processes related to compilation. Use this to filter and view relevant logs from the command line.

```sh
log show --last 2m --info --debug \
  --predicate 'process == "MTLCompilerService" OR (process BEGINSWITH "julia" AND eventMessage CONTAINS "Compil")'
```

--------------------------------

### Access Device Name Property

Source: https://github.com/juliagpu/metal.jl/blob/main/README.md

Accesses the 'name' property of a Metal device object.

```julia
julia> dev.name
NSString("Apple M1 Pro")
```

--------------------------------

### Element-wise Operations with map

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/array.md

Perform element-wise operations on MtlArrays using `map` or `broadcast` (`.=`). This avoids manual kernel writing for simple transformations.

```julia
a = MtlArray{Float32}(undef, (1,2))
```

```julia
a .= 5
```

```julia
map(sin, a)
```

--------------------------------

### Error: Malformed LLVM IR during Assembly

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

This error message indicates that the LLVM IR provided to `metallib-as` is malformed and fails LLVM's parsing stage. It typically occurs with incorrect intrinsic signatures or syntax errors.

```text
ERROR: LoadError: LLVM error: error: expected top-level entity
```

--------------------------------

### Matrix and Vector Types

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/mps.md

Exposes the MPSMatrix and MPSVector types from Metal Performance Shaders.

```APIDOC
## Matrix and Vector Types

### Description
Provides access to the `MPSMatrix` and `MPSVector` types.

### Types
- `MPS.MPSMatrix`
- `MPS.MPSVector`
```

--------------------------------

### Construct MtlArray

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/array.md

Construct MtlArrays similarly to Julia's native Arrays, specifying element type and dimensions. Use `undef` for uninitialized arrays.

```julia
MtlArray{Int}(undef, 2)
```

```julia
MtlArray{Int}(undef, (1,2))
```

```julia
similar(ans)
```

--------------------------------

### Kernel Indexing Functions

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Provides access to thread and grid indexing information within Metal kernel execution. These functions are essential for parallel computation and data access patterns on the GPU.

```julia
thread_index_in_quadgroup
thread_index_in_simdgroup
thread_index_in_threadgroup
thread_position_in_grid
thread_position_in_threadgroup
threadgroup_position_in_grid
threadgroups_per_grid
threads_per_grid
threads_per_simdgroup
threads_per_threadgroup
simdgroups_per_threadgroup
simdgroup_index_in_threadgroup
quadgroup_index_in_threadgroup
quadgroups_per_threadgroup
grid_size
grid_origin
thread_execution_width
```

--------------------------------

### Matrix Arithmetic Operators

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/mps.md

Details the matrix arithmetic operators available for Metal Performance Shaders.

```APIDOC
## Matrix Arithmetic Operators

### Description
Provides matrix arithmetic operations for Metal Performance Shaders.

### Functions
- `MPS.matmul!`
- `MPS.matvecmul!`
- `MPS.topk`
- `MPS.topk!`
```

--------------------------------

### Storage Mode Check Functions

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/array.md

Utility functions to check the storage mode of an `MtlArray`.

```APIDOC
## Storage Mode Check Functions

### Description
Convenience functions to determine if an `MtlArray` is using a specific storage mode.

- `is_private(array)`: Checks if the array uses `PrivateStorage`.
- `is_shared(array)`: Checks if the array uses `SharedStorage`.
- `is_managed(array)`: Checks if the array uses `ManagedStorage`.
```

--------------------------------

### MtlArray Types

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/array.md

Documentation for the core Metal array types: `MtlArray`, `MtlVector`, `MtlMatrix`, and `MtlVecOrMat`.

```APIDOC
## MtlArray Types

### Description
Core types for representing arrays on Metal.

- `MtlArray`: The general Metal array type.
- `MtlVector`: Represents a Metal vector.
- `MtlMatrix`: Represents a Metal matrix.
- `MtlVecOrMat`: A type alias for either `MtlVector` or `MtlMatrix`.
```

--------------------------------

### Metal Storage Modes

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/array.md

Information about Metal API storage modes: `PrivateStorage`, `SharedStorage`, and `ManagedStorage`.

```APIDOC
## Metal Storage Modes

### Description
Details the different storage modes available for Metal resources, affecting how they can be accessed.

- `Metal.PrivateStorage`: Default storage mode for `MtlArray`s.
- `Metal.SharedStorage`: Resource can be accessed by both the CPU and the GPU.
- `Metal.ManagedStorage`: Resource can be accessed by both the CPU and the GPU, with Metal managing data transfers.

Refer to the [official Metal documentation](https://developer.apple.com/documentation/metal/resource_fundamentals/setting_resource_storage_modes) for more details.
```

--------------------------------

### Metal Pipeline Descriptor JSON

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/faq/hacking.md

A JSON file specifying compute pipelines, including the compute function name and threadgroup configuration. This is used by `metal-tt` to define the pipeline state.

```json
{
  "pipelines": {
    "compute_pipelines": [
      {
        "compute_function": "_Z1g",
        "threadgroup_size_is_multiple_of_thread_execution_width": true,
        "max_total_threads_per_threadgroup": 704
      }
    ]
  }
}
```

--------------------------------

### Dimensional Reduction with reduce

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/array.md

Utilize `reduce`, `mapreduce`, and `mapreducedim!` for dimensionality reduction operations on MtlArrays. Specify the reduction dimension with `dims`.

```julia
a = Metal.ones(2,3)
```

```julia
reduce(+, a)
```

```julia
mapreduce(sin, *, a; dims=2)
```

```julia
b = Metal.zeros(1)
```

```julia
Base.mapreducedim!(identity, +, b, a)
```

--------------------------------

### GPU Random Number Generation

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/usage/array.md

Generate random numbers on the GPU using Metal.jl's convenience functions like `Metal.rand` and `Metal.randn`. These utilize `Metal.default_rng()`.

```julia
Metal.rand(2)
```

```julia
Metal.randn(Float32, 2, 1)
```

--------------------------------

### mtl Function

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/array.md

The `mtl` function is used to transfer an array to the GPU. It automatically converts `Float64` arrays to `Float32` for compatibility.

```APIDOC
## mtl

### Description
Transfers an array to the GPU. Automatically converts `Float64` to `Float32`.

### Usage
```julia
mtl(array)
```
```

--------------------------------

### Shared Memory

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Introduces `MtlThreadGroupArray` for managing shared memory within a threadgroup.

```APIDOC
### Shared Memory

- `MtlThreadGroupArray`
```

--------------------------------

### Shared Memory Array

Source: https://github.com/juliagpu/metal.jl/blob/main/docs/src/api/kernel.md

Defines `MtlThreadGroupArray` for managing shared memory within a thread group on the GPU. This is crucial for efficient data sharing and communication between threads in the same group.

```julia
MtlThreadGroupArray
```

--------------------------------

### Iterate Undocumented MTLDataType Values

Source: https://github.com/juliagpu/metal.jl/blob/main/res/wrap/README.md

This code iterates through possible undocumented values for the MTLDataType enum using the MTLTypeInternal type. It prints the description and its corresponding integer value.

```julia
load_framework("Metal")

@objcwrapper immutable=false MTLType <: NSObject

@objcwrapper immutable=false MTLTypeInternal <: MTLType

@objcproperties MTLTypeInternal begin
    @autoproperty dataType::UInt64
    @autoproperty description::id{NSString}
end

function MTLTypeInternal(dataType::Integer)
    obj = MTLTypeInternal(@objc [MTLTypeInternal alloc]::id{MTLTypeInternal})
    finalizer(dealloc, obj)
    @objc [obj::id{MTLTypeInternal} initWithDataType:dataType::UInt64]::id{MTLTypeInternal}
    return obj
end

dealloc(obj::MTLTypeInternal) = @objc [obj::id{MTLTypeInternal} dealloc]::Cvoid

for i in 1:200
    typ = MTLTypeInternal(i)
    name = string(typ.description)
    if name != "Unknown"
        println("    $name = $i")
    end
end
```

=== COMPLETE CONTENT === This response contains all available snippets from this library. No additional content exists. Do not make further requests.