What Metal Shading Language Is

It's C++11 with attribute brackets — and one architectural twist

Metal Shading Language (MSL) is what you write inside .metal files. To a CUDA developer it looks immediately familiar: function qualifiers mark GPU entry points, attribute brackets expose SIMT indices, and address-space keywords name memory regions. To a graphics developer it looks like a leaner, more typed cousin of GLSL.

Three extensions matter most:

kernel function qualifier — the GPU entry point launched by CPU. CUDA's __global__ twin.
Attribute brackets — [[thread_position_in_grid]], [[threadgroup_position_in_grid]], [[thread_index_in_threadgroup]]. Metal's answer to threadIdx + blockIdx.
Address-space keywords — device, threadgroup, constant, thread — explicitly mark which memory region a pointer lives in. CUDA infers this from __shared__/__constant__; Metal makes it part of the type.

Beyond that, you get standard C++11 + Apple extensions: packed vectors (float4, half3), threadgroup barriers, math built-ins. The whole MSL spec is one PDF — readable in an afternoon.

The architectural twist that hides in plain sight: on Apple Silicon, device memory and CPU-visible memory are the same physical RAM. We'll dedicate a whole lesson to that next.

Code

Metal kernel — same shape as CUDA, different vocabulary·metal

#include <metal_stdlib>
using namespace metal;

// 'kernel' = CUDA's __global__
// device   = pointer in GPU-visible memory
// constant = read-only uniform passed once per dispatch
// [[thread_position_in_grid]] = threadIdx + blockIdx fused
kernel void vector_add(
    const device float *A     [[buffer(0)]],
    const device float *B     [[buffer(1)]],
    device       float *C     [[buffer(2)]],
    constant     uint  &N     [[buffer(3)]],
    uint gid                  [[thread_position_in_grid]])
{
    if (gid < N) C[gid] = A[gid] + B[gid];
}

Common MSL built-ins worth memorizing·metal

// Per-thread identity
uint gid     [[thread_position_in_grid]];        // global thread index
uint tid     [[thread_index_in_threadgroup]];    // local index inside threadgroup
uint tg_idx  [[threadgroup_position_in_grid]];   // which threadgroup am I in
uint tg_size [[threads_per_threadgroup]];        // size of my threadgroup

// SIMD-group (Apple's warp, 32 threads)
uint simd_lane [[thread_index_in_simdgroup]];
uint simd_id   [[simdgroup_index_in_threadgroup]];

It's C++11 with attribute brackets — and one architectural twist

Code

External links

Exercise

Progress

Comments 0